From 54fba515b3720545fa30e4e4cf936d77aab1f9bd Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Mon, 9 Aug 2021 17:33:32 +0000 Subject: [PATCH] tidy --- ..._forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp | 6 ------ ..._forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp | 3 --- .../include/tensor_operation/blockwise_gemm_dlops_v3.hpp | 1 - .../include/tensor_operation/blockwise_gemm_xdlops.hpp | 3 --- composable_kernel/include/utility/sequence.hpp | 2 -- ...n_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.cpp | 1 - ...forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nhwk.hpp | 8 -------- .../include/driver_dynamic_gemm_xdlops_v2r3.hpp | 3 --- 8 files changed, 27 deletions(-) diff --git a/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp b/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp index 404129365f..4378314108 100644 --- a/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp +++ b/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp @@ -126,9 +126,6 @@ transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw_no_pad( const auto C = in_n_c_hi_wi_global_desc.GetLength(I1); const auto K = out_n_k_ho_wo_global_desc.GetLength(I1); - const auto Hi = in_n_c_hi_wi_global_desc.GetLength(I2); - const auto Wi = in_n_c_hi_wi_global_desc.GetLength(I3); - const auto Ho = out_n_k_ho_wo_global_desc.GetLength(I2); const auto Wo = out_n_k_ho_wo_global_desc.GetLength(I3); @@ -209,9 +206,6 @@ __host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_ const auto C = in_n_c_hi_wi_global_desc.GetLength(I1); const auto K = out_n_k_ho_wo_global_desc.GetLength(I1); - const auto Hi = in_n_c_hi_wi_global_desc.GetLength(I2); - const auto Wi = in_n_c_hi_wi_global_desc.GetLength(I3); - const auto Ho = out_n_k_ho_wo_global_desc.GetLength(I2); const auto Wo = out_n_k_ho_wo_global_desc.GetLength(I3); diff --git a/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp b/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp index 79051d9512..4764f02787 100644 --- a/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp +++ b/composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp @@ -125,9 +125,6 @@ __host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_ const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3); const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3); - const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1); - const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2); - const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1); const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2); diff --git a/composable_kernel/include/tensor_operation/blockwise_gemm_dlops_v3.hpp b/composable_kernel/include/tensor_operation/blockwise_gemm_dlops_v3.hpp index 074d519b76..25b2ba7ce8 100644 --- a/composable_kernel/include/tensor_operation/blockwise_gemm_dlops_v3.hpp +++ b/composable_kernel/include/tensor_operation/blockwise_gemm_dlops_v3.hpp @@ -69,7 +69,6 @@ struct BlockwiseGemmDlops_km_kn_m0m1n0n1_v3 "wrong! K dimension not consistent\n"); constexpr index_t K = BlockMatrixA{}.GetLength(I1); // A is transposed - constexpr index_t N = BlockMatrixB{}.GetLength(I1); constexpr index_t H = BlockMatrixB{}.GetLength(I2); constexpr index_t W = BlockMatrixB{}.GetLength(I3); diff --git a/composable_kernel/include/tensor_operation/blockwise_gemm_xdlops.hpp b/composable_kernel/include/tensor_operation/blockwise_gemm_xdlops.hpp index 98407ab7fc..74c9dc1547 100644 --- a/composable_kernel/include/tensor_operation/blockwise_gemm_xdlops.hpp +++ b/composable_kernel/include/tensor_operation/blockwise_gemm_xdlops.hpp @@ -52,7 +52,6 @@ struct BlockwiseGemmXdlops_km_kn_m0m1m2n_v1 const index_t waveId = thread_id / WaveSize; const index_t laneId = thread_id % WaveSize; const index_t waveId_m = waveId / NWaves; - const index_t waveId_n = waveId % NWaves; if constexpr(xdlops_gemm.IsKReduction) { @@ -73,7 +72,6 @@ struct BlockwiseGemmXdlops_km_kn_m0m1m2n_v1 const index_t thread_id = get_thread_local_1d_id(); const index_t waveId = thread_id / WaveSize; const index_t laneId = thread_id % WaveSize; - const index_t waveId_m = waveId / NWaves; const index_t waveId_n = waveId % NWaves; if constexpr(xdlops_gemm.IsKReduction) @@ -293,7 +291,6 @@ struct BlockwiseGemmXdlops_km_kn_m0m1m2n_v1_2x2pipeline const index_t thread_id = get_thread_local_1d_id(); const index_t waveId = thread_id / WaveSize; const index_t laneId = thread_id % WaveSize; - const index_t waveId_m = waveId / NWaves; const index_t waveId_n = waveId % NWaves; if constexpr(xdlops_gemm.IsKReduction) diff --git a/composable_kernel/include/utility/sequence.hpp b/composable_kernel/include/utility/sequence.hpp index 81eb488715..b35999d56f 100644 --- a/composable_kernel/include/utility/sequence.hpp +++ b/composable_kernel/include/utility/sequence.hpp @@ -685,8 +685,6 @@ __host__ __device__ constexpr auto operator+(Number, Sequence) template __host__ __device__ constexpr auto operator-(Number, Sequence) { - constexpr auto seq_x = Sequence{}; - return Sequence<(Y - Xs)...>{}; } diff --git a/composable_kernel/src/kernel_wrapper/dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.cpp b/composable_kernel/src/kernel_wrapper/dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.cpp index d49693b511..d946bc63ee 100644 --- a/composable_kernel/src/kernel_wrapper/dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.cpp +++ b/composable_kernel/src/kernel_wrapper/dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.cpp @@ -225,7 +225,6 @@ extern "C" __global__ void constexpr auto I0 = Number<0>{}; constexpr auto I1 = Number<1>{}; constexpr auto I2 = Number<2>{}; - constexpr auto I3 = Number<3>{}; constexpr auto in_n_hi_wi_c_desc = make_dynamic_naive_tensor_descriptor_packed_v2(make_tuple(256, 28, 28, 256)); diff --git a/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nhwk.hpp b/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nhwk.hpp index 10284b48f3..5310503318 100644 --- a/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nhwk.hpp +++ b/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nhwk.hpp @@ -35,11 +35,6 @@ void device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nh constexpr auto I1 = Number<1>{}; constexpr auto I2 = Number<2>{}; constexpr auto I3 = Number<3>{}; - constexpr auto I4 = Number<4>{}; - constexpr auto I5 = Number<5>{}; - constexpr auto I6 = Number<6>{}; - constexpr auto I7 = Number<7>{}; - constexpr auto I8 = Number<8>{}; DeviceMem in_n_hi_wi_c_device_buf(sizeof(TInWei) * in_n_hi_wi_c.mDesc.GetElementSpace()); DeviceMem wei_k_y_x_c_device_buf(sizeof(TInWei) * wei_k_y_x_c.mDesc.GetElementSpace()); @@ -218,9 +213,6 @@ void device_dynamic_convolution_forward_implicit_gemm_v4r4r2_xdlops_nhwc_kyxc_nh const auto K = out_n_ho_wo_k_lengths[I3]; const auto C = wei_k_y_x_c_lengths[I3]; - const auto Hi = in_n_hi_wi_c_lengths[I1]; - const auto Wi = in_n_hi_wi_c_lengths[I2]; - const auto Ho = out_n_ho_wo_k_lengths[I1]; const auto Wo = out_n_ho_wo_k_lengths[I2]; diff --git a/host/driver_offline/include/driver_dynamic_gemm_xdlops_v2r3.hpp b/host/driver_offline/include/driver_dynamic_gemm_xdlops_v2r3.hpp index 481d08188d..a2f4e28c54 100644 --- a/host/driver_offline/include/driver_dynamic_gemm_xdlops_v2r3.hpp +++ b/host/driver_offline/include/driver_dynamic_gemm_xdlops_v2r3.hpp @@ -66,9 +66,6 @@ __host__ float driver_dynamic_gemm_xdlops_v2r3(const FloatAB* p_a_grid, constexpr auto I0 = Number<0>{}; constexpr auto I1 = Number<1>{}; constexpr auto I2 = Number<2>{}; - constexpr auto I3 = Number<3>{}; - constexpr auto I4 = Number<4>{}; - constexpr auto I5 = Number<5>{}; using GridwiseGemm = GridwiseDynamicGemm_k0mk1_k0nk1_mn_xdlops_v2r3