每个 CUDA 线程的本地内存量
Amount of local memory per CUDA thread
我在 NVIDIA 文档(http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications、table #12)中读到,对于我的 GPU(GTX 580,计算能力 2.0),每个线程的本地内存量是 512 Ko。
我尝试用 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]
当我达到每个线程 65000 字节时,我在 GTX 580 上的运行时遇到 "out of memory" 错误。这是控制台中程序的确切输出:
MEMSIZE=65000 bytes.
GPUassert: out of memory cuda_test_private_memory.cu 48
我还使用 GTX 770 GPU 进行了测试(在 Linux 上使用 CUDA 6.5)。它 运行 对于 MEMSIZE=200000 没有错误,但是对于 MEMSIZE=250000 "out of memory error" 在运行时发生。
如何解释这种行为?难道我做错了什么 ?
看来您 运行 不是本地内存限制而是堆栈大小限制:
ptxas info : Function properties for _Z19kernel_test_privatePc
65000 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
在本例中,您原本打算成为本地变量的变量位于(GPU 线程)堆栈上。
根据@njuffa here 提供的信息,可用堆栈大小限制为以下较小者:
- 最大本地内存大小(cc2.x 及更高版本为 512KB)
- GPU 内存/(#of SM)/(每个 SM 的最大线程数)
显然,第一个限制不是问题。我假设你有一个 "standard" GTX580,它有 1.5GB 内存和 16 个 SM。 cc2.x 设备的每个多处理器最多有 1536 个常驻线程。这意味着我们有 1536MB/16/1536 = 1MB/16 = 65536 字节堆栈。从总可用内存中减去一些开销和其他内存使用量,因此堆栈大小限制低于 65536,显然在您的情况下介于 60000 和 65000 之间。
我怀疑在您的 GTX770 上进行类似的计算会产生类似的结果,即最大堆栈大小在 200000 到 250000 之间。
我在 NVIDIA 文档(http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications、table #12)中读到,对于我的 GPU(GTX 580,计算能力 2.0),每个线程的本地内存量是 512 Ko。
我尝试用 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]
当我达到每个线程 65000 字节时,我在 GTX 580 上的运行时遇到 "out of memory" 错误。这是控制台中程序的确切输出:
MEMSIZE=65000 bytes.
GPUassert: out of memory cuda_test_private_memory.cu 48
我还使用 GTX 770 GPU 进行了测试(在 Linux 上使用 CUDA 6.5)。它 运行 对于 MEMSIZE=200000 没有错误,但是对于 MEMSIZE=250000 "out of memory error" 在运行时发生。
如何解释这种行为?难道我做错了什么 ?
看来您 运行 不是本地内存限制而是堆栈大小限制:
ptxas info : Function properties for _Z19kernel_test_privatePc
65000 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
在本例中,您原本打算成为本地变量的变量位于(GPU 线程)堆栈上。
根据@njuffa here 提供的信息,可用堆栈大小限制为以下较小者:
- 最大本地内存大小(cc2.x 及更高版本为 512KB)
- GPU 内存/(#of SM)/(每个 SM 的最大线程数)
显然,第一个限制不是问题。我假设你有一个 "standard" GTX580,它有 1.5GB 内存和 16 个 SM。 cc2.x 设备的每个多处理器最多有 1536 个常驻线程。这意味着我们有 1536MB/16/1536 = 1MB/16 = 65536 字节堆栈。从总可用内存中减去一些开销和其他内存使用量,因此堆栈大小限制低于 65536,显然在您的情况下介于 60000 和 65000 之间。
我怀疑在您的 GTX770 上进行类似的计算会产生类似的结果,即最大堆栈大小在 200000 到 250000 之间。