2

为了练习使用 CUDA 编码，我做了一个小测试场景，其中包含三个文件：

  • memory.c持有纯C代码
  • memory_kernels.hCUDA 内核的声明和启动内核的函数
  • 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 上运行。计算能力 2.0 需要printf()在内核中使用。

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

经过试验，我感觉错误是在复制intArr到设备时引起的。也许我把我的指针都搞混了？

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

4

2 回答 2

2

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

声明并定义内核为：

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

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

于 2013-01-06T18:33:37.097 回答
1

我想扩展@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);
  ...
  }

湾。使用静态定位（也可能是分配的）数组。代码顺序大致如下：

__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 提供的答案并不完全符合这两种方法中的任何一种，因此例如关于在主机代码中使用设备地址的警告仍然存在（即使它似乎有效）。

于 2013-01-06T21:58:20.170 回答