0

特定の入力に対して cuda-gdb を使用して実行すると、Device Illegal Address が原因でカーネルの起動に失敗するコードを書きました。cuda-memcheck を使って実行したところ、Invalid write of size 4 エラーが発生しました。コードが大きすぎるので、ここでシナリオを説明します。

スタックとして機能する配列ポインターを渡すメインカーネルがあります。メインカーネルから呼び出され、スタックを使用するデバイス関数があります。

__device__ void find(int v , int* p, int* pv,int n, int* d_stackContents)
{

    int d_stackTop;
    d_stackTop = -1;
    *pv = p[v];
    if(*pv == -1){
            *pv = v;

    }
    else{
    cuPrintf("Stack top is %d\n",d_stackTop);
    d_stackTop = d_stackTop + 1;
    d_stackContents[d_stackTop] = v;
    cuPrintf("Stack top is %d\n",d_stackTop);
    while(*pv != -1){
            d_stackTop = d_stackTop + 1;
            d_stackContents[d_stackTop] = *pv;
            cuPrintf("Stack top is %d\n",d_stackTop);
            *pv = p[*pv];
    }

}

d_stackContents[d_stackTop] = *pv; でエラーが発生しています。

次のように、メイン カーネルでデバイス関数を呼び出しています。

find(v[idx], p,&pv,n, d_stackContents);

ここで、idx = threadIdx.x + blockDim.x * blockIdx.x で、pv を int pv; として宣言しました。

また、d_stackContents 配列は cudaMalloc を使用して main に割り当てられ、メイン カーネルに引数として渡されます。

4

1 に答える 1

4

単一ブロック内の単一スレッドでカーネルを呼び出さない限り、これは機能しません。そうしないと、すべてのスレッドが互いのスタックに走り書きします。次に、破損したスタックに格納されたポインターを逆参照すると、コードが不正なアドレスにアクセスしようとする理由がすぐに説明されます。

スレッドごとに個別のスタックを使用するか、アトミック操作によってのみ操作されるグローバル メモリ内のスタック ポインターを持つ単一のスタックを使用する必要があります。

于 2012-10-26T10:07:38.070 に答える