例でCudaのドットプロダクトが動作しません

Question

「Cuda By Example」ブックを読み始めました。「共有メモリ」を使用したドットの例に問題がありました。私は本からの例をコピーして貼り付けて、次のように設定します：N = x * 1024; threadsPerBlock = 32; blocksPerGrid = 8. "x"の値を2、3、4、5でテストします。 x = 3に設定すると結果は悪くなりますが、x = 2,4,5を使用した場合はすべてOKです。私はどこに問題があるのか​​理解していない。コードは次のとおりです。~7 decimal digitsだけで十分な精度を持っていない

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <stdio.h> 

#define imin(a, b) (a<b?a:b) 
#define sum_squares(x) (x*(x+1)*(2*x+1)/6) 

const int x = 3; 
const int N = 3 * 1024; 
const int threadsPerBlock = 32; 
const int blocksPerGrid = 8; 

__global__ void dot(float *a, float *b, float *c) 
{ 
    __shared__ float cache[threadsPerBlock]; 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    int cacheIndex = threadIdx.x; 
    float temp = 0; 

    while (tid < N) 
    { 
     temp += a[tid] * b[tid]; 
     tid += blockDim.x * gridDim.x; 
    } 

    cache[cacheIndex] = temp; 

    __syncthreads(); 

    int i = blockDim.x/2; 
    while (i != 0) 
    { 
     if (cacheIndex < i) 
      cache[cacheIndex] += cache[cacheIndex + i]; 
     __syncthreads(); 
     i /= 2; 
    } 

    if (cacheIndex == 0) 
     c[blockIdx.x] = cache[0]; 
} 

int main() 
{ 
    float *a, *b, *partial_c, result; 
    float *d_a, *d_b, *d_partial_c; 

    a = (float *)malloc(N * sizeof(float)); 
    b = (float *)malloc(N * sizeof(float)); 
    partial_c = (float *)malloc(blocksPerGrid * sizeof(float)); 

    cudaMalloc((void **)&d_a, N * sizeof(float)); 
    cudaMalloc((void **)&d_b, N * sizeof(float)); 
    cudaMalloc((void **)&d_partial_c, blocksPerGrid * sizeof(float)); 

    for (int i = 0; i < N; i++) 
    { 
     a[i] = i; 
     b[i] = 2 * i; 
    } 

    cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_b, b, N * sizeof(float), cudaMemcpyHostToDevice); 

    dot << <blocksPerGrid, threadsPerBlock >> >(d_a, d_b, d_partial_c); 

    cudaMemcpy(partial_c, d_partial_c, blocksPerGrid * sizeof(float),  cudaMemcpyDeviceToHost); 

    result = 0; 
    for (int i = 0; i < blocksPerGrid; i++) 
     result += partial_c[i]; 

    if (2 * sum_squares((float)(N - 1)) == result) 
     printf(":)\n"); 
    else 
     printf(":(\n"); 

    cudaFree(d_a); 
    cudaFree(d_b); 
    cudaFree(d_partial_c); 

    free(a); 
    free(b); 
    free(partial_c); 

    getchar(); 
    return 0; 
} 

Answer 1

floatので。しかし、x=3;については、あなたの期待する結果は

19317916672 

11桁です。

x=4,5の場合、結果は自分のマシンでも悪いです。