[CUB] Fix silent sample loss in DeviceHistogram for signed bytes and wide histograms

c/parallel/src/histogram.cu

@@ -330,7 +330,7 @@ static_assert(device_histogram_policy()(detail::current_tuning_cc()) == {4}, "Ho
   const int privatized_smem_bins =
     num_output_levels_val - 1 > cub::detail::histogram::max_privatized_smem_bins ? 0 : 256;
 
-  const bool is_byte_sample = d_samples.value_type.size == 1;
+  const bool is_byte_sample = d_samples.value_type.size == 1 && d_samples.value_type.type != CCCL_INT8;
 
   std::string init_kernel_name  = histogram::get_init_kernel_name(num_active_channels, counter_cpp, offset_cpp);
   std::string sweep_kernel_name = histogram::get_sweep_kernel_name(
@@ -626,8 +626,9 @@ CUresult cccl_device_histogram_even(
   int64_t row_stride_samples,
   CUstream stream)
 {
-  auto histogram_impl = d_samples.value_type.size == 1 ? cccl_device_histogram_even_impl<::cuda::std::true_type>
-                                                       : cccl_device_histogram_even_impl<::cuda::std::false_type>;
+  const bool is_byte_sample = d_samples.value_type.size == 1 && d_samples.value_type.type != CCCL_INT8;
+  auto histogram_impl       = is_byte_sample ? cccl_device_histogram_even_impl<::cuda::std::true_type>
+                                             : cccl_device_histogram_even_impl<::cuda::std::false_type>;
 
   return histogram_impl(
     build,

cub/cub/agent/agent_histogram.cuh

@@ -293,7 +293,7 @@ struct AgentHistogram
         int output_bin       = -1;
         const CounterT count = privatized_histograms[ch][bin];
         const bool is_valid  = count > 0;
-        output_decode_op[ch].template BinSelect<load_modifier>(static_cast<SampleT>(bin), output_bin, is_valid);
+        output_decode_op[ch].template BinSelect<load_modifier>(bin, output_bin, is_valid);
 
         if (output_bin >= 0)
         {

cub/cub/device/device_histogram.cuh

@@ -31,6 +31,7 @@
 #include <cuda/__execution/require.h>
 #include <cuda/std/__algorithm/copy.h>
 #include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_signed.h>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/array>
 #include <cuda/std/limits>
@@ -774,7 +775,8 @@ public:
     _CCCL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceHistogram::MultiHistogramEven");
 
     using SampleT = cub::detail::it_value_t<SampleIteratorT>;
-    ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
+    // Signed byte samples must not use the pass-thru path: negative values would yield negative privatized bins.
+    using is_byte_sample_t = ::cuda::std::bool_constant<sizeof(SampleT) == 1 && !::cuda::std::is_signed_v<SampleT>>;
 
     using default_policy_selector =
       detail::histogram::policy_selector_from_types<SampleT, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, true>;
@@ -799,7 +801,7 @@ public:
               (int) num_rows,
               (int) (row_stride_bytes / sizeof(SampleT)),
               stream,
-              is_byte_sample,
+              is_byte_sample_t{},
               policy_selector);
           }
         }
@@ -816,7 +818,7 @@ public:
           num_rows,
           (OffsetT) (row_stride_bytes / sizeof(SampleT)),
           stream,
-          is_byte_sample,
+          is_byte_sample_t{},
           policy_selector);
       });
   }
@@ -1500,7 +1502,8 @@ public:
     _CCCL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceHistogram::MultiHistogramRange");
 
     using SampleT = cub::detail::it_value_t<SampleIteratorT>;
-    ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
+    // Signed byte samples must not use the pass-thru path: negative values would yield negative privatized bins.
+    using is_byte_sample_t = ::cuda::std::bool_constant<sizeof(SampleT) == 1 && !::cuda::std::is_signed_v<SampleT>>;
 
     using default_policy_selector =
       detail::histogram::policy_selector_from_types<SampleT, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, false>;
@@ -1524,7 +1527,7 @@ public:
               (int) num_rows,
               (int) (row_stride_bytes / sizeof(SampleT)),
               stream,
-              is_byte_sample,
+              is_byte_sample_t{},
               policy_selector);
           }
         }
@@ -1540,7 +1543,7 @@ public:
           num_rows,
           (OffsetT) (row_stride_bytes / sizeof(SampleT)),
           stream,
-          is_byte_sample,
+          is_byte_sample_t{},
           policy_selector);
       });
   }
@@ -2062,7 +2065,8 @@ public:
     _CCCL_NVTX_RANGE_SCOPE("cub::DeviceHistogram::MultiHistogramEven");
 
     using SampleT = cub::detail::it_value_t<SampleIteratorT>;
