Clang/master does not compile - please rebase or merge

[daniel-schuermann]

Hello thewilsonator,
thank you for this work! I think it will be very beneficial and I am trying to make use of it.
I only figured that your branch is out of sync with master for to long to compile with clang/master.
Could you please rebase or merge the changes (I don't care, but maybe others do)?
I did so locally, so it's not urgent for me.
Thanks alot and keep up the work!

Daniel

Edit: Forgot to mention, this issue relates to your llvm/compute fork, not this repository.

[thewilsonator]

Already? I updated it 2 days ago, use the compute branch. Oh whoops that was this repo.

[thewilsonator]

Done. Be aware that I will soon (I hope approx. three weeks-ish) be making a breaking change w.r.t. the OpenCL functions moving from Itanium mangled C++ to proper LLVM intrinsics. See the Enums.td and Instructions.td, any input on those is very much appreciated.

[daniel-schuermann]

Thank you very much. The intrinsics change seems to have much impact on the clang OpenCL Codegen. For my purpose I only have to know how to insert proper get_global_id() and such function calls in Codegen. Am I right that OpenCL address spaces in LLVM get transformed to storage classes in SPIRV? What happens to the kernel arg metadata? Is there also a transformation or does it only use the qualifiers from the parameters directly?

[thewilsonator]

All that should change is the function will go from ::get_global_id(size_t) to __builtin_spirv_opencl_get_global_id(size_t)(GCC like builtin)/llvm.spirv.opencl.get_global_id.i32/64(LLVM intrinsic).

Am I right that OpenCL address spaces in LLVM get transformed to storage classes in SPIRV?

Correct.

What happens to the kernel arg metadata?

The write/read attributes of images and pipes get translated into AccessQualifiers. I think they're the only arg metadata.

[thewilsonator]

Also you may be interested in the llvm-dev mailing list thread [SPIR-V] SPIR-V in LLVM which discusses getting this backend into LLVM trunk.

[daniel-schuermann]

Hi!
Wanted to ask, how far you got with the logical address space layout for your mentioned spirl aka Vulkan back-end. Due to lack of existing drivers with OpenCL 2.1 support, I'm thinking about switching to Vulkan API as it seems, i don't really need the OpenCL specific stuff. Currently, I'm developing against the Intel OpenCL SDK (CPU only), which is the only existing OpenCL 2.1 implementation...

[thewilsonator]

None yet, first goal is to get this upstreamed.

[daniel-schuermann]

Could you please rebase or merge the changes into llvm/compute? Clang refuses to build again.
Maybe we find a fixed period for this? Because my project heavily depends on your backend.
I am writing an OpenMP - SPIR-V backend, which I hope to get upstream as well. Tell me, if I can support your pull request somehow.
Sidenote: Did you see https://github.com/google/clspv ? Maybe the llvm-passes can be reused for the lspir backend. (However, it's not urgent)

[thewilsonator]

Just synced everything, sorry about that. Yeah I did see that, I have to take a closer look. I'll be publishing a blog post in about an hour on my project that uses this repo and I'll try to gather some efforts to get this into shape for upstreaming, I'll link to it and the /r/programming discussion when thats done.

Currently I'm working on writing a tablegen backend for the tablegen descriptors I've written, the .td files in the top level directory of this repo. After that is done I'll do the intrinsics .td for OpenCL in /include/llvm/IR and then gradually change this repo to use them instead of the Itanium mangled C++ and utilise the tablegen description of the binary format.
After all that's done hopefully it will be accepted into LLVM.

[daniel-schuermann]

Sorry for disturbing again. I get a lot of undefined references when compiling your llvm compute backend with clang in-tree. Specifically in lib/CodeGen/MachineDominators.cpp and lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp. I even tried to clone everything again, but no success. Does it work for you?

[thewilsonator]

Hmmm. I just git pulled LLVM and rebuilt and it worked for me. I've pushed that pull so you can try it again. Did you ensure you updated the submodule?

If you want this to not break too often I can notify you in this thread when I going to pull the latest LLVM changes. Chances are that that will coincide with when clang breaks and you can stagger updating clang.

[daniel-schuermann]

Yes, this seems a good solution. I also put your llvm fork on my watch list, but don't know if I get notified.
I figured that updating this (Target/SPIRV) subproject brought the linking errors:
SPIRVUtil has unresolved dependencies to Mangler. Please check the CMakefiles (also rebuild the corresponding Makefiles). I think this came when moving libSPIRV, because the commit HEAD~10 (Make stuff compile) still works fine.
Kind regards!

[thewilsonator]

Hmm. Could you post the missing symbols?

All I've done since then was move stuff from the main target folder to ./Passes and update the includes and CMakeLists.txts accordingly. I'm not sure how that could have broken things.

[daniel-schuermann]

I don't know well about all the CMake stuff, but here is what I get when I try to compile master:
[100%] Linking CXX executable ../../../../bin/clang
../../../../lib/libLLVMSPIRVlibSPIRV.a(SPIRVUtil.cpp.o): In function SPIRV::transTypeDesc(llvm::Type*, SPIRV::BuiltinArgTypeMangleInfo const&)': /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:980: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:983: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:988: undefined reference to SPIR::AtomicType::AtomicType(SPIR::RefCountSPIR::ParamType)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:993: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:996: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:999: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1002: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1005: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' ../../../../lib/libLLVMSPIRVlibSPIRV.a(SPIRVUtil.cpp.o):/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1011: more undefined references to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' follow
../../../../lib/libLLVMSPIRVlibSPIRV.a(SPIRVUtil.cpp.o): In function SPIRV::transTypeDesc(llvm::Type*, SPIRV::BuiltinArgTypeMangleInfo const&)': /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1021: undefined reference to SPIR::VectorType::VectorType(SPIR::RefCountSPIR::ParamType, int)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1036: undefined reference to SPIR::UserDefinedType::UserDefinedType(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1044: undefined reference to SPIR::BlockType::BlockType()'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1060: undefined reference to SPIR::BlockType::BlockType()' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1063: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1064: undefined reference to SPIR::PointerType::PointerType(SPIR::RefCount<SPIR::ParamType>)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1065: undefined reference to SPIR::PointerType::setAddressSpace(SPIR::TypeAttributeEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1070: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1075: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1076: undefined reference to SPIR::PointerType::PointerType(SPIR::RefCount<SPIR::ParamType>)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1077: undefined reference to SPIR::PointerType::setAddressSpace(SPIR::TypeAttributeEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1081: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1086: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1093: undefined reference to SPIR::PointerType::PointerType(SPIR::RefCount<SPIR::ParamType>)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1095: undefined reference to SPIR::PointerType::setAddressSpace(SPIR::TypeAttributeEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1098: undefined reference to SPIR::PointerType::setQualifier(SPIR::TypeAttributeEnum, bool)' ../../../../lib/libLLVMSPIRVlibSPIRV.a(SPIRVUtil.cpp.o): In function SPIRV::mangleBuiltin(std::__cxx11::basic_string<char, std::char_traits, std::allocator > const&, llvm::ArrayRefllvm::Type*, SPIRV::BuiltinFuncMangleInfo*)':
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1358: undefined reference to SPIR::NameMangler::NameMangler(SPIR::SPIRversion)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1368: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)'
/home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1383: undefined reference to SPIR::PrimitiveType::PrimitiveType(SPIR::TypePrimitiveEnum)' /home/daniel/Workspace/OpenMP/SPIRV/lib/Target/SPIRV/libSPIRV/SPIRVUtil.cpp:1385: undefined reference to SPIR::NameMangler::mangle(SPIR::FunctionDescriptor const&, std::__cxx11::basic_string<char, std::char_traits, std::allocator >&)'
collect2: error: ld returned 1 exit status
tools/clang/tools/driver/CMakeFiles/clang.dir/build.make:349: recipe for target 'bin/clang-5.0' failed
make[3]: *** [bin/clang-5.0] Error 1
CMakeFiles/Makefile2:20045: recipe for target 'tools/clang/tools/driver/CMakeFiles/clang.dir/all' failed
make[2]: *** [tools/clang/tools/driver/CMakeFiles/clang.dir/all] Error 2
CMakeFiles/Makefile2:20057: recipe for target 'tools/clang/tools/driver/CMakeFiles/clang.dir/rule' failed
make[1]: *** [tools/clang/tools/driver/CMakeFiles/clang.dir/rule] Error 2
Makefile:5332: recipe for target 'clang' failed
make: *** [clang] Error 2

