tweak

composable_kernel/include/utility/config.amd.hpp.in

@@ -11,9 +11,9 @@
 #define CK_DEVICE_BACKEND_AMD 1
 
 // GPU ID
-#define CK_AMD_GPU_GFX906 0
+#define CK_AMD_GPU_GFX906 1
 #define CK_AMD_GPU_GFX908 0
-#define CK_AMD_GPU_GFX1030 1
+#define CK_AMD_GPU_GFX1030 0
 
 // HIP version
 #ifndef CK_HIP_VERSION_FLAT

driver/include/device_dynamic_convolution_forward_implicit_gemm_v4r4_nhwc_kyxc_nhwk.hpp

@@ -112,7 +112,7 @@ void device_dynamic_convolution_forward_implicit_gemm_v4r4_nhwc_kyxc_nhwk(
     wei_k_y_x_c_device_buf.ToDevice(wei_k_y_x_c.mData.data());
     out_n_ho_wo_k_device_buf.ToDevice(out_n_ho_wo_k.mData.data());
 
-#if 0
+#if 1
     // cdata = 16, BlockSize = 64, 16x64x4
     constexpr index_t BlockSize = 64;
 
@@ -417,9 +417,7 @@ void device_dynamic_convolution_forward_implicit_gemm_v4r4_nhwc_kyxc_nhwk(
                         in_n_hi_wi_c_device_buf.GetDeviceBuffer()),
                     static_cast<TOut*>(out_n_ho_wo_k_device_buf.GetDeviceBuffer()));
 
-#if 1
     out_n_ho_wo_k_device_buf.FromDevice(out_n_ho_wo_k.mData.data());
-#endif
 
     auto f_nhwk2nkhw = [&](auto n, auto k, auto ho, auto wo) {
         out_n_k_ho_wo(n, k, ho, wo) = out_n_ho_wo_k(n, ho, wo, k);

driver/src/conv_driver.cpp

@@ -64,7 +64,7 @@ int main(int argc, char* argv[])
 
     using LeftPads  = Sequence<0, 0>;
     using RightPads = Sequence<0, 0>;
-#elif 0
+#elif 1
     constexpr index_t N  = 1;
     constexpr index_t C  = 4;
     constexpr index_t HI = 1080;
@@ -78,7 +78,7 @@ int main(int argc, char* argv[])
 
     using LeftPads  = Sequence<1, 1>;
     using RightPads = Sequence<1, 1>;
-#elif 1
+#elif 0
     constexpr index_t N  = 1;
     constexpr index_t C  = 1;
     constexpr index_t HI = 1024;
@@ -630,11 +630,16 @@ int main(int argc, char* argv[])
     print_array("ConvStrides", to_multi_index(ConvStrides{}));
     print_array("ConvDilations", to_multi_index(ConvDilations{}));
 
-#if 0
+#if 1
     using in_data_t                  = float;
     constexpr index_t in_vector_size = 1;
     using acc_data_t                 = float;
     using out_data_t                 = float;
+#elif 1
+    using in_data_t                  = float;
+    constexpr index_t in_vector_size = 1;
+    using acc_data_t                 = float;
+    using out_data_t                 = int8_t;
 #elif 1
     using in_data_t                  = int8_t;
     constexpr index_t in_vector_size = 4;
@@ -736,7 +741,7 @@ int main(int argc, char* argv[])
          LeftPads{},
          RightPads{},
          nrepeat);
-#elif 0
+#elif 1
     device_dynamic_convolution_forward_implicit_gemm_v4r4_nhwc_kyxc_nhwk<in_data_t,
                                                                          in_vector_size,
                                                                          acc_data_t,