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

编译与done以下命令：

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 

在NVIDIA®（英伟达™）Tesla M2050上运行CUDA 4.0。计算能力2.0需要在内核中使用printf()。

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

经过实验，我有一种感觉，错误是由于将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)); 

为了保持全球效果，你可以通过指针作为参数到你的函数而不是声明全局数组。

Answer 2

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

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

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

湾使用静态定位（也许分配）数组。码序列大致如下：

__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提供的答案并不完全符合这些方法的执行，所以例如警告仍然存在有关在主机代码中使用的设备地址（即使它似乎工作）。