Optimize compilation time for the common case

[gevtushenko]

This PR contains compilation time improvement for the common case when agents fit into default shared memory size (48 KB).

[gevtushenko]

gpuCI: NVIDIA/thrust#1557
DVS: 30594457

[alliepiper]

Thanks @senior-zero, this is a huge improvement -- I'm seeing compilation time improve on both nvcc and nvc++. It isn't quite back to before, but it's a significant reduction!

Compile times for the thrust::sort test program:

Compiler	Old merge sort	Current merge sort	This PR
nvcc	18.79s	29.10s (+55%)	22.44s (+19%)
nvc++	61.81s	75.61s (+22%)	65.80s (+6%)

LGTM -- Let's get this merged once the tests are passing and see how this impacts the total build time.

[alliepiper]

Related to NVBugs 3418930 and 3419768.

[dongxiao92]

Could you help to explain why these changes can reduce compilation time?
From the code changes, I understand that 1) instantiation for DeviceMergeSortBlockSortKernel with use_vshmem=true may not be needed and 2) compilation for code in this if branch may not be needed.
Does the compilation time reduction come from these two changes?

[gevtushenko]

Could you help to explain why these changes can reduce compilation time? From the code changes, I understand that 1) instantiation for DeviceMergeSortBlockSortKernel with use_vshmem=true may not be needed and 2) compilation for code in this if branch may not be needed. Does the compilation time reduction come from these two changes?

Hello, @dongxiao92!

Two specializations of merge sort kernels exist:

kernel	use vshmem=false	use vshmem=true
DeviceMergeSortBlockSortKernel	+	+
DeviceMergeSortMergeKernel	+	+

Since the check for available shared memory is performed at runtime, we had to compile for both cases (in generic case). This patch relies on the thrust approach which consists of comparing kernel shared memory size requirements with the default available shared memory size (48KB). This check can be done at compile time. If we know that virtual shared memory is not required at compile-time, there's no need to compile merge sort kernels twice.