Могу ли я получить объем выделенной динамической разделяемой памяти из ядра?

Question

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

Answer 1

Да, есть специальный регистр, содержащий это значение. названный %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,