ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

version 3

Browse Source
This commit is contained in:
Lin, Qun
2025-05-29 08:25:18 -05:00
parent 5fd4aef2dc
commit 321dc5dd39
1 changed files with 9 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp
View File

@@ -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 <typename T>
__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];