diff --git a/benchmarks/bench/merge_sort/keys.cu b/benchmarks/bench/merge_sort/keys.cu index f26e6c6c7..76eb1d26e 100644 --- a/benchmarks/bench/merge_sort/keys.cu +++ b/benchmarks/bench/merge_sort/keys.cu @@ -2,6 +2,8 @@ #include +// %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 struct policy_hub_t { @@ -20,9 +39,9 @@ struct policy_hub_t using MergeSortPolicy = cub::AgentMergeSortPolicy(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 index 8b5e97faf..cd254574c 100644 --- a/benchmarks/bench/merge_sort/pairs.cu +++ b/benchmarks/bench/merge_sort/pairs.cu @@ -2,6 +2,8 @@ #include +// %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 struct policy_hub_t { @@ -18,9 +36,9 @@ struct policy_hub_t using MergeSortPolicy = cub::AgentMergeSortPolicy(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;