C/CUDA Nvidia Dotproduct 示例给出了不正确的结果

Question

我正在尝试在 C/CUDA 中实现点积。我主要从此处提供的 Nvidias 教程中复制了代码:http://www.nvidia.com/content/gtc-2010/pdfs/2131_gtc2010.pdf

我要的结果就是输出

*c     = 44870400
result = 44870400

但是我明白了

*c     = 44608256
result = 44870400

好像是“511*511 case”不是计算结果的一部分。我上下检查了代码，甚至找不到同步错误。我在这里做错了什么？

编译标志是:

cuda_dotp: ./cuda_dotp.cu
    nvcc -arch=sm_13 \
    -o cuda_dotp ./cuda_dotp.cu

和文件 cuda_dotp.cu 的内容

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

#define N 513
#define THREADS_PER_BLOCK 512

__global__ void dot(int *a, int *b, int *c) {
    __shared__ int temp[THREADS_PER_BLOCK];
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    temp[threadIdx.x] = a[index] * b[index];
    if (index >= N) return;

    __syncthreads();
    if(0 == threadIdx.x) {
        int sum = 0;
        int max = THREADS_PER_BLOCK;
        if (N < max) max = N;

        for (int i = 0; i < max; i++) {
            sum += temp[i];
        }
        c[0] = sum;
    }
}

void random_ints(int *a, int size)
{
    int i;
    for (i=0; i<size; i++)
        a[i] = i;
    return;
}

int main(void) {
    int i;
    int result;
    int *a, *b, *c; // host copies of a, b, c
    int *dev_a, *dev_b, *dev_c; // device copies of a, b, c
    int size = N * sizeof(int); // we need space for N ints
    // allocate device copies of a, b, c
    cudaMalloc( (void**)&dev_a, size );
    cudaMalloc( (void**)&dev_b, size );
    cudaMalloc( (void**)&dev_c, sizeof(int) );
    a = (int*)malloc( size );
    b = (int*)malloc( size );
    c = (int*)malloc( sizeof(int) );

    random_ints( a, N );
    random_ints( b, N );
    /*
    printf("a = ");
    for (i=0; i<N; i++) printf("%d, ", a[i]);
    printf("\n");
    printf("b = ");
    for (i=0; i<N; i++) printf("%d, ", b[i]);
    printf("\n");
    */
    result = 0;
    for (i=0; i<N; i++) result += a[i] * b[i];
    *c = 0;

    // copy inputs to device
    cudaMemcpy( dev_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy( dev_b, b, size, cudaMemcpyHostToDevice);

    int blocks = N/THREADS_PER_BLOCK;
    if(blocks<1) blocks=1;

    // launch dot() kernel
    dot <<< blocks, THREADS_PER_BLOCK >>> (dev_a, dev_b, dev_c);

    // copy device result back to host copy of c
    cudaMemcpy(c, dev_c, sizeof(int) , cudaMemcpyDeviceToHost);

    printf("*c     = %d\n", *c);
    printf("result = %d\n", result);

    free(a); free(b); free(c);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

Answer 1

有不少错误。访问分配空间之外的数组，运行 0 个 block ((int)10/(int)512 = 0)，在将 c 添加到内核之前不初始化 c。

将您的代码与以下内容进行比较。

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

#define N 10
#define THREADS_PER_BLOCK 512

__global__ void dot(int *a, int *b, int *c) {
    __shared__ int temp[THREADS_PER_BLOCK];
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    temp[threadIdx.x] = a[index] * b[index];
    if(index>=N) return;

    __syncthreads();
    if(0 == threadIdx.x) {
        int sum = 0;
        int max= THREADS_PER_BLOCK;
        if(N<max)max=N;

        for(int i = 0; i < max; i++){
            sum += temp[i];
        }
        c[0]=sum;
    }
}

void random_ints(int *a, int size)
{
    int i;
    for (i=0; i<size; i++)
        a[i] = i;
    return;
}

int main(void) {
    int i;
    int *a, *b, *c; // host copies of a, b, c
    int *dev_a, *dev_b, *dev_c; // device copies of a, b, c
    int size = N * sizeof(int); // we need space for N ints
    // allocate device copies of a, b, c
    cudaMalloc( (void**)&dev_a, size );
    cudaMalloc( (void**)&dev_b, size );
    cudaMalloc( (void**)&dev_c, sizeof(int) );
    a = (int*)malloc( size );
    b = (int*)malloc( size );
    c = (int*)malloc( sizeof(int) );

    random_ints( a, N );
    random_ints( b, N );
    printf("a = ");
    for (i=0; i<N; i++) printf("%d, ", a[i]);
    printf("\n");
    printf("b = ");
    for (i=0; i<N; i++) printf("%d, ", b[i]);
    printf("\n");
    *c = 0;

    // copy inputs to device
    cudaMemcpy( dev_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy( dev_b, b, size, cudaMemcpyHostToDevice);

    int blocks = N/THREADS_PER_BLOCK;
    if(blocks<1) blocks=1;

    // launch dot() kernel
    dot<<< blocks,THREADS_PER_BLOCK>>>( dev_a, dev_b, dev_c);


    // copy device result back to host copy of c
    cudaMemcpy(c, dev_c, sizeof(int) , cudaMemcpyDeviceToHost);

    printf("*c = %d\n", *c);

    free(a); free(b); free(c);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);


    return 0;
}