编译 CUDA 文件时在 Clang 编译器中从 _attribute_((shared)) 转换为 addrspace(3)

Question

clang 编译器包含 CUDA 头文件 host_defines.h，其中 __shared__ 定义为 __attribute__((shared))。当使用 clang 将 CUDA 源文件编译为内部表示 (IR) 时，__shared__ 会转换为 addrspace(3)。这些地址空间可以在clang文件llvm/tools/clang/lib/Basic/Targets.cpp行号1601中观察为一个数组

static const unsigned NVPTXAddrSpaceMap[] = {
    1, // opencl_global
    3, // opencl_local
    4, // opencl_constant
    // FIXME: generic has to be added to the target
    0, // opencl_generic
    1, // cuda_device
    4, // cuda_constant
    3, // cuda_shared
};

所以具体问题是在转换的哪个阶段，__attribute__((shared)) 被转换为 addrspace(3)。查看 clang 的解析和词法分析部分并没有给出任何暗示。有人可以帮忙吗？

Answer 1

shared属性定义在clang的Attr.td文件中，在内部称为CUDAShared，表示为CUDASharedAttr。在任何属性的词法分析和解析阶段，对 Attr.td 中定义的所有属性进行词法分析和解析。在此阶段，您不会找到任何必要的见解。

您将看到 CUDASharedAttr 的有价值代码的第一个点位于 clang/lib/Sema/SemaDeclAttr.cpp。 Sema class 构建 AST 并在 SemaDeclAttr.cpp 中完成对每个属性的处理。对于特定的 CUDASharedAttr handleSimpleAttribute<CUDASharedAttr>(S, D, Attr); 被调用。此函数只是将属性插入给定声明 (Decl& D)。

现在属性已附加到 Decl，您可以使用以下方法查询声明是否具有属性：D.hasAttr<CUDASharedAttr>()。例如，在 SemaDecl.cpp 中强制执行对 CUDA 共享内存声明的限制，并且共享内存变量的存储 class 设置为静态。

您将再次找到发出实际 LLVM IR 的 CUDASharedAttr bin clang/lib/CodeGen/CodeGenModule.cpp。在 CodeGenModule.cpp 中，您具有以下功能：

unsigned CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D,
                                                 unsigned AddrSpace) {
  if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
    if (D->hasAttr<CUDAConstantAttr>())
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_constant);
    else if (D->hasAttr<CUDASharedAttr>())
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_shared);
    else
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_device);
  }

  return AddrSpace;
}

该函数从实际目标查询共享函数的地址 space，即对于 nvptx 目标，使用您发布的地址 space 地图：

static const unsigned NVPTXAddrSpaceMap[] = {
    1, // opencl_global
    3, // opencl_local
    4, // opencl_constant
    // FIXME: generic has to be added to the target
    0, // opencl_generic
    1, // cuda_device
    4, // cuda_constant
    3, // cuda_shared
};

LangAS::cuda_shared对应地址space 3.

完成所有这些步骤后，您将在最终的 IR 模块中获得地址为 space 3 的全局变量，如下所示：

 ; 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
  }

编译 CUDA 文件时在 Clang 编译器中从 _attribute_((shared)) 转换为 addrspace(3)

Conversion from _attribute_((shared)) to addrspace(3) in Clang compiler when compiling CUDA files

cuda

llvm

clang

llvm-clang

llvm-ir

编译 CUDA 文件时在 Clang 编译器中从 ___attribute___((shared)) 转换为 addrspace(3)

Conversion from ___attribute___((shared)) to addrspace(3) in Clang compiler when compiling CUDA files

cuda

llvm

clang

llvm-clang

llvm-ir

编译 CUDA 文件时在 Clang 编译器中从 _attribute_((shared)) 转换为 addrspace(3)

Conversion from _attribute_((shared)) to addrspace(3) in Clang compiler when compiling CUDA files