使用 clang-llvm 编译器在 CUDA 中添加对类似于 __shared__ 的内存类型的支持

Posted 2023-02-16

【中文标题】使用 clang-llvm 编译器在 CUDA 中添加对类似于 __shared__ 的内存类型的支持

【英文标题】：Adding support for a memory type similar to __shared__ in CUDA using clang-llvm compiler

【发布时间】：2016-02-20 06:05:49

【问题描述】：

我正在努力在 CUDA 中添加一种类似于 __shared__ 的新内存类型，称为 __noc__，它需要使用 clang-llvm 进行编译。以下是参考answer实现新内存类型解析的步骤：

第一步：在clangs的Attr.td文件（clang/include/clang/Basic/Attr.td）中，添加了类似于shared关键字的noc关键字。

def CUDAShared : InheritableAttr 
  let Spellings = [GNU<"shared">];
  let Subjects = SubjectList<[Var]>;
  let LangOpts = [CUDA];
  let Documentation = [Undocumented];


def CUDANoc : InheritableAttr 
  let Spellings = [Keyword<"noc">];
  let Subjects = SubjectList<[Var]>;
  let LangOpts = [CUDA];
  let Documentation = [Undocumented];

第二步：与CUDASharedAttr类似，CUDANocAttr被添加到clang/lib/Sema/SemaDeclAttr.cpp中。

  case AttributeList::AT_CUDAShared:
    handleSimpleAttribute<CUDASharedAttr>(S, D, Attr);
    break;
  case AttributeList::AT_CUDANoc:
    handleSimpleAttribute<CUDANocAttr>(S, D, Attr);
    printf("\n T1:SemaDeclAttr.cpp"); //testpoint 1 : for noc debugging
    break;

第 3 步：在 SemaDecl.cpp 文件中，添加 CUDANocAttr 以强制 noc 为静态存储（类似于 shared）

  if (getLangOpts().CUDA) 
    if (EmitTLSUnsupportedError && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD))
      Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
           diag::err_thread_unsupported);
    // CUDA B.2.5: "__shared__ and __constant__ variables have implied static
    // storage [duration]."
    if (SC == SC_None && S->getFnParent() != nullptr &&
        (NewVD->hasAttr<CUDASharedAttr>() ||
         NewVD->hasAttr<CUDANocAttr>()||
         NewVD->hasAttr<CUDAConstantAttr>())) 
      NewVD->setStorageClass(SC_Static);
    
  

第4步：在CodeGenModule（llvm/tools/clang/lib/CodeGen/CodeGenModule.cpp）中添加NOC，允许cuda_noc地址空间从NVPTXAddrSpaceMap访问

    else if (D->hasAttr<CUDASharedAttr>())
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_shared);
    else if (D->hasAttr<CUDANocAttr>())
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_noc);
    else
      AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_device);
  

  return AddrSpace;

第 5 步：将 cuda_noc 添加到 NVPTXAddrSpaceMap 数组以允许新类型的地址空间

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
    6, // cuda_noc
;

第 6 步：将宏 #define __noc__ __location__(noc) 添加到文件 clang/lib/Headers/__clang_cuda_runtime_wrapper.h 中，其中包含来自 CUDA 的 host_defines.h。

llvm 源代码编译安装成功。但是当试图编译一个内存类型为 __noc__ 的 CUDA 源文件时，它会给出以下警告：

warning: unknown attribute 'noc' ignored [-Wunknown-attributes]
        __noc__ int c_shared[5];
        ^

/usr/local/bin/../lib/clang/3.8.0/include/__clang_cuda_runtime_wrapper.h:69:30: note: expanded from macro '__noc__'
#define __noc__ __location__(noc)
                             ^
1 warning generated.

从警告中可以看出__noc__ 被忽略了。在生成的IR中，没有观察到与__noc__对应的addrspace(6)。

从放入文件 clang/lib/Sema/SemaDeclAttr.cpp 的调试 printf（步骤 2）中，可以观察到 AttributeList::AT_CUDANoc 的情况没有得到执行。

任何建议或直觉都会有很大帮助。在编译 *.td 文件中的输入的 llvm 源代码以显示为 C++ 源代码之前，是否有任何脚本要显式运行...

【问题讨论】：

您是否正在使用 -std=cuda 编译源代码并且正在编译 .cu 文件？

是的，我正在编译一个 cuda 文件 (.cu)。但没有使用 -std=cuda。现在尝试添加 -std=cuda 但即便如此也没有发现任何变化。用于编译的命令是：'sudo clang++ -emit-llvm -S -std=cuda -o noc05 -I/home/ginu001/NVIDIA_CUDA-7.5_Samples/common/inc -L/usr/local/cuda/lib64 -L /usr/local/cuda/lib64/stubs noc.cu -lcudart_static -lcuda -ldl -lrt -pthread'

@GinuJacob 好的（不确定Keyword 会做什么），但请尝试使用let Spellings = [GNU<"noc">]; 并重建

@GinuJacob 我不明白？在第 1 步中，您有 Keyword<"noc">。

请考虑写一个答案以获得未回答列表的问题。

【参考方案1】：

__location__(noc) 扩展为 __attribute__((noc))。这是 GNU 或 gcc 属性语法。所以问题在于这一行：

let Spellings = [Keyword<"noc">];

为了使noc 与__location__ 宏一起使用，您应该使用GNU<"noc"> 而不是Keyword<"noc">。