topic Re:DPCPP Generate bswap in Intel® oneAPI DPC++/C++ Compiler

DPCPP Generates bswap.i64

Michoumichmich — Thu, 06 Jan 2022 17:48:03 GMT

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

Re:DPCPP Generate bswap

HemanthCH_Intel — Fri, 07 Jan 2022 10:49:17 GMT

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.

Re: DPCPP Generates bswap.i64

Michoumichmich — Mon, 10 Jan 2022 10:43:40 GMT

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

Re:DPCPP Generate bswap

HemanthCH_Intel — Fri, 14 Jan 2022 12:19:24 GMT

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.

Re:DPCPP Generate bswap

Viet_H_Intel — Fri, 14 Jan 2022 15:38:28 GMT

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.

Re:DPCPP Generate bswap

Viet_H_Intel — Fri, 14 Jan 2022 16:02:25 GMT

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,

Re: DPCPP Generates bswap.i64

Michoumichmich — Fri, 14 Jan 2022 16:04:08 GMT

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

Best,