CUDA 断言 - 在 __host__ __device__ 上重载,为什么没有 warnings/errors?

CUDA assert - overload on __host__ __device__, why no warnings/errors?

我正在尝试了解 CUDA 断言在幕后是如何工作的。 assert 宏调用 __assert_fail 函数,我可以在其中找到以下签名:

/usr/include/assert.h:extern void __assert_fail (const char *__assertion, const char *__file,
/usr/local/cuda-10.2/targets/x86_64-linux/include/crt/common_functions.h:extern __host__ __device__ __cudart_builtin__ void __assert_fail(

我可以看到它们具有相同的签名,但是 CUDA 版本有 __host__ __device__ 个限定符。

通常不可能重载基于 __host__ __device__ 的函数,因为它们不是函数签名的一部分(除非使用 Clang 而不是 NVCC 编译 CUDA 代码):

extern void foo();
extern __host__ __device__ void foo();

使用 NVCC 编译并将警告作为错误:

main.cu(4): error: a __host__ function("foo") redeclared with __host__ __device__, hence treated as a __host__ __device__ function

那么,为什么这个 warning/error 没有和 assert 一起出现?是否有任何其他魔法可以使这项工作正常进行?

考虑以下示例程序 main.cu,它是通过 nvcc --keep main.cu -o main

编译的
#include <cassert>

__global__ void kernel(){
    assert(false);
}

int main(){
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
    assert(false);
}

CUDA 程序的编译是按工具包文档中所述的几个步骤执行的 https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#cuda-compilation-trajectory 使用标志 --keep 进行编译会保留所有可以搜索 __assert_fail.

的中间文件

主文件。cpp1.ii

# 66 "/usr/include/assert.h" 3 4
extern "C" {

extern void __assert_fail (const char *__assertion, const char *__file,
      unsigned int __line, const char *__function)
     throw () __attribute__ ((__noreturn__));
...
}
# 169 "/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/common_functions.h"
extern "C"
{
# 197 "/usr/local/cuda/bin/../targets/x86_64-linux/include/crt/common_functions.h"
extern __attribute__((host)) __attribute__((device)) __attribute__((cudart_builtin)) void __assert_fail(
  const char *, const char *, unsigned int, const char *)
  
...
}

注意 nvcc 注入的版本中的属性 cudart_builtin。有了这个属性,就没有重新声明的警告了。

对于您的情况,正如您观察到的那样,此程序会产生警告

extern void foo();
extern __host__ __device__ void foo();

int main(){
    return 0;
}

但是,下面的编译没有警告。

extern void foo();
extern __host__ __device__ __attribute__((cudart_builtin)) 
void foo();

int main(){
    return 0;
}

因此,您可以使用此属性。但是,它可能有未知的副作用,foo 实际上并不是 cuda 运行时的内置函数。您不应在您的函数中使用此属性。

已使用CUDA 11.2。