问 在cuda中使用静态锁定内存时，全局设备内存大小限制

我认为全局内存的最大大小应该仅受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 全局内存，我将在有12 on全局内存的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.