cuda - カーネル内から割り当てられた動的共有メモリの量を取得できますか?

Question

ホスト側では、カーネルを起動する予定の動的共有メモリの量を保存して使用できます。それを引数としてカーネルに渡すこともできます。しかし、ホスト側の助けなしに、デバイスコードから直接取得する方法はありますか? つまり、カーネルのコードで、実行時に使用可能な動的共有メモリの量を決定しますか?

Answer 1

はい、その値を保持する特別なレジスタがあります。という名前%dynamic_smem_sizeです。インラインPTXを getter 関数でラップすることにより、CUDA C/C++ コードでこのレジスタの値を取得できます。

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

同様に、割り当てられた共有メモリ (静的 + 動的) の合計サイズを register から取得できます%total_smem_size。