Tune load / store / caching methods

gevtushenko · May 5, 2023 · 0ed82d4 · 0ed82d4
1 parent 997723a
commit 0ed82d4
Show file tree

Hide file tree

Showing 2 changed files with 43 additions and 6 deletions.
diff --git a/benchmarks/bench/merge_sort/keys.cu b/benchmarks/bench/merge_sort/keys.cu
@@ -2,6 +2,8 @@
 
 #include <nvbench_helper.cuh>
 
+// %RANGE% TUNE_TRANSPOSE trp 0:1:1
+// %RANGE% TUNE_LOAD ld 0:2:1
 // %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
 // %RANGE% TUNE_THREADS_PER_BLOCK_POW2 tpb 6:10:1
 
@@ -12,6 +14,23 @@
 using value_t = cub::NullType;
 
 #if !TUNE_BASE
+
+#if TUNE_TRANSPOSE == 0
+#define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
+#define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_DIRECT
+#else
+#define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
+#define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_WARP_TRANSPOSE
+#endif
+
+#if TUNE_LOAD == 0
+#define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
+#elif TUNE_LOAD == 1
+#define TUNE_LOAD_MODIFIER cub::LOAD_LDG
+#else
+#define TUNE_LOAD_MODIFIER cub::LOAD_CA
+#endif
+
 template <typename KeyT>
 struct policy_hub_t
 {
@@ -20,9 +39,9 @@ struct policy_hub_t
     using MergeSortPolicy =
       cub::AgentMergeSortPolicy<TUNE_THREADS_PER_BLOCK,
                                 cub::Nominal4BItemsToItems<KeyT>(TUNE_ITEMS_PER_THREAD),
-                                cub::BLOCK_LOAD_WARP_TRANSPOSE,
-                                cub::LOAD_DEFAULT,
-                                cub::BLOCK_STORE_WARP_TRANSPOSE>;
+                                TUNE_LOAD_ALGORITHM,
+                                TUNE_LOAD_MODIFIER,
+                                TUNE_STORE_ALGORITHM>;
   };
 
   using MaxPolicy = policy_t;

diff --git a/benchmarks/bench/merge_sort/pairs.cu b/benchmarks/bench/merge_sort/pairs.cu
@@ -2,6 +2,8 @@
 
 #include <nvbench_helper.cuh>
 
+// %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
+// %RANGE% TUNE_THREADS_PER_BLOCK_POW2 tpb 6:10:1
 // %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
 // %RANGE% TUNE_THREADS_PER_BLOCK_POW2 tpb 6:10:1
 
@@ -10,6 +12,22 @@
 #endif
 
 #if !TUNE_BASE
+#if TUNE_TRANSPOSE == 0
+#define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
+#define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_DIRECT
+#else
+#define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
+#define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_WARP_TRANSPOSE
+#endif
+
+#if TUNE_LOAD == 0
+#define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
+#elif TUNE_LOAD == 1
+#define TUNE_LOAD_MODIFIER cub::LOAD_LDG
+#else
+#define TUNE_LOAD_MODIFIER cub::LOAD_CA
+#endif
+
 template <typename KeyT>
 struct policy_hub_t
 {
@@ -18,9 +36,9 @@ struct policy_hub_t
     using MergeSortPolicy =
       cub::AgentMergeSortPolicy<TUNE_THREADS_PER_BLOCK,
                                 cub::Nominal4BItemsToItems<KeyT>(TUNE_ITEMS_PER_THREAD),
-                                cub::BLOCK_LOAD_WARP_TRANSPOSE,
-                                cub::LOAD_DEFAULT,
-                                cub::BLOCK_STORE_WARP_TRANSPOSE>;
+                                TUNE_LOAD_ALGORITHM,
+                                TUNE_LOAD_MODIFIER,
+                                TUNE_STORE_ALGORITHM>;
   };
 
   using MaxPolicy = policy_t;