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

@@ -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:

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

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()
+endif()

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
             {

test/TestFrontends.cpp

@@ -78,13 +78,17 @@ void TestFrontends::testLinking()
     }
 
     std::unique_ptr<std::istream> file0(new std::ifstream("./testing/test_linking_0.cl"));
-    std::unique_ptr<std::istream> file1(new std::ifstream("./testing/test_linking_0.cl"));
+    std::unique_ptr<std::istream> file1(new std::ifstream("./testing/test_linking_1.cl"));
+    std::unique_ptr<std::istream> file2(new std::ifstream("./testing/test_linking_2.cl"));
     std::unordered_map<std::istream*, Optional<std::string>> inputs{
         {file0.get(), Optional<std::string>{"./testing/test_linking_0.cl"}},
-        {file1.get(), Optional<std::string>{"./testing/test_linking_1.cl"}}};
+        {file1.get(), Optional<std::string>{"./testing/test_linking_1.cl"}},
+        {file2.get(), Optional<std::string>{"./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;

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);
-}
+    out[get_global_id(0)] = get_value(in);
+}

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;
-}
+    return get_global_id(0) + get_value_inner(in) % 0xFFFF;
+}

testing/test_linking_2.cl

@@ -0,0 +1,5 @@
+
+int get_value_inner(const __global int* in)
+{
+    return get_global_id(0) + *in;
+}