DPCPP Generates bswap.i64

[Michoumichmich]

Hello, 

When compiling the following byte swap (with x a 64 bit unsigned):

   ((x & 0xff00000000000000ULL) >> 56)
 | ((x & 0x00ff000000000000ULL) >> 40) 
 | ((x & 0x0000ff0000000000ULL) >> 24)
 | ((x & 0x000000ff00000000ULL) >> 8) 
 | ((x & 0x0000000000ff0000ULL) << 24)
 | ((x & 0x000000000000ff00ULL) << 40)
 | ((x & 0x00000000000000ffULL) << 56);

 DPCPP (llvm-spirv) fails with:

InvalidFunctionCall: Unexpected llvm intrinsic:
llvm.bswap.i64

The version of DPCPP is  2022.0.0 (2022.0.0.20211123). It happens when compiling SYCL kernels for spir64 as well as when I use AOT. 

The upstream version of the open source oneAPI compiles the code fine. 

To use Intel(R) oneAPI you can "defeat" the optimiser by replacing the `|` by `+` (which is still correct as the bits do not overlap) and it compiles (but is not optimal haha) . 

Best,

Michel

[HemanthCH_Intel]

Hi,

Thanks for reaching out to us.

Could you please provide the sample reproducer code and steps to reproduce the issue at our end?

Thanks & Regards,

Hemanth.

[Michoumichmich]

Hi, 

You may compile the following: 

#include <CL/sycl.hpp>

static inline uint64_t bswap(uint64_t x){
 return 
   ((x & 0xff00000000000000ULL) >> 56) 
 | ((x & 0x00ff000000000000ULL) >> 40) 
 | ((x & 0x0000ff0000000000ULL) >> 24) 
 | ((x & 0x000000ff00000000ULL) >> 8 )  
 | ((x & 0x0000000000ff0000ULL) << 24) 
 | ((x & 0x000000000000ff00ULL) << 40) 
 | ((x & 0x00000000000000ffULL) << 56);
}

int main(){
 sycl::queue q{};
 auto ptr = sycl::malloc_device<uint64_t>(1, q);
 q.single_task([=](){
     *ptr = bswap(*ptr);
 }).wait();
}

With the flags `dpcpp -fsycl -O3`. On Intel(R) DevCloud it crashes with the error I mentioned previously too.

Best regards, 

Michel

[HemanthCH_Intel]

Hi,

We are able to able to reproduce your issue at our end and working on this internally. We will get back to you soon.

Thanks & Regards,

Hemanth.

[Viet_H_Intel · Accepted solution]

Thank you for reporting this problem to us. This is a known issue and will be fixed in the next update.

Can we close this thread for now?

Regards.

[Viet_H_Intel]

Thanks for accepting this as a solution. This thread is now closed and won't be monitored by Intel. If you need anything else, please create a new thread.

Regards,

[Michoumichmich]

Okay, thanks for looking into this and have a nice day!

Best,