Move load modifier check to policy

NVIDIA · Jan 30, 2025 · 60df0e4 · 60df0e4
1 parent e7c7248
commit 60df0e4
Show file tree

Hide file tree

Showing 3 changed files with 18 additions and 5 deletions.
diff --git a/c/parallel/src/scan.cu b/c/parallel/src/scan.cu
@@ -60,6 +60,15 @@ struct scan_runtime_tuning_policy
   {
     return load_modifier;
   }
+
+  void CheckLoadModifier() const
+  {
+    if (LoadModifier() == cub::CacheLoadModifier::LOAD_LDG)
+    {
+      throw std::runtime_error("The memory consistency model does not apply to texture "
+                               "accesses");
+    }
+  }
 };
 
 template <typename Tuning, int N>

diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh
@@ -253,12 +253,9 @@ struct DispatchScan
   CUB_RUNTIME_FUNCTION _CCCL_HOST _CCCL_FORCEINLINE cudaError_t
   Invoke(InitKernelT init_kernel, ScanKernelT scan_kernel, ActivePolicyT policy = {})
   {
-    // TODO(ashwin): Does this now need to be a runtime check?
     // `LOAD_LDG` makes in-place execution UB and doesn't lead to better
     // performance.
-    // static_assert(policy.LoadModifier() != CacheLoadModifier::LOAD_LDG,
-    //               "The memory consistency model does not apply to texture "
-    //               "accesses");
+    policy.CheckLoadModifier();
 
     cudaError error = cudaSuccess;
     do

diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh
@@ -251,10 +251,17 @@ struct ScanPolicyWrapper<StaticPolicyT, ::cuda::std::void_t<decltype(StaticPolic
     return MakePolicyWrapper(typename StaticPolicyT::ScanPolicyT());
   }
 
-  CUB_RUNTIME_FUNCTION constexpr CacheLoadModifier LoadModifier()
+  CUB_RUNTIME_FUNCTION static constexpr CacheLoadModifier LoadModifier()
   {
     return StaticPolicyT::ScanPolicyT::LOAD_MODIFIER;
   }
+
+  CUB_RUNTIME_FUNCTION constexpr void CheckLoadModifier()
+  {
+    static_assert(LoadModifier() != CacheLoadModifier::LOAD_LDG,
+                  "The memory consistency model does not apply to texture "
+                  "accesses");
+  }
 };
 
 template <typename PolicyT>