From fed1474e4f7d8109aba6c7ff06c80e357709b4b1 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Thu, 17 Jul 2025 09:00:41 +0000 Subject: [PATCH] Revert "Disable support of hdim64 amnd hdim256 for quick compiling and testing" This reverts commit ed062f93e2ad835d15bdc3d862213286e2c63b7b. --- .../18_hstu_attention/generate_instances.py | 2 +- .../hstu_attention_hdim_switch.hpp | 12 +- ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...ion_batched_forward_bf16_instances_ref.hpp | 128 ++++++++++++++++++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...ion_batched_forward_fp16_instances_ref.hpp | 128 ++++++++++++++++++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...tion_jagged_forward_bf16_instances_ref.hpp | 128 ++++++++++++++++++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...tion_jagged_forward_fp16_instances_ref.hpp | 128 ++++++++++++++++++ ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 +++ ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 +++ ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 +++ ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 +++ ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 +++ ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 +++ ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 +++ ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 +++ 134 files changed, 3340 insertions(+), 2 deletions(-) create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp create mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index 487499faf0..0f6f6191fd 100644 --- a/example/ck_tile/18_hstu_attention/generate_instances.py +++ b/example/ck_tile/18_hstu_attention/generate_instances.py @@ -161,7 +161,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None: file.write(forward_instance) if __name__ == "__main__": - headdims_fwd = [128] + headdims_fwd = [64, 128, 256] this_dir = os.path.dirname(__file__) output_dir = Path(this_dir) / "instances" diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp index 08ef2acc19..ab8cbe31ed 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp @@ -8,11 +8,21 @@ #define HDIM_SWITCH(HDIM_1, HDIM_2, CONST_NAME, ...) \ [&] { \ - if(HDIM_1 <= 128 && HDIM_2 <= 128) \ + if(HDIM_1 <= 64 && HDIM_2 <= 64) \ + { \ + constexpr ck_tile::index_t CONST_NAME = 64; \ + __VA_ARGS__(); \ + } \ + else if(HDIM_1 <= 128 && HDIM_2 <= 128) \ { \ constexpr ck_tile::index_t CONST_NAME = 128; \ __VA_ARGS__(); \ } \ + else if(HDIM_1 <= 256 && HDIM_2 <= 256) \ + { \ + constexpr ck_tile::index_t CONST_NAME = 256; \ + __VA_ARGS__(); \ + } \ else \ { \ throw std::runtime_error("Head-dim sizes not supported!"); \ diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..1240e37c28 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..dc61eb8ce7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..fecc66f9ef --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..ee6bc64f14 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..e567c295f4 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..4a9fdb3052 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..a75ff8e76a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..a4f6adc2eb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..9bbdd62a9f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..4a65433223 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..bee6cb5b7b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..e0ece6745a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..a0fa6eaf31 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..49b1192f4f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..a5da3f40ab --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..2ba44c444e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp index c869807db9..c4a377f9ed 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp @@ -13,6 +13,70 @@ #include #include "hstu_attention_batched_forward_dispatch.hpp" +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + extern template void run_batched_forward_causal_local_bias_dropout_dispatch< ck_tile::bf16_t, true, @@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..450a70774d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..bdc14f4caa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..2bb09b84ca --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..24db78fdab --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..7788e7f89b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..6227be13ac --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..566da83708 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..f39d11a2aa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..f723f0fb38 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..09be671b6c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..6b2de8cc19 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..fe1e8b3509 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..fd41687481 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..89160eb715 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..1f4508cd6b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..dd0c7978f2 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..2bba03d1e9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..a10e926a29 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..dda7dc07aa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..ce48b9a401 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..6e700fce4b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..5c0206156c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..f2f1bdc3cd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..cd9979a4ee --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..693fe022f6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..73cb991466 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..70d9a31ae0 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..51acf1b238 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..295001b585 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..3658825391 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..4078670fd9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..314f577c91 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp index 38fd44a3e2..5c350254dc 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp @@ -13,6 +13,70 @@ #include #include "hstu_attention_batched_forward_dispatch.hpp" +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + extern template void run_batched_forward_causal_local_bias_dropout_dispatch< ck_tile::fp16_t, true, @@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..ab436fc646 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..d2cc4b579e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..71bb7a409d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..0f209d16e6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..cbefee339d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..84fe9cf731 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..d0f8a1131a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..a1eb2ed6f6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..d399cf6dbb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..91875c4deb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..653c378852 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..c926f2e9b7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..2d60c18f02 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..bbf5f6817e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..411afb0ec1 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..49bb0b3ce4 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..c1e80c9666 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..795c8cc4ed --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..0f3ee275fb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..c875587af3 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..8e73023e0c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..71050ae3d7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..bf40833a7f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..cef802a33e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..58bc729e42 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..6e8ba77772 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..f717ee4563 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..258b1a576e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..d9b4631acd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..7796f67298 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..a442eda92b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..d79f8ffacd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp index 44a80d98a8..0bdbdd91fc 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp @@ -13,6 +13,70 @@ #include #include "hstu_attention_jagged_forward_dispatch.hpp" +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< ck_tile::bf16_t, true, @@ -76,3 +140,67 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..cbccdcdf0a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..ea329e073f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..614eae2111 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..428e08cc13 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..aa35336a75 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..3070211e3f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..44abb0207f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..b08d8d7012 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..4f330cd73b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..22ca2c84ad --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..6a06c0e0d3 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..3bbbe402d2 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..8337a1559f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..3c09e8416f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..ab45759547 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..2843a7a1c5 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..10b05c8a77 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..83d43f43ef --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..5093f4742f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..60440c4916 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..1d55f69805 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..ea8e8895e9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..88464a5075 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..5be82db604 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..68af8bdfaf --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..0bedd49441 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..3679db4777 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..dd9ad9c2bf --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..33d24204a9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..12f1013041 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..4693f6ea05 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..c4427d8e6e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp index 1eac6f27e7..c4ae5f851b 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp @@ -13,6 +13,70 @@ #include #include "hstu_attention_jagged_forward_dispatch.hpp" +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); + extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< ck_tile::fp16_t, true, @@ -76,3 +140,67 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..00be5745df --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..2cb1cd7e83 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..25b0dc246e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..a4b8bb18e8 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..af606140ed --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..4b05b4e9b7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..b14466864c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..04e311ca20 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..23de813bbe --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..8aead96641 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..a36b542f6d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..59d20c38e7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..c78ee824fb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp new file mode 100644 index 0000000000..195fa6cc0a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp new file mode 100644 index 0000000000..0a7a51034d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp new file mode 100644 index 0000000000..31fbe9459c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream);