Skip to content

Commit dd76dec

Browse files
keryellhwu36
andauthored
[Doc] Make C++ code more plausible (#2156)
Co-authored-by: Haicheng Wu <[email protected]>
1 parent 19cc2a5 commit dd76dec

File tree

1 file changed

+7
-4
lines changed

1 file changed

+7
-4
lines changed

media/docs/cpp/blackwell_cluster_launch_control.md

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -43,17 +43,18 @@ __device__ non_persistent_kernel(...) {
4343
setup_common_data_structures();
4444
dim3 workCoordinates = blockIdx;
4545
coordinate_specific_compute(workCoordinates);
46-
}
46+
}
4747
```
4848
#### Static Persistent Kernel
4949
``` c++
5050
// Static Persistent Kernel
5151
__device__ static_persistent_kernel(...) {
5252
setup_common_data_structures(...);
5353
dim3 workCoordinates = blockIdx;
54+
bool isValidId;
5455
do {
5556
coordinate_specific_compute(workCoordinates);
56-
isValidId, workCoordinates = staticTileScheduler.fetch_next_work();
57+
std::tie(isValidId, workCoordinates) = staticTileScheduler.fetch_next_work();
5758
} while (isValidId);
5859
}
5960
```
@@ -65,9 +66,11 @@ __device__ static_persistent_kernel(...) {
6566
__device__ clc_dynamic_persistent_kernel(...) {
6667
setup_common_data_structures(...);
6768
dim3 workCoordinates = blockIdx;
69+
dim3 newClcID;
70+
bool isValidId;
6871
do {
6972
coordinate_specific_compute(workCoordinates);
70-
isValidId, newClcID = clcTileScheduler.fetch_next_work();
73+
std::tie(isValidId, newClcID) = clcTileScheduler.fetch_next_work();
7174
workCoordinates = newClcID;
7275
} while (isValidId);
7376
}
@@ -76,7 +79,7 @@ __device__ clc_dynamic_persistent_kernel(...) {
7679
7780
### Cluster Launch Control Pipeline Class
7881
79-
Please refer to the `PipelineCLCFetchAsync` pipeline class defined in [Cluster launch control pipeline class](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/pipeline/sm100_pipeline.hpp). Cluster launch control queries can be pipelined and mananged by an asynchronous pipeline with producer-consumer relationship (See
82+
Please refer to the `PipelineCLCFetchAsync` pipeline class defined in [Cluster launch control pipeline class](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/pipeline/sm100_pipeline.hpp). Cluster launch control queries can be pipelined and managed by an asynchronous pipeline with producer-consumer relationship (See
8083
[pipeline](pipeline.md) document). The producer is the scheduler warp of the 0th CTA in the cluster and the consumers are all warps that need `ClcID`s.
8184
8285
To setup a CLC pipeline correctly, we need to make sure the params are set to the right values:

0 commit comments

Comments
 (0)