From a405b9412035c1219f3278105e528c1de562ed1c Mon Sep 17 00:00:00 2001 From: doe300 Date: Sat, 7 Dec 2019 09:45:49 +0100 Subject: [PATCH] Small fixes - Reverts removal of SPIRV-Tools-opt target - Handles ODR violations when linking multiple LLVM modules with conflicting VC4CL std-lib function definitions - Improves linking test - Improves documentation, closes #131 --- .circleci/config.yml | 2 +- Readme.md | 2 +- cmake/spirv.cmake | 6 +++--- src/precompilation/FrontendCompiler.cpp | 12 ++++++++++-- test/TestFrontends.cpp | 10 +++++++--- testing/test_linking_0.cl | 4 ++-- testing/test_linking_1.cl | 6 ++++-- testing/test_linking_2.cl | 5 +++++ 8 files changed, 33 insertions(+), 14 deletions(-) create mode 100644 testing/test_linking_2.cl diff --git a/.circleci/config.yml b/.circleci/config.yml index dcb149a8..03bbcc4e 100644 --- a/.circleci/config.yml +++ b/.circleci/config.yml @@ -138,7 +138,7 @@ jobs: - run: dpkg -i build/vc4c-0.4-Linux.deb - run: ln -s `pwd`/build/cpptest-lite/src/cpptest-lite-project-build/libcpptest-lite.so.0.9 /usr/lib/libcpptest-lite.so.1.1.2 - run: ldconfig - - run: build/test/TestVC4C --output=plain --mode=verbose --test-instructions --test-operators --emulate-common --emulate-geometric --emulate-memory --emulate-relational --emulate-vector --test-patterns + - run: build/test/TestVC4C --output=plain --mode=verbose --test-instructions --test-operators --emulate-common --emulate-geometric --emulate-memory --emulate-relational --emulate-vector --test-patterns --test-frontend - run: dpkg -r vc4c - run: dpkg -r vc4cl-stdlib test-spirv: diff --git a/Readme.md b/Readme.md index 96de8d31..e6cab74b 100644 --- a/Readme.md +++ b/Readme.md @@ -11,7 +11,7 @@ This compiler supports OpenCL C (via LLVM or [SPIRV-LLVM](https://github.com/Khr - A C++14 capable compiler (Raspbian stretch ships with GCC 6.x, which is supported. GCC 4.9 used by Raspbian jessie is no longer supported! The clang from the Raspbian repository is also okay and can double as OpenCL compiler, see below) - CMake in version >= 3.1 -- A suitable OpenCL 1.2 compiler. Supported versions are the "original" LLVM/CLang (version 3.9 and up, **using this one is recommended**), which can be found in the Raspbian repositories, and Khronos [SPIRV-LLVM](https://github.com/KhronosGroup/SPIRV-LLVM) with the Khronos [SPIR-V compiler front-end](https://github.com/KhronosGroup/SPIR/tree/spirv-1.0) (only the tools clang and llvm-spirv need to be built). +- A suitable OpenCL 1.2 compiler. Supported versions are the "original" LLVM/CLang (version 3.9 and up, **using this one is recommended**), which can be found in the Raspbian repositories, the Khronos [SPIRV-LLVM Translator](https://github.com/KhronosGroup/SPIRV-LLVM-Translator) built on top of the "original" LLVM/CLang or the deprecated Khronos [SPIRV-LLVM](https://github.com/KhronosGroup/SPIRV-LLVM) with the Khronos [SPIR-V compiler front-end](https://github.com/KhronosGroup/SPIR/tree/spirv-1.0) (only the tools clang and llvm-spirv need to be built). - The source-code for [VC4CLStdLib](https://github.com/doe300/VC4CLStdLib) for the GPU-side standard-library ## Build diff --git a/cmake/spirv.cmake b/cmake/spirv.cmake index b695fbeb..480ed517 100644 --- a/cmake/spirv.cmake +++ b/cmake/spirv.cmake @@ -78,7 +78,7 @@ if(SPIRV_LLVM_SPIR_FOUND AND SPIRV_FRONTEND) set(SPIRV-Headers_SOURCE_DIR ${SPIRV_HEADERS_SOURCE_DIR}) FetchContent_MakeAvailable(spirv-tools-project) # Set variables expected by the VC4CC library to the SPIRV-Tools libraries. - set(SPIRV_Tools_LIBS SPIRV-Tools SPIRV-Tools-link) + set(SPIRV_Tools_LIBS SPIRV-Tools SPIRV-Tools-opt SPIRV-Tools-link) # the headers are already included by linking the targets set(SPIRV_Tools_HEADERS "") # we need to export these targets, since they are required by VC4CC which we export @@ -116,7 +116,7 @@ if(SPIRV_LLVM_SPIR_FOUND AND SPIRV_FRONTEND) ) ExternalProject_Get_Property(spirv-tools-project BINARY_DIR) ExternalProject_Get_Property(spirv-tools-project SOURCE_DIR) - set(SPIRV_Tools_LIBS "-Wl,--whole-archive ${BINARY_DIR}/source/libSPIRV-Tools.a ${BINARY_DIR}/source/link/libSPIRV-Tools-link.a -Wl,--no-whole-archive") + set(SPIRV_Tools_LIBS "-Wl,--whole-archive ${BINARY_DIR}/source/libSPIRV-Tools.a ${BINARY_DIR}/source/opt/libSPIRV-Tools-opt.a ${BINARY_DIR}/source/link/libSPIRV-Tools-link.a -Wl,--no-whole-archive") set(SPIRV_Tools_HEADERS ${SOURCE_DIR}/include) # This target is used to collect the dependencies on all SPIR-V library build steps @@ -125,4 +125,4 @@ if(SPIRV_LLVM_SPIR_FOUND AND SPIRV_FRONTEND) add_dependencies(SPIRV-Dependencies spirv-tools-project-build) endif() set(VC4C_ENABLE_SPIRV_FRONTEND ON) -endif() \ No newline at end of file +endif() diff --git a/src/precompilation/FrontendCompiler.cpp b/src/precompilation/FrontendCompiler.cpp index a94fdc8d..88679963 100644 --- a/src/precompilation/FrontendCompiler.cpp +++ b/src/precompilation/FrontendCompiler.cpp @@ -318,8 +318,16 @@ void precompilation::linkLLVMModules( const std::string out = result.file ? std::string("-o=") + result.file.value() : ""; std::string inputs = std::accumulate( sources.begin(), sources.end(), std::string{}, [&](const std::string& a, const LLVMIRSource& b) -> std::string { + /* + * If we have multiple input files compiled with the VC4CC compiler, then they might all contain the + * definition/implementation of one or more VC4CL std-lib functions (e.g. get_global_id()). + * To not fail on ODR violations, we allow all but the first linked in modules to simply override already + * defined symbols from the previous modules. + * TODO is there a better solution? + */ + auto separator = a.empty() ? " " : " -override="; if(b.file) - return (a + " ") + b.file.value(); + return a + separator + b.file.value(); if(inputStream != nullptr) { // there already is a stream input, need to move this input to temporary file @@ -328,7 +336,7 @@ void precompilation::linkLLVMModules( auto& file = tempFiles.back(); file->openOutputStream(s); (*s) << b.stream->rdbuf(); - return (a + " ") + file->fileName; + return a + separator + file->fileName; } else { diff --git a/test/TestFrontends.cpp b/test/TestFrontends.cpp index 0b336422..66f9c3d4 100644 --- a/test/TestFrontends.cpp +++ b/test/TestFrontends.cpp @@ -78,13 +78,17 @@ void TestFrontends::testLinking() } std::unique_ptr file0(new std::ifstream("./testing/test_linking_0.cl")); - std::unique_ptr file1(new std::ifstream("./testing/test_linking_0.cl")); + std::unique_ptr file1(new std::ifstream("./testing/test_linking_1.cl")); + std::unique_ptr file2(new std::ifstream("./testing/test_linking_2.cl")); std::unordered_map> inputs{ {file0.get(), Optional{"./testing/test_linking_0.cl"}}, - {file1.get(), Optional{"./testing/test_linking_1.cl"}}}; + {file1.get(), Optional{"./testing/test_linking_1.cl"}}, + {file2.get(), Optional{"./testing/test_linking_2.cl"}}}; std::stringstream tmp; - auto type = Precompiler::linkSourceCode(inputs, tmp); + // extra linking in std-lib tests handling of multiple times linking std-lib, since it will also be linked in for + // the OpenCL C -> LLVM IR compilation + auto type = Precompiler::linkSourceCode(inputs, tmp, true); TEST_ASSERT(type == SourceType::LLVM_IR_BIN || type == SourceType::SPIRV_BIN) std::stringstream out; diff --git a/testing/test_linking_0.cl b/testing/test_linking_0.cl index b33d5fdf..fecbc7b0 100644 --- a/testing/test_linking_0.cl +++ b/testing/test_linking_0.cl @@ -3,5 +3,5 @@ extern int get_value(const __global int* in); __kernel void test_linker(__global int* out, const __global int* in) { - *out = get_value(in); -} \ No newline at end of file + out[get_global_id(0)] = get_value(in); +} diff --git a/testing/test_linking_1.cl b/testing/test_linking_1.cl index 4297d3b0..514d6424 100644 --- a/testing/test_linking_1.cl +++ b/testing/test_linking_1.cl @@ -1,5 +1,7 @@ +extern int get_value_inner(const __global int* in); + int get_value(const __global int* in) { - return *in; -} \ No newline at end of file + return get_global_id(0) + get_value_inner(in) % 0xFFFF; +} diff --git a/testing/test_linking_2.cl b/testing/test_linking_2.cl new file mode 100644 index 00000000..4b0ee588 --- /dev/null +++ b/testing/test_linking_2.cl @@ -0,0 +1,5 @@ + +int get_value_inner(const __global int* in) +{ + return get_global_id(0) + *in; +}