diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp index ea63042735..cf9d6bab52 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_kernel.hpp @@ -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::max(); }; diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_splitkv_kernel.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_splitkv_kernel.hpp index 3b38b3e07c..f1a03a6a8d 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_splitkv_kernel.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_splitkv_kernel.hpp @@ -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::max(); };