トップページ

よおこそ

備忘録的な何かということで始めてみましたが,全然違うことも書いています.

• 記事一覧

• 自己紹介

• ホストサーバーについて

アクセスの多い記事

1. Failed to initialize NVML: Driver/library version mismatch
2. WICTextureLoaderで画像の読み込み
3. VisualStudioで__syncthreads()が使えない
4. thrustを使ってみた
5. xmobarで入力方式を表示

最新記事

CUDA Dynamic Parallelism API

子カーネル関数立ち上げ

引数アドレス制約

__global__ void parent_kernel(){
    int a = 0;
    child_kernel<<<1, 128>>>(&a);
    // ...

__device__ int a;
__global__ void parent_kernel(){
    child_kernel<<<1, 128>>>(&a);
    // ...

PTXで見るCUDA Dynamic Parallelism API

子カーネル立ち上げのcuコードからptxを出力すると.weakディレクティブ付きでcudaLaunchDeviceV2関数やcudaGetDevice関数、cudaMalloc関数などのダミー関数がPTX内に定義されています．
.weakですのでリンク時により強い宣言がある場合上書きされます．(そりゃdevice側のcudaMallocがあったとして，そのPTXをNVIDIAが簡単に我々下々の開発者に公開するわけがないですよね)

肝心の子カーネル関数の立ち上げ部分の流れは次のようになります．

1. cudaGetParameterBufferV2関数で子カーネル関数のアドレス，ブロック数，スレッド数などを指定します．
2. cudaLaunchDeviceV2関数で子カーネル関数の引数を格納した領域のアドレスなどを渡し，子カーネル関数を立ち上げます．

実験

デバイス側からのカーネル関数立ち上げが高速だった場合，<<<1, 1>>>で立ち上げ用のカーネル関数を立ち上げ，内部で目的のカーネル関数を立ち上げた方が速く立ち上げるかもと思い試してみましたが遅かったです．

あとがき