Replace volatile by proper memory ordering in HIP

[upsj]

rocm-clang supports the GCC __atomic intrinsics, which we can use to implement the atomic operations instead of using volatile and memory fences.

This is a follow-up to #1344, so I requested reviews from the same reviewers

[yhmtsai]

do you have the documentation about rocm-clang support gcc atomic on GPU?
Is there only one compiler rocm-clang from hipcc on AMD GPU?

[upsj]

This is based on communications with an AMD engineer, and the fact that clang has to be able to compile libstdc++, which relies on these intrinsics. We are not using any other compiler, and I'm not aware of any that we should be looking at. But thanks for the hint, I looked at what rocSPARSE is doing, and they are using __hip_atomic_load, which gives us the same kind of control over scope and ordering, but isn't documented, so I'm not sure we can rely on it too much

[upsj]

Relevant: https://github.com/ROCm-Developer-Tools/llvm-project/blob/6b9c186b2d4c8bd315034a9655a28d32bcf745ab/clang/test/SemaCUDA/atomic-ops.cu#L6

[upsj]

Looks like the intrinsics are not supported by older HIP versions yet. I'll fall back on the GCC versions then

[yhmtsai]

For the code self, LGTM.
I am still not sure whether gcc intrinsics work on GPU address. compiling to libstdc++ only mean for CPU side not GPU, right?

[yhmtsai]

maybe only for thrust::complex<double>?
thrust::complex<float> can be done with int64

[upsj]

I don't follow, we only use thrust::complex with float value types Formatting issues with templates. We don't need full atomicity with thrust::complex, only element-wise, so this allows a more efficient code generation. Maybe the compiler even combines them? Doesn't matter much

[yhmtsai]

is it used?

[upsj]

memcpy is defined in string.h, not sure if this is actually necessary though

[MarcelKoch]

I think this needs quite a bit of documentation, since this relies on functionality that is not in the official documentation. It needs at least links to where you found the__hip_atomic_load|store, and the different scopes. If the references themselves don't have much documentation, then also properly documenting them would be necessary.

[upsj]

@yhmtsai The tests run fine, and this was a suggestion by an AMD engineer, so I'm confident we can use them. If the intrinsics weren't supported, it would fail to compile instead.

[sonarcloud]

Kudos, SonarCloud Quality Gate passed!   

0 Bugs
0 Vulnerabilities
0 Security Hotspots
0 Code Smells

No Coverage information
0.0% Duplication

The version of Java (11.0.3) you have used to run this analysis is deprecated and we will stop accepting it soon. Please update to at least Java 17.
Read more here