當我改變矩陣大小時，Cuda內核結果出錯

Question

非常簡單的計算內核：tmp = X * y; tmp = sigmoid（temp）-L; Y =轉置（X）* TMP;但是，有時會返回正確的結果，有時會返回錯誤的結果，有時對於1000 * 1000大小的問題，它會返回正確的結果，但是當我增加問題大小時，它會返回錯誤的結果。看起來它有一些競爭條件。但所有數據都受到tid的限制。你能幫我找出錯誤！

謝謝！

#include <stdio.h> 
 
#include <stdlib.h> 
 
#include <time.h> 
 

 
#define BLOCK_ROWS 512 
 

 
__global__ void MVM(int trows, int tcols, float *d_x, float *d_y, float *d_l, float *d_out) 
 
{ 
 
     int tid = blockIdx.x*blockDim.x+threadIdx.x; 
 
     if(tid < trows) { 
 
       d_out[tid]=0; 
 
       for(int i=0; i<tcols; i++) 
 
       { 
 
         d_out[tid] = d_out[tid] + d_x[i*trows+tid]*d_y[i]; 
 
       } 
 
       d_out[tid] = 1.0/(exp(-d_out[tid])+1.0)-d_l[tid]; 
 
      
 
     } 
 

 
     __syncthreads; 
 

 
     if(tid < tcols) { 
 
       d_y[tid] =0; 
 
       for(int i=0; i<trows; i++) 
 
       { 
 
         d_y[tid] = d_y[tid] + d_x[tid*trows+i]*d_out[i]; 
 
       } 
 
     } 
 
} 
 
int main(void) 
 
{ 
 
    int trows = 100; int tcols = 100; 
 
    float *x, *y, *out, *l, *d_x, *d_y, *d_out, *d_l, *check, *check1; 
 
    x = (float*)malloc(trows*tcols*sizeof(float)); 
 
    y = (float*)malloc(tcols*sizeof(float)); 
 
    l = (float*)malloc(trows*sizeof(float)); 
 
    out = (float*)malloc(tcols*sizeof(float)); 
 
    check = (float*)malloc(trows*sizeof(float)); 
 
    check1 = (float*)malloc(tcols*sizeof(float)); 
 

 
    int result=0; 
 
    result = cudaMalloc(&d_x, trows*tcols*sizeof(float)); 
 
    if(result!=cudaSuccess) printf("GPU allocation fail\n"); 
 
    result = cudaMalloc(&d_y, tcols*sizeof(float)); 
 
    if(result!=cudaSuccess) printf("GPU allocation fail\n"); 
 
    result = cudaMalloc(&d_out, trows*sizeof(float)); 
 
    if(result!=cudaSuccess) printf("GPU allocation fail\n"); 
 
    result = cudaMalloc(&d_l, trows*sizeof(float)); 
 
    if(result!=cudaSuccess) printf("GPU allocation fail\n"); 
 

 
    for(int j = 0; j < tcols; j++) { 
 
     for (int i = 0; i < trows; i++) 
 
       x[j*trows+i] = (float)(i%10); 
 
    } 
 

 
    for(int i=0; i<tcols; i++) y[i] = (float)(i%10); 
 

 
    for(int i=0; i<trows; i++) l[i] = (float)((trows-i)%10); 
 

 
    result = cudaMemcpy(d_x, x, trows*tcols*sizeof(float), cudaMemcpyHostToDevice); 
 
    if(result!=cudaSuccess) printf("cpying to GPU fail\n"); 
 
    result = cudaMemcpy(d_y, y, tcols*sizeof(float), cudaMemcpyHostToDevice); 
 
    if(result!=cudaSuccess) printf("cpying to GPU fail\n"); 
 
    result = cudaMemcpy(d_l, l, trows*sizeof(float), cudaMemcpyHostToDevice); 
 
    if(result!=cudaSuccess) printf("cpying to GPU fail\n"); 
 

 
    int grid=0; 
 
    if(trows>tcols) grid = (trows-1)/BLOCK_ROWS+1; else grid = (tcols-1)/BLOCK_ROWS+1; 
 
    dim3 dimGrid(grid,1,1); 
 
    dim3 dimBlock(BLOCK_ROWS,1,1); 
 

 
    clock_t t; 
 
    t = clock(); 
 
    MVM<<<dimGrid, dimBlock>>>(trows, tcols, d_x, d_y, d_l, d_out); 
 
    t = clock()-t; 
 
    double time = ((double)t)/CLOCKS_PER_SEC; 
 
    printf("time: %f\n", time); 
 
    
 
    for(int i=0; i<trows; i++) { 
 
     float tmp = 0; 
 
     for(int j=0; j<tcols; j++) 
 
       tmp += x[j*trows+i]*y[j]; 
 
     tmp = 1.0/(exp(-tmp)+1.0) - l[i]; 
 
     check[i] = tmp; 
 
    } 
 
    for(int i=0; i<tcols; i++) { 
 
     float tmp = 0; 
 
     for(int j=0; j<trows; j++) 
 
       tmp = tmp+ x[i*trows+j]*check[j]; 
 
     check1[i] = tmp; 
 
    } 
 

 
    result = cudaMemcpy(out, d_y, tcols*sizeof(float), cudaMemcpyDeviceToHost); 
 
    if(result!=cudaSuccess) printf("cpying to CPU fail, error=%d\n",result); 
 

 
    float error=0; 
 
    for(int i=0; i< tcols;i++) { 
 
     error += abs(check1[i]-out[i])/(abs(check1[i])+1e-6); 
 
    } 
 
    printf("error = %f\n", error); 
 

 
    cudaFree(d_x); 
 
    cudaFree(d_y); 
 
    cudaFree(d_out); 
 
    cudaFree(d_l); 
 
    free(x); 
 
    free(y); 
 
    free(l); 
 
    free(out); 
 
    free(check); 
 
    free(check1); 
 
}

Answer 1

嘗試改變 'BLOCK_ROWS' 的 '爲dim3 dimBlock'，像上述 'trows' 的值。

dim3 dimBlock(trows,1,1); // instead of dim3 dimBlock(BLOCK_ROWS,1,1); 

如果在各列元素被一起處理，設置blockDimX等於行通常可避免因問題的數量「__syncthreads（）;」