Only enable check when tuning

src/targets/gpu/jit/ck_gemm.cpp

@@ -184,6 +184,9 @@ struct ck_gemm_compiler : compiler<ck_gemm_compiler>
         options.kernel_name    = v.get("kernel", "ck_gemm_kernel");
         options.virtual_inputs = inputs;
 
+        if (v.get("check", false))
+            options.params += " -DMIGRAPHX_CK_CHECK=1";
+
         auto src = interpolate_string(ck_gemm_kernel,
                                       {{"instance", join_strings(instance, ",")},
                                        {"params", enum_params(inputs.size(), "void * private_p")},

src/targets/gpu/kernels/include/migraphx/kernels/ck.hpp

@@ -109,5 +109,11 @@ struct ck_passthrough
     }
 };
 
+#ifdef MIGRAPHX_CK_CHECK
+#define MIGRAPHX_CK_STATIC_ASSERT static_assert
+#else
+#define MIGRAPHX_CK_STATIC_ASSERT(...)
+#endif
+
 } // namespace migraphx
 #endif // MIGRAPHX_GUARD_KERNELS_CK_HPP

src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp

@@ -59,7 +59,7 @@ __device__ void ck_gemm(E e, A a, B b, Ds... ds)
     constexpr auto e_grid_desc_mblock_mperblock_nblock_nperblock =
         GridwiseGemm::MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(e_grid_desc_m_n);
 
-    static_assert(GridwiseGemm::CheckValidity(
+    MIGRAPHX_CK_STATIC_ASSERT(GridwiseGemm::CheckValidity(
         a_grid_desc_m_k, b_grid_desc_n_k, ds_grid_desc_m_n, e_grid_desc_m_n, block_2_etile_map));
 
     __shared__ char p_shared_block[GridwiseGemm::GetSharedMemoryNumberOfByte()];

tools/tune_ck.py

@@ -23,6 +23,7 @@ def run_driver(b):
     with tmp_file(lambda tf: json.dump(b, tf)) as tf:
         cp = subprocess.run('./bin/gpu-driver {}'.format(tf),
                             capture_output=True,
+                            check=True,
                             shell=True)
         for line in cp.stdout.decode().split("\n"):
             s = line.strip()
@@ -49,6 +50,7 @@ def benchmark_ck(config, tuning):
         },
         'compile_op': {
             'name': 'ck_gemm',
+            'check': True,
             'tuning_val': tuning,
             'inputs': config
         }