rework DefualtSemiRingConfiguration struct API for clarity

hpcgarage · Nov 1, 2021 · e0d0c17 · e0d0c17
1 parent a40f7bc
commit e0d0c17
Show file tree

Hide file tree

Showing 261 changed files with 35,740 additions and 35,735 deletions.
diff --git a/bench/device/gen_simt.py b/bench/device/gen_simt.py
@@ -95,9 +95,9 @@
   using WarpShape        = cutlass::gemm::GemmShape<{13}, {14}, {12}>;
   using InstructionShape = cutlass::gemm::GemmShape<1, 1, 1>;
 
-  using Config = typename cuasr::gemm::device::DefaultSemiRingConfiguration< //
-      precision, precision, precision, precision, OpClass,                   //
-      cuasr::{0}<precision>, cuasr::{1}<precision>, SmArch>;
+  using Config = typename cuasr::gemm::device::DefaultSemiRingConfiguration<
+      precision, precision, precision, precision,
+      cuasr::{0}<precision>, cuasr::{1}<precision>, OpClass, SmArch>;
 
   using AddOp            = Config::AdditionOp;
   using MultOp           = Config::MultiplicationOp;

diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nn_n.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nn_t.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nt_n.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nt_t.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tn_n.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tn_t.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tt_n.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tt_t.cu b/bench/device/sm50_simt_binary_or_binary_and_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nn_n.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nn_t.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nt_n.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nt_t.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tn_n.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tn_t.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tt_n.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tt_t.cu b/bench/device/sm50_simt_binary_or_binary_and_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_dsrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_minimum_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_minimum_ssrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_minimum_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_multiplies_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_multiplies_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_dsrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_plus_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_nn_n.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_nn_t.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_nt_n.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_nt_t.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_tn_n.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_tn_t.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_tt_n.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_maximum_plus_ssrgemm_tt_t.cu b/bench/device/sm50_simt_maximum_plus_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_dsrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_maximum_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_maximum_ssrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_maximum_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_multiplies_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_multiplies_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_dsrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_plus_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_nn_n.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_nn_t.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_nt_n.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_nt_t.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_tn_n.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_tn_t.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_tt_n.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_minimum_plus_ssrgemm_tt_t.cu b/bench/device/sm50_simt_minimum_plus_ssrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_nn_n.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_nn_t.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_nt_n.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_nt_t.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_tn_n.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_tn_t.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_tt_n.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_dsrgemm_tt_t.cu b/bench/device/sm50_simt_plus_multiplies_dsrgemm_tt_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_nn_n.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_nn_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_nn_t.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_nn_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_nt_n.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_nt_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_nt_t.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_nt_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_tn_n.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_tn_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_tn_t.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_tn_t.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_tt_n.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_tt_n.cu
diff --git a/bench/device/sm50_simt_plus_multiplies_ssrgemm_tt_t.cu b/bench/device/sm50_simt_plus_multiplies_ssrgemm_tt_t.cu
diff --git a/include/cuasr/gemm/device/default_srgemm_configuration.h b/include/cuasr/gemm/device/default_srgemm_configuration.h
@@ -30,13 +30,15 @@ template <
   typename ElementB,
   typename ElementC,
   typename ElementAccumulator,
-  typename OperatorClass,
   typename AdditionOp,
   typename MultiplicationOp,
+  typename OperatorClass,
   typename ArchTag
 >
 struct DefaultSemiRingConfiguration;
 
+////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////// SM 50 //////////////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////
 
 // Plus-Times semi-ring GEMM configuration
@@ -50,9 +52,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::plus<Element>,
   cuasr::multiplies<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -80,9 +82,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::minimum<Element>,
   cuasr::plus<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -110,9 +112,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::maximum<Element>,
   cuasr::plus<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -139,9 +141,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::maximum<Element>,
   cuasr::minimum<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -168,9 +170,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::minimum<Element>,
   cuasr::maximum<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -197,9 +199,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::minimum<Element>,
   cuasr::multiplies<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -226,9 +228,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::maximum<Element>,
   cuasr::multiplies<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -255,9 +257,9 @@ struct DefaultSemiRingConfiguration<
   Element,
   Element,
   Element,
-  cutlass::arch::OpClassSimt,
   cuasr::binary_or<Element>,
   cuasr::binary_and<Element>,
+  cutlass::arch::OpClassSimt,
   ArchTag> {
 
   static int constexpr kAlignmentA = 1;
@@ -275,6 +277,9 @@ struct DefaultSemiRingConfiguration<
 };
 
 ////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////// SM 80 //////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////
+
 
 } // namespace device
 } // namespace gemm

diff --git a/include/cuasr/gemm/device/srgemm.h b/include/cuasr/gemm/device/srgemm.h
@@ -51,34 +51,34 @@ template <
   /// Threadblock-level tile size (concept: GemmShape)
   typename ThreadblockShape_ = typename DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::ThreadblockShape,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::ThreadblockShape,
   /// Warp-level tile size (concept: GemmShape)
   typename WarpShape_ = typename DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::WarpShape,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::WarpShape,
   /// Instruction-level tile size (concept: GemmShape)
   typename InstructionShape_ = typename DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::InstructionShape,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::InstructionShape,
   /// Epilogue output operator
   typename EpilogueOutputOp_ = typename DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::EpilogueOutputOp,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::EpilogueOutputOp,
   /// Threadblock-level swizzling operator
   typename ThreadblockSwizzle_ =
       typename cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
   /// Number of stages used in the pipelined mainloop
   int Stages = DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::kStages,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::kStages,
   /// Access granularity of A matrix in units of elements
   int AlignmentA = DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::kAlignmentA,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::kAlignmentA,
   /// Access granularity of B matrix in units of elements
   int AlignmentB = DefaultSemiRingConfiguration<
       ElementA_, ElementB_, ElementC_, ElementAccumulator_,
-      OperatorClass_, AdditionOp_, MultiplicationOp_, ArchTag_>::kAlignmentB,
+      AdditionOp_, MultiplicationOp_, OperatorClass_, ArchTag_>::kAlignmentB,
   /// If true, kernel supports split-K with serial reduction
   bool SplitKSerial = false
 >

diff --git a/test/device/gen_simt.py b/test/device/gen_simt.py
@@ -103,9 +103,9 @@
   using WarpShape        = cutlass::gemm::GemmShape<{13}, {14}, {12}>;
   using InstructionShape = cutlass::gemm::GemmShape<1, 1, 1>;
 
-  using Config = typename cuasr::gemm::device::DefaultSemiRingConfiguration< //
-      precision, precision, precision, precision, OpClass,                   //
-      cuasr::{0}<precision>, cuasr::{1}<precision>, SmArch>;
+  using Config = typename cuasr::gemm::device::DefaultSemiRingConfiguration<
+      precision, precision, precision, precision,
+      cuasr::{0}<precision>, cuasr::{1}<precision>, OpClass, SmArch>;
 
   using AddOp            = Config::AdditionOp;
   using MultOp           = Config::MultiplicationOp;