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 编译您的代码时，我收到警告

警告：设备变量“d_intArr”不能直接读取 宿主函数

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

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

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

Answer 2

我想扩展@sgar91 提供的答案，以提供一些额外的观点（我自己的）。在我看来，至少有两种方法可以在全局内存中实例化可以从主机和设备访问的数组。

一个。使用在主机端创建的动态定位/分配的数组。代码顺序大致如下：

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.使用静态定位（也可能是分配的）数组。代码顺序大致如下：

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

使用第一种方法，我必须将d_arr 的地址作为参数传递给内核。使用第二种方法我不需要这样做，因为数组是静态定位的，因此编译器和运行时能够找到它并在加载时正确修复代码。使用第二种方法，我可以直接从内核访问d_arr，即使我没有将它作为参数传递给内核。

请注意，使用第二种方法可以拥有一个动态大小（但静态定位）的数组，但为了简洁起见，我没有在这里说明。

sgar91 提供的答案并不完全符合这两种方法中的任何一种，因此例如关于在主机代码中使用设备地址的警告仍然存在（尽管它似乎有效）。