cub::DeviceRadixSort uses Policy900 when compiled for an A100 GPU

[mfbalin]

Hi everyone,

In dmlc/dgl#6770, we utilize cub::DeviceRadixSort and we profiled the execution after compiling for sm_80 to run on A100 GPU. However, it can be seen that the Policy900 was used in the profile below. Is this normal? I would have expected Policy800 to be selected for the A100 GPU.

[gevtushenko]

Hello @mfbalin!

If you are referring to the symbol name having Policy900 in it, it's due to the fact that we pass maximal policy (DeviceRadixSortOnesweepKernel<MaxPolicyT...) to the kernel and then statically dispatch to active policy inside the kernel. The description of chained policy can be found in our docs.

[mfbalin]

I see, thank you for the quick reply, that is good to know.