使用 CUDA 占用计算器

Question

我正在使用占用计算器，但我无法理解如何获取每个线程的寄存器/每个块的共享内存。我阅读了文档。我使用 Visual Studio。所以在 CUDA 构建规则->命令行->附加选项下的项目属性中我添加--ptxas-options=-v。程序编译得很好。但是我没有看到任何输出。有人可以帮忙吗？ 谢谢

Answer 1

打开此开关后，编译器输出窗口上应该会出现一行，告诉您寄存器的数量和共享内存的数量。
您在编译器输出窗口上看到什么了吗？您可以将其复制并粘贴到问题中吗？
它应该看起来像

ptxas info : Used 3 registers, 2084+1060 bytes smem, 40 bytes cmem[0], 12 bytes cmem[1]

Answer 2

试试这个简单的规则：

内核中的所有局部变量，如 int a、float b 等都存储在寄存器中。仅当代码中的局部变量保持在多处理器中可用寄存器的限制范围内时，See Limits。但是，如果您声明一千个整数，例如 int a[1000]，则 a 不会存储在寄存器中，而是存储在本地内存 (DRAM) 中。 p>

内核代码中使用的共享内存量是共享内存/块。例如，如果您定义__shared__ float shMem[256]，那么您使用的是 256*4(浮点大小) = 1024 字节的共享内存。

以下示例代码（它不能正常工作，只是举例）每个线程使用 9 个 32 位寄存器，它们是：int xIndex、yIndex、Idx、shY、 shX、aLocX、aLocY 和 float t, temp。代码每个块使用 324 字节 的共享内存，因为 BLOCK_DIM = 16。

__global__ void averageFilter (unsigned char * outImage,
                           int imageWidth,
                           int imageHeight,
                           cuviPoint2 loc){


    unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
    unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
    unsigned int Idx = yIndex*imageWidth + xIndex;
    float t = INC;


    if(xIndex>= imageWidth|| yIndex>=imageHeight)
        return;


    else if(xIndex==0 || xIndex== imageWidth-1 || yIndex==0 || yIndex==imageHeight-1){

          for (int i=-1; i<=1; i++)
             for (int j=-1; j<=1; j++)
                 t+= tex1Dfetch(texMem,Idx+i*imageWidth+j);
                    outImage[Idx] = t/6;

          }


    __shared__ unsigned char shMem[BLOCK_DIM+2][BLOCK_DIM+2];


    unsigned int shY = threadIdx.y + 1;
    unsigned int shX = threadIdx.x + 1;


   if (threadIdx.x==0 || threadIdx.x==BLOCK_DIM-1 || threadIdx.y==0 || threadIdx.y==BLOCK_DIM-1){


 for (int i=-1; i<=1; i++)
      for (int j=-1; j<=1; j++)
        shMem[shY+i][shX+j]=  tex1Dfetch(texMem,Idx+i*imageWidth+j);

    }
    else
    shMem[shY][shX] =  tex1Dfetch(texMem,Idx);

     __syncthreads();     



if(xIndex==0 || xIndex== imageWidth-1 || yIndex==0 || yIndex==imageHeight-1)
        return;     

  int aLocX = loc.x, aLocY = loc.y;

    float temp=INC;

      for (int i=aLocY; i<=aLocY+2; i++)
         for (int j=aLocX; j<=aLocX+2; j++)
        temp+= shMem[shY+i][shX+j];

        outImage[Idx] = floor(temp/9);

}

Answer 3

shoosh's answer 可能是查找寄存器和共享内存使用情况的最简单方法。确保您首先查看输出窗格（在“查看”下拉菜单中选择“输出”），然后重新编译。编译器应该在输出窗格中为您提供所有内核的 ptxas 信息，如下图所示...

Answer 4

另一种查找此信息的方法是使用视觉分析器或 nvidia 的并行 nsight。