AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

device: optimize threadfence for ll64 protocol (#1858)

Browse Source 
* device: optimize threadfence for ll64 protocol

* device: use __atomic_signal_fence()

---------

Co-authored-by: Nusrat Islam <nusislam@useocpslog-003.amd.com>
This commit is contained in:
Nusrat Islam
2025-08-18 09:16:41 -05:00
committed by GitHub
parent 876f985e0f
commit 6ade5065b4
1 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/device/prims_ll128.h
+4 -2
View File
@@ -138,11 +138,13 @@ private:
if (recvConnHeadPtr) STORE(recvConnHeadPtr, recvConnHead += 1); if (recvConnHeadPtr) STORE(recvConnHeadPtr, recvConnHead += 1);
} }
inline __device__ void postSend() { inline __device__ void postSend() {
__atomic_signal_fence(__ATOMIC_SEQ_CST);
asm volatile("s_waitcnt lgkmcnt(0) vmcnt(0)");
__atomic_signal_fence(__ATOMIC_SEQ_CST);
if (sendConnTailPtr) { if (sendConnTailPtr) {
#if __CUDA_ARCH__ >= 900 #if __CUDA_ARCH__ >= 900
__threadfence_system(); __threadfence_system();
#else
__threadfence();
#endif #endif
STORE((unsigned long long *)sendConnTailPtr, sendConnTail += 1); STORE((unsigned long long *)sendConnTailPtr, sendConnTail += 1);
} }