CUDA 断言 - 在 host device 上重载，为什么没有 warnings/errors？

Question

我正在尝试了解 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 一起出现？是否有任何其他魔法可以使这项工作正常进行？

Answer 1

考虑以下示例程序 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。

CUDA 断言 - 在 host device 上重载，为什么没有 warnings/errors？

CUDA assert - overload on host device, why no warnings/errors?

c++

assert

cuda

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

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

c++

assert

cuda

CUDA 断言 - 在 host device 上重载，为什么没有 warnings/errors？

CUDA assert - overload on host device, why no warnings/errors?