llvm
llvm copied to clipboard
[SYCL] Enable SPV_KHR_bit_instructions for SYCL code
SPV_KHR_bit_instructions is required when using __builtin_bitreverse32(). SYCL code that was compiled without that extension enabled and would fail when __builtin_bitreverse32() were present in source.
Example of code that would not compile:
#include <sycl/sycl.hpp>
using namespace sycl;
int main(int argc, char **argv)
{
queue q;
int *out = malloc_shared<int>(1024, q);
q.submit([&](handler &h) {
h.parallel_for(range<1>(1024), [=](sycl::id<1> i){
out[i] = __builtin_bitreverse32(i);
});
});
free(out);
return 0;
}
The error message the compiler would give when called via clang++ -fsycl bitreverse.cpp is:
InvalidFunctionCall: Unexpected llvm intrinsic:
Translation of llvm.bitreverse intrinsic requires SPV_KHR_bit_instructions extension.
%9 = call i32 @llvm.bitreverse.i32(i32 %conv.i.i)
llvm-foreach:
clang++: error: llvm-spirv command failed with exit code 8 (use -v to see invocation)
This used to work until commit 85acfa24169fe47fc8b23ada866c42ba0a587647
@intel/llvm-gatekeepers, this should be ready for merge - thanks!