Weist CudaMallocManaged Speicher auf dem Gerät zu?

Question

Ich verwende Unified Memory, um den Zugriff auf Daten auf der CPU und der GPU zu vereinfachen. Soweit ich weiß, sollte cudaMallocManaged Speicher auf dem Gerät reservieren. Ich schrieb einen einfachen Code zu überprüfen, dass:

#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; 
} 

So bin ich nicht einmal die Daten auf der CPU-Zugriff auf alle Seitenfehler zu vermeiden, so dass der Speicher angeblich auf dem Gerätespeicher sein sollte. Allerdings, wenn ich nvprof auf diesem Code ausführen, erhalte ich folgende Ergebnisse:

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 

So offensichtlich das Array auf den Systemspeicher und nicht der Gerätespeicher zugeordnet ist. Was fehlt mir hier?

Answer 1

Verwalteter Speicher weist tatsächlich physischen Speicher auf der GPU zu. Sie können sich selbst bestätigen dies der Fall ist durch so etwas wie die folgenden, um Ihren Code zu tun:

#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(); 
} 

die jetzt 16Mb für jede der zwei gemanagter Allokation ist die Zuteilung und Freigabe sie dann. Es findet kein Host- oder Gerätezugriff statt, daher sollten keine ausgelösten Übertragungen oder Synchronisation ausgelöst werden. Die Größe ist groß genug, um die minimale Granularität des GPU-Speichermanagers zu überschreiten und Änderungen im sichtbaren freien Speicher auszulösen. Kompilieren und Ausführen es tut dies:

$ 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 

Die physikalische freie Speicherplatz auf der GPU deutlich erhöht wird und erniedrigte durch 16Mb an jedem Alloc/frei.