1

カーネル（内部または外部）に共有メモリを割り当てて、カーネルから呼び出される他のデバイス関数で使用することは可能ですか？私にとって特に興味深いのは、返されたパラメーター/配列としてそれを使用できるかどうか/どのようにできるかです。

デバイス機能の入力パラメータとして共有メモリを使用することは問題ないようです（少なくとも、問題、エラー、または予期しない結果は発生しません。

これを戻りパラメーターとして使用すると、いくつかの問題が発生します。

  • プログラムがデバッグ構成からビルドされたときに実行できます。

  • しかし、デバッグできません->共有メモリを使用するとデバイス機能でクラッシュします

  • また、アドレスが範囲外であり、共有アドレス空間から読み取られたため、 cuda-memchecker->無効な読み取りでエラーが発生します__global__

では、デバイス関数からカーネルに配列を返すために共有メモリを使用することは可能ですか？

編集：

私が行った他のエラーを除外するために、非常に簡単な例を書きました。

#define CUDA_CHECK_RETURN(value) {                                      \
    cudaError_t _m_cudaStat = (value);                                  \
    if (_m_cudaStat != cudaSuccess) {                                   \
        printf( "Error %s at line %d in file %s\n",                     \
                cudaGetErrorString(_m_cudaStat), __LINE__, __FILE__);   \
        exit(-1);                                                       \
    } }

__device__ void Function( const int *aInput, volatile int *aOutput )
{
    for( int i = 0; i < 10; i++ )
        aOutput[i] = aInput[i] * aInput[i];
}

__global__ void Kernel( int *aInOut )
{
     __shared__ int aShared[10];

    for(int i=0; i<10; i++)
        aShared[i] = i+1;

    Function( aShared, aInOut );
}

int main( int argc, char** argv )
{
    int *hArray = NULL;
    int *dArray = NULL;

    hArray = ( int* )malloc( 10*sizeof(int) );
    CUDA_CHECK_RETURN( cudaMalloc( (void**)&dArray, 10*sizeof(int) ) );

    for( int i = 0; i < 10; i++ )
            hArray[i] = i+1;

    CUDA_CHECK_RETURN( cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice ) );
    cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice );

    Kernel<<<1,1>>>( dArray );

    CUDA_CHECK_RETURN( cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost ) );
    cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost );

    free( hArray );
    CUDA_CHECK_RETURN( cudaFree( dArray ) );
    cudaFree( dArray );

    return 0;
}

ブロックごとに1つのスレッドブロックと1つのスレッドでカーネルを実行します。プログラムを作成して実行することは問題ありません。期待通りの結果が得られました。しかし、プログラムがcuda-memcheckerを使用してtestetである場合、カーネルを終了し、次のログが表示されます。

Error unspecified launch failure at line 49 in file ../CuTest.cu
========= Invalid __global__ read of size 4
=========     at 0x00000078 in /home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:14:Function(int const *, int volatile *)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x01000000 is out of bounds
=========     Device Frame:/home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:25:Kernel(int*) (Kernel(int*) : 0xd0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x34b) [0x55d0b]
=========     Host Frame:/usr/lib/libcudart.so.5.0 [0x8f6a]
=========
========= Program hit error 4 on CUDA API call to cudaMemcpy 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/libcuda.so [0x24e129]
=========     Host Frame:/usr/lib/libcudart.so.5.0 (cudaMemcpy + 0x2bc) [0x3772c]
=========     Host Frame:[0x5400000]
=========
========= ERROR SUMMARY: 2 errors

共有メモリを調整する必要がありますか、何か他のことをする必要がありますか、それとも無視できますか？そうは思いませんか？

4

2 に答える 2

1

ここで説明したように、それは単なるドライバーの問題でした。現在のものに更新した後、すべてが正常に機能しています。

于 2013-01-16T13:56:19.633 に答える
1

CUDA 5.0 インストール ファイル /usr/local/cuda-5.0/samples/6_Advanced/reduction/doc/reduction.ppt を参照してください。

sdatadevice function のローカル変数ですwarpReduce()。共有メモリのアドレスを格納します。共有メモリは、デバイス関数内の addr によって読み書きできます。最終的なリダクション結果は、外部の共有メモリから読み取られますwarpReduce()

template <unsigned int blockSize>
__device__ void warpReduce(volatile int *sdata, unsigned int tid) {
    if (blockSize >=  64) sdata[tid] += sdata[tid + 32];
    if (blockSize >=  32) sdata[tid] += sdata[tid + 16];
    if (blockSize >=  16) sdata[tid] += sdata[tid +  8];
    if (blockSize >=   8) sdata[tid] += sdata[tid +  4];
    if (blockSize >=   4) sdata[tid] += sdata[tid +  2];
    if (blockSize >=   2) sdata[tid] += sdata[tid +  1];
}
template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) {
    extern __shared__ int sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockSize*2) + tid;
    unsigned int gridSize = blockSize*2*gridDim.x;
    sdata[tid] = 0;

    while (i < n) { sdata[tid] += g_idata[i] + g_idata[i+blockSize];  i += gridSize;  }
    __syncthreads();

    if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
    if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
    if (blockSize >= 128) { if (tid <  64) { sdata[tid] += sdata[tid +  64]; } __syncthreads(); }

    if (tid < 32) warpReduce(sdata, tid);
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
于 2013-01-14T12:58:57.590 に答える