microsoft/mscclpp
1
0
Fork 0
mirror of https://github.com/microsoft/mscclpp.git synced 2026-05-13 09:46:00 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #30 from microsoft/saemal/consistency_bug_fix

Browse Source 
an important deadlock bug fix
This commit is contained in:
Saeed Maleki
2023-03-26 14:02:42 -07:00
committed by GitHub
parent 57b885c9ab 35ca25781a
commit f6a7962511
1 changed files with 1 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
src/include/mscclppfifo.h
View File

@@ -40,6 +40,7 @@ struct mscclppConcurrentFifo {
__forceinline__ __device__ uint64_t push(uint64_t type, uint64_t dstDataOffset, uint64_t srcDataOffset, uint64_t dataSize){
uint64_t curFifoHead = atomicAdd((unsigned long long int*)this->triggerFifoHead,1);
while (curFifoHead >= MSCCLPP_PROXY_FIFO_SIZE + *((volatile uint64_t*)this->triggerFifoTail));
while (*(volatile uint64_t*)&this->triggerFifo[curFifoHead % MSCCLPP_PROXY_FIFO_SIZE] != 0);
uint64_t* valptr = (uint64_t*)&(this->triggerFifo[curFifoHead % MSCCLPP_PROXY_FIFO_SIZE].value);
asm volatile(
"st.volatile.global.v2.u64 [%0], {%1,%2};" ::"l"(valptr),