【发布时间】：2013-02-17 01:41:37
【问题描述】：
我编写了一个 CUDA 代码，它基本上为我汇总了一个数组。数组大小N应该是2的幂，即2^x。但是，我的代码无法正常工作。例如，如果输出为150177410，我的代码输出150177408。在过去的5 小时内，我一直在尝试调试它。任何帮助将不胜感激。下面是代码：
//only for array size of 2^x and TPB of 2^y as godata is = num of blocks. But num of blocks 2^sth if previous satisfied
//Works for arbitrary size array of type 2^x
#include<stdio.h>
__global__ void computeAddShared(int *in , int *out, int sizeInput){
//not made parameters gidata and godata to emphasize that parameters get copy of address and are different from pointers in host code
extern __shared__ float temp[];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int ltid = threadIdx.x;
temp[ltid] = 0;
while(tid < sizeInput){
temp[ltid] += in[tid];
tid+=gridDim.x * blockDim.x; // to handle array of any size
}
__syncthreads();
int offset = 1;
while(offset < blockDim.x){
if(ltid % (offset * 2) == 0){
temp[ltid] = temp[ltid] + temp[ltid + offset];
}
__syncthreads();
offset*=2;
}
if(ltid == 0){
out[blockIdx.x] = temp[0];
}
}
int main(){
int N = 8192;//should be 2^sth
int size = N;
int *a = (int*)malloc(N * sizeof(int));
/* TO create random number
FILE *f;
f = fopen("invertedList.txt" , "w");
a[0] = 1 + (rand() % 8);
fprintf(f, "%d,",a[0]);
for( int i = 1 ; i< N; i++){
a[i] = a[i-1] + (rand() % 8) + 1;
fprintf(f, "%d,",a[i]);
}
fclose(f);
return 0;*/
FILE *f;
f = fopen("invertedList.txt","r");
if( f == NULL){
printf("File not found\n");
system("pause");
exit(1);
}
int count = 0 ;
long actualSum = 0;
for( int i =0 ; i < N ; i++){
fscanf(f, "%d,", &a[count]);
actualSum+=a[count];
count++;
}
fclose(f);
printf("The actual sum is %d\n",actualSum);
int* gidata;
int* godata;
cudaMalloc((void**)&gidata, N* sizeof(int));
cudaMemcpy(gidata,a, size * sizeof(int), cudaMemcpyHostToDevice);
int TPB = 256;
int blocks = 10; //to get things kicked off
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
while(blocks != 1 ){
if(size < TPB){
TPB = size; // size is 2^sth
}
blocks = (size+ TPB -1 ) / TPB;
cudaMalloc((void**)&godata, blocks * sizeof(int));
computeAddShared<<<blocks, TPB,TPB*sizeof(int)>>>(gidata, godata,size);
//cudaFree(gidata);
gidata = godata;
size = blocks;
}
//printf("The error by cuda is %s",cudaGetErrorString(cudaGetLastError()));
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime , start, stop);
printf("time is %f ms", elapsedTime);
int *output = (int*)malloc(sizeof(int));
cudaMemcpy(output, gidata,size * sizeof(int), cudaMemcpyDeviceToHost);
//Cant free either earlier as both point to same location
cudaError_t chk = cudaFree(godata);
if(chk!=0){
printf("First chk also printed error. Maybe error in my logic\n");
}
printf("The error by threadsyn is %s", cudaGetErrorString(cudaGetLastError()));
printf("The sum of the array is %d\n", output[0]);
getchar();
return 0;
}
【问题讨论】：
-
缩小输入域会发生什么？错误模式可能有助于诊断。
-
是否真的有必要发布依赖于我们无权访问的外部文件的代码，其中包括 10 行冗余注释掉的代码和 130 多行仅包含无意义 cmets 的列宽行?
-
我很惊讶它给出了任何答案！您正在检查
offset < blockDim.x，因为它应该是ltid+offset < blockDim.x。在第二个 while 循环中，您正在使用不属于您的共享内存。编辑：动态__shared__大小不会有任何好处，因为您使用的是temp[ltid] = 0;所以只有具有块大小的共享内存是有效的。 -
我认为归约内核本身虽然非正统，但实际上是可以的，就像调用内核的看起来相当奇怪的主机循环一样（你只需要像这样调用任何归约 两次 i>，那么为什么要使用循环？）。但是在
cudaMemcpy(output, gidata,size * sizeof(int), cudaMemcpyDeviceToHost);中至少存在主机端缓冲区溢出以及可能的其他错误，但是由于没有人可以实际运行此代码，我不确定您如何期望任何人能够真正回答这个问题...