add int8 buildins

composable_kernel/include/utility/amd_xdlops.hpp

@@ -307,7 +307,7 @@ struct intrin_mfma_i32_32x32x8i8<32, 32>
     __device__ static void Run(const int8x4_t& reg_a, const int8x4_t& reg_b, FloatC& reg_c)
     {
         reg_c.template AsType<int32x16_t>()(Number<0>{}) =
-            llvm_intrin_amdgcn_mfma_i32_32x32x8i8(bit_cast<int>(reg_a),
+            __builtin_amdgcn_mfma_i32_32x32x8i8(bit_cast<int>(reg_a),
                                                 bit_cast<int>(reg_b),
                                                 reg_c.template AsType<int32x16_t>()[Number<0>{}],
                                                 0,
@@ -326,7 +326,7 @@ struct intrin_mfma_i32_16x16x16i8<16, 16>
     __device__ static void Run(const int8x4_t& reg_a, const int8x4_t& reg_b, FloatC& reg_c)
     {
         reg_c.template AsType<int32x4_t>()(Number<0>{}) =
-            llvm_intrin_amdgcn_mfma_i32_16x16x16i8(bit_cast<int>(reg_a),
+            __builtin_amdgcn_mfma_i32_16x16x16i8(bit_cast<int>(reg_a),
                                                  bit_cast<int>(reg_b),
                                                  reg_c.template AsType<int32x4_t>()[Number<0>{}],
                                                  0,

composable_kernel/include/utility/dynamic_buffer.hpp

@@ -169,6 +169,8 @@ struct DynamicBuffer
                                        is_same<remove_cvref_t<X>, int8x2_t>::value) ||
                                       (is_same<remove_cvref_t<T>, int8_t>::value &&
                                        is_same<remove_cvref_t<X>, int8x4_t>::value) ||
+                                      (is_same<remove_cvref_t<T>, int8_t>::value &&
+                                       is_same<remove_cvref_t<X>, int8x8_t>::value) ||
                                       (is_same<remove_cvref_t<T>, int8x4_t>::value &&
                                        is_same<remove_cvref_t<X>, int8x4_t>::value) ||
                                       (is_same<remove_cvref_t<T>, int8x8_t>::value &&
@@ -202,6 +204,14 @@ struct DynamicBuffer
                         *c_style_pointer_cast<int32_t*>(&p_data_[i]) =
                             *c_style_pointer_cast<const int32_t*>(&x);
                     }
+                    else if constexpr(is_same<remove_cvref_t<T>, int8_t>::value &&
+                                      is_same<remove_cvref_t<X>, int8x8_t>::value)
+                    {
+                        // HACK: cast pointer of x is bad
+                        // TODO: remove this after compiler fix
+                        *c_style_pointer_cast<int32x2_t*>(&p_data_[i]) =
+                            *c_style_pointer_cast<const int32x2_t*>(&x);
+                    }
                     else if constexpr(is_same<remove_cvref_t<T>, int8x4_t>::value &&
                                       is_same<remove_cvref_t<X>, int8x4_t>::value)
                     {

host/driver_offline/src/gemm_driver_offline.cpp

@@ -239,7 +239,7 @@ int main(int argc, char* argv[])
     using ab_data_t  = float;
     using acc_data_t = float;
     using c_data_t   = float;
-#elif 1
+#elif 0
     using ab_data_t  = half_t;
     using acc_data_t = float;
     using c_data_t   = half_t;