CudaMallocManaged выделяет память на устройстве?

Question

Я использую Unified Memory для упрощения доступа к данным на CPU и GPU. Насколько мне известно, cudaMallocManaged должен выделять память на устройстве. Я написал простой код, чтобы проверить, что:

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

Так что я даже не доступ к данным на CPU, чтобы избежать каких-либо ошибок страницы, так что память должна быть предположительно в памяти устройства. Однако, когда я бегу nvprof на этом коде, я получаю следующие результаты:

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 

Таким образом, очевидно массив выделяется на системной памяти, а не в памяти устройства. Что мне здесь не хватает?

Answer 1

Управляемая память действительно выделяет физическую память на GPU. Вы можете подтвердить себя это дело, делая что-то вроде следующего к коду:

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

Который сейчас выделяющий 16Mb для каждого из двух управляемых распределений, а затем освободить их. Отсутствует доступ к хосту или устройству, поэтому не должно быть никаких переключений или синхронизации. Размер достаточно велик, чтобы он превышал минимальную гранулярность диспетчера памяти GPU и вызывал изменения в видимой свободной памяти. Компиляция и запуск его делает это:

$ 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 

Физический объем свободной памяти на GPU явно инкрементируется и уменьшается на 16Мб на каждой Alloc/свободный.