1

CUDAカーネル呼び出しを単純化するためにマクロを作成しました。

#define LAUNCH LAUNCH_ASYNC

#define LAUNCH_ASYNC(kernel_name, gridsize, blocksize, ...) \
    LOG("Async kernel launch: " #kernel_name);              \
    kernel_name <<< (gridsize), (blocksize) >>> (__VA_ARGS__);

#define LAUNCH_SYNC(kernel_name, gridsize, blocksize, ...)     \
    LOG("Sync kernel launch: " #kernel_name);                  \
    kernel_name <<< (gridsize), (blocksize) >>> (__VA_ARGS__); \
    cudaDeviceSynchronize();                                   \
    // error check, etc...

使用法：

LAUNCH(my_kernel, 32, 32, param1, param2)

LAUNCH(my_kernel<int>, 32, 32, param1, param2)

これは正常に機能します。最初の定義で、デバッグのための同期呼び出しとエラーチェックを有効にできます。

ただし、以下のような複数のテンプレート引数では機能しません。

LAUNCH(my_kernel<int,float>, 32, 32, param1, param3)

マクロを呼び出した行に表示されるエラーメッセージ：

error : expected a ">"

このマクロを複数のテンプレート引数で機能させることは可能ですか？

4

3 に答える 3

5

問題は、プリプロセッサが山かっこの入れ子について何も知らないため、それらの間のコンマをマクロ引数の区切り文字として解釈することです。

kernel-launch構文がカーネル名の前後の括弧をサポートしている場合（CUDAマシンではなく、今は確認できません）、次のように実行できます。

LAUNCH((my_kernel<int, float>), 32, 32, param1, param3)
于 2013-02-14T12:37:30.233 に答える
1

（投稿したマクロに基づいて）私が使用した他の方法は、カーネルブロックサイズとグリッドサイズの引数を独自のマクロにラップすることです。

#define KERNEL_ARGS2(grid, block) <<< grid, block >>>
#define KERNEL_ARGS3(grid, block, sh_mem) <<< grid, block, sh_mem >>>
#define KERNEL_ARGS4(grid, block, sh_mem, stream) <<< grid, block, sh_mem, stream >>>

これで、次のようにマクロを使用できるようになります。

#define CUDA_LAUNCH(kernel_name, gridsize, blocksize, ...) \
kernel_name KERNEL_ARGS2(gridsize, blocksize)(__VA_ARGS__);

次のように使用できます。

CUDA_LAUNCH(my_kernel, grid_size, block_size, float* input, float* output, int size);

これにより、指定されたグリッドとブロックサイズ、および入力引数を使用して、「my_kernal」というカーネルが起動します。

于 2017-02-15T18:32:09.573 に答える
0

エラーもスローするこのソリューションを検討してください

inline void echoError(cudaError_t e, const char *strs) {
    char a[255];
    if (e != cudaSuccess) {
        strncpy(a, strs, 255);
        fprintf(stderr, "Failed to %s,errorCode %s",
                a, cudaGetErrorString(e));
        exit(EXIT_FAILURE);
    }
}


#define CUDA_KERNEL_DYN(kernel, bpg, tpb, shd, ...){                     \
    kernel<<<bpg,tpb,shd>>>( __VA_ARGS__ );                              \
    cudaError_t err = cudaGetLastError();                                \
    echoError(err, #kernel);                                              \
}
于 2018-07-25T14:45:50.133 に答える