为什么简单的 CUDA 程序的结果每次都不一样？

Question

我想做一个简单的平铺卷积代码。来自Coursera：异构并行编程的讲座。讲座提供了一个简单的卷积代码，带有平铺的方法，但代码并不完整。因此，我填写了代码中的空白，下面是结果。

这段代码的目标是计算卷积。 输入尺寸：（24 x 24），
内核大小：（9乘9）， 输出尺寸：（16 x 16）。

此外，我主要将计算时间检查代码与 CPU 版本进行比较。

问题是，每当我运行这段代码时，结果都是不同的。找了几天的问题，但每次试都没有成功。我还在 Internet blog 中找到了类似的代码，但它与我的问题相同。我不知道为什么每次结果都不一样。有人说这是由于比赛条件，但我没有找到任何相关信息。

这是卷积的示例结果（16 x 16 大小）。

81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
0 0 0 0 0 0 0 0 0 0 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81

81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0

我的设备是 CUDA Driver = CUDART, CUDA Driver Version = 7.5, CUDA Runtime Version = 7.5, NumDevs = 1, Device0 = GeForce GT 630。 我使用的是 Ubuntu 14.04。

提前致谢！

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

#define O_TILE_WIDTH 10
#define MASK_WIDTH 9
#define I_TILE_WIDTH (O_TILE_WIDTH+MASK_WIDTH-1)

__global__ void Convolution2DBasicKernel(float *out, float *in, int in_height, int in_width, const float *__restrict__ mask, int output_dim)
{

    int tx=threadIdx.x;
    int ty=threadIdx.y;


    int row_o=blockIdx.y*O_TILE_WIDTH+ty;
    int col_o=blockIdx.x*O_TILE_WIDTH+tx; 


    int row_i=row_o;
    int col_i=col_o;
    __syncthreads();


    __shared__ float Ns[I_TILE_WIDTH][I_TILE_WIDTH];

///////////////////////////////////////////////////////////
//////////////////// reading input data ///////////////////
    if( (row_i>=0)&&(row_i<in_height)&&(col_i>=0)&&(col_i<in_width) )
    {
        Ns[ty][tx]=in[row_i*in_width + col_i];
    }
    else
    {
        Ns[ty][tx]=0.0f;
    }
    __syncthreads();    


///////////////////////////////////////////////////////////
//////////////////// calculating convol ///////////////////
    float output=0.0f;
    if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) )
    {
        for(int i=0; i<MASK_WIDTH; i++)
        {
            for(int j=0; j<MASK_WIDTH; j++)
            {
                output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];

            }
        }

    }
    __syncthreads();


    if( (row_o<output_dim)&&(col_o<output_dim) )
    {
        out[row_o*output_dim+col_o]=output;//in_width
    }
    __syncthreads();
}

int main() {

int input_dim=24;
    int kernel_dim=9;
    int output_dim=16;


float *input = new float[input_dim*input_dim];
float *kernel = new float[kernel_dim*kernel_dim];
float *output = new float[output_dim*output_dim];

float *d_input;
float *d_kernel;
float *d_output;
cudaMalloc(&d_input, sizeof(float)*input_dim*input_dim);
cudaMalloc(&d_kernel, sizeof(float)*kernel_dim*kernel_dim);
cudaMalloc(&d_output, sizeof(float)*output_dim*output_dim);



for(int i=0; i<input_dim*input_dim; i++)
{
    input[i]=1.0;
}
for(int i=0; i<kernel_dim*kernel_dim; i++)
{
    kernel[i]=1.0;
}


cudaMemcpy(d_input, input, sizeof(float)*input_dim*input_dim, cudaMemcpyHostToDevice);
cudaMemcpy(d_kernel, kernel, sizeof(float)*kernel_dim*kernel_dim, cudaMemcpyHostToDevice);



dim3 dimBlock (I_TILE_WIDTH, I_TILE_WIDTH, 1);
dim3 dimGrid ((output_dim-1)/O_TILE_WIDTH+1, (output_dim-1)/O_TILE_WIDTH+1, 1);


clock_t begin, end;
double time_spent;
begin = clock();

for(int iteration=0; iteration<1; iteration++)//100000
{
    Convolution2DBasicKernel<<<dimGrid, dimBlock>>>(d_output, d_input, input_dim, input_dim, d_kernel, output_dim);
}

end = clock();
time_spent = (double)(end - begin) / CLOCKS_PER_SEC;
printf("time: %f\n", time_spent);

cudaMemcpy(output, d_output, sizeof(float)*output_dim*output_dim, cudaMemcpyDeviceToHost);

for(int y=0; y<output_dim; y++)
{
    for(int x=0; x<output_dim; x++)
        printf("%d\t", int(output[y*16+x]));
    printf("\n");
}


}

Answer 1

这是一个竞争条件。
这是一个例子。
您正在每个块和 2x2 块启动 18x18 个线程。
示例：
ThreadA threadIdx.x = 10 threadIdx.y = 0 blockIdx.x = 0 blockIdx.y = 0
ThreadB threadIdx.x = 0 threadIdx.y = 0 blockIdx.x = 1 blockIdx.y = 0

在内核内部计算时：
int tx=threadIdx.x
int ty=threadIdx.y
int row_o=blockIdx.y*O_TILE_WIDTH+ty
int col_o=blockIdx.x*O_TILE_WIDTH+tx
using O_TILE_WIDTH = 10

ThreadA row_o = 0*10+10 = 10 col_o = 0
@ 987654330@

这意味着两个线程将在内存中的相同位置输出结果，但计算方式不同。

Answer 2

您的共享内存访问超出范围。

假设你确信你的程序或多或少是正确的，你需要确保你没有越界：

if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) ) {
    for(int i=0; i<MASK_WIDTH; i++) {
        if(ty +i < O_TILE_WIDTH) { // Changed here
            for(int j=0; j<MASK_WIDTH; j++) {
                if(tx +j < O_TILE_WIDTH) { // Changed here
                    output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];
                }
            }
        }
    }
}