在 cuda 中使用静态分配的内存时的全局设备内存大小限制

Posted 2023-03-23

【中文标题】在 cuda 中使用静态分配的内存时的全局设备内存大小限制

【英文标题】：Global device memory size limit when using statically alocated memory in cuda

【发布时间】：2016-10-16 07:42:42

【问题描述】：

我认为全局内存的最大大小应该只受 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 上运行编译和运行它。

编译器命令：

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：操作系统和 CUDA：CentOS 6.5 x64、CUDA-7.5。

【参考方案1】：

这似乎是 CUDA 运行时 API 的限制。根本原因是这个函数（在 CUDA 7.5 中）：

__cudaRegisterVar(
        void **fatCubinHandle,
        char  *hostVar,
        char  *deviceAddress,
  const char  *deviceName,
        int    ext,
        int    size,
        int    constant,
        int    global
);

它只接受任何静态声明的设备变量大小的有符号整数。这会将最大大小限制为 2^31 (2147483648) 字节。您看到的警告是因为 CUDA 前端正在发出样板代码，其中包含对 __cudaResgisterVar 的调用，如下所示：

__cudaRegisterManagedVariable(__T26, __shadow_var(A,::A), 0, 4294967296, 0, 0);
__cudaRegisterManagedVariable(__T26, __shadow_var(B,::B), 0, 4294967296, 0, 0);

问题的根源在于 4294967296。大小将溢出有符号整数并导致 API 调用崩溃。因此，目前您的每个静态变量似乎仅限于 2Gb。如果这对您的应用程序来说是一个严重的问题，我建议将此作为一个错误向 NVIDIA 提出。

【讨论】：

你是对的，静态内存限制为 2Gb。我会尝试将此作为错误提出。顺便说一句，我想使用静态声明的多维数组的原因是我发现它们对于我的多维物理建模代码效率更高。在 3D 模具中寻址数据元素时，使用 1D 动态分配的数组会产生大量整数运算。

我已经尝试了 CUDA 8.0 RC，错误仍然存​​在。 :(