intel / llvm

Intel staging area for llvm.org contribution. Home for Intel LLVM-based projects.
Other
1.25k stars 738 forks source link

sycl::complex, std::complex, or sycl::ext::oneapi::experimental::complex #11502

Open jinz2014 opened 1 year ago

jinz2014 commented 1 year ago

Is "sycl::complex" mentioned in the SYCL specification ? Which type is recommended for Intel, AMD, and NVIDIA GPUs ? Thanks.

no template named 'complex' in namespace 'sycl'; did you mean 'std::complex'?
  auto z1 = sycl::complex<double>(r1, r2);
            ^~~~~~~~~~~~~
            std::complex
bader commented 1 year ago

Sounds like a question for @intel/dpcpp-specification-reviewers team.

AlexeySachkov commented 1 year ago

Hi @jinz2014,

From 5.4. Language restrictions for device functions:

The rules for kernels apply to both the kernel function objects themselves and all functions, operators, member functions, constructors and destructors called by the kernel. This means that kernels can only use library functions that have been adapted to work with SYCL. Implementations are not required to support any library routines in kernels beyond those explicitly mentioned as usable in kernels in this spec. Developers should refer to the SYCL built-in functions in Section 4.17 to find functions that are specified to be usable in kernels.

complex is not a part of the SYCL core specification.

It seems like our intel/llvm implementation supports std::complex through devicelib, at least we have some tests for it: complex<float>, complex<double>

At the same time, at intel/llvm we have a couple of extensions: sycl_ext_oneapi_complex introduces sycl::ext::oneapi::experimental::complex. According to the description, sycl::complex is supposed to be better integrated with other SYCL stuff like vec or marray.

We also have an extension, which enables both std::complex and sycl::complex in SYCL group algorithms: sycl_ext_oneapi_complex_algorithms

gmlueck commented 1 year ago

Note that sycl_ext_oneapi_complex_algorithms is only a proposed extension, so application's shouldn't use it. However, sycl_ext_oneapi_complex is implemented as "experimental". Applications can use an experimental extension, but should be aware that the API might change in a future release of the compiler.

jinz2014 commented 1 year ago

Could the "experimental" API change soon so that "sycl::complex" will not cause compiler errors ?

AlexeySachkov commented 1 year ago

Could the "experimental" API change soon so that "sycl::complex" will not cause compiler errors ?

I don't think that we will do that. Whilst I don't see a strict restriction from vendors putting something extra things into sycl:: namespace, it goes against the guidelines for portable extensions. To make it sycl::complex, it should be a part of the core SYCL specification.

KornevNikita commented 5 months ago

@jinz2014 hi! Did you got an answer to your question? Can we close this?

jinz2014 commented 4 months ago

Does the icpx compiler support the extension ?

error: no member named 'complex' in namespace 'sycl::ext::oneapi::experimental'
   14 |   auto z1 = sycl::ext::oneapi::experimental::complex<float>(r1, r2);

There is an error when building a SYCL program using the clang++ compiler with CUDA support. The program is located in https://github.com/zjin-lcf/HeCBench/tree/master/src/complex-sycl

clang++  -std=c++17 -Wall -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 --gcc-toolchain= -O3 -DUSE_GPU main.o -o main
ptxas fatal   : Unresolved extern function 'cabsf'
llvm-foreach:
clang++: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 18.0.0git (https://github.com/intel/llvm.git 1e36e61ca14cae4663d0a0dab6b987a388ceab0e)
Target: x86_64-unknown-linux-gnu
Thread model: posix
abagusetty commented 3 months ago

@jinz2014 Looks like one would have to (a) enable the extension explicitly, (b) had to include the header manually, (c) Use the expt name space function APIs

For instance,

#define SYCL_EXT_ONEAPI_COMPLEX
#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/experimental/complex/complex.hpp>

return sycl::ext::oneapi::experimental::abs(static_cast<sycl::ext::oneapi::experimental::complex<float>>(a));
jinz2014 commented 3 months ago

@abagusetty Thanks for your answer. Could you reproduce the error ?

icpx -std=c++17 -Wall -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 --gcc-toolchain= -O3 -DUSE_GPU main.o -o main ptxas fatal : Unresolved extern function 'cabsf' llvm-foreach: icpx: error: ptxas command failed with exit code 255 (use -v to see invocation) Intel(R) oneAPI DPC++/C++ Compiler 2024.2.0 (2024.2.0.20240602)