Добавление поддержки типа памяти, аналогичного __shared__ в CUDA с использованием компилятора clang-llvm

Question

Я работаю над добавлением нового типа памяти, подобного __shared__ в CUDA под названием __noc__, который необходимо скомпилировать с помощью clang-llvm. Ниже приведены шаги, а затем достичь синтаксический для нового принимая обозначению типа памяти от answer:

Шаг 1: В clangs Attr.td файле игровая (лязг/включить/лязг/Basic/Attr.td), ключевое слово noc было добавлено аналогично ключевому слову shared.

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

Шаг 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; 

Шаг 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); 
    } 
    } 

Шаг 4: НОК добавлена ​​в CodeGenModule (LLVM/инструменты/лязг/Библиотека/CodeGen/CodeGenModule.cpp), чтобы разрешить доступ 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/Библиотека/Заголовки/__ clang_cuda_runtime_wrapper .h, где включен host_defines.h из CUDA.

Исходный код llvm скомпилирован и установлен успешно. Но при попытке скомпилировать исходный файл CUDA с памятью типа __noc__, он дает следующее предупреждение:

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__ игнорируется. В генерируемом ИК-диапазоне addrspace(6), который соответствует __noc__, не наблюдается.

Из отладочной печати, помещенной в файл clang/lib/Sema/SemaDeclAttr.cpp (шаг 2), можно заметить, что случай для AttributeList::AT_CUDANoc не выполняется.

Любые предложения или интуиции могут очень помочь. Есть ли какой-либо скрипт, который должен выполняться явно, прежде чем компилировать исходный код llvm для входов в файле .td, чтобы он отображался как исходный код C++ ...

Answer 1

__location__(noc) расширился до __attribute__((noc)). Это синтаксис атрибутов GNU или gcc. Таким образом, проблема с этой линии:

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

Для того, чтобы noc работать с __location__ макросъемки, вы должны использовать GNU<"noc"> вместо Keyword<"noc">.