关于CUDA宏__CUDA_ARCH__的问题

Question

我在ttt.cu 中有一个简单的 cuda 代码

#include <iostream>
__global__ void example(){
   printf("__CUDA_ARCH__: %d \n", __CUDA_ARCH__);
}
int main(){
example<<<1,1>>>();
}

使用 CMakeLists.txt：

cmake_minimum_required(VERSION 3.18)
project(Hello)
find_package(CUDA REQUIRED)

cuda_add_executable(sss ttt.cu)

然后我得到了错误：identifier "__CUDA_ARCH__" is undefined。我想知道为什么会发生这种情况，我应该怎么做才能使__CUDA_ARCH__ 有效？我们可以在标题.h 文件中的主机代码中使用有效的__CUDA_ARCH__ 吗？

更新：

我打算使用以下 cmake 来生成 750 cuda 拱门，但是，这总是导致 __CUDA_ARCH__ = 300（2080 ti 与 cuda 10.1）。我尝试了set_property 和target_compile_options，都失败了。

cmake_minimum_required(VERSION 3.18)
project(Hello)
find_package(CUDA REQUIRED)
cuda_add_executable(oounne ttt.cu)
set_property(TARGET oounne PROPERTY CUDA_ARCHITECTURES 75)
#target_compile_options(oounne PRIVATE  $<$<COMPILE_LANGUAGE:CUDA>:-gencode 
arch=compute_75,code=sm_75>)

Answer 1

__CUDA_ARCH__ 是一个编译器宏。

我们可以在主机代码中使用有效的__CUDA_ARCH__

不，是intended to be used in device code only：

主机代码（非 GPU 代码）不能依赖它。

您无法按照您想象的方式打印编译器宏。它不是 C++ 中定义的普通数值变量。您可以执行 this 之类的操作，但这会在编译时打印，而不是在运行时打印。

要在运行时打印，您可以执行以下操作：

$ cat t2.cu
#include <cstdio>
#define STR_HELPER(x) #x
#define STR(x) STR_HELPER(x)

__device__ void print_arch(){
  const char my_compile_time_arch[] = STR(__CUDA_ARCH__);
  printf("__CUDA_ARCH__: %s\n", my_compile_time_arch);
}
__global__ void example()
{
   print_arch();
}

int main(){

example<<<1,1>>>();
cudaDeviceSynchronize();
}
$ nvcc -o t2 t2.cu

$ ./t2
__CUDA_ARCH__: 520
$

请注意，cuda 标签上有很多问题here 讨论__CUDA_ARCH__，您不妨查看其中的一些问题。