Remove special handling for dialect in Thrust's build system.

.gitignore

@@ -16,3 +16,4 @@ CMakeUserPresets.json
 *.pyc
 __pycache__
 *.pyd
+wheelhouse/

CMakeLists.txt

@@ -9,7 +9,6 @@ endif()
 
 # Determine whether CCCL is the top-level project or included into
 # another project via add_subdirectory()
-set(CCCL_TOPLEVEL_PROJECT OFF)
 if ("${CMAKE_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_LIST_DIR}")
   set(CCCL_TOPLEVEL_PROJECT ON)
 endif()

CMakePresets.json

@@ -71,12 +71,8 @@
         "CUB_ENABLE_HEADER_TESTING": true,
         "CUB_ENABLE_TESTING": true,
         "CUB_ENABLE_EXAMPLES": true,
-        "CUB_ENABLE_DIALECT_CPP17": true,
-        "CUB_ENABLE_DIALECT_CPP20": true,
         "THRUST_ENABLE_MULTICONFIG": true,
         "THRUST_MULTICONFIG_WORKLOAD": "LARGE",
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP17": true,
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP20": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_CPP": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_CUDA": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_OMP": true,
@@ -200,34 +196,34 @@
       }
     },
     {
-      "name": "thrust-base",
-      "hidden": true,
+      "name": "thrust",
+      "displayName": "Thrust",
       "inherits": "base",
       "cacheVariables": {
         "CCCL_ENABLE_THRUST": true,
         "THRUST_ENABLE_MULTICONFIG": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_CPP": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_CUDA": true,
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_OMP": true,
-        "THRUST_MULTICONFIG_ENABLE_SYSTEM_TBB": true,
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP17": false,
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP20": false
+        "THRUST_MULTICONFIG_ENABLE_SYSTEM_TBB": true
       }
     },
     {
       "name": "thrust-cpp17",
       "displayName": "Thrust: C++17",
-      "inherits": "thrust-base",
+      "inherits": "thrust",
       "cacheVariables": {
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP17": true
+        "CMAKE_CXX_STANDARD": "17",
+        "CMAKE_CUDA_STANDARD": "17"
       }
     },
     {
       "name": "thrust-cpp20",
       "displayName": "Thrust: C++20",
-      "inherits": "thrust-base",
+      "inherits": "thrust",
       "cacheVariables": {
-        "THRUST_MULTICONFIG_ENABLE_DIALECT_CPP20": true
+        "CMAKE_CXX_STANDARD": "20",
+        "CMAKE_CUDA_STANDARD": "20"
       }
     },
     {
@@ -420,6 +416,10 @@
       "name": "cub-cpp20",
       "configurePreset": "cub-cpp20"
     },
