[SYCL][HIP] Add basic HIP atomics #8003

hdelan · 2023-01-13T12:00:42Z

Adding support for basic atomic operations for HIP AMD backend.

hdelan · 2023-01-13T12:11:01Z

/verify with intel/llvm-test-suite#1510

clang/lib/Frontend/InitPreprocessor.cpp

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_add.cl

sycl/include/sycl/atomic_ref.hpp

Co-authored-by: Steffen Larsen <[email protected]>

steffenlarsen

Good stuff! 🚀

premanandrao

A couple of comments needed, but otherwise FE changes LGTM.

clang/test/Preprocessor/opencl-macro-target-specific.cl

clang/test/Preprocessor/sycl-macro-target-specific.cpp

Fznamznon

Clang changes look ok to me.

hdelan · 2023-01-23T10:25:12Z

/verify with intel/llvm-test-suite#1510

zjin-lcf · 2023-01-24T20:28:03Z

Faster than hipcc ...

SYCL
./main 10000
Atomic min for data type U64 | Average execution time: 0.000025 (s)
Atomic min for data type S64 | Average execution time: 0.000021 (s)
Atomic min for data type F64 | Average execution time: 0.000022 (s)
Atomic max for data type U64 | Average execution time: 0.000052 (s)
Atomic max for data type S64 | Average execution time: 0.000054 (s)
Atomic max for data type F64 | Average execution time: 0.000052 (s)
Atomic add for data type U64 | Average execution time: 3.235254 (s)
Atomic add for data type S64 | Average execution time: 3.215109 (s)
Atomic add for data type F64 | Average execution time: 3.245661 (s)
atomic max results: 262144 -9223372036854775808 262144.000000
FAIL

=================================================
HIP
./main 10000
Atomic min for data type U64 | Average execution time: 0.000053 (s)
Atomic min for data type S64 | Average execution time: 0.000029 (s)
Atomic min for data type F64 | Average execution time: 0.000029 (s)
Atomic max for data type U64 | Average execution time: 0.000280 (s)
Atomic max for data type S64 | Average execution time: 0.000292 (s)
Atomic max for data type F64 | Average execution time: 0.000280 (s)
Atomic add for data type U64 | Average execution time: 4.652767 (s)
Atomic add for data type S64 | Average execution time: 4.517521 (s)
Atomic add for data type F64 | Average execution time: 4.601057 (s)
PASS

hdelan · 2023-01-25T10:26:44Z

@zjin-lcf wow nice. Note that AtomicFMin and AtomicFMax are not included in this PR, which is why the tests are failing.

hdelan · 2023-01-25T10:27:48Z

@bader the LLVM test suite failures are all XPASSes, so I think this can be merged

zjin-lcf · 2023-01-25T13:05:01Z

@zjin-lcf wow nice. Note that AtomicFMin and AtomicFMax are not included in this PR, which is why the tests are failing.

Is there something missing in ROCm for not including FMin/FMax ?

hdelan · 2023-01-25T14:02:31Z

No it is just lower priority than the FAdd etc. I will be adding FMin, FMax as well as some other atomics as soon as I have finished some other stuff I am working on

bader · 2023-01-25T18:46:07Z

@bader the LLVM test suite failures are all XPASSes, so I think this can be merged

I see these failures:

Failed Tests (4):
SYCL :: Reduction/reduction_big_data.cpp
SYCL :: Reduction/reduction_nd_ext_double.cpp
SYCL :: Reduction/reduction_usm.cpp
SYCL :: Reduction/reduction_usm_dw.cpp

lld: error: undefined hidden symbol: __spirv_AtomicFMaxEXT(float AS1*, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag, float)

referenced by lto.tmp:(typeinfo name for sycl::_V1::detail::reduction::MainKrn<main::A1, (sycl::_V1::detail::reduction::strategy)4>)
referenced by lto.tmp:(typeinfo name for sycl::_V1::detail::reduction::MainKrn<main::A1, (sycl::_V1::detail::reduction::strategy)4>)
referenced by lto.tmp:(_ZTSN4sycl3_V16detail9reduction7MainKrnIZ4mainE2A1LNS2_8strategyE4EJEEE_with_offset)
referenced 1 more times
llvm-foreach:
clang-16: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)

error: command failed with exit status: 1

These seem to be related to the change.

NOTE: Unfortunately, "/verify with" doesn't cover AMDGPU, so there is no safe way to test both changes (compiler + test) in CI.

hdelan · 2023-01-26T13:09:56Z

