Clang/CUDA 解析 CUDA 关键字 shared

Question

由于可以使用Clang进行CUDA编译，我有兴趣研究clang到中间表示（IR）的cuda代码（.cu文件）转换。

Clang 的 CUDA 编译需要某些 CUDA 库。那么CUDA程序中关键字__shared__的解析是由Clang还是由CUDA编译器完成的呢？从我最初的搜索来看，我相信转换是由 CUDA 而不是 Clang 完成的。这种理解是否正确？

Answer 1

当 clang 编译 CUDA 代码时，Nvidia NVCC 编译器不参与。

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

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

用 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 在模块内部具有静态（但内部）链接，位于地址 space 3。

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

通过将 PTX 代码与 CUDA API 链接，PTX 代码可用于对 GPU 进行编程。

编辑： 由于问题出现在此处定义 __shared__ 属性的位置是：在 clang 中 headers host_defines.h 包含在 CUDA 工具包中。在 host_defines.h（来自 CUDA 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)).

Clang/CUDA 解析 CUDA 关键字 shared

Parsing of CUDA keyword shared by Clang/CUDA

cuda

clang

llvm-clang

llvm-ir

Clang/CUDA 解析 CUDA 关键字 __shared__

Parsing of CUDA keyword __shared__ by Clang/CUDA

cuda

clang

llvm-clang

llvm-ir

Clang/CUDA 解析 CUDA 关键字 shared

Parsing of CUDA keyword shared by Clang/CUDA