Unspecified launch failure - parallel scan in CUDA

Question

I am using GeForce GT 520 (compute capablility v2.1) to run a program that performs the scan operation on an array of int elements. Here's the code:

/*
This is an implementation of the parallel scan algorithm.
Only a single block of threads is used. Maximum array size = 2048
*/

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

#define errorCheck(ans) { gpuAssert((ans), __FILE__, __LINE__); }

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess) 
{
    fprintf(stderr,"GPUassert: %s, file: %s line: %d\n", cudaGetErrorString(code), file, line);
    if (abort) exit(code);
}
}

__global__ void blelloch_scan(int* d_in, int* d_out, int n) 
{
    extern __shared__ int temp[];// allocated on invocation

    int thid = threadIdx.x;
    int offset = 1;

    temp[2*thid] = d_in[2*thid]; // load input into shared memory
    temp[2*thid+1] = d_in[2*thid+1];

    // build sum in place up the tree
    for (int d = n>>1; d > 0; d >>= 1)
    {
        __syncthreads();
        if (thid < d)
        {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }

    // clear the last element
    if (thid == 0)
    temp[n - 1] = 0; 
    __syncthreads();

    // traverse down tree & build scan
    for (int d = 1; d < n; d *= 2)
    {
        offset >>= 1;
        __syncthreads();
        if (thid < d)
        {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            int t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();  

    d_out[2*thid] = temp[2*thid]; // write results to device memory
    d_out[2*thid+1] = temp[2*thid+1];
}

int main(int argc, char **argv)
{
int ARRAY_SIZE;
if(argc != 2)
{
    printf("Input Syntax: ./a.out <number-of-elements>\nProgram terminated.\n");
    exit (1);        
}
else
ARRAY_SIZE = (int) atoi(*(argv+1));

int *h_in, *h_out, *d_in, *d_out, i;
h_in = (int *) malloc(sizeof(int) * ARRAY_SIZE);
h_out = (int *) malloc(sizeof(int) * ARRAY_SIZE);

cudaSetDevice(0);
cudaDeviceProp devProps;
if (cudaGetDeviceProperties(&devProps, 0) == 0)
{
    printf("Using device %d:\n", 0);
    printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
           devProps.name, (int)devProps.totalGlobalMem, 
           (int)devProps.major, (int)devProps.minor, 
           (int)devProps.clockRate);
}

for(i = 0; i < ARRAY_SIZE; i++)
{
    h_in[i] = i;    
}

errorCheck(cudaMalloc((void **) &d_in, sizeof(int) * ARRAY_SIZE));
errorCheck(cudaMalloc((void **) &d_out, sizeof(int) * ARRAY_SIZE));    
errorCheck(cudaMemcpy(d_in, h_in, ARRAY_SIZE * sizeof(int), cudaMemcpyHostToDevice));

blelloch_scan <<<1, ARRAY_SIZE / 2, sizeof(int) * ARRAY_SIZE>>> (d_in, d_out, ARRAY_SIZE);
cudaDeviceSynchronize();
errorCheck(cudaGetLastError());

errorCheck(cudaMemcpy(h_out, d_out, ARRAY_SIZE * sizeof(int), cudaMemcpyDeviceToHost));

printf("Results:\n");    
for(i = 0; i < ARRAY_SIZE; i++)
{
    printf("h_in[%d] = %d, h_out[%d] = %d\n", i, h_in[i], i, h_out[i]);    
}
return 0;
}

On compiling using nvcc -arch=sm_21 parallel-scan.cu -o parallel-scan, I get an error: GPUassert: unspecified launch failure, file: parallel-scan-single-block.cu line: 106

Line 106 is the line after kernel launch when we check for errors using errorCheck.

This is what I am planning to implement:

From the kernel, it can be seen that if a block has 1000 threads, it can operate on 2000 elements. Therefore, blockSize = ARRAY_SIZE / 2.
And, shared memory = sizeof(int) * ARRAY_SIZE
Everything is loaded into shared mem. Then, up sweep is done, with last element being set to 0. Finally, down sweep is done to give an exclusive scan of the elements.

I have used this file as the reference to write this code. I do not understand what's the mistake in my code. Any help would be greatly appreciated.

run your code with cuda-memcheck, if you get a read or write error, recompile your code with the -lineinfo switch. Then rerun the code with cuda-memcheck and it will point you to the kernel line that is generating the out-of-bounds access. — Robert Crovella

talonmies talonmies · Accepted Answer · 2014-06-07T10:11:21

You are launching the kernel like so

blelloch_scan <<<1, ARRAY_SIZE / 2, sizeof(int) * ARRAY_SIZE>>>

meaning that witihin then kernel 0 < thid < int(ARRAY_SIZE/2).

However, your kernel requires a minimum of (2 * int(ARRAY_SIZE/2)) + 1 words of available shared memory to work correctly, otherwise this:

temp[2*thid+1] = d_in[2*thid+1];

will produce an out-of-bounds shared memory access.

If my integer mathematical skillz are not too rusty, this should mean that the code will be safe if ARRAY_SIZE is odd, because ARRAY_SIZE == (2 * int(ARRAY_SIZE/2)) + 1 for any odd integer. However, if ARRAY_SIZE is even, then ARRAY_SIZE < (2 * int(ARRAY_SIZE/2)) + 1 and you have a problem.

It might be that shared memory page size granularity saves you for some even values of ARRAY_SIZE which should theoretically fail, because the hardware will always round up the dynamic shared memory allocation to the next page size larger than the request size. But there should be a number of even values of ARRAY_SIZE for which this fails.

I can't comment on whether the rest of the kernel is correct or not, but using a shared memory size of sizeof(int) * size_t(1 + ARRAY_SIZE) should make this particular problem go away.

Unspecified launch failure - parallel scan in CUDA

1 Answers