@bader that should be fixed now. @zjin-lcf remaining atomics are all added, so FMax, FMin etc should be working

zjin-lcf · 2023-01-26T16:23:54Z

I tried to run atomic tests. Can you please run them again ? Thanks.

https://github.com/zjin-lcf/HeCBench/tree/master/atomicCAS-sycl

hdelan · 2023-01-26T16:37:18Z

I tried to run atomic tests. Can you please run them again ? Thanks.

https://github.com/zjin-lcf/HeCBench/tree/master/atomicCAS-sycl
@zjin-lcf

$ sycl-ls
[ext_oneapi_hip:gpu:0] AMD HIP BACKEND, AMD Radeon VII 0.0 [HIP 40421.43]
$ ./main 10000
Atomic min for data type U64 | Average execution time: 0.000072 (s)
Atomic min for data type S64 | Average execution time: 0.000087 (s)
Atomic min for data type F64 | Average execution time: 0.000061 (s)
Atomic max for data type U64 | Average execution time: 0.000218 (s)
Atomic max for data type S64 | Average execution time: 0.000226 (s)
Atomic max for data type F64 | Average execution time: 0.000250 (s)
Atomic add for data type U64 | Average execution time: 2.484193 (s)
Atomic add for data type S64 | Average execution time: 2.413692 (s)
Atomic add for data type F64 | Average execution time: 2.370539 (s)
PASS

zjin-lcf · 2023-01-26T17:19:03Z

@hdelan
I pushed some changes to the sycl benchmark to fix some issues, and observed the same result (PASS) on a MI200 GPU. Congratulations!

I tried to run the benchmark
https://github.com/zjin-lcf/HeCBench/tree/master/sheath-hip
https://github.com/zjin-lcf/HeCBench/tree/master/sheath-sycl

The option -munsafe-fp-atomics is passed to the hip compiler and the option -DSYCL_USE_NATIVE_FP_ATOMICS to the sycl compiler. However, I think that the sycl compiler does not generate fast fp atomics due to the significant difference in execution time on a MI200 GPU. Do I pass the right compiler option to the sycl compiler ? Thanks again.

bader · 2023-01-26T17:22:08Z

so I think this can be merged

One more thing. We need to merge these changes together with test changes, but the test PR is not reviewed yet.

hdelan · 2023-01-26T22:03:18Z

@hdelan I pushed some changes to the sycl benchmark to fix some issues, and observed the same result (PASS) on a MI200 GPU. Congratulations!

I tried to run the benchmark https://github.com/zjin-lcf/HeCBench/tree/master/sheath-hip https://github.com/zjin-lcf/HeCBench/tree/master/sheath-sycl

The option -munsafe-fp-atomics is passed to the hip compiler and the option -DSYCL_USE_NATIVE_FP_ATOMICS to the sycl compiler. However, I think that the sycl compiler does not generate fast fp atomics due to the significant difference in execution time on a MI200 GPU. Do I pass the right compiler option to the sycl compiler ? Thanks again.

@zjin-lcf So the fast gfx90a+ atomics are not part of this PR. They will be enabled in a future PR, will keep you posted

Also SYCL_USE_NATIVE_FP_ATOMICS is now enabled by default for the AMD HIP backend, so you no longer need to pass this to the compiler

zjin-lcf · 2023-01-27T03:22:44Z