-    ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
+    // Signed byte samples must not use the pass-thru path: negative values would yield negative privatized bins.
+    using is_byte_sample_t = ::cuda::std::bool_constant<sizeof(SampleT) == 1 && !::cuda::std::is_signed_v<SampleT>>;
 
     using default_policy_selector =
       detail::histogram::policy_selector_from_types<SampleT, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, true>;
@@ -2084,7 +2088,7 @@ public:
               (int) num_rows,
               (int) (row_stride_bytes / sizeof(SampleT)),
               stream,
-              is_byte_sample,
+              is_byte_sample_t{},
               policy_selector);
           }
         }
@@ -2101,7 +2105,7 @@ public:
           num_rows,
           (OffsetT) (row_stride_bytes / sizeof(SampleT)),
           stream,
-          is_byte_sample,
+          is_byte_sample_t{},
           policy_selector);
       });
   }
@@ -2525,7 +2529,8 @@ public:
     _CCCL_NVTX_RANGE_SCOPE("cub::DeviceHistogram::MultiHistogramRange");
 
     using SampleT = cub::detail::it_value_t<SampleIteratorT>;
-    ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
+    // Signed byte samples must not use the pass-thru path: negative values would yield negative privatized bins.
+    using is_byte_sample_t = ::cuda::std::bool_constant<sizeof(SampleT) == 1 && !::cuda::std::is_signed_v<SampleT>>;
 
     using default_policy_selector =
       detail::histogram::policy_selector_from_types<SampleT, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, false>;
@@ -2546,7 +2551,7 @@ public:
               (int) num_rows,
               (int) (row_stride_bytes / sizeof(SampleT)),
               stream,
-              is_byte_sample,
+              is_byte_sample_t{},
               policy_selector);
           }
         }
@@ -2562,7 +2567,7 @@ public:
           num_rows,
           (OffsetT) (row_stride_bytes / sizeof(SampleT)),
           stream,
-          is_byte_sample,
+          is_byte_sample_t{},
           policy_selector);
       });
   }

cub/cub/device/dispatch/kernels/kernel_histogram.cuh

@@ -265,9 +265,10 @@ struct Transforms
       m_scale = this->ComputeScale(num_levels, m_max, m_min);
     }
 
-    // Method for converting samples to bin-ids
-    template <CacheLoadModifier LOAD_MODIFIER>
-    _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(SampleT sample, int& bin, bool valid) const
+    // Method for converting samples to bin-ids. The sample type is a template parameter because the
+    // agent also feeds privatized bin indices through this op, which must not round-trip through SampleT.
+    template <CacheLoadModifier LOAD_MODIFIER, typename _SampleT>
+    _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(_SampleT sample, int& bin, bool valid) const
     {
       const CommonT common_sample = static_cast<CommonT>(sample);
 

cub/test/catch2_test_device_histogram.cu

@@ -831,3 +831,175 @@ CUB_TEST("DeviceHistogram::HistogramEven bin calculation regression", "[histogra
     static_cast<int>(d_samples.size()));
   CHECK(h_histogram_ref == d_histogram);
 }
