CUB 1.15.0 compilation errors if included after Windows.h

[ptheywood]

If Windows.h is included prior to cub/cub.cuh under MSVC compilation will fail with the following:

D:/a/jitify-cub-1.15-MSVC-demo/jitify-cub-1.15-MSVC-demo/build/_deps/thrust-src/dependencies/cub\cub/device/dispatch/dispatch_segmented_sort.cuh(338): error : invalid combination of type specifiers [D:\a\jitify-cub-1.15-MSVC-demo\jitify-cub-1.15-MSVC-demo\build\jitify-before-cub.vcxproj]
D:/a/jitify-cub-1.15-MSVC-demo/jitify-cub-1.15-MSVC-demo/build/_deps/thrust-src/dependencies/cub\cub/device/dispatch/dispatch_segmented_sort.cuh(338): error : expected an identifier [D:\a\jitify-cub-1.15-MSVC-demo\jitify-cub-1.15-MSVC-demo\build\jitify-before-cub.vcxproj]
D:/a/jitify-cub-1.15-MSVC-demo/jitify-cub-1.15-MSVC-demo/build/_deps/thrust-src/dependencies/cub\cub/device/dispatch/dispatch_segmented_sort.cuh(379): error : expected a member name [D:\a\jitify-cub-1.15-MSVC-demo\jitify-cub-1.15-MSVC-demo\build\jitify-before-cub.vcxproj]

This impacts the use of NVIDIA/Jitify and NVRTC.
nvrtc.h conditionally includes Windows.h if NVRTC_GET_TYPE_NAME is defined to a non-zero value (as used in Jitify).

This only occurs since CUB 1.15.0 with multiple CUDA versions, and only if Windows.h is included prior to CUB.

It is not due to the min/max macros (checked via -DNOMINMAX).

I've created a repository to reproduce this, and demonstrate this through GitHub Actions:

https://github.com/ptheywood/cub-1.15-windows.h-demo

The minimal example to reproduce this is:

#if defined(_WIN32)
    #include <Windows.h>
#endif
#include <cub/cub.cuh>
int main(int argc, char* argv[]) {
    return EXIT_SUCCESS;
}

GitHub Action run with logs valid for 90 days (?): https://github.com/ptheywood/cub-1.15-windows.h-demo/runs/4851931940?check_suite_focus=true

The CUDA 11.6 + CUB 1.15.0 log has been archived at https://gist.github.com/ptheywood/87b821860931bf6a2c700bbe3803d3f4 for longer term storage.

[gevtushenko]

It would seem that windows.h defines a macro called small, which happens to be a name of a variable in the new segmented sort:

  __shared__ union
  {
    // ...
    typename SmallAgentWarpMergeSortT::TempStorage small[segments_per_small_block];
  } temp_storage;

Fix consists of renaming this member. I'll create a PR later today.