Skip to content

SYCL/SubGroup/shuffle_fp64.cpp test in intel/llvm-test-suite fails for CUDA backend #8516

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Closed
steffenlarsen opened this issue Mar 2, 2023 · 2 comments
Closed

SYCL/SubGroup/shuffle_fp64.cpp test in intel/llvm-test-suite fails for CUDA backend #8516

steffenlarsen opened this issue Mar 2, 2023 · 2 comments
Assignees
@maksimsab
Labels
bug Something isn't working confirmed cuda CUDA back-end

Comments

@steffenlarsen
Copy link
Contributor

The SYCL/SubGroup/shuffle_fp64.cpp test in intel/llvm-test-suite fails after intel/llvm-test-suite#1618 for the CUDA backend with errors like

  .../include/sycl/detail/spirv.hpp:680:17: error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
      result[s] = SubgroupShuffleDown(x[s], delta);
                  ^
  .../include/sycl/detail/spirv.hpp:611:27: note: 'SubgroupShuffleDown<double>' declared here
  EnableIfGenericShuffle<T> SubgroupShuffleDown(T x, id<1> local_id);
                            ^
  .../include/sycl/detail/spirv.hpp:677:26: note: called by 'SubgroupShuffleDown<sycl::vec<double, 16>>'
  EnableIfVectorShuffle<T> SubgroupShuffleDown(T x, uint32_t delta) {

See logs on #8425 for more information.
@steffenlarsen steffenlarsen added bug Something isn't working cuda CUDA back-end confirmed labels Mar 2, 2023
@steffenlarsen steffenlarsen mentioned this issue Mar 2, 2023
Merged
steffenlarsen added a commit to steffenlarsen/llvm-test-suite that referenced this issue Mar 2, 2023
@steffenlarsen
[SYCL][CUDA] Disable subgroup shuffle_fp64 test for CUDA
ccfc7f5 
The CUDA backend is failing the shuffle_fp64 after recent additions.
This has been reported in intel/llvm#8516 and
this commit disables the test for CUDA while it is fixed.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
This was referenced Mar 2, 2023
Merged
Merged
Merged
bader pushed a commit to intel/llvm-test-suite that referenced this issue Mar 4, 2023
@steffenlarsen
[SYCL][CUDA] Disable subgroup shuffle_fp64 test for CUDA (#1626)
a7917bf 
The CUDA backend is failing the shuffle_fp64 after recent additions.
This has been reported in intel/llvm#8516 and
this commit disables the test for CUDA while it is fixed.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
@bader bader mentioned this issue Mar 4, 2023
Merged
@steffenlarsen steffenlarsen mentioned this issue Mar 6, 2023
Merged
@KseniyaTikhomirova KseniyaTikhomirova mentioned this issue Mar 6, 2023
Merged
@maksimsab maksimsab self-assigned this Mar 6, 2023
@steffenlarsen steffenlarsen mentioned this issue Mar 8, 2023
Merged
@GeorgeWeb
Copy link
Contributor

GeorgeWeb commented Mar 10, 2023 •
edited
Loading

Hey @steffenlarsen 👋

#8614 fixes this.

georgi@devpc:~/workspaces/forks/llvm-test-suite/build$ /usr/bin/python3.10 \
/home/georgi/workspaces/dpcpp/llvm/build/./bin/llvm-lit \
-sv --param sycl_be=cuda --param target_devices=gpu \
~/workspaces/forks/llvm-test-suite/build/SYCL/SubGroup/shuffle_fp64.cpp
llvm-lit: /home/georgi/workspaces/forks/llvm-test-suite/SYCL/lit.cfg.py:235: note: Backend: ext_oneapi_cuda

Testing Time: 2.75s
  Passed: 1

2 warning(s) in tests

Will open a PR for re-enabling the test, once the fix gets merged.

@GeorgeWeb GeorgeWeb mentioned this issue Mar 10, 2023
Merged
steffenlarsen pushed a commit that referenced this issue Mar 15, 2023
@GeorgeWeb
[SYCL] Fix forward declarations to SubgroupShuffleUp/Down overloads (#…
ecb148b 
…8614)

This PR fixes up the template overload function forward declarations for (Generic) `SubroupShuffleDown` and `SubroupShuffleUp` to use the `uint32_t delta` parameter instead of `id<1> local_id`. Also removed two unnecessary forward declarations  `EnableIfBitcastShuffle SubroupShuffleUp/Down` with `id<1> local_id`.

These forward declarations are needed because the `EnableIfVector` imlpementations call these functions but their bodies are defined after vector versions, so the compiler can't identify them.

Without this fix users will get `SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute ...` compile errors. That error also manifested in issue [#8516](#8516), which will be fixed with the changes of the PR.

Side note: This also fixes compiling the `group_broadcast` tests for the `test_group_functions` target in the SYCL-CTS.
@GeorgeWeb GeorgeWeb mentioned this issue Mar 15, 2023
Merged
@steffenlarsen
Copy link
Contributor Author

Fixed by #8614.

aelovikov-intel pushed a commit to aelovikov-intel/llvm that referenced this issue Mar 27, 2023
@steffenlarsen
[SYCL][CUDA] Disable subgroup shuffle_fp64 test for CUDA (intel/llvm-…
7a12d60 
…test-suite#1626)

The CUDA backend is failing the shuffle_fp64 after recent additions.
This has been reported in intel#8516 and
this commit disables the test for CUDA while it is fixed.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@maksimsab maksimsab

Labels
bug Something isn't working confirmed cuda CUDA back-end
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@GeorgeWeb @steffenlarsen @maksimsab