+
+// Regression test for NVIDIA/cccl#10977: signed byte samples produced negative privatized bin
+// indices in the pass-thru byte-sample path and were never counted.
+CUB_TEST("DeviceHistogram::Histogram* negative signed byte samples", "[histogram][device]", CUB_SMALL)
+{
+  constexpr int num_bins         = 4;
+  constexpr int num_levels       = num_bins + 1;
+  const int8_t h_samples[]       = {-60, -1, 10, 63};
+  auto d_samples                 = c2h::device_vector<int8_t>(cs::begin(h_samples), cs::end(h_samples));
+  const auto* const d_sample_ptr = thrust::raw_pointer_cast(d_samples.data());
+  c2h::host_vector<int> h_expected{1, 1, 1, 1};
+
+  SECTION("HistogramEven")
+  {
+    auto d_histogram = c2h::device_vector<int>(num_bins);
+    histogram_even(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      num_levels,
+      static_cast<int8_t>(-60),
+      static_cast<int8_t>(64),
+      static_cast<int>(d_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+
+  SECTION("HistogramRange")
+  {
+    const int h_levels[] = {-60, -30, 0, 30, 64};
+    auto d_levels        = c2h::device_vector<int>(cs::begin(h_levels), cs::end(h_levels));
+    auto d_histogram     = c2h::device_vector<int>(num_bins);
+    histogram_range(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      num_levels,
+      thrust::raw_pointer_cast(d_levels.data()),
+      static_cast<int>(d_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+}
+
+// Regression test for NVIDIA/cccl#10977: with more than 127 bins, full-range int8_t histograms lost
+// bins through the SampleT round trip of privatized bin indices in StoreOutput.
+CUB_TEST("DeviceHistogram::Histogram* full 8-bit domain", "[histogram][device]", CUB_SMALL)
+{
+  c2h::host_vector<int8_t> h_all_samples(256);
+  for (int i = 0; i < 256; ++i)
+  {
+    h_all_samples[i] = static_cast<int8_t>(i - 128);
+  }
+  auto d_all_samples             = c2h::device_vector<int8_t>(h_all_samples.begin(), h_all_samples.end());
+  const auto* const d_sample_ptr = thrust::raw_pointer_cast(d_all_samples.data());
+  c2h::host_vector<int> h_expected(256, 1);
+
+  SECTION("HistogramEven")
+  {
+    auto d_histogram = c2h::device_vector<int>(256);
+    histogram_even(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      257,
+      -128,
+      128,
+      static_cast<int>(d_all_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+
+  SECTION("HistogramRange")
+  {
+    c2h::host_vector<int> h_levels(257);
+    for (int i = 0; i < 257; ++i)
+    {
+      h_levels[i] = i - 128;
+    }
+    auto d_levels    = c2h::device_vector<int>(h_levels.begin(), h_levels.end());
+    auto d_histogram = c2h::device_vector<int>(256);
+    histogram_range(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      257,
+      thrust::raw_pointer_cast(d_levels.data()),
+      static_cast<int>(d_all_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+}
+
+// Regression test for NVIDIA/cccl#10976: the privatized bin index was cast through the sample type
+// before the output decode op, so bin indices that are not exactly representable in the sample type
+// were merged or dropped (and, for float-like types, could write out of bounds).
+CUB_TEST("DeviceHistogram::Histogram* bin indices survive the output decode", "[histogram][device]", CUB_SMALL)
+{
+  constexpr int num_bins   = 1 << 16;
+  constexpr int num_levels = num_bins + 1;
+
+  c2h::host_vector<int16_t> h_i16_samples(num_bins);
+  for (int i = 0; i < num_bins; ++i)
+  {
+    h_i16_samples[i] = static_cast<int16_t>(static_cast<int>(i) - 32768);
+  }
+  auto d_samples                 = c2h::device_vector<int16_t>(h_i16_samples.begin(), h_i16_samples.end());
+  const auto* const d_sample_ptr = thrust::raw_pointer_cast(d_samples.data());
+
+  c2h::host_vector<int> h_expected(num_bins, 1);
+
+  SECTION("HistogramEven")
+  {
+    auto d_histogram = c2h::device_vector<int>(num_bins);
+    histogram_even(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      num_levels,
+      -32768,
+      32768,
+      static_cast<int>(d_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+
+  SECTION("HistogramRange")
+  {
+    c2h::host_vector<int> h_i16_levels(num_levels);
+    for (int i = 0; i < num_levels; ++i)
+    {
+      h_i16_levels[i] = i - 32768;
+    }
+    auto d_levels    = c2h::device_vector<int>(h_i16_levels.begin(), h_i16_levels.end());
+    auto d_histogram = c2h::device_vector<int>(num_bins);
+    histogram_range(
+      d_sample_ptr,
+      thrust::raw_pointer_cast(d_histogram.data()),
+      num_levels,
+      thrust::raw_pointer_cast(d_levels.data()),
+      static_cast<int>(d_samples.size()));
+    CHECK(d_histogram == h_expected);
+  }
+
+#if TEST_HALF_T()
+  SECTION("__half samples with more than 2048 bins")
+  {
+    // Strictly increasing exactly representable half values: k/1024 below 2048, then every second
+    // integer step above, since half needs 11 significand bits and ulp doubles at 2.0.
+    constexpr int half_num_bins   = 3000;
+    constexpr int half_num_levels = half_num_bins + 1;
+    const auto exact_half         = [](int i) {
+      const int k = i < 2048 ? i : 2048 + (i - 2048) * 2;
+      return half_t(static_cast<float>(k) / 1024.0f);
+    };
+
+    c2h::host_vector<half_t> h_half_levels(half_num_levels);
+    for (int i = 0; i < half_num_levels; ++i)
+    {
+      h_half_levels[i] = exact_half(i);
+    }
+    c2h::host_vector<half_t> h_half_samples(half_num_bins);
+    for (int i = 0; i < half_num_bins; ++i)
+    {
+      h_half_samples[i] = exact_half(i);
+    }
+    auto d_half_levels  = c2h::device_vector<half_t>(h_half_levels.begin(), h_half_levels.end());
+    auto d_half_samples = c2h::device_vector<half_t>(h_half_samples.begin(), h_half_samples.end());
+
+    auto d_histogram = c2h::device_vector<int>(half_num_bins);
+    histogram_range(
+      cast_if_half_pointer(thrust::raw_pointer_cast(d_half_samples.data())),
+      thrust::raw_pointer_cast(d_histogram.data()),
+      half_num_levels,
+      cast_if_half_pointer(thrust::raw_pointer_cast(d_half_levels.data())),
+      half_num_bins);
+
+    c2h::host_vector<int> h_half_expected(half_num_bins, 1);
+    CHECK(d_histogram == h_half_expected);
+  }
+#endif // TEST_HALF_T()
+}