Potential deadlock in DeviceSelect?

[fkallen]

The decoupled look-back step to calculate the prefix sum requires results from all previous tiles.
In agent_select_if.cuh, thread blocks are mapped to tiles using the build-in blockId.
int tile_idx = (blockIdx.x * gridDim.y) + blockIdx.y;

Can this result in a dead-lock because the scheduling order of thread blocks is unspecified?

[alliepiper]

@dumerrill @canonizer I think you two are most familiar with the decoupled-lookback algorithm used by CUB. Can either of you comment on this?

[alliepiper]

The tl;dr is "here be dragons".

This is not strictly safe within the CUDA programming model, but we take advantage of the fact that our current HW behaves this way. We may change this in the future if needed.