Remote int4 related

include/ck/utility/amd_wmma.hpp

@@ -98,27 +98,5 @@ struct intrin_wmma_i32_16x16x16_iu8_w32<16, 16, neg_a, neg_b, clamp>
     }
 };
 
-#ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
-// src: iu4, dst: i32
-template <index_t MPerWave, index_t NPerWave, bool neg_a, bool neg_b, bool clamp>
-struct intrin_wmma_i32_16x16x16_iu4_w32;
-
-template <bool neg_a, bool neg_b, bool clamp>
-struct intrin_wmma_i32_16x16x16_iu4_w32<16, 16, neg_a, neg_b, clamp>
-{
-    template <class FloatC>
-    __device__ static void Run(const int4x16_t& reg_a, const int4x16_t& reg_b, FloatC& reg_c)
-    {
-        reg_c.template AsType<int32x8_t>()(Number<0>{}) =
-            __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(
-                neg_a,
-                bit_cast<int32x4_t>(reg_a),
-                neg_b,
-                bit_cast<int32x4_t>(reg_b),
-                reg_c.template AsType<int32x8_t>()[Number<0>{}],
-                clamp);
-    }
-};
-#endif
 } // namespace ck
 #endif

include/ck/utility/data_type.hpp

@@ -942,11 +942,6 @@ using int8x16_t = typename vector_type<int8_t, 16>::type;
 using int8x32_t = typename vector_type<int8_t, 32>::type;
 using int8x64_t = typename vector_type<int8_t, 64>::type;
 
-#ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
-// i4
-using int4x16_t = typename vector_type<int4_t, 16>::type;
-#endif
-
 // Convert X to Y
 template <typename Y, typename X>
 __host__ __device__ constexpr Y type_convert(X x)

test/wmma_op/wmma_op.cpp

@@ -60,9 +60,6 @@ int main(int, char*[])
     pass &= run_test<ck::half_t,  ck::half_t,  ck::half_t,  ck::half_t, 16    >();
     pass &= run_test<ck::bhalf_t, ck::bhalf_t, ck::bhalf_t, float,      16    >();
     pass &= run_test<int8_t,      int8_t,      int32_t,     int32_t,    8     >();
-#ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
-    pass &= run_test<int4_t,      int4_t,      int32_t,     int32_t,    8     >();
-#endif
     // clang-format on
 
     std::cout << "TestGemm ..... " << (pass ? "SUCCESS" : "FAILURE") << std::endl;