llvm icon indicating copy to clipboard operation
llvm copied to clipboard

[SYCL] Enable SPV_KHR_bit_instructions for SYCL code

Open stefanatwork opened this issue 1 year ago • 2 comments

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.

stefanatwork avatar Feb 19 '24 07:02 stefanatwork

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)

stefanatwork avatar Feb 19 '24 07:02 stefanatwork

This used to work until commit 85acfa24169fe47fc8b23ada866c42ba0a587647

stefanatwork avatar Feb 19 '24 07:02 stefanatwork

@intel/llvm-gatekeepers, this should be ready for merge - thanks!

mdtoguchi avatar Feb 29 '24 15:02 mdtoguchi