CUDA Error - unspecified launch failure

Question

To practice coding with CUDA, I made a little test scenario where I have three files:

• memory.c Holds pure C code
• memory_kernels.h Declarations for CUDA kernels and functions to launch kernels
• memory_kernels.cu Definitions of kernels

What the program should do is create an integer array on the host, copy it to the device and query the elements. The kernel will print out a few details.

However, I am getting the error:

Error in memory_kernels.cu at line 43 with error code "unspecified launch failure"

The source code for the three files are given below:

/** 
 * memory.c
 *
 * Test copying large arrays to device
 * and printing from kernel
 */

/* Include standard libraries */
#include <stdlib.h>
#include <stdio.h>

/* Include local header files */
#include "memory_kernels.h"

int main() {

  /* Size of array */
  int i, N = 1024;

  /* Array */
  int *intArr = (int *) malloc( N * sizeof(int) );

  /* Fill array */
  for( i = 0; i < N; i++ ) {
    intArr[i] = i;
  }

  /* Run CUDA code */
  cuda_mem( &intArr );

  /* Clean up device */
  cudaDeviceReset();

  /* Everything done */
  exit(EXIT_SUCCESS);
}

/** 
 * memory_kernels.h
 *
 * Declarations for CUDA kernels
 */

/* Determine compiler */
#ifdef __cplusplus
#define EXTCFUNC extern "C"
#else
#define EXTCFUNC extern
#endif

#ifndef KERNELS_H
#define KERNELS_H

/* Standard libraries (only needed for debugging) */
#include <stdio.h>

/* Include CUDA header files */
#include <cuda.h>
#include <cuda_runtime.h>

#define CUDA_CALL(x) do { if((x) != cudaSuccess) {                                                         \
  printf("Error in %s at line %d with error code \"%s\"\n",__FILE__,__LINE__,cudaGetErrorString(x));       \
  exit(x);}} while(0)

/* Device globals */
__device__ int *d_intArr;

/* Device kernels */
__global__ void mem();

/* Host access functions */
EXTCFUNC void cuda_mem( int **intArr );

#endif

/** 
 * memory_kernels.cu
 *
 * CUDA kernel implementations
 */

/* Include header file */
#include "memory_kernels.h"

__global__ void mem() {
  int i = threadIdx.x;
  int a = d_intArr[i];

  printf("i = %d    a = %d\n",i,a);
}

/* Determine compiler */
#ifdef __cplusplus
#define EXTCFUNC extern "C"
#else
#define EXTCFUNC extern
#endif

/** 
 * cuda_mem()
 *
 * Test copying large array to device 
 * and printing from kernel
 */
EXTCFUNC void cuda_mem( int **intArr ) {
  /* Local variables */
  int N = 1024;

  /* Initialise device variables */
  CUDA_CALL( cudaMalloc( (void **) &d_intArr, sizeof(int) * N ) );

  /* Copy to device initial values */
  CUDA_CALL( cudaMemcpy( d_intArr, *intArr, sizeof(int) * N, cudaMemcpyHostToDevice ) );

  /* Run kernel */
  mem <<< 1,N >>> ();
  CUDA_CALL( cudaPeekAtLastError() );
  CUDA_CALL( cudaDeviceSynchronize() );

  /* Free local scoped dynamically allocated memory */
  CUDA_CALL( cudaFree( d_intArr ) );
}

Compilation is done with the following commands:

nvcc -c -o memory.o memory.c -arch=sm_20
nvcc -c -o memory_kernels.o memory_kernels.cu -arch=sm_20
nvcc -o memory memory.o memory_kernels.o -arch=sm_20

and ran on an NVIDIA Tesla M2050 with CUDA 4.0. Compute capability 2.0 is required to use printf() in the kernel.

Having searched around for a solution, the error code suggests I have a segmentation fault in the kernel, when reading from global memory. However, I am launching the same number of threads as the size of the array.

Having experimented around, I have a feeling the error is caused when copying intArr to the device. Maybe I am getting my pointers all mixed up?

I understand if the file structure is a bit strange, but it's all part of a larger program, but I have reduced the error to this smaller case.

Answer 1

The error raised because global arrays cannot be read/written directly by the kernel. The correct approach is to pass a pointer of the global array as an argument to the kernel.

Declare and define the kernel as:

__global__ void mem(int *dArr);

__global__ void mem(int *dArr) 
{
  int i = threadIdx.x;
  int a = dArr[i];

  printf("i = %d    a = %d\n",i,a);
}

Call the kernel as:

mem <<< 1,N >>> (d_intArr);

The above approached solved the problem for me and the program works perfectly.

Additional consideration:

You cannot use a variable declared with a __device__ modifier directly in the host code. When I compiled your code with CUDA 5, I get warning that

warning: a device variable "d_intArr" cannot be directly read in a host function

The following function call generates the warning:

CUDA_CALL( cudaMemcpy( d_intArr, *intArr, sizeof(int) * N, cudaMemcpyHostToDevice ) );

To keep the global effect, you can pass the pointers as arguments to your functions instead of declaring global arrays.

Answer 2

I'd like to expand on the answer provided by @sgar91 to provide some additional perspective (my own). As I see it there are at least two ways to instantiate arrays that are accessible from both the host and the device, in global memory.

a. use dynamically located/allocated arrays created on the host side. The code sequence is roughly as follows:

int main(){
  int *arr, *d_arr;
  arr = (int *)malloc(N*sizeof(int));
  cudaMalloc((void **) &d_arr, N*sizeof(int));
  cudaMemcpy(d_arr, arr, N*sizeof(int), cudaMemcpyHostToDevice);
  ...
  }

b. use statically located (and perhaps allocated) arrays. The code sequence is roughly as follows:

__device__ int d_arr[N];
...
int main(){
  int *arr;
  arr = (int *)malloc(N*sizeof(int));
  cudaMemcpyToSymbol(d_arr, arr, N*sizeof(int));
  ...
  }

With the first method, I must pass the address of the d_arr to the kernel as a parameter. With the second method I need not do so, as the array is statically located and therefore the compiler and runtime are able to locate it and properly fix up the code at load time. With the second method, I can directly access the d_arr from the kernel, even though I have not passed it as a parameter to the kernel.

Note that it's possible to have a dynamically sized (but statically located) array using the second method, but for the sake of brevity I'm not illustrating that here.

The answer provided by sgar91 didn't completely conform to either of these approaches, and so for example the warning is still there about using a device address in host code (even though it seems to work).