Fixes to CCCL.C

Author: bernhardmgruber

c/parallel/src/radix_sort.cu

@@ -204,10 +204,9 @@ CUresult cccl_device_radix_sort_build_ex(
     const char* name = "test";
 
     const auto key_cpp = cccl_type_enum_to_name(input_keys_it.value_type.type);
-    const auto value_cpp =
-      input_values_it.type == cccl_iterator_kind_t::CCCL_POINTER && input_values_it.state == nullptr
-        ? "cub::NullType"
-        : cccl_type_enum_to_name(input_values_it.value_type.type);
+    const auto keys_only =
+      input_values_it.type == cccl_iterator_kind_t::CCCL_POINTER && input_values_it.state == nullptr;
+    const auto value_cpp = keys_only ? "cub::NullType" : cccl_type_enum_to_name(input_values_it.value_type.type);
     const std::string op_src =
       (decomposer.name == nullptr || (decomposer.name != nullptr && decomposer.name[0] == '\0'))
         ? "using op_wrapper = cub::detail::identity_decomposer_t;"
@@ -224,15 +223,16 @@ CUresult cccl_device_radix_sort_build_ex(
         case CCCL_FLOAT32:
           return cub::detail::type_t::float32;
         case CCCL_FLOAT64:
-          return cub::detail::type_t::float32;
+          return cub::detail::type_t::float64;
         default:
           return cub::detail::type_t::other;
       }
     }();
 
     const auto cub_arch_policies = cub::detail::radix_sort::arch_policies{
       static_cast<int>(input_keys_it.value_type.size),
-      static_cast<int>(input_values_it.value_type.size),
+      // FIXME(bgruber): input_values_it.value_type.size is 4 when it represents cub::NullType, which is very odd
+      keys_only ? 0 : static_cast<int>(input_values_it.value_type.size),
       int{sizeof(OffsetT)},
       key_type};