0

私は次のコードを持っていますcuda_computation.cu

#include <iostream>
#include <stdio.h>
#include <cuda.h>
#include <assert.h>

void checkCUDAError(const char *msg);

__global__ void euclid_kernel(float *x, float* y, float* f)
{
  int idx = blockIdx.x*blockDim.x + threadIdx.x;
  int i = blockIdx.x;
  int j = threadIdx.x;
  f[idx] = sqrt((x[i]-x[j])*(x[i]-x[j]) + (y[i]-y[j])*(y[i]-y[j]));
}
int main()
{
  float *xh;
  float *yh;
  float *fh;
  float *xd;
  float *yd;
  float *fd;

  size_t n = 256;
  size_t numBlocks = n;
  size_t numThreadsPerBlock = n;

  size_t memSize = numBlocks * numThreadsPerBlock * sizeof(float);
  xh = (float *) malloc(n * sizeof(float));
  yh = (float *) malloc(n * sizeof(float));
  fh = (float *) malloc(memSize);

  for(int ii(0); ii!=n; ++ii)
    {
      xh[ii] = ii;
      yh[ii] = ii;
    }

  cudaMalloc( (void **) &xd, n * sizeof(float) );
  cudaMalloc( (void **) &yd, n * sizeof(float) );
  cudaMalloc( (void **) &fd, memSize );
  for(int run(0); run!=10000; ++run)
    {
      //change value to avoid optimizations
      xh[0] = ((float)run)/10000.0;
      cudaMemcpy( xd, xh, n * sizeof(float), cudaMemcpyHostToDevice );
      checkCUDAError("cudaMemcpy");
      cudaMemcpy( yd, yh, n * sizeof(float), cudaMemcpyHostToDevice );
      checkCUDAError("cudaMemcpy");
      dim3 dimGrid(numBlocks);
      dim3 dimBlock(numThreadsPerBlock);
      euclid_kernel<<< dimGrid, dimBlock >>>( xd, yd, fd );
      cudaThreadSynchronize();
      checkCUDAError("kernel execution");
      cudaMemcpy( fh, fd, memSize, cudaMemcpyDeviceToHost );
      checkCUDAError("cudaMemcpy");
    }
  cudaFree(xd);
  cudaFree(yd);
  cudaFree(fd);
  free(xh);
  free(yh);
  free(fh);
  return 0;
}

void checkCUDAError(const char *msg)
{
  cudaError_t err = cudaGetLastError();
  if( cudaSuccess != err) 
    {
      fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
      exit(-1);
    }                         
}

FX QUADRO 380で実行するには約6インチかかりますが、i7-870コアを1つだけ使用する対応するシリアルバージョンは約3インチかかります。私は何かが恋しいですか？コードはいくつかの点で最適化されていませんか？それとも、単純な計算（このすべてのペアのユークリッド距離など）の場合、メモリを移動するために必要なオーバーヘッドが計算ゲインを超えると予想される動作ですか？

4

3 に答える 3

2

問題を分割して、各ブロックが単一のiとすべての256jを担当するようにします。これらの256jはブロックごとにリロードする必要があり、合計2 * 256 *（256 + 1）のロードになるため、これは悪い局所性です。代わりに、グリッドを分割して、各ブロックがたとえば16iと16jの範囲を担当するようにします。これは、256ブロック*256スレッドのままです。ただし、各ブロックは2 *（16 + 16）の値のみをロードするようになり、合計または2 * 256*32の合計ロードになります。アイデアは、ロードされた各値を可能な限り何度も再利用することです。これは256x256では大きな影響はないかもしれませんが、サイズが大きくなるにつれてますます重要になります。

この最適化は、同様の局所性の問題がある効率的な行列乗算に使用されます。詳細については、 http：//en.wikipedia.org/wiki/Loop_tiling、または「最適化された行列の乗算」についてはgoogleを参照してください。そして、おそらくNVIDIA SDKの行列乗算カーネルは、いくつかの詳細とアイデアを提供します。

于 2012-05-23T13:53:10.593 に答える
2

私はあなたがデータを移動する時までに殺されていると思います。特に、個々の値を使用してCUDAカーネルを呼び出しているため、大量の値のセットを1D配列としてアップロードし、それらを操作する方が速い場合があります。

また、sqrtはCudaのHWでは実行されません（少なくとも私のGPUでは実行されません）が、CPUはこれのためにFPU HWを最適化し、おそらくGPUよりも10倍高速であり、このような小さなジョブでは、おそらくすべての結果をタイミング実行間のキャッシュ。

于 2012-05-22T22:29:02.493 に答える
2

グローバルメモリの読み取りは高価であるため、減らしてください。スレッドごとに4つのグローバルメモリ読み取りがありますが、共有メモリを使用すると2つに減らすことができます。

__global__ void euclid_kernel(const float * inX_g, const float* inY_g, float * outF_g)
{
    const unsigned int threadId = blockIdx.x * blockDim.x + threadIdx.x;

    __shared__ float xBlock_s;
    __shared__ float yBlock_s;

    if(threadIdx.x == 0)
    {
        xBlock_s = inX_g[blockIdx.x];
        yBlock_s = inY_g[blockIdx.x];
    }
    __syncthreads();

    float xSub = xBlock_s - inX_g[threadIdx.x];
    float ySub = yBlock_s - inY_g[threadIdx.x];

    outF_g[threadId] = sqrt(xSub * xSub + ySub * ySub);
}

また、さまざまなブロックサイズでテストする必要があります（100％の占有率がある場合）。

于 2012-05-23T00:10:43.847 に答える