0

我的主程序中有一个大型 char 数组，我将其分块复制到设备内存中。我在我的程序中运行了大约 500,000 个线程，每个线程访问 2000 个字符。所以我使用代码一次传输 500,000 * 2000 = 1GB 字节

err = cudaMemcpy (dev_database, adjusted_database[k], JOBS * 2000 * sizeof(char), cudaMemcpyHostToDevice);
if(err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); exit(EXIT_FAILURE); }

在我的内核中，我还定义了三个共享数组

//__shared__ char dev_query[200];
__shared__ float dev_scores[200*5];
__shared__ int dev_index[26];

并用

if(threadIdx.x == 0) { 
  //for(i = 0; i < 200; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 200 * 5; i++){ dev_scores[i] = dev_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_index[i] = dev_index_constant[i]; }
}
__syncthreads();

如果我用两行注释运行我的程序，我的内核返回奇怪的值，当我复制 char 数组的第二块时，我得到错误

CUDA 错误：未指定的启动失败

如果我取消注释上面代码中的行，一切正常。如果我复制较小的数组块，例如 100MB 而不是 1GB，它可以正常工作，直到我得到与上述相同的错误的第 6 个块。

这是非常奇怪的行为，我想了解为什么会这样。某处是否存在导致此问题的错误？很难确定它，因为如果我传输一小块（例如 100MB）而忽略其他块，程序可以正常工作。如果我取消注释与共享变量相关的行或将共享变量更改为常量，它也可以正常工作。任何帮助将不胜感激。谢谢！

编辑：这是我的内核。总而言之，我通过比较它们的第 i 个字符来计算两个字符串的相似度分数，这些字符在 0 和它们的长度之间。下面的代码将产生上述错误，除非您立即取消注释 if(threadIdx.x == 0) {. 或者，如果您将下面的共享数组替换为常量数组，那么它也可以正常工作。

__global__ void assign7(int jobs_todo, char* database, float* results, int flag) {
unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;

if(id < jobs_todo) {
__shared__ char dev_query[200];
__shared__ float dev_pos_specific_scores[200*5];
__shared__ int dev_subst_index[26];

int j_, i, p, stop, k; //stop2;
float score=0, max=0;
char ch; //ch1, ch2;

if(threadIdx.x == 0) {
//for(i = 0; i < 51; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 5 * 200; i++){ dev_pos_specific_scores[i] = dev_pos_specific_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_subst_index[i] = dev_subst_index_constant[i]; }
}
__syncthreads(); 

for(i = 1; i <= 2000 - 51; i += 1){
  p = jobs_todo*(i-1);
  score = 0;
  stop = 51/1; stop = stop*1;
  for(j_ = 1; j_ <= stop; j_ += 1){
    k = (j_-1)*5;
    ch = database[p + id];
    score += dev_pos_specific_scores[k + dev_subst_index[ch - 'A']];
    if(score < 0) score = 0;
    if(score > max) max = score;                                      
    p += jobs_todo;
  }
}
results[id] = max;
}
}
4

2 回答 2

2

下面的部分在k没有初始化的情况下使用：

ch = database[p + id];
score += dev_scores[k + dev_index[ch - 'A']];

这无关紧要，但这部分：

if(threadIdx.x == 0) { 
  //for(i = 0; i < 200; i++){ dev_query[i] = dev_query_constant[i]; }
  for(i = 0; i < 200 * 5; i++){ dev_scores[i] = dev_scores_constant[i]; }
  for(i = 0; i < 26; i++){ dev_index[i] = dev_index_constant[i]; }
}

可以更改为：

if(threadIdx.x < 200) {
  // dev_query[i] = dev_query_constant[i];
}

if(threadIdx.x < 200 * 5) { // or iterate whole block 5 times..
  dev_scores[i] = dev_scores_constant[i];
}
...
于 2012-02-27T22:06:42.453 回答
0

数据中有一些字符导致dev_index[ch-'A']返回 -1。当 k = 0 时，这使得索引为dev_scores-1。我相信这是我的代码中内存错误的根源。我评论了所有内容并逐渐取消了评论。它现在工作正常。感谢@talonmies、@harrism 和@perreal 的评论！

于 2012-03-01T14:05:31.103 回答