atomicAdd causing error Unable to Launch/Execute Kernel

Question

I have the following CUDA C code:

 int i = threadIdx.x + blockIdx.x*blockDim.x;
 int stride = blockDim.x*gridDim.x;
 while(i < size)
 {
     atomicAdd(&(histo_private[buffer[i]]),1);
     i+=stride;
 }

which causes my program to crash with the error: "unable to launch/execute kernel"

Here buffer is an input array of integers to this function of size elements and histo_private is an array of integers in shared memory of histo_size elements. I know this isn't an index out of bounds error because when I use the code:

int i = threadIdx.x + blockIdx.x*blockDim.x;
int stride = blockDim.x*gridDim.x;
while(i < size)
{
     int a = histo_private[buffer[i]];
     i+=stride;
}

So I gather that there is something wrong with the atomicAdd function and/or the memory address of this 32-bit int array.

The kernel.cu file contains the following code:

// Define your kernels in this file you may use more than one kernel if you
// need to

// INSERT KERNEL(S) HERE

__global__ void histo_kernel(unsigned int* buffer, unsigned int size, int* histo, unsigned int histo_size)
{
    extern __shared__ int histo_private[];
    if(threadIdx.x < histo_size)
        histo_private[threadIdx.x] = 0;
    __syncthreads();

    // compute block's histogram
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    int stride = blockDim.x*gridDim.x;
    while(i < size)
    {
        //int a = histo_private[buffer[i]];
        atomicAdd(&(histo_private[buffer[i]]),1);
        i+=stride;
    }

    // store to global histogram
    __syncthreads();
    //if(threadIdx.x < histo_size)
    //  atomicAdd(&(histo[threadIdx.x]),histo_private[threadIdx.x]);
}   

// ensures that no bins contains more than 255 elements
__global__ void enforce_saturation(int* histo, unsigned int histo_size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < histo_size)
    {
        if(histo[i] > 255)  // this will be necessary to prevent data loss
            histo[i] = 255; // when converting from int to uint8_t                          
    }
}

__global__ void construct_histo(uint8_t* histo_unpacked, int* histo, unsigned int histo_size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < histo_size)
        histo_unpacked[i] = histo[i];
}

// unpacks the input array into an output array with 'spaces'
__global__ void unpack(uint8_t* in, uint8_t* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[4*i] = in[i];
        out[4*i+1] = 0;
        out[4*i+2] = 0;
        out[4*i+3] = 0;
    }
}

// converts the input uint8_t array to an int array
__global__ void convert(uint8_t* in, int* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[i] = (int) in[4*i];
    }
}

// converts the input int array to a uint8_t array 
__global__ void convert_back(int* in, uint8_t* out, unsigned int size)
{
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < size)
    {
        out[i] = (uint8_t) in[i];
    }
}



void histogram(unsigned int* input, uint8_t* bins, unsigned int num_elements, unsigned int num_bins) 
{

    int BLOCK_SIZE = (int) num_bins;
    BLOCK_SIZE = 512;
    dim3 dim_grid, dim_block;
    dim_block.x = BLOCK_SIZE; dim_block.y = dim_block.z = 1;
        dim_grid.x = 1+(num_elements-1)/BLOCK_SIZE; dim_grid.y = dim_grid.z = 1;

    // create an array of uint8_t to be converted into an array of int
    uint8_t* bins_unpacked;
    cudaMalloc((void**)&bins_unpacked, 4 * num_bins * sizeof(uint8_t));

    // unpack the input uint8_t array
    unpack<<<dim_grid,dim_block>>>(bins, bins_unpacked, num_bins);

    // need an int version of bins_d
    int* bins_int_d;
    cudaMalloc((void**)&bins_int_d, num_bins * sizeof(int));

    // convert the uint8_t array to an int array
    convert<<<dim_grid,dim_block>>>(bins_unpacked, bins_int_d, num_bins);   

    // run kernel and enforce saturation requirements
    int histo_private_size = num_bins;
    histo_kernel<<<dim_grid,dim_block,histo_private_size>>>(input, num_elements, bins_int_d, num_bins);
    enforce_saturation<<<dim_grid,dim_block>>>(bins_int_d,num_bins);

    // convert the int array back to uint8_t
    convert_back<<<dim_grid,dim_block>>>(bins_int_d, bins, num_bins);
}       

While the function that calls this last histogram function is in main.cu (I did NOT make this second file--it was provided to me--also, I have been testing this on consistent data by compiling via make test-mode):

#include <stdio.h>
#include <stdint.h>

#include "support.h"
#include "kernel.cu"

