Dot product in Cuda by example does not work for me

Question

I'm starting to read "Cuda By Example" Book and I've been a problem with the dot example using "shared memory". I copy-paste the example from the book and I set: N = x * 1024; threadsPerBlock = 32; blocksPerGrid = 8. Where I test the "x" values with 2, 3, 4, 5. If I set x = 3, the result is bad, but when I used x = 2,4,5 all is ok. I don't understand where is the problem. The code is:

#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

Because float does not have enough precision, which is ~7 decimal digits only. But for x=3; your expected result is

19317916672

containing 11 digits.

for x=4,5, the results are bad on my machine too.