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

Fix performance downgrade issue & update doc (#229)

Browse Source 
For push function, we only need to make sure the instruction `st.global`
will be executed after the while loop. Since there is a Write-After-Read
hazard for `trigger.fst` (Check `this->triggers[curFifoHead % size].fst
!= 0` first then write value to `triggers[curFifoHead % size]`), we can
expect the compiler and hardware can handle this situation correctly.
Remove the `release.sys` there.

BTW, `st.global.release.sys.v2.u64` will cause perf regression issue.
Previous we use `st.global.release.cta.v2.u64`, but seems not necessary.
This commit is contained in:
Binyang Li
2023-12-05 02:20:10 +08:00
committed by GitHub
parent 351b95b926
commit f1b2c9df12
2 changed files with 6 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
docs/quickstart.md
View File

@@ -106,7 +106,7 @@ $ make -j allgather_test_perf allreduce_test_perf
For example, the following command runs the `allreduce5` algorithm with 8 GPUs starting from 3MB to 48MB messages, by doubling the message size in between. You can try different algorithms by changing the `-k 5` option to another value (e.g., `-k 3` runs `allreduce3`). Check all algorithms from the code: [allreduce_test.cu](https://github.com/microsoft/mscclpp/blob/main/test/mscclpp-test/allreduce_test.cu) and [allgather_test.cu](https://github.com/microsoft/mscclpp/blob/main/test/mscclpp-test/allgather_test.cu).
```bash
$ mpirun --bind-to-numa -np 8 ./test/mscclpp-test/allreduce_test_perf -b 3m -e 48m -G 100 -n 100 -w 20 -f 2 -k 5
$ mpirun --bind-to numa -np 8 ./test/mscclpp-test/allreduce_test_perf -b 3m -e 48m -G 100 -n 100 -w 20 -f 2 -k 5
```
*NOTE: a few algorithms set a condition on the total data size, such as to be a multiple of 3. If the condition is unmet, the command will throw a regarding error.*

8
include/mscclpp/fifo_device.hpp
View File

@@ -25,7 +25,8 @@ struct alignas(16) ProxyTrigger {
uint64_t fst, snd;
};
/// A concurrent FIFO where multiple device threads can push work elements and a single host proxy thread consumes them.
/// A concurrent FIFO where multiple device threads (the number of threads should not exceed the fifo size) can push
/// work elements and a single host proxy thread consumes them.
///
/// The FIFO has a head pointer allocated on the device which starts at 0 and goes up to 2^64-1, which is almost
/// infinity. There are two copies of the tail, one on the device, @ref FifoDeviceHandle::tailReplica, and another on
@@ -64,9 +65,10 @@ struct FifoDeviceHandle {
ProxyTrigger* triggerPtr = &(this->triggers[curFifoHead % size]);
// store with memory order release so that the while loop does not go pass this.
// There is a Write-After-Read hazard for the triggerPtr->fst. So the st instruction will not be executed
// before the loop.
#if defined(MSCCLPP_DEVICE_CUDA)
asm volatile("st.global.release.sys.v2.u64 [%0], {%1,%2};" ::"l"(triggerPtr), "l"(trigger.fst), "l"(trigger.snd));
asm volatile("st.global.relaxed.sys.v2.u64 [%0], {%1,%2};" ::"l"(triggerPtr), "l"(trigger.fst), "l"(trigger.snd));
#else // !defined(MSCCLPP_DEVICE_CUDA)
// TODO: both atomic and clang built-ins are buggy here
triggerPtr->fst = trigger.fst;