Skip to content
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

rocPRIM test_warp_load failing with Intel CPU driver #480

Open
pjaaskel opened this issue Jun 7, 2023 · 2 comments · Fixed by #533
Open

rocPRIM test_warp_load failing with Intel CPU driver #480

pjaaskel opened this issue Jun 7, 2023 · 2 comments · Fixed by #533
Assignees
@pjaaskel
Labels
Intel CPU OpenCL

Comments

@pjaaskel
Copy link
Collaborator

pjaaskel commented Jun 7, 2023

Both Intel's and PoCL's CPU drivers produce wrong results with this quite simple test. As it's very short test it's something to investigate and could affect other kernels too. Maybe assumes warp lock-step semantics or multiple writers to same destination to succeed or such.
@pjaaskel pjaaskel added this to the 1.0 milestone Jun 7, 2023
@pjaaskel pjaaskel self-assigned this Jul 7, 2023
@pjaaskel
Copy link
Collaborator Author

pjaaskel commented Jul 7, 2023

Still fails with PoCL-CPU (v4.0) and Intel CPU.

[  PASSED  ] 16 tests.
[  FAILED  ] 14 tests, listed below:
[  FAILED  ] WarpLoadTest/8.WarpLoad, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)0>
[  FAILED  ] WarpLoadTest/8.WarpLoadGuarded, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)0>
[  FAILED  ] WarpLoadTest/9.WarpLoad, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)1>
[  FAILED  ] WarpLoadTest/9.WarpLoadGuarded, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)1>
[  FAILED  ] WarpLoadTest/10.WarpLoad, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)2>
[  FAILED  ] WarpLoadTest/10.WarpLoadGuarded, where TypeParam = Params<int, 4u, 32u, (rocprim::warp_load_method)2>
[  FAILED  ] WarpLoadTest/11.WarpLoad, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)0>
[  FAILED  ] WarpLoadTest/11.WarpLoadGuarded, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)0>
[  FAILED  ] WarpLoadTest/12.WarpLoad, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)1>
[  FAILED  ] WarpLoadTest/12.WarpLoadGuarded, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)1>
[  FAILED  ] WarpLoadTest/13.WarpLoad, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)2>
[  FAILED  ] WarpLoadTest/13.WarpLoadGuarded, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)2>
[  FAILED  ] WarpLoadTest/14.WarpLoad, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)3>
[  FAILED  ] WarpLoadTest/14.WarpLoadGuarded, where TypeParam = Params<HIP_vector_type<float, 2u>, 4u, 32u, (rocprim::warp_load_method)3>

Works with Intel iGPU via OpenCL, Intel iGPU via Level Zero and PoCL via the Level0 driver, so not blocking the 1.0 due to this.

@pjaaskel pjaaskel removed this from the 1.0 milestone Jul 7, 2023
pjaaskel added a commit to pjaaskel/chip-spv that referenced this issue Jul 10, 2023
@pjaaskel
Add get_sub_group_load_id() to the warpfuncs
0aa5cc2 
This is done to trigger required subgroup size to
warp size when there are warp-id-dependent memory accesses
in the kernel. This fixes CHIP-SPV#480 for PoCL-CPU.
@pjaaskel pjaaskel mentioned this issue Jul 10, 2023
Merged
pvelesko added a commit that referenced this issue Jul 11, 2023
@pjaaskel @pvelesko
Add get_sub_group_load_id() to the warpfuncs (#533)
0404178 
This is done to trigger required subgroup size to
warp size when there are warp-id-dependent memory accesses
in the kernel. This fixes #480 for PoCL-CPU.

Co-authored-by: Paulius Velesko <[email protected]>
@pjaaskel
Copy link
Collaborator Author

This still fails with Intel OpenCL CPU, so keeping it open.

@pjaaskel pjaaskel reopened this Jul 11, 2023
@pjaaskel pjaaskel changed the title rocPRIM test_warp_load failing with CPU drivers rocPRIM test_warp_load failing with Intel CPU driver Jul 11, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@pjaaskel pjaaskel

Labels
Intel CPU OpenCL
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Add get_sub_group_load_id() to the warpfuncs pjaaskel/chip-spv
1 participant
@pjaaskel