From 1088f84dddc9840286a40960d75613937b1c3152 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Henry=20Linjam=C3=A4ki?= Date: Thu, 14 Dec 2023 11:51:38 +0200 Subject: [PATCH 1/4] Link __ballot() implementation at runtime Both the level zero and the OpenCL implementations may support SPIR-V counterpart of __ballot() below SPIR-V v1.3 with appropriate extension. However, there does not seem to be a way to produce SPIR-V with ballot instructions on version v1.2 and below without using chipStar's fork of the llvm-spirv. To get around this limitation the ballot() is put into separate source, compiled to SPIR-V v1.3 and linked at runtime if the device supports it. Consequently, this patch should make the llvm-spirv fork unneeded. --- bitcode/CMakeLists.txt | 45 +++++++++++++++++++------ bitcode/ballot_native.cl | 37 ++++++++++++++++++++ bitcode/devicelib.cl | 10 ++---- src/CHIPBackend.hh | 2 ++ src/backend/Level0/CHIPBackendLevel0.cc | 4 +++ src/backend/Level0/CHIPBackendLevel0.hh | 2 ++ src/backend/OpenCL/CHIPBackendOpenCL.cc | 23 +++++++++---- src/backend/OpenCL/CHIPBackendOpenCL.hh | 3 ++ src/spirv.cc | 8 +++-- tests/runtime/CMakeLists.txt | 2 ++ tests/runtime/TestBallot.hip | 42 +++++++++++++++++++++++ 11 files changed, 151 insertions(+), 27 deletions(-) create mode 100644 bitcode/ballot_native.cl create mode 100644 tests/runtime/TestBallot.hip diff --git a/bitcode/CMakeLists.txt b/bitcode/CMakeLists.txt index 90d31db39..7bb2277cd 100644 --- a/bitcode/CMakeLists.txt +++ b/bitcode/CMakeLists.txt @@ -53,7 +53,8 @@ endif() set(BITCODE_CL_COMPILE_FLAGS -Xclang -finclude-default-header -O2 -x cl -cl-std=CL2.0 - --target=${BC_TRIPLE} -DDEFAULT_WARP_SIZE=${DEFAULT_WARP_SIZE} -emit-llvm + -cl-ext=+cl_khr_subgroup --target=${BC_TRIPLE} + -DDEFAULT_WARP_SIZE=${DEFAULT_WARP_SIZE} -emit-llvm ${EXTRA_FLAGS}) set(BITCODE_C_COMPILE_FLAGS --target=${BC_TRIPLE} -O2 -x c -DDEFAULT_WARP_SIZE=${DEFAULT_WARP_SIZE} @@ -153,19 +154,25 @@ install(FILES "${CMAKE_BINARY_DIR}/${BC_DESTINATION}/${BC_FILE}" # with the extension's atomic operations. Otherwise, the runtime links # in a slower, emulated version. # -# RTDEVLIB_SOURCES defines OpenCL C sources for the rtdevlib. They are +# RTDEVLIB_SOURCES* defines OpenCL C sources for the rtdevlib. They are # compiled to SPIR-V binary and embedded into the CHIP # library. /bitcode/rtdevlib-modules.h declares the # embedded modules as 'std::array '. - +# # Use only characters allowed in C/C++/OpenCL-C language for the # source names. -set(RTDEVLIB_SOURCES +# +# Sources requiring SPIR-V 1.2 at most. +set(RTDEVLIB_SOURCES_v1_2 atomicAddFloat_native atomicAddFloat_emulation atomicAddDouble_native atomicAddDouble_emulation) +# Sources requiring SPIR-V 1.3 at most. +set(RTDEVLIB_SOURCES_v1_3 + ballot_native) + # Alias for add_opencl_bitcode. -foreach(SOURCE IN LISTS RTDEVLIB_SOURCES) +foreach(SOURCE IN LISTS RTDEVLIB_SOURCES_v1_2 RTDEVLIB_SOURCES_v1_3) add_opencl_bitcode( "${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE}.cl" "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.bc") @@ -176,7 +183,8 @@ add_custom_target("rtdevlib-bitcodes" DEPENDS ${RTDEVLIB_BITCODES}) # Compile LLVM bitcode to SPIR-V binary which is then embedded into # std::array . -function(embed_spirv_in_cpp ARRAY_NAME BC_SOURCE OUTPUT_SOURCE OUTPUT_HEADER) +function(embed_spirv_in_cpp + ARRAY_NAME BC_SOURCE OUTPUT_SOURCE OUTPUT_HEADER MAX_SPIRV_VERSION) set(SPIRV_EXTENSIONS "+SPV_EXT_shader_atomic_float_add") get_filename_component(SOURCE_BASENAME "${BC_SOURCE}" NAME_WLE) # Name of the intermediate SPIR-V binary. The name of the C array will be @@ -186,7 +194,9 @@ function(embed_spirv_in_cpp ARRAY_NAME BC_SOURCE OUTPUT_SOURCE OUTPUT_HEADER) OUTPUT "${OUTPUT_SOURCE}" "${OUTPUT_HEADER}" DEPENDS "${BC_SOURCE}" BYPRODUCTS "${SPIR_BINARY}" - COMMAND "${LLVM_SPIRV}" --spirv-ext=${SPIRV_EXTENSIONS} + COMMAND "${LLVM_SPIRV}" + --spirv-ext=${SPIRV_EXTENSIONS} + --spirv-max-version=${MAX_SPIRV_VERSION} "${BC_SOURCE}" -o "${SPIR_BINARY}" COMMAND ${CMAKE_SOURCE_DIR}/scripts/embed-binary-in-cpp.bash ${ARRAY_NAME} ${SPIR_BINARY} ${OUTPUT_SOURCE} ${OUTPUT_HEADER} @@ -195,13 +205,27 @@ function(embed_spirv_in_cpp ARRAY_NAME BC_SOURCE OUTPUT_SOURCE OUTPUT_HEADER) ) endfunction() -foreach(SOURCE IN LISTS RTDEVLIB_SOURCES) +foreach(SOURCE IN LISTS RTDEVLIB_SOURCES_v1_2) embed_spirv_in_cpp( ${SOURCE} "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.bc" "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.cc" + "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.h" + "1.2") + list(APPEND RTDEVLIB_SOURCES_v1_2 + "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.cc") + list(APPEND RTDEVLIB_HEADERS "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.h") - list(APPEND RTDEVLIB_SOURCES +endforeach() + +foreach(SOURCE IN LISTS RTDEVLIB_SOURCES_v1_3) + embed_spirv_in_cpp( + ${SOURCE} + "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.bc" + "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.cc" + "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.h" + "1.3") + list(APPEND RTDEVLIB_SOURCES_v1_3 "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.cc") list(APPEND RTDEVLIB_HEADERS "${CMAKE_CURRENT_BINARY_DIR}/${SOURCE}.h") @@ -220,4 +244,5 @@ add_custom_command( add_custom_target(rtdevlib-header DEPENDS ${RTDEVLIB_HEADER}) -add_library(rtdevlib OBJECT ${RTDEVLIB_SOURCES} ${RTDEVLIB_HEADER}) +add_library(rtdevlib OBJECT + ${RTDEVLIB_SOURCES_v1_2} ${RTDEVLIB_SOURCES_v1_3} ${RTDEVLIB_HEADER}) diff --git a/bitcode/ballot_native.cl b/bitcode/ballot_native.cl new file mode 100644 index 000000000..3b2634dc3 --- /dev/null +++ b/bitcode/ballot_native.cl @@ -0,0 +1,37 @@ +/* + * Copyright (c) 2023 chipStar developers + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS IN THE SOFTWARE. + */ +// __ballot implementation with sub_group_ballot(). + +// SPIR-V requirements: 1.3+ unless environment supports +// cl_khr_subgroup_ballot or ZE_extension_subgroups. +// OpenCL requirements: cl_khr_subgroup_ballot. +// Level0 requirements: ZE_extension_subgroups. + +#define OVERLOADED __attribute__((overloadable)) + +OVERLOADED ulong __chip_ballot(int predicate) { +#if DEFAULT_WARP_SIZE <= 32 + return sub_group_ballot(predicate).x; +#else + return sub_group_ballot(predicate).x | (sub_group_ballot(predicate).y << 32); +#endif +} diff --git a/bitcode/devicelib.cl b/bitcode/devicelib.cl index d94ac9663..eb66f4305 100644 --- a/bitcode/devicelib.cl +++ b/bitcode/devicelib.cl @@ -820,14 +820,8 @@ __SHFL_DOWN(ulong); __SHFL_DOWN(float); __SHFL_DOWN(double); -__attribute__((overloadable)) uint4 sub_group_ballot(int predicate); -EXPORT OVLD ulong __chip_ballot(int predicate) { -#if DEFAULT_WARP_SIZE <= 32 - return sub_group_ballot(predicate).x; -#else - return sub_group_ballot(predicate).x | (sub_group_ballot(predicate).y << 32); -#endif -} +// The definition is linked at runtime from one of the ballot*.cl files. +EXPORT OVLD ulong __chip_ballot(int predicate); EXPORT OVLD int __chip_all(int predicate) { return __chip_ballot(predicate) == ~0; diff --git a/src/CHIPBackend.hh b/src/CHIPBackend.hh index 333c8e272..187801427 100644 --- a/src/CHIPBackend.hh +++ b/src/CHIPBackend.hh @@ -1596,6 +1596,8 @@ public: void invalidateDeviceVariables(); void deallocateDeviceVariables(); + virtual bool hasBallot() const noexcept { return false; } + protected: /** * @brief The backend hook for reset(). diff --git a/src/backend/Level0/CHIPBackendLevel0.cc b/src/backend/Level0/CHIPBackendLevel0.cc index a10edebd6..984b3e303 100644 --- a/src/backend/Level0/CHIPBackendLevel0.cc +++ b/src/backend/Level0/CHIPBackendLevel0.cc @@ -2501,6 +2501,10 @@ static void appendDeviceLibrarySources( else AppendSource(chipstar::atomicAddDouble_emulation); + // OpGroupNonUniformBallot instructions seems to compile and work + // despite not having ZE_extension_subgroups. + AppendSource(chipstar::ballot_native); + assert(SrcSizes.size() == Sources.size() && Sources.size() == BuildFlags.size()); } diff --git a/src/backend/Level0/CHIPBackendLevel0.hh b/src/backend/Level0/CHIPBackendLevel0.hh index 97b918c37..f170ae327 100644 --- a/src/backend/Level0/CHIPBackendLevel0.hh +++ b/src/backend/Level0/CHIPBackendLevel0.hh @@ -620,6 +620,8 @@ public: const ze_float_atomic_ext_properties_t &getFpAtomicProps() const noexcept { return FpAtomicProps_; } + + bool hasBallot() const noexcept override { return true; } }; class CHIPBackendLevel0 : public chipstar::Backend { diff --git a/src/backend/OpenCL/CHIPBackendOpenCL.cc b/src/backend/OpenCL/CHIPBackendOpenCL.cc index a04671e33..e3a8ff66d 100644 --- a/src/backend/OpenCL/CHIPBackendOpenCL.cc +++ b/src/backend/OpenCL/CHIPBackendOpenCL.cc @@ -407,6 +407,10 @@ CHIPDeviceOpenCL::CHIPDeviceOpenCL(CHIPContextOpenCL *ChipCtx, *DevIn, CL_DEVICE_SINGLE_FP_ATOMIC_CAPABILITIES_EXT); Fp64AtomicAddCapabilities_ = getFPAtomicCapabilities( *DevIn, CL_DEVICE_DOUBLE_FP_ATOMIC_CAPABILITIES_EXT); + + auto DevExts = DevIn->getInfo(); + HasSubgroupBallot_ = + DevExts.find("cl_khr_subgroup_ballot") != std::string::npos; } CHIPDeviceOpenCL *CHIPDeviceOpenCL::create(cl::Device *ClDevice, @@ -789,18 +793,23 @@ static void appendRuntimeObjects(cl::Context Ctx, CHIPDeviceOpenCL &ChipDev, // TODO: Reuse already compiled modules. + auto AppendSource = [&](auto &Source) -> void { + Objects.push_back(compileIL(Ctx, ChipDev, Source)); + }; + if (ChipDev.hasFP32AtomicAdd()) - Objects.push_back(compileIL(Ctx, ChipDev, chipstar::atomicAddFloat_native)); + AppendSource(chipstar::atomicAddFloat_native); else - Objects.push_back( - compileIL(Ctx, ChipDev, chipstar::atomicAddFloat_emulation)); + AppendSource(chipstar::atomicAddFloat_emulation); if (ChipDev.hasFP64AtomicAdd()) - Objects.push_back( - compileIL(Ctx, ChipDev, chipstar::atomicAddDouble_native)); + AppendSource(chipstar::atomicAddDouble_native); else - Objects.push_back( - compileIL(Ctx, ChipDev, chipstar::atomicAddDouble_emulation)); + AppendSource(chipstar::atomicAddDouble_emulation); + + if (ChipDev.hasBallot()) + AppendSource(chipstar::ballot_native); + // No fall-back implementation for ballot - let linker raise an error. } void CHIPModuleOpenCL::compile(chipstar::Device *ChipDev) { diff --git a/src/backend/OpenCL/CHIPBackendOpenCL.hh b/src/backend/OpenCL/CHIPBackendOpenCL.hh index b11331895..6001eba7f 100644 --- a/src/backend/OpenCL/CHIPBackendOpenCL.hh +++ b/src/backend/OpenCL/CHIPBackendOpenCL.hh @@ -199,6 +199,7 @@ private: cl_device_fp_atomic_capabilities_ext Fp32AtomicAddCapabilities_; cl_device_fp_atomic_capabilities_ext Fp64AtomicAddCapabilities_; + bool HasSubgroupBallot_ = false; public: ~CHIPDeviceOpenCL() override { @@ -243,6 +244,8 @@ public: return (Fp64AtomicAddCapabilities_ & CL_DEVICE_GLOBAL_FP_ATOMIC_ADD_EXT) && (Fp64AtomicAddCapabilities_ & CL_DEVICE_LOCAL_FP_ATOMIC_ADD_EXT); } + + bool hasBallot() const noexcept override { return HasSubgroupBallot_; } }; class CHIPQueueOpenCL : public chipstar::Queue { diff --git a/src/spirv.cc b/src/spirv.cc index 6b9e30c46..2bda0004f 100644 --- a/src/spirv.cc +++ b/src/spirv.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include "common.hh" #include "spirv.hh" @@ -810,6 +811,9 @@ bool filterSPIRV(const char *Bytes, size_t NumBytes, std::string &Dst) { Dst.reserve(NumBytes); Dst.append(Bytes, (const char *)WordsPtr); // Copy the header. + // Matches chipStar device library and SPIR-V translator symbols. + auto CompilerMagicSymbol = + std::regex(R"RE((__spirv_|__chip_|_Z\d*__chip_).*)RE"); std::set EntryPoints; std::unordered_set BuiltIns; std::unordered_map MissingDefs; @@ -858,8 +862,8 @@ bool filterSPIRV(const char *Bytes, size_t NumBytes, std::string &Dst) { // // Issue warning unless it's a builtin, magic chipStar or // llvm-spirv symbol. - if (!startsWith(LinkName, "spirv_") && - !startsWith(LinkName, "__chip_") && + if (!std::regex_match(LinkName.begin(), LinkName.end(), + CompilerMagicSymbol) && !BuiltIns.count(Insn.getWord(1))) MissingDefs[Insn.getWord(1)] = LinkName; } diff --git a/tests/runtime/CMakeLists.txt b/tests/runtime/CMakeLists.txt index b5870174d..568a4fc05 100644 --- a/tests/runtime/CMakeLists.txt +++ b/tests/runtime/CMakeLists.txt @@ -93,3 +93,5 @@ add_hip_runtime_test(TestBitInsert.hip) # Imports __bitinsert_*. target_compile_definitions(TestBitInsert PRIVATE CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE) + +add_hip_runtime_test(TestBallot.hip) diff --git a/tests/runtime/TestBallot.hip b/tests/runtime/TestBallot.hip new file mode 100644 index 000000000..13c09775f --- /dev/null +++ b/tests/runtime/TestBallot.hip @@ -0,0 +1,42 @@ +#include "CHIPDriver.hh" +#include +#include +#include + +__global__ void k(uint64_t *Out, uint64_t In) { + auto Votes = __ballot((In >> threadIdx.x) & 1u); + if (threadIdx.x == 0) + *Out = Votes; +} + +void checkBallot(unsigned BlockSize, uint64_t In) { + uint64_t *OutD, OutH = ~In; + (void)hipMalloc(&OutD, sizeof(uint64_t)); + k<<<1, BlockSize>>>(OutD, In); + (void)hipMemcpy(&OutH, OutD, sizeof(uint64_t), hipMemcpyHostToDevice); + + hipDeviceProp_t Props{}; + (void)hipGetDeviceProperties(&Props, 0); + + if (Props.warpSize < 64 && OutH >> Props.warpSize) { + printf("Error: Garbage bits in the __ballot result!\n"); + exit(1); + } + + if (OutH != In) { + printf("BlockSize=%u:\nError: Expected '%lu'. Got '%lu'\n", BlockSize, In, + OutH); + exit(1); + } + (void)hipFree(OutD); +} + +int main() { + if (!Backend->getActiveDevice()->hasBallot()) { + printf("SKIP: device does not support __ballot()\n"); + return 2; + } + + checkBallot(32, 0xBADF00D1); + return 0; +} From a036a35ade8d46f4f7e231b1e14cdfe6f152620b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Henry=20Linjam=C3=A4ki?= Date: Tue, 19 Dec 2023 09:52:29 +0200 Subject: [PATCH 2/4] README.md: Remove mention about chipStar's fork of llvm-spirv --- README.md | 1 - 1 file changed, 1 deletion(-) diff --git a/README.md b/README.md index 09d8fb401..8e426b290 100644 --- a/README.md +++ b/README.md @@ -31,7 +31,6 @@ Release notes for [1.1](docs/release_notes/chipStar_1.1.rst), [1.0](docs/release (e.g. llvm\_release\_170 for LLVM 17) [llvm-spirv](https://github.com/KhronosGroup/SPIRV-LLVM-Translator). * Make sure the built llvm-spirv binary is installed into the same path as clang binary, otherwise clang might find and use a different llvm-spirv, leading to errors. - * For the best results, install it from a chipStar [branch](https://github.com/CHIP-SPV/SPIRV-LLVM-Translator/tree/chipStar-llvm-17) which has fixes that are not yet upstreamed. ### Compiling Clang, LLVM and SPIRV-LLVM-Translator From f722e0b0e0b15a2366dd02c332bddbd5051da8d2 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Henry=20Linjam=C3=A4ki?= Date: Tue, 19 Dec 2023 13:46:21 +0200 Subject: [PATCH 3/4] Update CI to use upstream llvm-spirv --- .github/actions/get-shas/action.yml | 2 +- .github/workflows/presubmit.yml | 2 +- scripts/configure_llvm.sh | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/.github/actions/get-shas/action.yml b/.github/actions/get-shas/action.yml index c96a1042c..8ab1dccce 100644 --- a/.github/actions/get-shas/action.yml +++ b/.github/actions/get-shas/action.yml @@ -32,7 +32,7 @@ runs: - id: spirv-llvm-sha run: scripts/get_sha.sh >> $GITHUB_OUTPUT env: - URL: https://api.github.com/repos/CHIP-SPV/SPIRV-LLVM-Translator/git/ref/heads/chipStar-llvm-${{ inputs.version }} + URL: https://api.github.com/repos/KhronosGroup/SPIRV-LLVM-Translator/git/ref/heads/llvm_release_${{ inputs.version }}0 TOKEN: u:${{github.token}} shell: bash - id: pocl-sha diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index 1faa5c402..a585339de 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -99,7 +99,7 @@ jobs: if: steps.llvm.outputs.cache-hit != 'true' - uses: actions/checkout@v3 with: - repository: CHIP-SPV/SPIRV-LLVM-Translator + repository: KhronosGroup/SPIRV-LLVM-Translator ref: ${{ steps.get-shas.outputs.spirv-llvm-sha }} path: llvm-${{ matrix.version }}/llvm/projects/SPIRV-LLVM-Translator if: steps.llvm.outputs.cache-hit != 'true' diff --git a/scripts/configure_llvm.sh b/scripts/configure_llvm.sh index 8b37d5e9a..c1450ddbc 100755 --- a/scripts/configure_llvm.sh +++ b/scripts/configure_llvm.sh @@ -37,10 +37,10 @@ LINK_TYPE=$3 # set the brach name for checkuot based on only-necessary-spirv-exts if [ "$4" == "on" ]; then LLVM_BRANCH="spirv-ext-fixes-${VERSION}" - TRANSLATOR_BRANCH="chipStar-llvm-${VERSION}" + TRANSLATOR_BRANCH="llvm_release_${VERSION}0" else LLVM_BRANCH="chipStar-llvm-${VERSION}" - TRANSLATOR_BRANCH="chipStar-llvm-${VERSION}" + TRANSLATOR_BRANCH="llvm_release_${VERSION}0" fi export LLVM_DIR=`pwd`/llvm-project/llvm @@ -49,7 +49,7 @@ export LLVM_DIR=`pwd`/llvm-project/llvm if [ ! -d llvm-project ]; then git clone https://github.com/CHIP-SPV/llvm-project.git -b ${LLVM_BRANCH} --depth 1 cd ${LLVM_DIR}/projects - git clone https://github.com/CHIP-SPV/SPIRV-LLVM-Translator.git -b ${TRANSLATOR_BRANCH} --depth 1 + git clone https://github.com/KhronosGroup/SPIRV-LLVM-Translator.git -b ${TRANSLATOR_BRANCH} --depth 1 cd ${LLVM_DIR} else # Warn the user, error out From 64fd61ea79933ee1cda5cc628ae02020fa46b005 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Tue, 19 Dec 2023 09:07:11 -0600 Subject: [PATCH 4/4] switched to unpatched llvm-spirv --- scripts/unit_tests.sh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/scripts/unit_tests.sh b/scripts/unit_tests.sh index d1e8e0dc1..8ecba8d58 100755 --- a/scripts/unit_tests.sh +++ b/scripts/unit_tests.sh @@ -30,13 +30,13 @@ build_type=$(echo "$1" | tr '[:lower:]' '[:upper:]') if [ "$2" == "llvm-15" ]; then LLVM=llvm-15 - CLANG=llvm/15.0/dynamic + CLANG=llvm/15.0-unpatched-spirv elif [ "$2" == "llvm-16" ]; then LLVM=llvm-16 - CLANG=llvm/16.0/dynamic + CLANG=llvm/16.0-unpatched-spirv elif [ "$2" == "llvm-17" ]; then LLVM=llvm-17 - CLANG=llvm/17.0/dynamic + CLANG=llvm/17.0-unpatched-spirv else echo "$2" echo "Invalid 2nd argument. Use either 'llvm-15', 'llvm-16' or 'llvm-17'."