CUDA Advent Calendar 2019 15日目の記事です.
今日も小ネタ
逆も然りで,__device__修飾したconsexpr関数をhost側から呼ぶこともできます.
誰かがCUDAのカーネル関数内で使うことを想定せずに書いたconstexpr関数をkernel関数内で使えるようになります.
nvccがC++17に対応してくれれば,例えばVecPP/cste_mathとかが使えるようになり便利そうですね.
今日も小ネタ
例
たとえばこんなコードがあったとしましょう. これを
nvcc text.cuとするとhost用の関数であるadd_2をdeviceから呼んでいるわけですから当然怒られます.
error: calling a constexpr __host__ function("add_2") from a __global__ function("kernel") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.しかし,ここにも書かれている通り試験的にではありますが--expt-relaxed-constexprをつけることでconstexprのhostの関数をdeviceで使えるようになります.
逆も然りで,__device__修飾したconsexpr関数をhost側から呼ぶこともできます.
何が嬉しいの?
世の中はconstexpr関数で満ち溢れています.誰かがCUDAのカーネル関数内で使うことを想定せずに書いたconstexpr関数をkernel関数内で使えるようになります.
nvccがC++17に対応してくれれば,例えばVecPP/cste_mathとかが使えるようになり便利そうですね.
参考
Constexpr functions and function templatesカテゴリー:CUDA
記事作成日:2019-12-15