Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel ICX Compiler , Intel® DPC++ Compatibility Tool, and GDB*

DPCPP Generates bswap.i64

Michoumichmich
Beginner
899 Views

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

 

 

 

0 Kudos
1 Solution
Viet_H_Intel
Moderator
779 Views

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.


View solution in original post

6 Replies
HemanthCH_Intel
Moderator
869 Views

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.


0 Kudos
Michoumichmich
Beginner
844 Views

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

0 Kudos
HemanthCH_Intel
Moderator
783 Views

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.


0 Kudos
Viet_H_Intel
Moderator
780 Views

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
Moderator
770 Views

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,


0 Kudos
Michoumichmich
Beginner
768 Views

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

Best, 

 

0 Kudos
Reply