用Clang/CUDA解析CUDA关键字__shared__

Question

由于可以使用Clang进行CUDA编译，所以我对clang转换为中间表示(IR)的cuda代码(.cu文件)很感兴趣。

Clang的CUDA汇编需要某些CUDA图书馆。那么，对CUDA程序中关键字__shared__的解析是由Clang还是由CUDA编译器完成的呢？根据我最初的搜索，我相信转换是由数据自动化系统而不是Clang完成的。这种理解是正确的吗？

Answer 1

当clang编译CUDA代码时，Nvidia NVCC编译器不涉及。

__shared__或更准确地说，__attribute__((shared))是clang知道的一个属性。如果clang遇到一个使用共享属性标记的变量，它将做两件事：

1. 该变量将具有静态链接。这意味着变量的定义从内核函数移动到模块范围。
2. 变量将放置在地址空间3中，该地址空间被定义为共享内存地址空间。

用clang编译这个小程序：

__global__ void foo(int* tmp)
{
  __shared__ int vec[32];
  vec[threadIdx.x] = tmp[threadIdx.x];
  tmp[threadIdx.y] = vec[threadIdx.y];
}

int main()
{
  int* tmp;
  foo<<<1, 1>>>(tmp);
  return tmp[0];
}

下列IR的结果：

  ; ModuleID = 'sm.cu'
  target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
  target triple = "nvptx64-unknown-unknown"

  @vec= internal unnamed_addr addrspace(3) global [32 x i32] zeroinitializer, align 4

  ; Function Attrs: nounwind readnone
  declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0

  ; Function Attrs: nounwind readnone
  declare i32 @llvm.nvvm.read.ptx.sreg.tid.y() #0

  define ptx_kernel void @__pacxx_kernel0(i32 addrspace(1)* %tmp) {
    %1 = tail call spir_func i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1
    %2 = zext i32 %1 to i64
    %3 = getelementptr i32, i32 addrspace(1)* %tmp, i64 %2
    %4 = load i32, i32 addrspace(1)* %3, align 4
    %5 = getelementptr [32 x i32], [32 x i32] addrspace(3)* @vec, i64 0, i64 %2
    store i32 %4, i32 addrspace(3)* %5, align 4
    %6 = tail call spir_func i32 @llvm.nvvm.read.ptx.sreg.tid.y() #1
    %7 = zext i32 %6 to i64
    %8 = getelementptr [32 x i32], [32 x i32] addrspace(3)* @vec, i64 0, i64 %7
    %9 = load i32, i32 addrspace(3)* %8, align 4
    %10 = getelementptr i32, i32 addrspace(1)* %tmp, i64 %7
    store i32 %9, i32 addrspace(1)* %10, align 4
    ret void
  }

您可以看到变量vec在模块内部具有静态(但内部)链接，并且驻留在地址空间3中。

Clang遵循NVVM规范，可以找到这里。但是，NVVM IR是为LLVM 3.4指定的，如果使用更新的LLVM/Clang版本生成的IR，可能会遇到问题。然而，来自LLVM的NVPTX后端没有这种限制，并且可以生成PTX代码而没有问题。Clang (在新版本中)将像NVCC一样构建一个胖垃圾箱。在旧版本的Clang中，您必须自己构建可执行文件，并使用CUDAIsDevice命令行标志编译程序的设备部分。

PTX代码可以通过链接到CUDA API来编程GPU。

编辑：，既然问题在哪里定义，__shared__属性在哪里定义:在clang中，host_defines.h是从CUDA工具包中包含的。在host_defines.h (来自数据自动化系统7.5)中，您可以看到：

  192 #define __shared__ \
  193         __location__(shared)

__location__ (这是另一个宏定义)扩展到__annotate__

   85 #define __annotate__(a) \
   86         __attribute__((a))
   87 #define __location__(a) \
   88         __annotate__(a)

正如我在答案的第一部分中所写的那样，它被扩展到__attribute__。因此，__shared__被扩展到__attribute__((shared))。