CUDA 错误 - 未指定的启动失败

Question

为了练习使用 CUDA 进行编码，我制作了一个小测试场景，其中包含三个文件:

memory.c持有纯C代码

memory_kernels.h用于启动内核的 CUDA 内核和函数声明

memory_kernels.cu内核定义

程序应该做的是在主机上创建一个整数数组，将其复制到设备并查询元素。内核会打印出一些细节。

但是，我收到错误:
Error in memory_kernels.cu at line 43 with error code "unspecified launch failure"
这三个文件的源代码如下:

/** 
 * 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 ) );
}

编译通过以下命令完成:

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

并在带有 CUDA 4.0 的 NVIDIA Tesla M2050 上运行。使用 printf()需要计算能力2.0在内核中。

四处寻找解决方案后，错误代码表明我在从全局内存读取时内核中存在段错误。但是，我启动的线程数与数组的大小相同。

试了下，感觉是复制 intArr时报错到设备。也许我把我的指针搞混了？

我知道文件结构是否有点奇怪，但它都是更大程序的一部分，但我已将错误减少到这个较小的情况。

Answer 1

引发错误是因为内核无法直接读取/写入全局数组。
正确的做法是将全局数组的指针作为参数传递给内核。

声明并定义内核为:

__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);
}

调用内核为:

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

以上方法为我解决了问题，程序运行良好。

额外考虑:

不能使用用 __device__ 声明的变量修饰符直接在宿主代码中。当我使用 CUDA 5 编译您的代码时，我收到警告

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

以下函数调用会生成警告:

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

保留 全局效果，您可以将指针作为参数传递给您的函数，而不是声明全局数组。