3

メモリへの標準アクセスと 1Dtexture アクセスの違いを測定しています。そのために、2 つのカーネルを作成しました。

__global__ void texture1D(float* doarray,int size)
{
  int index;
  //calculate each thread global index
  index=blockIdx.x*blockDim.x+threadIdx.x;
  //fetch global memory through texture reference
  doarray[index]=tex1Dfetch(texreference,index);
  return;
}
__global__ void standard1D(float* diarray, float* doarray, int size)
{
  int index;
  //calculate each thread global index
  index=blockIdx.x*blockDim.x+threadIdx.x;
  //fetch global memory through texture reference
  doarray[index]= diarray[index];
  return;
}

次に、各カーネルを呼び出して、所要時間を測定します。

//copy array from host to device memory
  cudaMemcpy(diarray,harray,sizeof(float)*size,cudaMemcpyHostToDevice);

  checkCuda( cudaEventCreate(&startEvent) );
  checkCuda( cudaEventCreate(&stopEvent) );
  checkCuda( cudaEventRecord(startEvent, 0) );

  //bind texture reference with linear memory
  cudaBindTexture(0,texreference,diarray,sizeof(float)*size);

  //execute device kernel
  texture1D<<<(int)ceil((float)size/threadSize),threadSize>>>(doarray,size);

  //unbind texture reference to free resource
  cudaUnbindTexture(texreference);

  checkCuda( cudaEventRecord(stopEvent, 0) );
  checkCuda( cudaEventSynchronize(stopEvent) );

  //copy result array from device to host memory
  cudaMemcpy(horray,doarray,sizeof(float)*size,cudaMemcpyDeviceToHost);

  //check result
  checkResutl(horray, harray, size);

  cudaEvent_t startEvent2, stopEvent2;
  checkCuda( cudaEventCreate(&startEvent2) );
  checkCuda( cudaEventCreate(&stopEvent2) );
  checkCuda( cudaEventRecord(startEvent2, 0) );
  standard1D<<<(int)ceil((float)size/threadSize),threadSize>>>(diarray,doarray,size);
  checkCuda( cudaEventRecord(stopEvent2, 0) );
  checkCuda( cudaEventSynchronize(stopEvent2) );

  //copy back to CPU
  cudaMemcpy(horray,doarray,sizeof(float)*size,cudaMemcpyDeviceToHost);

結果を出力します:

  float time,time2;
  checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
  checkCuda( cudaEventElapsedTime(&time2, startEvent2, stopEvent2) );
  printf("Texture  bandwidth (GB/s): %f\n",bytes * 1e-6 / time);
  printf("Standard bandwidth (GB/s): %f\n",bytes * 1e-6 / time2);

割り当てている配列のサイズ ( size) に関係なく、標準の帯域幅は常にはるかに高いことがわかります。それが想定されている方法ですか、それともある時点でそれを台無しにしていますか? テクスチャ メモリ アクセスについての私の理解では、グローバル メモリ アクセスを高速化できるということでした。

4

1 に答える 1

9

1D 複素数値関数の補間について、グローバル メモリとテクスチャ メモリ (キャッシュ目的のみに使用され、フィルタリングには使用されません) を比較しました。

私が比較しているカーネルは4、2グローバル メモリを使用するカーネルと2テクスチャ メモリを使用するカーネルです。1 float2これらは、複雑な値にアクセスする方法 (または)に従って区別され、2 floats以下に報告されます。誰かが批判をしたり、独自のテストを実行したりする場合に備えて、完全な Visual Studio 2010 をどこかに投稿します。

__global__ void linear_interpolation_kernel_function_GPU(float* __restrict__ result_d, const float* __restrict__ data_d, const float* __restrict__ x_out_d, const int M, const int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        float reg_x_out = x_out_d[j/2]+M/2;
        int k = __float2int_rz(reg_x_out);
        float a = reg_x_out - __int2float_rz(k);
        float dk = data_d[2*k+(j&1)];
        float dkp1 = data_d[2*k+2+(j&1)];
        result_d[j] = a * dkp1 + (-dk * a + dk);
    } 
}

__global__ void linear_interpolation_kernel_function_GPU_alternative(float2* __restrict__ result_d, const float2* __restrict__ data_d, const float* __restrict__ x_out_d, const int M, const int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        float reg_x_out = x_out_d[j]+M/2;
        int k = __float2int_rz(reg_x_out);
        float a = reg_x_out - __int2float_rz(k);
        float2 dk = data_d[k];
        float2 dkp1 = data_d[k+1];
        result_d[j].x = a * dkp1.x + (-dk.x * a + dk.x);
        result_d[j].y = a * dkp1.y + (-dk.y * a + dk.y);
    } 
}

__global__ void linear_interpolation_kernel_function_GPU_texture(float2* __restrict__ result_d, const float* __restrict__ x_out_d, const int M, const int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        float reg_x_out = x_out_d[j]+M/2;
        int k = __float2int_rz(reg_x_out);
        float a = reg_x_out - __int2float_rz(k);
        float2 dk = tex1Dfetch(data_d_texture,k);
        float2 dkp1 = tex1Dfetch(data_d_texture,k+1);
        result_d[j].x = a * dkp1.x + (-dk.x * a + dk.x);
        result_d[j].y = a * dkp1.y + (-dk.y * a + dk.y);
    } 
}

__global__ void linear_interpolation_kernel_function_GPU_texture_alternative(float* __restrict__ result_d, const float* __restrict__ x_out_d, const int M, const int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        float reg_x_out = x_out_d[j/2]+M/4;
        int k = __float2int_rz(reg_x_out);
        float a = reg_x_out - __int2float_rz(k);
        float dk = tex1Dfetch(data_d_texture2,2*k+(j&1));
        float dkp1 = tex1Dfetch(data_d_texture2,2*k+2+(j&1));
        result_d[j] = a * dkp1 + (-dk * a + dk);
    } 
}

GeForce GT540M (cc 2.1)、Tesla C2050 (cc 2.0)、Kepler K20c (cc 3.5)、および GT210 (cc 1.2) の 4 つの異なる GPU を検討しました。結果は以下の図で報告されています。おわかりのように、古い計算機能でテクスチャをキャッシュとして使用すると、グローバル メモリの使用よりも改善されますが、2 つのソリューションは最新のアーキテクチャではほぼ同等です。

もちろん、この例はすべてを網羅しているわけではなく、実際には、特定のアプリケーションで前者または後者を優先する必要がある他のケースが存在する可能性があります。

ps 処理時間の単位は [ms] であり、図のラベルに示されている [s] ではありません。

GT210 テスラ C2050 GeForce GT540M ケプラー K20c

于 2013-11-08T23:38:06.503 に答える