For a 2-D grid kernel launch, does blockIdx.y change faster than blockIdx.x?

[zhaolianshuizls]

When AgentSelectIf is launched through a 2-D grid configuration,
https://github.com/NVIDIA/cub/blob/b2e8bccb8c0cd15279974fe4b9b8d6fcd1842b57/cub/device/dispatch/dispatch_select_if.cuh#L452,
larger-index block should wait for the previous blocks finish. My intuition tells that blockIdx.x changes faster, but the following code seems to say blockIdx.y changes faster
https://github.com/NVIDIA/cub/blob/b2e8bccb8c0cd15279974fe4b9b8d6fcd1842b57/cub/agent/agent_select_if.cuh#L686

I know in the above case, only tile_idx value matters, however, I'm eager to know in what order the thread blocks are launched. Along the blockIdx.x or blockIdx.y direction?

I know thread blocks can execute independently, but are they scheduled to run randomly?

Does select_if kernel assumes all thread blocks can run concurrently on the device? So the number of grid blocks should not be larger than a certain value, i.e. the device capability.

[gevtushenko]

Hello, @zhaolianshuizls! There isn't a guarantee on the order. The thread blocks are scheduled independently. I think due to the current num_items restriction, we haven't actually tested the scenario where blockIdx.y is larger than 1. We should take a look once we get to supporting 64-bit offsets in select / partition. If this is critical to you, please, feel free to open an issue against https://github.com/NVIDIA/cccl.