diff options
| author | Johannes Doerfert <johannes@jdoerfert.de> | 2023-03-20 18:55:05 -0700 |
|---|---|---|
| committer | Johannes Doerfert <johannes@jdoerfert.de> | 2023-04-17 15:27:17 -0700 |
| commit | 67fed132f39c81e8006c4463ab1f173fea5e4e4b (patch) | |
| tree | 82e7a0b1fdbbc21b7da948e3e182b162bda7dc4e /lldb/source/Plugins/Process/gdb-remote/ProcessGDBRemote.cpp | |
| parent | c8bb7c234c6814b80e2be27eba9718de7ab1ad79 (diff) | |
| download | llvm-67fed132f39c81e8006c4463ab1f173fea5e4e4b.zip llvm-67fed132f39c81e8006c4463ab1f173fea5e4e4b.tar.gz llvm-67fed132f39c81e8006c4463ab1f173fea5e4e4b.tar.bz2 | |
[OpenMP] Ensure memory fences are created with barriers for AMDGPUs
It turns out that the __builtin_amdgcn_s_barrier() alone does not emit
a fence. We somehow got away with this and assumed it would work as it
(hopefully) is correct on the NVIDIA path where we just emit a
__syncthreads. After talking to @arsenm we now (mostly) align with the
OpenCL barrier implementation [1] and emit explicit fences for AMDGPUs.
It seems this was the underlying cause for #59759, but I am not 100%
certain. There is a chance this simply hides the problem.
Fixes: https://github.com/llvm/llvm-project/issues/59759
[1] https://github.com/RadeonOpenCompute/ROCm-Device-Libs/blob/07b347366eb2c6ebc3414af323c623cbbbafc854/opencl/src/workgroup/wgbarrier.cl#L21
Diffstat (limited to 'lldb/source/Plugins/Process/gdb-remote/ProcessGDBRemote.cpp')
0 files changed, 0 insertions, 0 deletions