Блог

На стороне хоста я могу сохранить объем динамической общей памяти, с которой я собираюсь запустить ядро, и использовать его. Я даже могу передать это в качестве аргумента ядру. Но — есть ли способ получить его непосредственно из кода устройства, без помощи со стороны хоста? То есть, должен ли код для ядра определять во время выполнения, сколько динамической общей памяти у него есть?

Да, есть специальный регистр, содержащий это значение. имя %dynamic_smem_size . Вы можете получить значение этого регистра в своем коде CUDA C / C , обернув некоторые встроенные PTX функцией получения:

 __device__ unsigned dynamic_smem_size()
{
    unsigned ret; 
    asm volatile ("mov.u32 %0, %dynamic_smem_size;" : "=r"(ret));
    return ret;
}
 

Аналогичным образом вы можете получить общий размер выделенной общей памяти (статической динамической) из регистра %total_smem_size .