cudaで静的に割り当てられたメモリを使用する場合のグローバルデバイスのメモリサイズ制限

Question

グローバルメモリの最大サイズは、__device__ __manged__を使用して静的に割り当てられても、cudaMallocを使用して動的に割り当てられる場合でもGPUデバイスによってのみ制限されるべきだと考えました。

しかし、私は__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ギガバイトグローバルメモリを持つNVIDIAのテスラ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）のグローバルメモリを使用する代わりに、com警告とランタイムエラーが発生しました。

CUDAの静的グローバルデバイスメモリの割り当て可能なサイズに関する特別な制限がある場合、私は思ったんだけど。

ありがとうございます！

PS：OSとCUDA：CentOSの6.5 x64の、CUDA-7.5。

Answer 1

これはCUDAランタイムAPIの制限であるように思われます。唯一任意静的に宣言されたデバイス変数のサイズのために署名された整数を受け付ける

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

：根本的な原因は、（CUDA 7.5）この関数です。これにより、最大サイズは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のバグとして提起することをお勧めします。