CudaMallocManaged alloue-t-il de la mémoire sur le périphérique?

Question

J'utilise la mémoire unifiée pour simplifier l'accès aux données sur le processeur et le GPU. Pour autant que je sache, cudaMallocManaged devrait allouer de la mémoire sur le périphérique. J'ai écrit un code simple pour vérifier que:

#define TYPE float 
#define BDIMX 16 
#define BDIMY 16 
#include <cuda.h> 
#include <cstdio> 
#include <iostream> 
__global__ void kernel(TYPE *g_output, TYPE *g_input, const int dimx, const int dimy) 
{ 
__shared__ float s_data[BDIMY][BDIMX]; 
    int ix = blockIdx.x * blockDim.x + threadIdx.x; 
    int iy = blockIdx.y * blockDim.y + threadIdx.y; 
    int in_idx = iy * dimx + ix; // index for reading input 
    int tx = threadIdx.x; // thread’s x-index into corresponding shared memory tile 
    int ty = threadIdx.y; // thread’s y-index into corresponding shared memory tile 
    s_data[ty][tx] = g_input[in_idx]; 
    __syncthreads(); 
    g_output[in_idx] = s_data[ty][tx] * 1.3; 
    } 


int main(){ 
    int size_x = 16, size_y = 16; 
    dim3 numTB; 
    numTB.x = (int)ceil((double)(size_x)/(double)BDIMX) ; 
    numTB.y = (int)ceil((double)(size_y)/(double)BDIMY) ; 
    dim3 tbSize; 
    tbSize.x = BDIMX; 
    tbSize.y = BDIMY; 
    float* a,* a_out; 
    cudaMallocManaged((void**)&a,  size_x * size_y * sizeof(TYPE)); 
    cudaMallocManaged((void**)&a_out, size_x * size_y * sizeof(TYPE)); 

    kernel <<<numTB, tbSize>>>(a_out, a, size_x, size_y); 
    cudaDeviceSynchronize(); 
    return 0; 
} 

Je ne suis même pas accéder aux données sur la CPU pour éviter les défauts de page de sorte que la mémoire doit être soi-disant sur la mémoire de l'appareil. Cependant quand je lance nvprof sur ce code, je reçois les résultats suivants:

invocations        Metric Name      Metric Description   Min   Max   Avg 
Device "Tesla K40c (0)" 
Kernel: kernel(float*, float*, int, int) 
     1     local_load_transactions     Local Load Transactions   0   0   0 
     1     local_store_transactions     Local Store Transactions   0   0   0 
     1     shared_load_transactions     Shared Load Transactions   8   8   8 
     1     shared_store_transactions     Shared Store Transactions   8   8   8 
     1       gld_transactions     Global Load Transactions   8   8   8 
     1       gst_transactions     Global Store Transactions   8   8   8 
     1     sysmem_read_transactions   System Memory Read Transactions   32   32   32 
     1     sysmem_write_transactions   System Memory Write Transactions   34   34   34 
     1     tex_cache_transactions    Texture Cache Transactions   0   0   0 
     1     dram_read_transactions   Device Memory Read Transactions   0   0   0 
     1     dram_write_transactions   Device Memory Write Transactions   0   0   0 

donc apparemment le tableau est alloué sur la mémoire du système et non la mémoire de l'appareil. Qu'est-ce que j'oublie ici?

Answer 1

La mémoire gérée alloue réellement de la mémoire physique sur le GPU. Vous pouvez vous confirmer cela est le cas en faisant quelque chose comme ce qui suit à votre code:

#include <iostream> 

void report_gpu_mem() 
{ 
    size_t free, total; 
    cudaMemGetInfo(&free, &total); 
    std::cout << "Free = " << free << " Total = " << total <<std::endl; 
} 

int main() 
{ 
    float* a,* a_out; 
    size_t sz = 1 << 24; // 16Mb 
    report_gpu_mem(); 
    cudaMallocManaged((void**)&a, sz); 
    report_gpu_mem(); 
    cudaMallocManaged((void**)&a_out, sz); 
    report_gpu_mem(); 
    cudaFree(a); 
    report_gpu_mem(); 
    cudaFree(a_out); 
    report_gpu_mem(); 
    return cudaDeviceReset(); 
} 

qui affecte à présent 16Mo pour chacune des deux allocations gérées, puis de les libérer. Aucun accès hôte ou périphérique ne se produit, il ne doit donc pas y avoir de transferts ou de synchronisation déclenchés. La taille est suffisamment grande pour dépasser la granularité minimale du gestionnaire de mémoire GPU et déclencher des changements dans la mémoire libre visible. Et l'exécuter Compiler fait ceci:

$ nvcc -arch=sm_52 sleepy.cu 
$ CUDA_VISIBLE_DEVICES="0" ./a.out 
Free = 4211929088 Total = 4294770688 
Free = 4194869248 Total = 4294770688 
Free = 4178092032 Total = 4294770688 
Free = 4194869248 Total = 4294770688 
Free = 4211654656 Total = 4294770688 

La mémoire physique libre sur le GPU est clairement incrémenté et décrémenté par 16Mo à chaque alloc/gratuit.