SYCL builtin support for marray is incomplete

[han-minhee]

Bug summary
I'm actually not sure if it's a bug in AdaptiveCPP or it's an issue with clang
I'm using RX 7900 GRE, which is gfx1100, but acpp doesn't recognize gfx1100 as the target
and acpp --acpp-targets=hip would set the target to gfx906, so that might be the issue.

When using "cos" function in the code, it yields an error during the compilation.

vec.cpp:43:21: error: no matching function for call to 'cos'
   43 |         re_acc[i] = cos(in_acc[0]);
      |                     ^~~

To Reproduce
Try to compile the code below(it's from Data Parallel C++ book)

#include <cassert>
#include <sycl/sycl.hpp>
#define M 32

using namespace sycl;

int main()
{
    queue q;
    marray<float, 4> input{1.0004f, 1e-4f, 1.4f, 14.0f};
    marray<float, 4> res[M];
    for (int i = 0; i < M; i++)
        res[i] = {-(i + 1), -(i + 1), -(i + 1), -(i + 1)};
    {
        buffer in_buf(&input, range{1});
        buffer re_buf(res, range{M});

        q.submit([&](handler &cgh)
                 {
      accessor re_acc{re_buf, cgh, read_write};
      accessor in_acc{in_buf, cgh, read_only};

      cgh.parallel_for(range<1>(M), [=](id<1> idx) {
        int i = idx[0];
        re_acc[i] = cos(in_acc[0]);
      }); });
    }
    return 0;
}

If possible, provide a minimal reproducer. The shorter the reproducing code snippet is, the easier will it be for us to debug and understand the issue.

Expected behavior
A clear and concise description of what you expected to happen.

Describe your setup
Ubuntu 24.04

Ubuntu clang version 18.1.3 (1ubuntu1)
Target: x86_64-pc-linux-gnu
Thread model: posix

ROCm-6.2

acpp vec.cpp -o vec -O3 --acpp-targets=generic

Optional additional diagnostic information
`acpp --hipsycl-version
acpp [AdaptiveCpp compilation driver], Copyright (C) 2018-2024 Aksel Alpay and the AdaptiveCpp project
AdaptiveCpp version: 24.06.0+git.fca48591.20240810.branch.develop
Installation root: /home/minhee/env/adaptiveCPP
Plugin LLVM version: 18, can accelerate CPU: True
Available runtime backends:
librt-backend-omp.so
librt-backend-hip.so

Full configuration [can be overridden using environment variables or command line arguments]:
default-rocm-path: /opt/rocm
default-rocm-link-line: -Wl,-rpath=$HIPSYCL_ROCM_PATH/lib -Wl,-rpath=$HIPSYCL_ROCM_PATH/hip/lib -L/opt/rocm/lib -L/opt/rocm/hip/lib -lamdhip64
default-rocm-cxx-flags: -isystem $HIPSYCL_PATH/include/AdaptiveCpp/hipSYCL/std/hiplike -isystem /usr/lib/llvm-18/lib/clang/18/include/.. -U__FLOAT128__ -U__SIZEOF_FLOAT128__ -I$HIPSYCL_ROCM_PATH/include -I$HIPSYCL_ROCM_PATH/include --rocm-device-lib-path=$HIPSYCL_ROCM_PATH/amdgcn/bitcode --rocm-path=$HIPSYCL_ROCM_PATH -fhip-new-launch-api -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -D__HIP_ROCclr__
version-major: 24
version-minor: 06
version-patch: 0
version-suffix: +git.fca48591.20240810.branch.develop
plugin-llvm-version-major: 18
plugin-with-cpu-acceleration: true
default-clang: /usr/lib/llvm-18/bin/clang++
default-targets: generic
default-cpu-cxx: /usr/bin/c++
default-use-bootstrap-mode: false
default-is-dryrun: false
default-use-accelerated-cpu: true
default-clang-include-path: /usr/lib/llvm-18/lib/clang/18/include/..
default-sequential-link-line: -L/usr/lib/x86_64-linux-gnu -lboost_context -lboost_fiber -Wl,-rpath=/usr/lib/x86_64-linux-gnu
default-sequential-cxx-flags: -I/usr/include -D_ENABLE_EXTENDED_ALIGNED_STORAGE
default-omp-link-line: -L/usr/lib/x86_64-linux-gnu -lboost_context -lboost_fiber -Wl,-rpath=/usr/lib/x86_64-linux-gnu -fopenmp
default-omp-cxx-flags: -I/usr/include -fopenmp -D_ENABLE_EXTENDED_ALIGNED_STORAGE
default-is-explicit-multipass: false
default-save-temps: false`

[illuhad]

I'm using RX 7900 GRE, which is gfx1100, but acpp doesn't recognize gfx1100 as the target
and acpp --acpp-targets=hip would set the target to gfx906, so that might be the issue.

No, it's not.

If gfx1100 is not recognized then your GPU might be detected as something else by ROCm (see acpp-info).
Unless you have very good reason and know what you are doing, the recommendation is also to use --acpp-targets=generic (which is the default in a standard installation) since it usually generates faster binaries, has more features and is more portable.

If you use --acpp-targets=hip, you need to follow up with the exact device architecture, otherwise it might just fall back to some default (e.g. gfx906): --acpp-targets=hip:gfx1100.

But this is all not related to your problem. Your problem is with the code. If you call the unqualified cos and have using namespace sycl in your code, then you have an ambiguous function call because there is both std::cos and sycl::cos. This code is not guaranteed to work in the SYCL specification.
It's probably generally bad practice to use using namespace sycl, especially when you use types or functions that also occur in std::. queue might also be affected here since there is also an std::queue.
Qualify the cos call as sycl::cos and optionally remove using namespace sycl.
I don't know to what extent this issue is explained in the DPC++ book, so tagging some of the authors for awareness:

EDIT: Sorry, this is also not the actual problem - you don't have using namespace std, so there's not actually an ambiguous overload. But it seems our builtin definitions for marray type are incomplete.

[han-minhee]

Yeah I tried it with generic at the beginning, but thought that might be an issue; I thought that was the right way to pass the target device code name. Just a question: why would compiling for generic target would yield binaries faster than specifying the target device?

[illuhad]

Just a question: why would compiling for generic target would yield binaries faster than specifying the target device?

generic and hip are two very different compilation flows. hip goes through the clang HIP toolchain, which AdaptiveCpp augments to also understand SYCL code. The primary use case for this is source code interoperability or code migration from CUDA/HIP.
generic utilizes AdaptiveCpp's own generic single-pass compiler. This is our standalone AdaptiveCpp-specific compiler. There are a number of reasons why this is faster and more feature-rich than going through the clang HIP toolchain:

• The vast majority of our optimization effort focuses on this compiler since it is our primary compiler.
• It is a JIT-compiler and actively takes into account additional information that is only available at runtime during code generation/optimization. --acpp-targets=hip cannot do that.
• Since it is our own, AdaptiveCpp-specific compiler, it has much better integration into our runtime infrastructure and into SYCL and thus is also more feature rich.

generic as a name is perhaps a bit misleading - it does not mean "compile and run a generic target similary to gcc -mtune=generic on CPU". It means: Embed kernels as a generic, backend-independent code representation, and lower to target-specific code at runtime once the target device is known. It's a just-in-time compiler.

[han-minhee]

Ah I see! Thank you for the explanation!

[nilsfriess]

I have started to map the builtins for marray. Just wanted to make a quick note here so that the work is not done twice.