ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 11:34:54 +00:00
Code Issues Packages Projects Releases Wiki Activity

Removing the class derivation to simplify struct HstuAttentionFwdCommonDropoutKargs

Browse Source
This commit is contained in:
Qianfeng Zhang
2026-06-24 09:03:39 +00:00
parent db5cb5d19d
commit bdc300591b
2 changed files with 10 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp
View File

@@ -203,12 +203,6 @@ struct HstuAttentionFwdKernel
ck_tile::index_t nhead_stride_bias;
};
struct HstuAttentionFwdDropoutSeedOffset
{
uint64_t drop_seed;
uint64_t drop_offset;
};
struct HstuAttentionFwdBatchedLSEKargs
{
void* lse_ptr;
@@ -224,7 +218,7 @@ struct HstuAttentionFwdKernel
ck_tile::index_t nhead_stride_lse;
};
struct HstuAttentionFwdCommonDropoutKargs : HstuAttentionFwdDropoutSeedOffset
struct HstuAttentionFwdCommonDropoutKargs
{
void init_dropout(float p_drop, uint64_t seed, uint64_t offset)
{
@@ -237,6 +231,10 @@ struct HstuAttentionFwdKernel
this->drop_offset = offset;
}
// used for generating random numbers
uint64_t drop_seed;
uint64_t drop_offset;
float rp_undrop = 1;
uint8_t p_undrop_in_uint8_t = std::numeric_limits<uint8_t>::max();
};

12
example/ck_tile/18_hstu_attention/hstu_attention_fwd_splitkv_kernel.hpp
View File

@@ -199,18 +199,12 @@ struct HstuAttentionFwdSplitKVKernel
ck_tile::index_t nhead_stride_bias;
};
struct HstuAttentionFwdDropoutSeedOffset
{
uint64_t drop_seed;
uint64_t drop_offset;
};
struct HstuAttentionFwdSoftmaxKargs
{
void* lse_acc_ptr;
};
struct HstuAttentionFwdCommonDropoutKargs : HstuAttentionFwdDropoutSeedOffset
struct HstuAttentionFwdCommonDropoutKargs
{
void init_dropout(float p_drop, uint64_t seed, uint64_t offset)
{
@@ -223,6 +217,10 @@ struct HstuAttentionFwdSplitKVKernel
this->drop_offset = offset;
}
// used for generating random numbers
uint64_t drop_seed;
uint64_t drop_offset;
float rp_undrop = 1;
uint8_t p_undrop_in_uint8_t = std::numeric_limits<uint8_t>::max();
};