int main(int argc, char* argv[])
{
    Timer timer;

    // Initialize host variables ----------------------------------------------

    #if TEST_MODE
    printf("\n***Running in test mode***\n"); fflush(stdout);
    #endif

    printf("\nSetting up the problem..."); fflush(stdout);
    startTime(&timer);

    unsigned int *in_h;
    uint8_t* bins_h;
    unsigned int *in_d;
    uint8_t* bins_d;
    unsigned int num_elements, num_bins;
    cudaError_t cuda_ret;

    if(argc == 1) {
        num_elements = 1000000;
        num_bins = 4096;
    } else if(argc == 2) {
        num_elements = atoi(argv[1]);
        num_bins = 4096;
    } else if(argc == 3) {
        num_elements = atoi(argv[1]);
        num_bins = atoi(argv[2]);
    } else {
        printf("\n    Invalid input parameters!"
           "\n    Usage: ./histogram            # Input: 1,000,000, Bins: 4,096"
           "\n    Usage: ./histogram <m>        # Input: m, Bins: 4,096"
           "\n    Usage: ./histogram <m> <n>    # Input: m, Bins: n"
           "\n");
        exit(0);
    }
    initVector(&in_h, num_elements, num_bins);
    bins_h = (uint8_t*) malloc(num_bins*sizeof(uint8_t));

    // TESTING
    for(unsigned int i = 0; i < num_bins; ++i) 
    {
        bins_h[i] = i;
        //printf("uint8_t Element %u: is %u \n", i, bins_h[i]);
    }



    stopTime(&timer); printf("%f s\n", elapsedTime(timer));
    printf("    Input size = %u\n    Number of bins = %u\n", num_elements,
        num_bins);

    // Allocate device variables ----------------------------------------------

    printf("Allocating device variables..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMalloc((void**)&in_d, num_elements * sizeof(unsigned int));
    if(cuda_ret != cudaSuccess) FATAL("Unable to allocate device memory");
    cuda_ret = cudaMalloc((void**)&bins_d, num_bins * sizeof(uint8_t));
    if(cuda_ret != cudaSuccess) FATAL("Unable to allocate device memory");

    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Copy host variables to device ------------------------------------------

    printf("Copying data from host to device..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMemcpy(in_d, in_h, num_elements * sizeof(unsigned int),
        cudaMemcpyHostToDevice);
    if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to the device");

    cuda_ret = cudaMemset(bins_d, 0, num_bins * sizeof(uint8_t));
    if(cuda_ret != cudaSuccess) FATAL("Unable to set device memory");

    // TESTING
    //cuda_ret = cudaMemcpy(bins_d, bins_h, num_bins * sizeof(uint8_t),
    //    cudaMemcpyHostToDevice);
    //if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to the device");



    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Launch kernel ----------------------------------------------------------
    printf("Launching kernel..."); fflush(stdout);
    startTime(&timer);

    histogram(in_d, bins_d, num_elements, num_bins);
    cuda_ret = cudaDeviceSynchronize();
    if(cuda_ret != cudaSuccess) FATAL("Unable to launch/execute kernel");

    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    // Copy device variables from host ----------------------------------------

    printf("Copying data from device to host..."); fflush(stdout);
    startTime(&timer);

    cuda_ret = cudaMemcpy(bins_h, bins_d, num_bins * sizeof(uint8_t),
        cudaMemcpyDeviceToHost);
    if(cuda_ret != cudaSuccess) FATAL("Unable to copy memory to host");

    cudaDeviceSynchronize();
    stopTime(&timer); printf("%f s\n", elapsedTime(timer));

    #if TEST_MODE
    printf("\nResult:\n");
    for(unsigned int binIdx = 0; binIdx < num_bins; ++binIdx) {
       printf("Bin %u: %u elements\n", binIdx, bins_h[binIdx]);
    }

    printf("\nElements Vec:\n");
    for(unsigned int i = 0; i < num_elements; ++i) {
        printf("Element %u: %u  is \n", i, in_h[i]);
    }



    #endif

    // Verify correctness -----------------------------------------------------

    printf("Verifying results..."); fflush(stdout);

    verify(in_h, bins_h, num_elements, num_bins);

    // Free memory ------------------------------------------------------------

    cudaFree(in_d); cudaFree(bins_d);
    free(in_h); free(bins_h);

    return 0;
}

Answer 1

Turns out that this was just an index out of bounds error. The element buffer[i] was greater than the length of histo_private. As another poster mentioned, this was not obvious due to the following artifact of the c compiler:

The compiler is permitted to assume every access is within bounds. That line of my test code did nothing if the access is within bounds and therefore the compiler is permitted to assume that line of code does nothing. Thus it didn't require an access so the successful run of the test code was misleading. Once that line was changed to where the variable hist_private was modified at buffer[i], runtime errors came about.