I see. I am also puzzled by the issue (#7252).
The atomic add works with "shared" USM, but other atomics (e.g. atomic sub, atomic xor) don't.

hdelan · 2023-01-30T16:56:27Z

@bader Can we get the test suite changes reviewed?

hdelan · 2023-01-31T16:12:38Z

@bader the test suite changes have been reviewed so I think this is safe to merge

Hugh Delaney added 3 commits January 11, 2023 11:20

First commit adding some HIP atomics

75abcef

Adding initial support for HIP atomics

8953f41

Removing cmpxchg

41af655

hdelan requested review from a team as code owners January 13, 2023 12:00

hdelan requested a review from sergey-semenov January 13, 2023 12:00

hdelan mentioned this pull request Jan 13, 2023

[SYCL][HIP] Adding basic atomics for HIP intel/llvm-test-suite#1510

Merged

This was referenced Jan 13, 2023

SYCL atomics are slow on MI250 #7179

Closed

[SYCL][ROCM] lld: error: undefined hidden symbol: __spirv_AtomicFAddEXT(float AS1*, __spv::Scope::Flag, __spv::MemorySemanticsMask::Flag, float) > #4471

Closed

hdelan temporarily deployed to aws January 13, 2023 12:49 — with GitHub Actions Inactive

hdelan temporarily deployed to aws January 13, 2023 13:34 — with GitHub Actions Inactive

premanandrao reviewed Jan 13, 2023

View reviewed changes

clang/lib/Frontend/InitPreprocessor.cpp Show resolved Hide resolved

steffenlarsen reviewed Jan 16, 2023

View reviewed changes

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_add.cl Show resolved Hide resolved

sycl/include/sycl/atomic_ref.hpp Outdated Show resolved Hide resolved

Adding compare exchange as well as generic AS

263c7a7

hdelan temporarily deployed to aws January 18, 2023 10:16 — with GitHub Actions Inactive

Update sycl/include/sycl/atomic_ref.hpp

7289005

Co-authored-by: Steffen Larsen <[email protected]>

steffenlarsen approved these changes Jan 18, 2023

View reviewed changes

hdelan temporarily deployed to aws January 18, 2023 10:51 — with GitHub Actions Inactive

Add AtomicAdd for FP64

ad575d4

hdelan temporarily deployed to aws January 18, 2023 12:50 — with GitHub Actions Inactive

Updating preprocessor tests

89fb7e6

hdelan temporarily deployed to aws January 18, 2023 15:19 — with GitHub Actions Inactive

premanandrao reviewed Jan 18, 2023

View reviewed changes

clang/test/Preprocessor/opencl-macro-target-specific.cl Show resolved Hide resolved

clang/test/Preprocessor/sycl-macro-target-specific.cpp Show resolved Hide resolved

Add comments to test and remove empty line in CMakeLists

4dfac50

hdelan temporarily deployed to aws January 19, 2023 10:24 — with GitHub Actions Inactive

premanandrao approved these changes Jan 19, 2023

View reviewed changes

Making test cl-std=2.0

85e96b9

hdelan temporarily deployed to aws January 19, 2023 12:25 — with GitHub Actions Inactive

hdelan temporarily deployed to aws January 19, 2023 12:56 — with GitHub Actions Inactive

smanna12 approved these changes Jan 19, 2023

View reviewed changes

Fznamznon approved these changes Jan 19, 2023

View reviewed changes

Add remaining atomic ops

33ba670

hdelan temporarily deployed to aws January 26, 2023 13:31 — with GitHub Actions Inactive

hdelan temporarily deployed to aws January 26, 2023 14:02 — with GitHub Actions Inactive

hdelan mentioned this pull request Jan 27, 2023

[SYCL][HIP] sycl::atomic_ref::fetch_sub does not work with USM #7252

Closed

bader merged commit c3c5e92 into intel:sycl Jan 31, 2023

hdelan mentioned this pull request Feb 1, 2023

[SYCL][HIP] HIP backend timing out on max_local in CI #7604

Closed

JackAKirk mentioned this pull request Sep 25, 2024

AtomicRef tests enabled by #534 fail on HIP back-end. #8804

Closed

[SYCL][HIP] Add basic HIP atomics #8003

[SYCL][HIP] Add basic HIP atomics #8003

Uh oh!

Conversation

hdelan commented Jan 13, 2023

Uh oh!

hdelan commented Jan 13, 2023

Uh oh!

Uh oh!

Uh oh!

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

premanandrao left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Fznamznon left a comment

Choose a reason for hiding this comment

Uh oh!

hdelan commented Jan 23, 2023

Uh oh!

zjin-lcf commented Jan 24, 2023

Uh oh!

hdelan commented Jan 25, 2023

Uh oh!

hdelan commented Jan 25, 2023

Uh oh!

zjin-lcf commented Jan 25, 2023

Uh oh!

hdelan commented Jan 25, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bader commented Jan 25, 2023

Uh oh!

hdelan commented Jan 26, 2023

Uh oh!

zjin-lcf commented Jan 26, 2023

Uh oh!

hdelan commented Jan 26, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

zjin-lcf commented Jan 26, 2023

Uh oh!

bader commented Jan 26, 2023

Uh oh!

hdelan commented Jan 26, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

zjin-lcf commented Jan 27, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

hdelan commented Jan 30, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

hdelan commented Jan 31, 2023

Uh oh!

Uh oh!

hdelan commented Jan 25, 2023 •

edited

Loading

hdelan commented Jan 26, 2023 •

edited

Loading

hdelan commented Jan 26, 2023 •

edited

Loading

zjin-lcf commented Jan 27, 2023 •

edited

Loading

hdelan commented Jan 30, 2023 •

edited

Loading