我正在努力添加类似于的新内存类型__shared__
在 CUDA 中称为__noc__
需要使用 clang-llvm 进行编译。以下是实现对新内存类型的解析的步骤,引用来自answer https://stackoverflow.com/questions/35356840/conversion-from-attribute-shared-to-addrspace3-in-clang-compiler-when:
Step 1:在 clangs 的 Attr.td 文件(clang/include/clang/Basic/Attr.td)中,添加了与共享关键字类似的 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];
}
Step 2: 如同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;
Step 3:在SemaDecl.cpp文件中,CUDANocAttr
添加以强制 noc 为静态存储(类似于共享)
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);
}
}
Step 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;
}
Step 5: The 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
};
Step 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中,addrspace(6)
这对应于__noc__
没有观察到。
从放入文件 clang/lib/Sema/SemaDeclAttr.cpp (步骤 2)的调试 printf 中,可以观察到以下情况:AttributeList::AT_CUDANoc
没有被处决。
任何建议或直觉都会有很大帮助。在编译 llvm 源代码之前是否有任何脚本需要显式运行,以使 *.td 文件中的输入显示为 C++ 源代码...