fixed c initial

composable_kernel/include/tensor_operation/gridwise_gemm_xdlops_fp16_bfp16.hpp

@@ -51,8 +51,6 @@ struct make_block_work_sequence<MBlockWork, NBlockWork, NBlock1MBlock0>
     __device__ constexpr auto get() { return Sequence<NBlockWork, MBlockWork>{}; }
 };
 
-#define ACCVGPR_ZERO(acc_reg_id) asm volatile("v_accvgpr_write_b32 a[" #acc_reg_id "], 0" : :);
-
 template <index_t GridSize,
           index_t BlockSize,
           class ABFloat,
@@ -214,11 +212,6 @@ struct GridwiseBatchGemmXdlops_gkmkpack_gknkpack_gmn_v2_org
         constexpr index_t c_thread_size = MPerBlock * NPerBlock / BlockSize;
         auto c_thread_vec               = GetRegBuffer<AccFloat, c_thread_size>();
 
-        ACCVGPR_ZERO(0)
-        ACCVGPR_ZERO(1)
-        ACCVGPR_ZERO(2)
-        ACCVGPR_ZERO(3)
-
         // preload data into LDS
         {
             a_blockwise_copy.Run(p_a_global, p_a_block);
@@ -503,11 +496,6 @@ struct GridwiseBatchGemmXdlops_gkmkpack_gknkpack_gmn_v2
         constexpr index_t c_thread_size = MPerBlock * NPerBlock / BlockSize;
         auto c_thread_vec               = GetRegBuffer<AccFloat, c_thread_size>();
 
-        ACCVGPR_ZERO(0)
-        ACCVGPR_ZERO(1)
-        ACCVGPR_ZERO(2)
-        ACCVGPR_ZERO(3)
-
         // preload data into LDS
         {
             a_blockwise_copy.Run(p_a_global, p_a_block);

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

@@ -37,7 +37,7 @@ union float_vec4_t
     StaticallyIndexedArray<float2_t, 2> s2;
     StaticallyIndexedArray<float4_t, 1> s4;
     float n[4];
-    __host__ __device__ constexpr float_vec4_t() {}
+    __host__ __device__ constexpr float_vec4_t() { s4(Number<0>{}) = 0; }
 
     template <index_t vs>
     __host__ __device__ auto& GetVector(Number<vs>);