From 321dc5dd394e993766ff86136fb71926c1e6f02f Mon Sep 17 00:00:00 2001 From: "Lin, Qun" Date: Thu, 29 May 2025 08:25:18 -0500 Subject: [PATCH] version 3 --- .../grouped_conv_bwd_weight_xdl_fp16.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp index 5212d1db12..b24125ed3a 100644 --- a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp +++ b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp @@ -35,7 +35,7 @@ static constexpr index_t Tile_Align_W = Tile_W + 1; static constexpr index_t ShareMemSize = Tile_H * Tile_Align_W * N_Pack * SizeOfType; static constexpr index_t ScratchSize = ShareMemSize / 64 / 4; static constexpr index_t Num_Wave = 1; - +#define MergeShareMem 1 template __device__ T warp_shuffle_up(const T& v_local, uint32_t lane_delta) { @@ -179,11 +179,18 @@ __global__ void const index_t lane_id = __lane_id(); //index_t n_idx = 0; +#if MergeShareMem + constexpr index_t ShaderMemSizePerWave =ShareMemSize * 2 - Pad_H * Tile_Align_W * sizeof(uint32_t); + __shared__ uint32_t p_share_mem[ShaderMemSizePerWave/sizeof(uint32_t) * Num_Wave]; + auto* p_share_in = &p_share_mem[0]; + //__shared__ char p_input_1[ShareMemSize]; + auto* p_share_out = &p_share_mem[(ShareMemSize - Pad_H * Tile_Align_W * sizeof(uint32_t))/ sizeof(uint32_t)]; +#else __shared__ uint32_t p_share_in[ShareMemSize/sizeof(uint32_t) * Num_Wave]; //__shared__ char p_input_1[ShareMemSize]; __shared__ uint32_t p_share_out[ShareMemSize/sizeof(uint32_t) * Num_Wave]; //__shared__ char p_output_1[ShareMemSize]; - +#endif uint32_t p_input_0_scratch[ScratchSize]; uint32_t p_input_1_scratch[ScratchSize]; uint32_t p_output_0_scratch[ScratchSize];