[SYCL] Added tests for atomics with various memory orders and scopes #534

t4c1 · 2021-10-26T13:58:20Z

Added tests for atomics with various memory orders and scopes. Reductions tests also have updated sm requirements, as they call work group atomics, which are now implemented and have higher sm requirements than device scoped ones.

This adds tests for changes introduced in intel/llvm#4820 and intel/llvm#5192.

steffenlarsen

Generally speaking I think the tests are good but as mentioned in one of my comments, the new tests may generate kernels with code that may be unsupported on a device and even if the tests intend to skip these kernels they may still be built together with the others which may cause failure at runtime.

This is not unique to 64-bit atomic tests however, as the same could be the case for the various memory orders and memory scopes. This is also why there are so many variants of the atomic_memory_order_*.cpp tests.

SYCL/AtomicRef/add.h

SYCL/AtomicRef/compare_exchange.h

SYCL/AtomicRef/exchange.h

SYCL/AtomicRef/min.h

SYCL/AtomicRef/xor.h

SYCL/AtomicRef/and.h

steffenlarsen · 2021-11-05T10:16:22Z

SYCL/AtomicRef/add_orders_scopes.cpp

+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    add_test_orders_scopes<long long>(q, N);
+    add_test_orders_scopes<unsigned long long>(q, N);
+  }


The device associated with the queue must support aspect::atomic64 for these tests to be valid. However, even if you check for that here the compiler may still try to generate the kernels for the 64-bit tests which it will try to build with the other kernels, even though it never intends to run them. This is the reason why many other tests have a _atomic64 test variant.

To handle all possible combinations of supported aspects each of these tests would have to be split into many files (for every combination of supported atomic64, 3 memory scopes and 4 memory orders). That would be a bit too many files in my opinion.

In practice memory orders are currently only supported by sm_70 or higher NVidia devices. These also support all memory scopes and atomic64. That is why I decided to put all new tests into a single file that requires sm_70.

If you have a better suggestion how to handle this, I am happy to change it.

I agree, the number of combinations becomes quite a headache, but the limitations of order/scope/atomic64 may apply to other devices than just CUDA ones.

An option would be to guard different combinations with macros you then define in different //RUN commands of the tests. This should make sure that the kernels are only generated for the combination test cases, avoiding compilation of invalid kernels by only running them if the device supports the combination of that test case. Does that make sense?

That would still result in 24 RUN commands per file, but it does allow to merge all test files for each operation into just one.

Although I have to say I do not completely understand what mechanism skips the RUN lines for aspects that are unsupported in the hardware the test is running.

That is a drawback of this approach; you cannot have e.g. XFAIL for each, but since you have to check for the features at runtime anyway the tests would have to dispatch.

The benefit of splitting them up (even through the macro alternative) is that the tests will not load binaries with unsupported instructions for a given device. That is, the individual RUNs would start out by checking that the corresponding order/scope/atomic64 is supported for the given device, and if not it should skip the test (with a print but not an error.)

I still have to do runtime checks? That would result in these tests not being run until intel/llvm#4853 is merged

You have to do the runtime checks to make sure the device actually supports the features. CUDA is special in that regard as the SM version controls whether or not a feature is supported, but that's not the case for all devices.

That would result in these tests not being run until intel/llvm#4853 is merged

Since the patch is up I don't see why we wouldn't wait for it. It would make the tests more robust from the get-go.

I applied this kind of changes to tests for add. Can you check this is what you had in mind before I do the rest?

That is exactly what I had in mind! It is quite a lot of RUN lines, but I suppose I got what was coming. Maybe it would make sense, just for good measure, to have a comment on each RUN grouping to have a short note specifying which test-configuration it is running.

# Conflicts: # SYCL/AtomicRef/add.cpp # SYCL/AtomicRef/add.h # SYCL/AtomicRef/add_atomic64.cpp # SYCL/AtomicRef/compare_exchange.h # SYCL/AtomicRef/exchange.h # SYCL/AtomicRef/max.h # SYCL/AtomicRef/min.h

…test-suite into atomic_orders_scopes2

bader · 2021-12-25T11:52:59Z

@t4c1, could you resolve merge conflicts, please?

Adds implementations for atomic operations that do not have direct PTX equivalents. Sub is implemented by using add. Floating point min and max use compare exchange loops. I pushed the tests for this into: intel/llvm-test-suite#534

t4c1 · 2022-02-09T13:08:59Z

Are any of the test failures related to the changes in this PR?

steffenlarsen · 2022-02-14T14:41:50Z

SYCL/AtomicRef/add_generic.cpp

 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

 // CUDA backend has had no support for the generic address space yet
-// XFAIL: cuda || hip
+// XFAIL: cuda || hip || level_zero || opencl


Comment only mentions why CUDA is disabled. Why are the other backends disabled? (Same applies to a handful of other tests)

It is the same reason as for CUDA. Will update the comments.

Actually ... I might be wrong about that. I need to look into this a bit more.

I have similar concern. Today atomic test are split into two parts: + generic pointers part not supported on HIP and CUDA backends.
This separation allows us:

