Как работает CUDA cudaMemcpyFromSymbol?

Question

Я понимаю концепцию передачи символа, но задавался вопросом, что именно происходит за кулисами. Если это не адрес переменной, то что это?

Answer 1

Я считаю, что детали для каждой переменной __device__, cudafe создает нормальную глобальную переменную, как в C, а также переменную PTX, специфичную для CUDA. Глобальная переменная C используется так, что хост-программа может ссылаться на переменную по ее адресу, а переменная PTX используется для фактического хранения переменной. Наличие переменной хоста также позволяет компилятору хоста успешно анализировать программу. Когда программа устройства выполняется, она работает с переменной PTX, когда она манипулирует переменной по имени.

Если вы написали программу для печати адреса __device__ переменных, адрес будет отличаться в зависимости от того, печатать ли вы его от хоста или устройств:

#include <cstdio> 

__device__ int device_variable = 13; 

__global__ void kernel() 
{ 
    printf("device_variable address from device: %p\n", &device_variable); 
} 

int main() 
{ 
    printf("device_variable address from host: %p\n", &device_variable); 

    kernel<<<1,1>>>(); 
    cudaDeviceSynchronize(); 

    return 0; 
} 

$ nvcc test_device.cu -run 
device_variable address from host: 0x65f3e8 
device_variable address from device: 0x403ee0000 

Поскольку ни процессор соглашается на адрес переменные, что делает копирование на него проблематично, да и __host__ функции не имеет доступа к __device__ переменным непосредственно:

__device__ int device_variable; 

int main() 
{ 
    device_variable = 13; 

    return 0; 
} 

$ nvcc warning.cu 
error.cu(5): warning: a __device__ variable "device_variable" cannot be directly written in a host function 

cudaMemcpyFromSymbol позволяет копировать данные обратно из переменной __device__, если программист случайно знает (искаженное) имя переменной в исходной программе.

cudafe облегчает это, создавая сопоставление от искаженных имен к адресам устройств в переменных во время инициализации программы. Программа обнаруживает адрес устройства каждой переменной, запрашивая драйвер CUDA для токена драйвера, учитывая его искаженное имя.

Так что реализация cudaMemcpyFromSymbol будет выглядеть примерно так в псевдокоде:

std::map<const char*, void*> names_to_addresses; 

cudaError_t cudaMemcpyFromSymbol(void* dst, const char* symbol, size_t count, size_t offset, cudaMemcpyKind kind) 
{ 
    void* ptr = names_to_addresses[symbol]; 

    return cudaMemcpy(dst, ptr + offset, count, kind); 
} 

Если посмотреть на выходе nvcc --keep, вы можете увидеть для себя так, что программа взаимодействует со специальным API, CUDART, которые как правило, не доступны для создания отображения:

$ nvcc --keep test_device.cu 
$ grep device_variable test_device.cudafe1.stub.c 
static void __nv_cudaEntityRegisterCallback(void **__T22) { __nv_dummy_param_ref(__T22); __nv_save_fatbinhandle_for_managed_rt(__T22); __cudaRegisterEntry(__T22, ((void (*)(void))kernel), _Z6kernelv, (-1)); __cudaRegisterVariable(__T22, __shadow_var(device_variable,::device_variable), 0, 4, 0, 0); } 

Если вы проверяете выход, вы можете увидеть, что cudafe вставил вызов __cudaRegisterVariable создать отображение для device_variable. Пользователи не должны пытаться использовать этот API самостоятельно.