CUDA C: Kernel outputs bad results

Question

first of all I want to say that this is not a homework assignment and I'm just beginning with CUDA.

I'm trying to run the following code to add 2 vectors... The thing is that after every run the result vector (c_device) stays the same and doesn't get the result of the addition of the two vectors.

I have tried changing the length of the vector and working with ints and unsigned ints and also trying to move between x64 and win32 in visual studio.

I have attached the code here:

This is the .h file

#ifndef ODINN_CUDA_MAIN_H
#define ODINN_CUDA_MAIN_H

#define ARR_SIZE 100
#define ITER_AMOUNT 1

typedef enum cudaError cudaError_t;

static void HandleError(cudaError_t err, const char *file, int line) {
    if (err != CUDA_SUCCESS) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

#define GET_CURRENT_CLOCKS(var) (var = clock())
#define GET_CLOCK_INTERVAL_SEC(start, end, result) (result = ((double)((double)end - (double)start) / (double)CLOCKS_PER_SEC))

__host__ dim3 requestBlockSize(int x, int y=0, int z=0);
__host__ dim3 requestNumBlocks(int x, int y=0, int z=0);
__host__ void allocateVectors(unsigned int **a_host, unsigned int **b_host, unsigned int **c_host, unsigned int **a_device, unsigned int **b_device, unsigned int **c_device);
__global__ void addVectors(unsigned int* a, unsigned int* b, unsigned int* result, int n);
__host__ void cleanUp(unsigned int *a_host, unsigned int *b_host, unsigned int *c_host, unsigned int *a_device, unsigned int *b_device, unsigned int *c_device);

#endif

And this is the .cu file:

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

#include "main.h"

static cudaDeviceProp prop;

int main(void) {
        // Start lazy init now so first cudaMallow will run faster.
        cudaSetDevice(0);
        cudaFree(0);

        unsigned int *a_host, *b_host, *c_host;
        unsigned int *a_device, *b_device, *c_device;
        double delta_in_sec;
        size_t size = sizeof(unsigned int) * ARR_SIZE;
        clock_t start_clock, end_clock;

        HANDLE_ERROR(cudaGetDeviceProperties(&prop, 0));

        dim3 block_size = requestBlockSize(1024);
        int blocks_requested = floor((double)(ARR_SIZE / block_size.x));
    dim3 n_blocks = requestNumBlocks(blocks_requested > 0 ? blocks_requested : 1);

        fprintf(stdout, "Allocating vectors ...\n");
        allocateVectors(&a_host, &b_host, &c_host, &a_device, &b_device, &c_device);

        fprintf(stdout, "Copying to device ...\n");
        HANDLE_ERROR(cudaMemcpy(a_device, a_host, size, cudaMemcpyHostToDevice));
        HANDLE_ERROR(cudaMemcpy(b_device, b_host, size, cudaMemcpyHostToDevice));

        fprintf(stdout, "Running kernel ...\n");
        GET_CURRENT_CLOCKS(start_clock);

        for(int i=0; i<ITER_AMOUNT; i++) {\
                addVectors<<<n_blocks, block_size>>>(a_device, b_device, c_device, ARR_SIZE);
        }

        GET_CURRENT_CLOCKS(end_clock);
        GET_CLOCK_INTERVAL_SEC(start_clock, end_clock, delta_in_sec);

        fprintf(stdout, "Runtime of kernel %d times on arrays in length %d took %f seconds\n"
                "Copying results back to host ...\n", ITER_AMOUNT, ARR_SIZE, delta_in_sec);

        HANDLE_ERROR(cudaMemcpy(c_host, c_device, size, cudaMemcpyDeviceToHost));;
        fprintf(stdout, "%u + %u != %u\n", a_host[0], b_host[0], c_host[0]);

        fprintf(stdout, "Cleaning up ...\n");
        cleanUp(a_host, b_host, c_host, a_device, b_device, c_device);

        fprintf(stdout, "Done!\n");
}

__host__ dim3 requestBlockSize(int x, int y, int z) {
        dim3 blocksize(
                x <= prop.maxThreadsDim[0] ? x : prop.maxThreadsDim[0],
                y <= prop.maxThreadsDim[1] ? y : prop.maxThreadsDim[1],
                z <= prop.maxThreadsDim[2] ? z : prop.maxThreadsDim[2]
        );

        return blocksize;
}

__host__ dim3 requestNumBlocks(int x, int y, int z) {
        dim3 numblocks(x, y, z);

        return numblocks;
}

__host__ void allocateVectors(unsigned int **a_host, unsigned int **b_host, unsigned int **c_host, unsigned int **a_device, unsigned int **b_device, unsigned int **c_device) {
        size_t size = sizeof(unsigned int) * ARR_SIZE;

        *a_host = (unsigned int *)malloc(size);
        *b_host = (unsigned int *)malloc(size);
        *c_host = (unsigned int *)malloc(size);

        HANDLE_ERROR(cudaMalloc((void **)a_device, size));
        HANDLE_ERROR(cudaMalloc((void **)b_device, size));
        HANDLE_ERROR(cudaMalloc((void **)c_device, size));

        srand(time(NULL));

        for(int i=0; i<ARR_SIZE; i++) {
                (*a_host)[i] = rand() % ARR_SIZE;
                (*b_host)[i] = rand() % ARR_SIZE;
        }
}

__global__ void addVectors(unsigned int* a, unsigned int* b, unsigned int* result, int n) {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;

        if(idx >= 0 && idx < n)
                result[idx] = a[idx] + b[idx];
}

__host__ void cleanUp(unsigned int *a_host, unsigned int *b_host, unsigned int *c_host, unsigned int *a_device, unsigned int *b_device, unsigned int *c_device) {
        free(a_host);
        free(b_host);
        free(c_host);

        HANDLE_ERROR(cudaFree(a_device));
        HANDLE_ERROR(cudaFree(b_device));
        HANDLE_ERROR(cudaFree(c_device));
}

Here is the link to pastebin if you prefer looking at the code there: http://pastebin.com/04jy1CaB

I would like to mention that when copying a_device to c_host it works. I have also tried to copy c_host to c_device and see what happens and the same result happened.

Any suggestions?

You haven't asked a question here ("any suggestions?" isn't a valid question), and "bad results" is so vague a problem description that I am not sure what kind of answer you are expecting. Your API error checking is incomplete, it is likely that the kernel is never running, but you are not detecting the error. See [this answer](http://stackoverflow.com/a/14038590/681865) for details of how to correctly check for runtime errors when launching a kernel. — talonmies, Jun 16 '13 at 06:52
Ok, I see what you mean... I did some more error checking like in the link you gave me and it said that I'm passing invalid arguments to the kernel call... So that probably is the reason... Now the question is what's wrong with the parameters I'm passing... They are regular DIM3 parameters that in both of the x>0 and y=z=0. — Odinn, Jun 16 '13 at 07:48

Odinn · Accepted Answer · 2013-06-18T08:26:18.570

OK, so thanks to talonmies's comment on my question I have realized that I wasn't doing enough error checking and when I did the additional checks I found out that I was passing bad arguments to the kernel call.

I was passing invalid dim3 y and z values to the kernel thread amount and block amount parameter. If you notice my default values were 0 and they should be 1.

Long live the debugger :)

Thanks for your help talonmies.

CUDA C: Kernel outputs bad results

1 Answers1