每個CUDA線程的本地內存量

我在NVIDIA文檔（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications，表＃12）中讀到我的GPU（GTX 580，計算能力2.0）的每個線程的本地內存量爲512 Ko。每個CUDA線程的本地內存量

我嘗試使用CUDA 6.5在Linux上檢查此限制失敗。

下面是我使用的代碼（其唯一目的是測試本地內存的限制，它沒有做任何有用的計算）：

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

#define MEMSIZE 65000 // 65000 -> out of memory, 60000 -> ok 

inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if(abort) 
      exit(code); 
    } 
} 

inline void gpuCheckKernelExecutionError(const char *file, int line) 
{ 
    gpuAssert(cudaPeekAtLastError(), file, line); 
    gpuAssert(cudaDeviceSynchronize(), file, line);  
} 


__global__ void kernel_test_private(char *output) 
{ 
    int c = blockIdx.x*blockDim.x + threadIdx.x; // absolute col 
    int r = blockIdx.y*blockDim.y + threadIdx.y; // absolute row 

    char tmp[MEMSIZE]; 
    for(int i = 0; i < MEMSIZE; i++) 
     tmp[i] = 4*r + c; // dummy computation in local mem 
    for(int i = 0; i < MEMSIZE; i++) 
     output[i] = tmp[i]; 
} 

int main(void) 
{ 
    printf("MEMSIZE=%d bytes.\n", MEMSIZE); 

    // allocate memory 
    char output[MEMSIZE]; 
    char *gpuOutput; 
    cudaMalloc((void**) &gpuOutput, MEMSIZE); 

    // run kernel 
    dim3 dimBlock(1, 1); 
    dim3 dimGrid(1, 1); 
    kernel_test_private<<<dimGrid, dimBlock>>>(gpuOutput); 
    gpuCheckKernelExecutionError(__FILE__, __LINE__); 

    // transfer data from GPU memory to CPU memory 
    cudaMemcpy(output, gpuOutput, MEMSIZE, cudaMemcpyDeviceToHost); 

    // release resources 
    cudaFree(gpuOutput); 
    cudaDeviceReset(); 

    return 0; 
}

而且編譯命令行：

nvcc -o cuda_test_private_memory -Xptxas -v -O2 --compiler-options -Wall cuda_test_private_memory.cu

編譯就可以了，並報告：

ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z19kernel_test_privatePc' for 'sm_20' 
ptxas info : Function properties for _Z19kernel_test_privatePc 
    65000 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 21 registers, 40 bytes cmem[0]

我在運行時就得到了一個「內存不足」的錯誤GTX 580，當我達到每線程65000字節。下面是該程序的控制檯的確切輸出：

MEMSIZE=65000 bytes. 
GPUassert: out of memory cuda_test_private_memory.cu 48

我也做了測試與GTX 770 GPU（Linux上，CUDA 6.5）。它在沒有錯誤的情況下運行於MEMSIZE = 200000，但是在運行時出現「內存不足錯誤」，出現在MEMSIZE = 250000的位置。

如何解釋這種行爲？難道我做錯了什麼？

來源

2015-03-02 devel484

其CUDA版本之間的最大堆棧大小您使用的？這是Linux還是Windows？當你編譯**代碼或者當你運行**代碼時，你會得到「內存不足錯誤」嗎？（將精確的錯誤文本粘貼到問題中）您用於編譯代碼的命令行是什麼？我的猜測是你正在編譯一個pre-cc2.0架構。如果我爲cc1.1體系結構編譯此代碼，編譯時會出現「內存不足錯誤」，因爲cc1.x設備對每個線程（16KB）的本地內存有更小的限制。如果我爲cc2.0體系結構進行編譯，您的代碼將爲我編譯並正常運行。 – 2015-03-02 14:54:21

您的問題也可能來自以下代碼行：'char output [MEMSIZE];'This（host code）創建基於堆棧的分配，並且這些類型的分配可能會受限於平臺。將確切的錯誤文本粘貼到問題中會有所幫助。（你可以編輯你自己的問題。） – 2015-03-02 15:00:38

@RobertCrovella謝謝你對我的問題感興趣。我編輯了我的問題以添加缺少的信息。運行時由cudaGetErrorString（）報告的確切錯誤文本是「內存不足」。 – devel484 2015-03-03 09:14:03

看來你正在運行到沒有本地存儲限制，但堆棧大小限制：

ptxas信息：對於_Z19kernel_test_privatePc

65000字節的功能特性堆棧幀，0字節溢出店， 0字節溢出加載

在這種情況下，您想要本地化的變量位於（GPU線程）堆棧上。

基於由@njuffa here提供的信息中，可用的堆棧大小限制爲中的較小者：

的最大本地存儲器大小（512KB爲cc2.x和更高）
GPU存儲器/ （＃SMs）/（每SM最大線程數）

顯然，第一個限制不是問題。我假設你有一個「標準」的GTX580，它有1.5GB的內存和16個SM。 cc2.x設備每個多處理器最多有1536個駐留線程。這意味着我們有1536MB/16/1536 = 1MB/16 = 65536字節的堆棧。從總可用內存中減去一些開銷和其他內存使用量，因此堆棧大小限制在65536以下，某種程度上介於60000和65000之間，顯然是這樣。

我懷疑你的GTX770類似的計算將產生類似的結果，即200000及250000

來源

2015-03-03 20:33:47

謝謝你的解釋。我對GTX 770進行了相同的計算（4 GB RAM，8個SM，每個SM最多2048個線程），堆棧大小爲262144字節。這與GTX 770在200000字節和250000字節之間（更精確地說，在242000字節和245000字節之間）遇到的限制是一致的。我還使用malloc（）替換了內核中的靜態分配：然後我可以分配512 K的本地內存（甚至更多，最多7 Mo！）。 – devel484 2015-03-04 09:45:08

如果你需要比in-kernel'malloc'更多的內存，請查看[文檔]（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-內存分配和操作）來提高堆內存分配限制。 – 2015-03-04 14:16:25

好的。我認爲對'malloc'的內核調用會分配**本地**內存。現在我明白了內核中的'malloc'分配了默認大小爲8M的全局內存堆。 – devel484 2015-03-05 13:09:59

每個CUDA線程的本地內存量

回答

相關問題