to have full validation of all atomic features on OpenCL and Level Zero backends

make sure that everything except atomics for generic pointers works on HIP and CUDA

(minor) test that atomics for generic pointers fail on HIP and CUDA

New version doesn't have such clear separation. Let's look at "add" operation.

Test Current New

add.cpp everything works everywhere XFAIL on host, opencl and hip

add_generic.cpp XFAIL on hip and cuda XFAIL on hip, cuda, opencl and level_zero

add_atomic64.cpp everything works everywhere moved to add.cpp

add_atomic64_generic.cpp XFAIL on hip and cuda moved to add_generic.cpp

With new test logic distribution, we don't know if anything is working on OpenCL backends, whereas all test cases are passing. It's because new tests for floating point atomics were added to all test files and they are not supported by OpenCL. In general having multiple checks in a single test file, where at least one of the checks fails, XFAIL hides the status of all successful checks. That was the reason to move tests for "generic" pointer to a separate test.
I think it make sense to do for a new test cases which do not pass on all backends.

@t4c1, what do you think?

I believe I might be able to resolve some of these XFAILs. Ones for floating point atomics (on add, sub, min and max) and for missing barrier on host will stay.

This PR is adding many more different subsets of tests that could be put into separate files. Previously there were no tests for atomics in local memory, atomics with non-relaxed memory orders, atomics with non-device scopes. If I followed the current convention and put every combination in a separate file it would get us a huge number of files (global/local generic/AS, 32b/64b, 3 scopes, 4 orders, totaling 64 files per operation).

We already discussed it here: #534 (comment) and as far as I understood this was deemed an acceptable solution. If, however, you have a better suggestion how to handle this, I am happy to change it.

I'm okay to add new test cases to existing test (e.g. add.cpp) as long as they doesn't change the test status i.e. the test should pass everywhere.

If I followed the current convention and put every combination in a separate file it would get us a huge number of files (global/local generic/AS, 32b/64b, 3 scopes, 4 orders, totaling 64 files per operation).

I think we can put test for features with the same validation status into the same file. Following this logic, I don't think we need a huge number of files. I expect to have 1-2 files per operation. Right?

Hmm, we would still need to split the 3 address spaces and floating point operations, leaving us with 6 files (or 3 for operations without floating point). I guess that is still manageable.

t4c1 · 2022-02-21T08:44:27Z

Can somebody copy the tests results for me again?

bader · 2022-02-21T08:48:43Z

/verify with intel/llvm#5620

bader · 2022-02-21T14:50:20Z

Take a look at "Jenkins/verify-test-suite-logs".

steffenlarsen

Woaw! This has become quite the PR. It looks fine to me.

@alexbatashev - Ping for awareness. This adds a solid handful of tests. Do you have any CI-related concerns?

bader · 2022-02-22T12:22:33Z

@t4c1, FYI.

