diff --git a/src/cunumeric/stat/histogram.cu b/src/cunumeric/stat/histogram.cu
index 7a812cf66..354da3723 100644
--- a/src/cunumeric/stat/histogram.cu
+++ b/src/cunumeric/stat/histogram.cu
@@ -126,15 +126,36 @@ struct HistogramImplBody<VariantKind::GPU, CODE> {
 
     CHECK_CUDA(cudaStreamSynchronize(stream));
 
-    detail::histogram_weights(thrust::device,  // DEFAULT_POLICY.on(stream),
-                              src_ptr,
+#if 0
+    {
+      // VAL* src_ptr                 = nullptr;
+      // size_t src_size              = 0;
+      // WeightType* bins_ptr         = nullptr;
+      // size_t num_intervals         = 0;
+      // WeightType* local_result_ptr = nullptr;
+      // WeightType* weights_ptr = nullptr;
+
+      detail::histogram_proto(DEFAULT_POLICY.on(stream),
+                              src_copy.ptr(0),  // src_ptr compiled...why?
                               src_size,
                               bins_ptr,
                               num_intervals,
                               local_result_ptr,
-                              weights_ptr,
-                              false,
+                              weights_copy.ptr(0),  // weights_ptr,  // <- PROBLEM here
                               stream_);
+    }
+#endif
+
+    // #if 0
+    detail::histogram_weights(DEFAULT_POLICY.on(stream),
+                              src_copy.ptr(0),
+                              src_size,
+                              bins_ptr,
+                              num_intervals,
+                              local_result_ptr,
+                              weights_copy.ptr(0),
+                              stream_);
+    // #endif
 
     CHECK_CUDA(cudaStreamSynchronize(stream));
 
diff --git a/src/cunumeric/stat/histogram_impl.h b/src/cunumeric/stat/histogram_impl.h
index d22c7bb5c..251e9c952 100644
--- a/src/cunumeric/stat/histogram_impl.h
+++ b/src/cunumeric/stat/histogram_impl.h
@@ -68,6 +68,18 @@ struct transform_op_t {
   weight_t Sw_{0};
 };
 
+template <typename exe_policy_t, typename array_t, typename bin_t, typename weight_t>
+void histogram_proto(exe_policy_t exe_pol,
+                     array_t* ptr_src,  // source array, a
+                     size_t n_samples,
+                     bin_t const* ptr_over,          // bins array,
+                     size_t n_intervals,             // |bins| - 1
+                     weight_t* ptr_hist,             // result; pre-allocated, sz = n_intervals
+                     weight_t* ptr_w     = nullptr,  // weights array, w
+                     cudaStream_t stream = nullptr)
+{
+}
+
 template <typename exe_policy_t,
           typename array_t,
           typename bin_t,