我认为全局内存的最大大小应仅受GPU设备的限制,无论是使用 __device__ __manged__
静态分配还是使用 cudaMalloc
动态分配 .
但我发现如果使用 __device__ manged__
方式,我可以声明的最大数组大小远小于GPU设备限制 .
最小的工作示例如下:
#include <stdio.h>
#include <cuda_runtime.h>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
#define MX 64
#define MY 64
#define MZ 64
#define NX 64
#define NY 64
#define M (MX * MY * MZ)
__device__ __managed__ float A[NY][NX][M];
__device__ __managed__ float B[NY][NX][M];
__global__ void swapAB()
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
for(int j = 0; j < NY; j++)
for(int i = 0; i < NX; i++)
A[j][i][tid] = B[j][i][tid];
}
int main()
{
swapAB<<<M/256,256>>>();
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
return 0;
}
它使用 64 ^5 * 2 * 4 / 2^30 GB = 8 GB
全局内存,我将运行编译并在具有12GB全局内存的Nvidia Telsa K40c GPU上运行它 .
编译器cmd:
nvcc test.cu -gencode arch=compute_30,code=sm_30
输出警告:
warning: overflow in implicit constant conversion.
当我运行生成的可执行文件时,错误说:
GPUassert: an illegal memory access was encountered test.cu
令人惊讶的是,如果我通过 cudaMalloc
API使用相同大小(8GB)的动态分配的全局内存,则不会出现编译警告和运行时错误 .
我想知道CUDA中静态全局设备内存的可分配大小是否有任何特殊限制 .
谢谢!
PS:OS和CUDA:CentOS 6.5 x64,CUDA-7.5 .
1 回答
这似乎是CUDA运行时API的限制 . 根本原因是这个功能(在CUDA 7.5中):
它只接受任何静态声明的设备变量大小的signed int . 这会将最大大小限制为2 ^ 31(2147483648)个字节 . 您看到的警告是因为CUDA前端正在发出包含对
__cudaResgisterVar
的调用的样板代码,如下所示:这是问题的根源是4294967296 . 大小将溢出有符号整数并导致API调用爆炸 . 因此,目前看来你的每个静态变量仅限于2Gb . 如果这对您的应用程序来说是一个严重的问题,我建议将其作为NVIDIA的一个错误 .