基本的に配列を合計する 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;
}