2

ヒストグラムを使用して画像の合計を計算する CUDA 関数を作成しました。

複数の計算機能用にカーネルとラッパー関数をコンパイルしようとしています。

カーネル：

__global__ void calc_hist(unsigned char* pSrc, int* hist, int width, int height, int pitch)
{
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;

#if __CUDA_ARCH__ > 110   //Shared Memory For Devices Above Compute 1.1
    __shared__ int shared_hist[256];
#endif

    int global_tid = yIndex * pitch + xIndex;

    int block_tid = threadIdx.y * blockDim.x + threadIdx.x;

    if(xIndex>=width || yIndex>=height) return;

#if __CUDA_ARCH__ == 110 //Calculate Histogram In Global Memory For Compute 1.1

    atomicAdd(&hist[pSrc[global_tid]],1);   /*< Atomic Add In Global Memory */

#elif __CUDA_ARCH__ > 110   //Calculate Histogram In Shared Memory For Compute Above 1.1

    shared_hist[block_tid] = 0;   /*< Clear Shared Memory */
    __syncthreads();

    atomicAdd(&shared_hist[pSrc[global_tid]],1);    /*< Atomic Add In Shared Memory */
    __syncthreads();

    if(shared_hist[block_tid] > 0)  /* Only Write Non Zero Bins Into Global Memory */
        atomicAdd(&(hist[block_tid]),shared_hist[block_tid]);
#else 
    return;     //Do Nothing For Devices Of Compute Capabilty 1.0
#endif
}

ラッパー機能:

int sum_8u_c1(unsigned char* pSrc, double* sum, int width, int height, int pitch, cudaStream_t stream = NULL)
{

#if __CUDA_ARCH__ == 100
    printf("Compute Capability Not Supported\n");
    return 0;

#else
    int *hHist,*dHist;
    cudaMalloc(&dHist,256*sizeof(int));
    cudaHostAlloc(&hHist,256 * sizeof(int),cudaHostAllocDefault);

    cudaMemsetAsync(dHist,0,256 * sizeof(int),stream);

    dim3 Block(16,16);
    dim3 Grid;

    Grid.x = (width + Block.x - 1)/Block.x;
    Grid.y = (height + Block.y - 1)/Block.y;

    calc_hist<<<Grid,Block,0,stream>>>(pSrc,dHist,width,height,pitch);

    cudaMemcpyAsync(hHist,dHist,256 * sizeof(int),cudaMemcpyDeviceToHost,stream);

    cudaStreamSynchronize(stream);

    (*sum) = 0.0;
    for(int i=1; i<256; i++)
        (*sum) += (hHist[i] * i);

    printf("sum = %f\n",(*sum));

    cudaFree(dHist);
    cudaFreeHost(hHist);

    return 1;
#endif

}

質問1：

用にコンパイルするsm_10場合、ラッパーとカーネルは実行されません。しかし、それは起こりません。ラッパー関数全体が実行されます。出力は を示していますsum = 0.0。

ラッパー関数の先頭にステートメントをCompute Capability Not Supported追加したため、出力が期待されていました。printf

でラッパー関数が実行されないようにするにはどうすればよいsm_10ですか? if ステートメントなどの実行時チェックを追加したくありません。テンプレート メタ プログラミングによって実現できますか?

質問2：

greater than でコンパイルする場合、カーネル呼び出しの後sm_10に追加した場合にのみ、プログラムが正しく実行されます。cudaStreamSynchronizeしかし、同期しないと、出力はsum = 0.0. なぜそれが起こっているのですか？関数を可能な限りホストに対して非同期にしたい。カーネル内の唯一のループをシフトすることは可能ですか?

Windows 8 で GTX460M、CUDA 5.0、Visual Studio 2008 を使用しています。

4

1 に答える 1

2

広告。質問1

ロバートがコメントで説明したように、__CUDA_ARCH__デバイスコードをコンパイルするときにのみ定義されます。明確にするために: nvcc を呼び出すと、コードは 2 回 (CPU 用に 1 回、GPU 用に 1 回) 解析およびコンパイルされます。の存在を__CUDA_ARCH__使用して、これら 2 つのパスのどちらが発生したかを確認できます。次に、カーネルで行うように、デバイス コードについて、どの GPU をターゲットにしているかを確認できます。

ただし、ホスト側にとっては、すべてが失われるわけではありません。がなくても、 GPU に関する多くの情報を返す__CUDA_ARCH__API 関数cudaGetDevicePropertiesを呼び出すことができます。特に、 Compute Capability を示すフィールドmajorに関心を持つことができます。minor注 - これは前処理段階ではなく実行時に行われるため、同じ CPU コードがすべての GPU で機能します。

広告。質問2

カーネル呼び出しcudaMemoryAsyncは非同期です。これは、呼び出しcudaStreamSynchronize(または類似) を行わない場合、GPU が作業を完了していなくても、フォローアップの CPU コードが引き続き実行されることを意味します。これは、ループ内で操作を開始したときに、コピー元のデータがまだ存在dHistしhHistない可能性があることを意味します。hHistカーネルからの出力で作業したい場合は、カーネルが終了するまで待つ必要があります。

cudaMemcpy( without Async) には暗黙的な同期が含まれていることに注意してください。

于 2012-11-18T08:22:55.503 に答える