Multiple simd tuple tests failing on GPU

lbushi25 commented 1 month ago

Describe the bug

The tests

  SYCL :: InvokeSimd/Spec/ImplicitSubgroup/tuple.cpp,
  SYCL :: InvokeSimd/Spec/ImplicitSubgroup/tuple_return.cpp,
  SYCL :: InvokeSimd/Spec/ImplicitSubgroup/tuple_vadd.cpp,
  SYCL :: InvokeSimd/Spec/tuple.cpp,
  SYCL :: InvokeSimd/Spec/tuple_return.cpp,
  SYCL :: InvokeSimd/Spec/tuple_vadd.cpp

are failing on GPU. Upon resolution, please remove XFAIL from test sources. Snippet taken from: https://github.com/intel/llvm/actions/runs/9716365508/job/26824198678?pr=14354

To reproduce

FAIL: SYCL :: InvokeSimd/Spec/ImplicitSubgroup/tuple.cpp (532 of 668)
******************** TEST 'SYCL :: InvokeSimd/Spec/ImplicitSubgroup/tuple.cpp' FAILED ********************
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 4
/__w/llvm/llvm/toolchain/bin//clang++   -DIMPL_SUBGROUP -fsycl -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr /__w/llvm/llvm/llvm/sycl/test-e2e/InvokeSimd/Spec/ImplicitSubgroup/../tuple.cpp -o /__w/llvm/llvm/build-e2e/InvokeSimd/Spec/ImplicitSubgroup/Output/tuple.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -DIMPL_SUBGROUP -fsycl -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr /__w/llvm/llvm/llvm/sycl/test-e2e/InvokeSimd/Spec/ImplicitSubgroup/../tuple.cpp -o /__w/llvm/llvm/build-e2e/InvokeSimd/Spec/ImplicitSubgroup/Output/tuple.cpp.tmp.out
# .---command stderr------------
# | In file included from /__w/llvm/llvm/llvm/sycl/test-e2e/InvokeSimd/Spec/ImplicitSubgroup/../tuple.cpp:25:
# | In file included from /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/intel/esimd.hpp:83:
# | In file included from /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/intel/esimd/alt_ui.hpp:13:
# | In file included from /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/intel/esimd/simd.hpp:22:
# | /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/oneapi/experimental/invoke_simd.hpp:427:17: error: static assertion failed due to requirement '!has_non_uniform_struct_arg': Structure arguments must be uniform
# |   427 |   static_assert(!has_non_uniform_struct_arg,
# |       |                 ^~~~~~~~~~~~~~~~~~~~~~~~~~~
# | /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/oneapi/experimental/invoke_simd.hpp:436:3: note: in instantiation of function template specialization 'sycl::ext::oneapi::experimental::detail::verify_no_non_uniform_struct_args<std::tuple<float, int>, float>' requested here
# |   436 |   verify_no_non_uniform_struct_args<Ts...>();
# |       |   ^
# | /__w/llvm/llvm/toolchain/bin/../include/sycl/ext/oneapi/experimental/invoke_simd.hpp:470:11: note: in instantiation of function template specialization 'sycl::ext::oneapi::experimental::detail::verify_valid_args_and_ret<std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>> (&)(std::tuple<std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>, std::experimental::simd<int, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>>, std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>) __attribute__((regcall)), std::tuple<float, int>, float>' requested here
# |   470 |   detail::verify_valid_args_and_ret<Callable, T...>();
# |       |           ^
# | /__w/llvm/llvm/llvm/sycl/test-e2e/InvokeSimd/Spec/ImplicitSubgroup/../tuple.cpp:102:21: note: in instantiation of function template specialization 'sycl::ext::oneapi::experimental::invoke_simd<std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>> (&)(std::tuple<std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>, std::experimental::simd<int, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>>, std::experimental::simd<float, std::experimental::__simd_abi<std::experimental::_StorageKind::_VecExt, 16>>) __attribute__((regcall)), std::tuple<float, int>, float>' requested here
# |   102 |         float res = invoke_simd(sg, SIMD_CALLEE, tup, A[wi_id]);
# |       |                     ^
# | 1 error generated.
# `-----------------------------
# error: command failed with exit status: 1

Environment

OS: Agnostic
Target device and vendor: GPU
DPC++ version: [e.g. commit hash or output of clang++ --version]
Dependencies version: [e.g. the output of sycl-ls --verbose]

Additional context

No response

sarnex commented 1 month ago

@lbushi25 This is expected, it's not implemented yet. Should I just close this issue?

lbushi25 commented 1 month ago

@lbushi25 This is expected, it's not implemented yet. Should I just close this issue?

We'd like to keep these open so that people can periodically check whether the issue is closed and take that to mean that the XFAIL can be removed from the test. Unless there is already a github issue tracking this and I've missed it?

sarnex commented 1 month ago

Nope, we can use this one, thanks.

intel / llvm