+    {
+      "name": "thrust",
+      "configurePreset": "thrust"
+    },
     {
       "name": "thrust-cpp17",
       "configurePreset": "thrust-cpp17"
@@ -640,64 +640,59 @@
       "inherits": "cub-lid2"
     },
     {
-      "name": "thrust-base",
-      "hidden": true,
-      "inherits": "base",
-      "filter": {
-        "exclude": {
-          "name": "^thrust.*\\.test\\.async[._].*$"
-        }
-      }
+      "name": "thrust",
+      "configurePreset": "thrust",
+      "inherits": "base"
     },
     {
-      "name": "thrust-gpu-base",
-      "hidden": true,
-      "inherits": "thrust-base",
-      "filter": {
-        "include": {
-          "name": "^thrust.*\\.cuda\\..*$"
-        }
-      }
+      "name": "thrust-cpp17",
+      "configurePreset": "thrust-cpp17",
+      "inherits": "thrust"
     },
     {
-      "name": "thrust-cpu-base",
-      "hidden": true,
-      "inherits": "thrust-base",
+      "name": "thrust-cpp20",
+      "configurePreset": "thrust-cpp20",
+      "inherits": "thrust"
+    },
+    {
+      "name": "thrust-gpu",
+      "configurePreset": "thrust",
+      "inherits": "thrust",
       "filter": {
-        "exclude": {
+        "include": {
           "name": "^thrust.*\\.cuda\\..*$"
         }
       }
     },
     {
       "name": "thrust-gpu-cpp17",
       "configurePreset": "thrust-cpp17",
-      "inherits": "thrust-gpu-base"
+      "inherits": "thrust-gpu"
     },
     {
       "name": "thrust-gpu-cpp20",
       "configurePreset": "thrust-cpp20",
-      "inherits": "thrust-gpu-base"
-    },
-    {
-      "name": "thrust-cpu-cpp17",
-      "configurePreset": "thrust-cpp17",
-      "inherits": "thrust-cpu-base"
+      "inherits": "thrust-gpu"
     },
     {
-      "name": "thrust-cpu-cpp20",
-      "configurePreset": "thrust-cpp20",
-      "inherits": "thrust-cpu-base"
+      "name": "thrust-cpu",
+      "configurePreset": "thrust",
+      "inherits": "thrust",
+      "filter": {
+        "exclude": {
+          "name": "^thrust.*\\.cuda\\..*$"
+        }
+      }
     },
     {
-      "name": "thrust-cpp17",
+      "name": "thrust-cpu-cpp17",
       "configurePreset": "thrust-cpp17",
-      "inherits": "thrust-base"
+      "inherits": "thrust-cpu"
     },
     {
-      "name": "thrust-cpp20",
+      "name": "thrust-cpu-cpp20",
       "configurePreset": "thrust-cpp20",
-      "inherits": "thrust-base"
+      "inherits": "thrust-cpu"
     },
     {
       "name": "cudax",

c/experimental/stf/test/CMakeLists.txt

@@ -17,11 +17,11 @@ function(cccl_c_experimental_stf_add_test target_name_var source)
   target_link_libraries(
     ${target_name}
     PRIVATE
+      cccl.compiler_interface
       cccl.c.experimental.stf
       CUDA::cudart_static
       CUDA::nvrtc
       cccl.c2h.main
-      cccl.compiler_interface_cpp20
       CUDA::cuda_driver
       CCCL::cudax
   )

c/parallel/CMakeLists.txt

@@ -51,11 +51,11 @@ set_target_properties(cccl.c.parallel PROPERTIES CUDA_RUNTIME_LIBRARY STATIC)
 target_link_libraries(
   cccl.c.parallel
   PRIVATE
+    cccl.compiler_interface
     CUDA::cudart_static
     CUDA::nvrtc
     CUDA::nvJitLink
     CUDA::cuda_driver
-    cccl.compiler_interface_cpp20
     cccl.c.parallel.jit_template
     CUB::CUB
     Thrust::Thrust

c/parallel/src/merge_sort.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/device_merge_sort.cuh>
 
 #include <format>
@@ -324,7 +324,7 @@ struct device_merge_sort_vsmem_helper {{
   }};
 }};
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_merge_sort_policy")>()
     = cub::detail::merge_sort::MergeSortPolicyWrapper<device_merge_sort_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/radix_sort.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/device_radix_sort.cuh>
 
 #include <format>
@@ -323,7 +323,7 @@ struct __align__({3}) values_storage_t {{
 {4}
 using {5} = {6}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_radix_sort_policy")>()
     = cub::detail::radix::RadixSortPolicyWrapper<{5}::ActivePolicy>::EncodedPolicy();

c/parallel/src/reduce.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/device_reduce.cuh>
 #include <cub/grid/grid_even_share.cuh>
 #include <cub/util_device.cuh>
@@ -208,7 +208,7 @@ struct __align__({2}) storage_t {{
 {5}
 using device_reduce_policy = {6}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_reduce_policy")>()
     = cub::detail::reduce::ReducePolicyWrapper<device_reduce_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/scan.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/dispatch/dispatch_scan.cuh>
 #include <cub/thread/thread_load.cuh>
 #include <cub/util_arch.cuh>
@@ -270,7 +270,7 @@ struct __align__({1}) storage_t {{
 {4}
 using device_scan_policy = {5}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_scan_policy")>()
     = cub::detail::scan::ScanPolicyWrapper<device_scan_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/segmented_reduce.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh> // cub::detail::choose_offset_t
 #include <cub/detail/launcher/cuda_driver.cuh> // cub::detail::CudaDriverLauncherFactory
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/dispatch/dispatch_fixed_size_segmented_reduce.cuh>
 #include <cub/device/dispatch/dispatch_segmented_reduce.cuh> // cub::DispatchSegmentedReduce
 #include <cub/thread/thread_load.cuh> // cub::LoadModifier
@@ -200,7 +200,7 @@ struct __align__({2}) storage_t {{
 {7}
 using device_segmented_reduce_policy = {8}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_segmented_reduce_policy")>()
     = cub::detail::reduce::ReducePolicyWrapper<device_segmented_reduce_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/segmented_sort.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh> // cub::detail::choose_offset_t
 #include <cub/detail/launcher/cuda_driver.cuh> // cub::detail::CudaDriverLauncherFactory
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/dispatch/dispatch_segmented_sort.cuh> // cub::DispatchSegmentedSort
 #include <cub/device/dispatch/kernels/kernel_segmented_sort.cuh> // DeviceSegmentedSort kernels
 #include <cub/device/dispatch/tuning/tuning_segmented_sort.cuh> // policy_hub
@@ -664,7 +664,7 @@ struct __align__({4}) items_storage_t {{
 using device_segmented_sort_policy = {11}::MaxPolicy;
 using device_three_way_partition_policy = {12}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& segmented_sort_policy_generator() {{
   return ptx_json::id<ptx_json::string("device_segmented_sort_policy")>()
     = cub::detail::segmented_sort::SegmentedSortPolicyWrapper<device_segmented_sort_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/three_way_partition.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh> // cub::detail::choose_offset_t
 #include <cub/detail/launcher/cuda_driver.cuh> // cub::detail::CudaDriverLauncherFactory
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/dispatch/dispatch_three_way_partition.cuh> // cub::DispatchThreeWayPartitionIf
 #include <cub/device/dispatch/kernels/kernel_three_way_partition.cuh> // DeviceThreeWayPartition kernels
 #include <cub/device/dispatch/tuning/tuning_three_way_partition.cuh> // policy_hub
@@ -209,7 +209,7 @@ struct __align__({2}) storage_t {{
 {9}
 using device_three_way_partition_policy = {10}::MaxPolicy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_three_way_partition_policy")>()
     = cub::detail::three_way_partition::ThreeWayPartitionPolicyWrapper<device_three_way_partition_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/transform.cu

@@ -10,7 +10,7 @@
 
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/dispatch/dispatch_transform.cuh>
 #include <cub/device/dispatch/tuning/tuning_transform.cuh>
 #include <cub/util_arch.cuh>
@@ -289,7 +289,7 @@ struct __align__({3}) output_storage_t {{
 {6}
 using device_transform_policy = {7}::max_policy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_transform_policy")>()
     = cub::detail::transform::TransformPolicyWrapper<device_transform_policy::ActivePolicy>::EncodedPolicy();
@@ -498,7 +498,7 @@ struct __align__({5}) output_storage_t {{
 {9}
 using device_transform_policy = {10}::max_policy;
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_transform_policy")>()
     = cub::detail::transform::TransformPolicyWrapper<device_transform_policy::ActivePolicy>::EncodedPolicy();

c/parallel/src/unique_by_key.cu

@@ -11,7 +11,7 @@
 #include <cub/block/block_scan.cuh>
 #include <cub/detail/choose_offset.cuh>
 #include <cub/detail/launcher/cuda_driver.cuh>
-#include <cub/detail/ptx-json-parser.h>
+#include <cub/detail/ptx-json-parser.cuh>
 #include <cub/device/device_select.cuh>
 
 #include <format>
@@ -298,7 +298,7 @@ struct device_unique_by_key_vsmem_helper {{
   }};
 }};
 
-#include <cub/detail/ptx-json/json.h>
+#include <cub/detail/ptx-json/json.cuh>
 __device__ consteval auto& policy_generator() {{
   return ptx_json::id<ptx_json::string("device_unique_by_key_policy")>()
     = cub::detail::unique_by_key::UniqueByKeyPolicyWrapper<device_unique_by_key_policy::ActivePolicy>::EncodedPolicy();

c/parallel/test/CMakeLists.txt

@@ -18,11 +18,11 @@ function(cccl_c_parallel_add_test target_name_var source)
   target_link_libraries(
     ${target_name}
     PRIVATE
+      cccl.compiler_interface
       cccl.c.parallel
       CUDA::cudart_static
       CUDA::nvrtc
       cccl.c2h.main
-      cccl.compiler_interface_cpp20
   )
 
   # Get the first CUDA include directory only

c2h/CMakeLists.txt

@@ -33,9 +33,9 @@ target_include_directories(cccl.c2h PUBLIC "${C2H_SOURCE_DIR}/include")
 target_link_libraries(
   cccl.c2h
   PUBLIC #
+    cccl.compiler_interface
     CCCL::CCCL
     Catch2::Catch2
-    cccl.compiler_interface_cpp17
 )
 cccl_configure_target(cccl.c2h DIALECT 17)