2

CUDA でのコーディングを練習するために、3 つのファイルがある小さなテスト シナリオを作成しました。

  • memory.c純粋なCコードを保持
  • memory_kernels.hカーネルを起動するための CUDA カーネルと関数の宣言
  • memory_kernels.cuカーネルの定義

プログラムがすべきことは、ホスト上に整数配列を作成し、それをデバイスにコピーして要素をクエリすることです。カーネルはいくつかの詳細を表示します。

ただし、次のエラーが表示されます。

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

3 つのファイルのソース コードは次のとおりです。

/** 
 * 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ます。多分私は私のポインタをすべて混同していますか？

ファイル構造が少し変わっているかどうかは理解していますが、それはすべて大きなプログラムの一部ですが、エラーをこの小さなケースに減らしました。

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によって提供された回答を拡張して、いくつかの追加の視点を提供したいと思います（私自身）。私が見ているように、グローバルメモリでホストとデバイスの両方からアクセス可能なアレイをインスタンス化する方法は少なくとも2つあります。

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

b。静的に配置された（そしておそらく割り当てられた）アレイを使用します。コードシーケンスはおおまかに次のとおりです。

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

最初の方法では、のアドレスをパラメータとしてカーネルに渡す必要があります。d_arr2番目の方法では、配列が静的に配置されているため、これを行う必要はありません。したがって、コンパイラとランタイムは、配列を検索し、ロード時にコードを適切に修正できます。d_arr2番目の方法では、パラメーターとしてカーネルに渡していない場合でも、カーネルから直接アクセスできます。

2番目の方法を使用して、動的にサイズ設定された（ただし静的に配置された）配列を作成することは可能ですが、簡潔にするために、ここでは説明しません。

sgar91によって提供された回答は、これらのアプローチのいずれにも完全には準拠していなかったため、たとえば、ホストコードでデバイスアドレスを使用することについての警告がまだあります（機能しているように見えますが）。

于 2013-01-06T21:58:20.170 に答える