2

我正在尝试将一个featWidth * featHeight * 31立方体与另一个modelWidth * modelHeight * 31立方体“卷积”。问题是这个内核非常慢（嗯，我设法比顺序 CPU 代码快，但与 OpenMP 版本一样慢）。我正在使用 Quadro FX 1800（是的，64 个 CUDA 核心......）。

__constant__ float d_model[31*22*22];
#define IMUL(a,b) ( __mul24((a), (b)) )
#define IMAD(a,b,c) ( __mul24((a), (b)) + (c) )
__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight)
{
  const int x = IMAD(blockIdx.x, blockDim.x, threadIdx.x);
  const int y = IMAD(blockIdx.y, blockDim.y, threadIdx.y);
  if(x < scoreWidth && y < scoreHeight)
  {
   const int scoreIdx = IMAD(x, scoreHeight, y);
   score[scoreIdx] = 0.f;
   const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y);
   for(int z = 0; z < 31; ++z)
   {
     // Index positionning
     int featIdx =  IMAD(z, IMUL(featWidth,featHeight), baseFeatIdx);
     int modelIdx = IMUL(z, IMUL(modelWidth,modelHeight));

     float value = 0.f;

     // filter
     for(int xx=0; xx<modelWidth; xx++)
     {
       const int xxmodelIdx = IMAD(xx, modelHeight, modelIdx);
       const int xxfeatIdx = IMAD(xx, featHeight, featIdx);
       for(int yy=0; yy<modelHeight; yy++)
       {
         value += d_model[xxmodelIdx+yy] * tex1Dfetch(texFeatures,xxfeatIdx+yy);
       }
     }
     score[scoreIdx] += value;
  }
 }
}

无论如何，我用8*8块中的线程启动这个内核，网格大小为(scoreWidth/8)*(scoreHeight/8)(scoreWidth 和 scoreHeight 是生成的矩阵大小)。我想知道你是否知道我的代码出了什么问题或者什么地方比较慢。

编辑：

一个更快的版本（480 毫秒的进程下降了 150 毫秒！）感谢 tera：

__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight)
{
    const int y = IMUL(4,IMAD(blockIdx.x, blockDim.x, threadIdx.x));
    const int x = IMAD(blockIdx.y, blockDim.y, threadIdx.y);
    if(x < scoreWidth && y < scoreHeight)
    {
    const int scoreIdx = IMAD(x, scoreHeight, y);
    const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y);
    float value=0.f, value1 = 0.f, value2 = 0.f, value3 = 0.f;
    float feat,feat1,feat2,feat3;

    // Index positionning
    int featIdx =  0;
    int modelIdx = 0;
    int xxmodelIdx;
    int xxfeatIdx; 
    float val;
    for(int z = 0; z < 31; ++z)
    {
        featIdx = IMAD(z,IMUL(featWidth,featHeight),baseFeatIdx);
        modelIdx = IMUL(z,IMUL(modelWidth,modelHeight));

        // filter
        for(int xx=0; xx<modelWidth; xx++)
        {
            xxmodelIdx  = IMAD(xx, modelHeight, modelIdx);
            xxfeatIdx = IMAD(xx, featHeight, featIdx);
            feat=tex1Dfetch(texFeatures,xxfeatIdx+0);
            feat1=tex1Dfetch(texFeatures,xxfeatIdx+1);
            feat2=tex1Dfetch(texFeatures,xxfeatIdx+2);
            feat3=tex1Dfetch(texFeatures,xxfeatIdx+3);
            for(int yy=0; yy<modelHeight; yy++)
            {
                val = d_model[xxmodelIdx+yy];
                value += val * feat;
                value1 += val * feat1;
                value2 += val * feat2;
                value3 += val * feat3;
                feat = feat1;
                feat1 = feat2;
                feat2 = feat3;
                feat3 = tex1Dfetch(texFeatures,xxfeatIdx+yy+4);
            }
        }
    }
    score[scoreIdx] = value;
    if(y+1 < scoreHeight)
        score[scoreIdx+1] = value1;
    if(y+2 < scoreHeight)
        score[scoreIdx+2] = value2;
    if(y+3 < scoreHeight)
        score[scoreIdx+3] = value3;
}

以此启动dim3 threads(16,16); dim3 grid(divup(scoreHeight,64), divup(scoreWidth,16));。

4

1 回答 1

1

探查器说什么？NVidia NSight（适用于 Windows 上的 Visual Studio 和 Linux 上的Eclipse的插件）允许您两个查看停顿位置并提供各种提示以优化性能。

我的猜测（没有查看分析器）是你拥有的块太小了。warp 内部有 32 个线程，这是基本的调度元素。NVIDIA GPU 速度很快，因为它可以通过在当前线程执行上一条指令时在其他线程上操作来隐藏延迟。虽然每个 SM（在特斯拉和费米上）或 16 个（在开普勒上）可能有 8 个块，但在峰值处仍然有 16-32 个扭曲，这可能非常小（我可能错了，但启动块有一定的延迟）。我会考虑使用更大的块。

如果我正确理解代码，则纹理获取不是最佳的 - warp 中的线程与 in 不同modelHeight - 1，baseFeatId因此在featIdxand中xxfeatIdx。因此，纹理提取是完全随机的，它不利用数据局部性。倒车x会y让它更有效率。

然而，好的规则是检查分析器 - 如果您的问题是 GPU 上的计算限制，那么您应该专注于计算方面。如果您的问题是内存限制，您应该查看内存访问模式。可能还有其他几个部分看起来像是要优化的地方，但直到你看到瓶颈是什么你才会知道。一旦你知道了，你可能想阅读关于最佳实践指南的特定章节。

于 2013-04-19T11:51:54.370 回答