內存要求CUDA

Question

我最近寫了一個非常簡單的內核：

__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'中。

我的第一個測試不能正常工作。所以，爲了弄清楚發生了什麼，我開始評論這筆錢的一些內容。一旦我減少了操作次數，它就開始工作了。

我的理論是，它可能與可用內存有關，以存儲表達式的中間結果。因此，減少每個塊的線程數量，它開始工作完美，無需減少操作數量。

基於這樣的經驗，我想知道如何更好地估計每個塊的線程數量，以避免內存要求超過我的可用空間。我怎麼知道上面的操作需要多少內存？ （當我們處於這種狀態時，它是什麼樣的內存？緩存，共享內存等）。

謝謝！

Answer 1

這大多數情況下可能是寄存器，您可以通過將-Xptxas="-v"選項添加到編譯內核的nvcc調用中來找出每個線程的寄存器消耗。彙編器將返回每個線程的寄存器數量，靜態共享內存，本地內存和編譯代碼使用的常量內存。

NVIDIA製作佔位計算器電子表格（available here），您可以在其中插入彙編器的輸出以查看塊大小的可行範圍及其對GPU佔用率的影響。 CUDA編程指南的第3章還詳細討論了佔用概念以及塊大小和內核資源需求如何相互作用。