[thewilsonator]

Very weird. Did llvm-spirv link successfully? Might be a misconfiguration on the clang side or a stale cmake.

[daniel-schuermann]

llvm-spirv was successful, yes. But meanwhile I found the issue and this way clang builds again.
In libSPIRV/CMakeList.txt please add the last line:

add_definitions(-D_SPIRV_LLVM_API)
include_directories(..)
add_llvm_library(LLVMSPIRVlibSPIRV
SPIRVBasicBlock.cpp
SPIRVDebug.cpp
SPIRVDecorate.cpp
SPIRVEntry.cpp
SPIRVFunction.cpp
SPIRVInstruction.cpp
SPIRVModule.cpp
SPIRVStream.cpp
SPIRVType.cpp
SPIRVUtil.cpp
SPIRVValue.cpp
LINK_LIBS LLVMSPIRVMangler
)

[thewilsonator]

d3f75f1

[daniel-schuermann]

Hello Nicholas,
seems, a change to the target registry introduced a conflict, when I merged upstream locally:
TargetInfo/SPIRVTargetInfo.cpp:32:52: error: no matching function for call to ‘llvm::RegisterTarget<(llvm::Triple::ArchType)42u>::RegisterTarget(llvm::Target&, const char [8], const char [14])’
"32-bit SPIR-V");
include/llvm/Support/TargetRegistry.h:898:3: note: candidate: llvm::RegisterTarget<TargetArchType, HasJIT>::RegisterTarget(llvm::Target&, const char*, const char*, const char*) [with llvm::Triple::ArchType TargetArchType = (llvm::Triple::ArchType)42u; bool HasJIT = false]
RegisterTarget(Target &T, const char *Name, const char *Desc,
include/llvm/Support/TargetRegistry.h:898:3: note: candidate expects 4 arguments, 3 provided

[thewilsonator]

74687b7623f5bec69b9ed9d3497cab2db77ff78e