SYCL/Basic/event_profiling_info.cpp:28: bool verifyProfiling(sycl::event): Assertion `Submit <= Start' failed.

t4c1 · 2022-02-22T12:30:39Z

This PR does not modify SYCL/Basic/event_profiling_info.cpp test, so I doubt the failure is related to this PR.

steffenlarsen · 2022-02-22T12:38:03Z

Event profiling in the CUDA backend has historically been very flaky. I agree with @t4c1 that this is unlikely to be related.

@t4c1 - Would you mind disabling it for CUDA and open a ticket for it?

SYCL/Basic/event_profiling_info.cpp

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

alexbatashev · 2022-02-23T14:03:05Z

Woaw! This has become quite the PR. It looks fine to me.

@alexbatashev - Ping for awareness. This adds a solid handful of tests. Do you have any CI-related concerns?

That's a lot of testing, but hopefully it'll be fine

bader · 2022-02-24T14:28:53Z

@t4c1, it looks like some tests are not working on AMDGPU. See https://github.com/intel/llvm/runs/5319895067?check_suite_focus=true.
Could you fix them as soon as possible, please? We run these tests in pre-commit validation.

bader · 2022-02-24T18:17:33Z

I've filed intel/llvm#8804 and disabled the failing tests with 08100b0.

…ntel#534) Added tests for atomics with various memory orders and scopes. Reductions tests also have updated sm requirements, as they call work group atomics, which are now implemented and have higher sm requirements than device scoped ones. This adds tests for changes introduced in intel/llvm#4820 and intel/llvm#5192.

…ntel/llvm-test-suite#534) Added tests for atomics with various memory orders and scopes. Reductions tests also have updated sm requirements, as they call work group atomics, which are now implemented and have higher sm requirements than device scoped ones. This adds tests for changes introduced in intel#4820 and intel#5192.

[SYCL] Added tests for atomics with various memory orders and scopes

6260333

t4c1 requested review from AGindinson and v-klochkov as code owners October 26, 2021 13:58

t4c1 mentioned this pull request Oct 26, 2021

[SYCL][CUDA][libclc] Added atomics with scopes and memory orders intel/llvm#4820

Merged

steffenlarsen reviewed Nov 5, 2021

View reviewed changes

t4c1 added 10 commits November 18, 2021 15:01

[SYCL] merged add tests into one file

a900c8f

[SYCL] merged tests for other operations

d7f7e34

Merge branch 'intel' into atomic_orders_scopes2

0e272ec

# Conflicts: # SYCL/AtomicRef/add.cpp # SYCL/AtomicRef/add.h # SYCL/AtomicRef/add_atomic64.cpp # SYCL/AtomicRef/compare_exchange.h # SYCL/AtomicRef/exchange.h # SYCL/AtomicRef/max.h # SYCL/AtomicRef/min.h

[SYCL] format

0375249

[SYCL] add testing for both AtomicRef implementations

c215e68

Merge branch 'atomic_orders_scopes2' of https://github.com/t4c1/llvm-…

8a185e1

…test-suite into atomic_orders_scopes2

[SYCL] format

d18ca34

[SYCL] fixed add test

f3e6079

[SYCL] format

94b90b7

tests for remaining atomics

0ff5fe0

This was referenced Dec 21, 2021

[SYCL][CUDA] Added tests for atomic loads and stores for various orders and scopes #648

Merged

[SYCL][CUDA][libclc] Add remaining atomics intel/llvm#5192

Merged

format

5351b6d

enabled add test for pointers with orders and scopes

a8fb5f8

t4c1 requested a review from a team as a code owner January 4, 2022 11:56

t4c1 added 2 commits January 6, 2022 07:40

Merge branch 'intel' into atomic_orders_scopes2

db06775

Merge branch 'intel' into atomic_orders_scopes2

840c89d

This was referenced Jan 13, 2022

[CUDA] Enable tests in AtomicRef for CUDA after adding support for generic address space in DPC++ intel/llvm#8807

Closed

[SYCL][CUDA][libclc] Add atomic loads and stores with various memory orders and scopes intel/llvm#5191

Merged

t4c1 added 3 commits January 27, 2022 08:08

Merge branch 'intel' into atomic_orders_scopes2

138a98a

fixed missing newlines at the end of files

996581a

Merge branch 'intel' into atomic_orders_scopes2

d4af22d

t4c1 requested a review from steffenlarsen February 14, 2022 07:40

steffenlarsen reviewed Feb 14, 2022

View reviewed changes

t4c1 added 4 commits February 15, 2022 09:55

fixed generic tests

d1e7591

split tests

9708521

fix copy paste error in sub

55795d3

split native floating point tests

9a23a34

t4c1 added 2 commits February 21, 2022 15:43

fixed sub_generic tests and XFAILS for native fp tests

ad77010

format

8bae70f

bader requested a review from steffenlarsen February 21, 2022 17:59

steffenlarsen previously approved these changes Feb 22, 2022

View reviewed changes

disable event_profiling_info for CUDA

19c3856

t4c1 dismissed steffenlarsen’s stale review via 19c3856 February 22, 2022 12:41

t4c1 mentioned this pull request Mar 28, 2023

SYCL/Basic/event_profiling_info.cpp fails sporadicaly with CUDA backend intel/llvm#8805

Closed

bader reviewed Feb 22, 2022

View reviewed changes

SYCL/Basic/event_profiling_info.cpp Outdated Show resolved Hide resolved

Update SYCL/Basic/event_profiling_info.cpp

6a4fa77

Co-authored-by: Alexey Bader <alexey.bader@intel.com>

bader merged commit 88ee9d1 into intel:intel Feb 24, 2022

mlychkov mentioned this pull request Mar 3, 2022

[SYCL][CUDA] llvm-test-suite : AtomicRef/sub.cpp test fails for DPC++ with new Pass Manager turned on intel/llvm#5716

Closed

t4c1 deleted the atomic_orders_scopes2 branch March 15, 2022 08:52

Test	Current	New
add.cpp	everything works everywhere	XFAIL on host, opencl and hip
add_generic.cpp	XFAIL on hip and cuda	XFAIL on hip, cuda, opencl and level_zero
add_atomic64.cpp	everything works everywhere	moved to add.cpp
add_atomic64_generic.cpp	XFAIL on hip and cuda	moved to add_generic.cpp

[SYCL] Added tests for atomics with various memory orders and scopes #534

[SYCL] Added tests for atomics with various memory orders and scopes #534

Uh oh!

Conversation

t4c1 commented Oct 26, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

steffenlarsen left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bader commented Dec 25, 2021

Uh oh!

t4c1 commented Feb 9, 2022

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

t4c1 Feb 15, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

t4c1 commented Feb 21, 2022

Uh oh!

bader commented Feb 21, 2022

Uh oh!

bader commented Feb 21, 2022

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

bader commented Feb 22, 2022

Uh oh!

t4c1 commented Feb 22, 2022

Uh oh!

steffenlarsen commented Feb 22, 2022

Uh oh!

Uh oh!

alexbatashev commented Feb 23, 2022

Uh oh!

bader commented Feb 24, 2022

Uh oh!

bader commented Feb 24, 2022

Uh oh!

Uh oh!

t4c1 commented Oct 26, 2021 •

edited

Loading

steffenlarsen left a comment •

edited

Loading

t4c1 Feb 15, 2022 •

edited

Loading