内存要求 CUDA答案

【问题标题】：Memory requirements CUDA内存要求 CUDA
【发布时间】：2011-07-02 04:52:10
【问题描述】：

我最近写了一个很简单的内核：

__device__ uchar elem(const Matrix m, int row, int col) {
    if(row == -1) {
        row = 0;
    } else if(row > m.rows-1) {
        row = m.rows-1;
    }

    if(col == -1) {
        col = 0;
    } else if(col > m.cols-1) {
        col = m.cols-1;
    }
    return *((uchar*)(m.data + row*m.step + col));
}

/**
* Each thread will calculate the value of one pixel of the image 'res'
*/
__global__ void resizeKernel(const Matrix img, Matrix res) {
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;

    if(row < res.rows && col < res.cols) {
        uchar* e = res.data + row * res.step + col;

        *e = (elem(img, 2*row, 2*col) >> 2) +
             ((elem(img, 2*row, 2*col-1) + elem(img, 2*row, 2*col+1) 
             + elem(img, 2*row-1, 2*col) + elem(img, 2*row+1, 2*col)) >> 3) +
             ((elem(img, 2*row-1, 2*col-1) + elem(img, 2*row+1, 2*col+1)
             + elem(img, 2*row+1, 2*col-1) + elem(img, 2*row-1, 2*col+1)) >> 4);
    }
}

基本上，它的作用是使用较大图像的值来计算缩小图像的像素值。在 resizeKernel 中的 'if' 内。

我的第一次测试没有正常工作。所以，为了找出发生了什么，我开始评论这个总和的一些行。一旦我减少了操作次数，它就开始工作了。

当时我的理论是，它可能与存储表达式中间结果的可用内存有关。因此，减少每个块的线程数，它开始完美地工作，无需减少操作数。

基于此经验，我想知道如何更好地估计每个块的线程数，以避免内存需求超过我可用的内存。我怎么知道上面的操作需要多少内存？（当我们讨论它时，它是什么类型的内存？缓存、共享内存等）。

谢谢！

【问题讨论】：

标签： memory cuda

【解决方案1】：

它很可能是寄存器，您可以通过在编译内核的 nvcc 调用中添加-Xptxas="-v" 选项来找出每个线程的寄存器消耗。汇编器将返回每个线程的寄存器数量、静态共享内存、本地内存和编译代码使用的常量内存。

NVIDIA 制作了一个占用计算器电子表格 (available here)，您可以在其中插入汇编程序的输出，以查看块大小的可行范围及其对 GPU 占用率的影响。 CUDA 编程指南的第 3 章还详细讨论了占用的概念以及块大小和内核资源需求如何相互作用。

【讨论】：