diff --git a/example/ck_tile/18_hstu_attention/example_helper.hpp b/example/ck_tile/18_hstu_attention/example_helper.hpp new file mode 100644 index 0000000000..3322d8538e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/example_helper.hpp @@ -0,0 +1,120 @@ + +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include + +template +static void dumpBufferToFile(const char* fileName, T* data, size_t dataNumItems) +{ + std::ofstream outFile(fileName, std::ios::binary); + if(outFile) + { + outFile.write(reinterpret_cast(data), dataNumItems * sizeof(T)); + outFile.close(); + printf("Write output to file %s\n", fileName); + } + else + { + printf("Could not open file %s for writing\n", fileName); + } +} + +template +static void readDataToBufferFromFile(T* data, size_t dataNumItems, const std::string& fileName) +{ + std::ifstream infile(fileName, std::ios::binary); + if(infile) + { + try + { + infile.read(reinterpret_cast(data), dataNumItems * sizeof(T)); + infile.close(); + } + catch(const std::runtime_error& e) + { + throw e; + }; + } + else + { + throw std::runtime_error("could not open the file for reading"); + } +} + +static std::vector get_integers_from_string(std::string srcStr) +{ + std::vector integers; + std::size_t pos = 0; + std::size_t new_pos; + + new_pos = srcStr.find(',', pos); + while(new_pos != std::string::npos) + { + std::string sliceStr = srcStr.substr(pos, new_pos - pos); + + int len = std::stoi(sliceStr); + + integers.push_back(len); + + pos = new_pos + 1; + new_pos = srcStr.find(',', pos); + }; + + std::string sliceStr = srcStr.substr(pos); + + if(!sliceStr.empty()) + { + int len = std::stoi(sliceStr); + + integers.push_back(len); + }; + + return (integers); +}; + +static std::vector get_floats_from_string(std::string srcStr) +{ + std::vector values; + std::size_t pos = 0; + std::size_t new_pos; + + new_pos = srcStr.find(',', pos); + while(new_pos != std::string::npos) + { + std::string sliceStr = srcStr.substr(pos, new_pos - pos); + + float val = std::stof(sliceStr); + + values.push_back(val); + + pos = new_pos + 1; + new_pos = srcStr.find(',', pos); + }; + + std::string sliceStr = srcStr.substr(pos); + + if(!sliceStr.empty()) + { + float val = std::stof(sliceStr); + + values.push_back(val); + }; + + return (values); +}; + +template +static void supplement_array_by_last_element(std::vector& arr, int target_num_elements) +{ + if(static_cast(arr.size()) < target_num_elements) + { + T last_val = arr.back(); + + for(int i = arr.size(); i < target_num_elements; i++) + arr.push_back(last_val); + }; +}; diff --git a/example/ck_tile/18_hstu_attention/example_hstu_attention_fwd.cpp b/example/ck_tile/18_hstu_attention/example_hstu_attention_fwd.cpp index acff2ecfd6..b512f45b0d 100644 --- a/example/ck_tile/18_hstu_attention/example_hstu_attention_fwd.cpp +++ b/example/ck_tile/18_hstu_attention/example_hstu_attention_fwd.cpp @@ -2,9 +2,7 @@ // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #include -#include -#include -#include +#include #include #include #include @@ -30,43 +28,7 @@ #include "hstu_attention_host_util.hpp" #include "hstu_attention_api.hpp" -template -void dumpBufferToFile(const char* fileName, T* data, size_t dataNumItems) -{ - std::ofstream outFile(fileName, std::ios::binary); - if(outFile) - { - outFile.write(reinterpret_cast(data), dataNumItems * sizeof(T)); - outFile.close(); - printf("Write output to file %s\n", fileName); - } - else - { - printf("Could not open file %s for writing\n", fileName); - } -} - -template -void readDataToBufferFromFile(T* data, size_t dataNumItems, const std::string& fileName) -{ - std::ifstream infile(fileName, std::ios::binary); - if(infile) - { - try - { - infile.read(reinterpret_cast(data), dataNumItems * sizeof(T)); - infile.close(); - } - catch(const std::runtime_error& e) - { - throw e; - }; - } - else - { - throw std::runtime_error("could not open the file for reading"); - } -} +#include "example_helper.hpp" template std::ostream& operator<<(std::ostream& os, const std::vector& v) @@ -130,80 +92,6 @@ auto create_args(int argc, char* argv[]) return std::make_tuple(result, arg_parser); } -static std::vector get_integers_from_string(std::string srcStr) -{ - std::vector integers; - std::size_t pos = 0; - std::size_t new_pos; - - new_pos = srcStr.find(',', pos); - while(new_pos != std::string::npos) - { - std::string sliceStr = srcStr.substr(pos, new_pos - pos); - - int len = std::stoi(sliceStr); - - integers.push_back(len); - - pos = new_pos + 1; - new_pos = srcStr.find(',', pos); - }; - - std::string sliceStr = srcStr.substr(pos); - - if(!sliceStr.empty()) - { - int len = std::stoi(sliceStr); - - integers.push_back(len); - }; - - return (integers); -}; - -static std::vector get_floats_from_string(std::string srcStr) -{ - std::vector values; - std::size_t pos = 0; - std::size_t new_pos; - - new_pos = srcStr.find(',', pos); - while(new_pos != std::string::npos) - { - std::string sliceStr = srcStr.substr(pos, new_pos - pos); - - float val = std::stof(sliceStr); - - values.push_back(val); - - pos = new_pos + 1; - new_pos = srcStr.find(',', pos); - }; - - std::string sliceStr = srcStr.substr(pos); - - if(!sliceStr.empty()) - { - float val = std::stof(sliceStr); - - values.push_back(val); - }; - - return (values); -}; - -template -void supplement_array_by_last_element(std::vector& arr, int target_num_elements) -{ - if(static_cast(arr.size()) < target_num_elements) - { - T last_val = arr.back(); - - for(int i = arr.size(); i < target_num_elements; i++) - arr.push_back(last_val); - }; -}; - // threshold for different dtypes template auto get_elimit() @@ -223,7 +111,7 @@ auto get_elimit() } template -bool run_no_group_hstu(const ck_tile::ArgParser& arg_parser, bool is_jagged) +bool run_no_group_hstu_forward(const ck_tile::ArgParser& arg_parser, bool is_jagged) { using CompDataType = typename HstuAttentionFwdTypeConfig::CompDataType; @@ -690,7 +578,7 @@ bool run_no_group_hstu(const ck_tile::ArgParser& arg_parser, bool is_jagged) } template -bool run_group_hstu(const ck_tile::ArgParser& arg_parser, int num_group) +bool run_group_hstu_forward(const ck_tile::ArgParser& arg_parser, int num_group) { using CompDataType = typename HstuAttentionFwdTypeConfig::CompDataType; @@ -1180,11 +1068,11 @@ int main(int argc, char* argv[]) { if(data_type == "fp16") { - return run_group_hstu(arg_parser, num_group) ? 0 : -2; + return run_group_hstu_forward(arg_parser, num_group) ? 0 : -2; } else if(data_type == "bf16") { - return run_group_hstu(arg_parser, num_group) ? 0 : -2; + return run_group_hstu_forward(arg_parser, num_group) ? 0 : -2; } } else @@ -1193,11 +1081,11 @@ int main(int argc, char* argv[]) if(data_type == "fp16") { - return run_no_group_hstu(arg_parser, is_jagged) ? 0 : -2; + return run_no_group_hstu_forward(arg_parser, is_jagged) ? 0 : -2; } else if(data_type == "bf16") { - return run_no_group_hstu(arg_parser, is_jagged) ? 0 : -2; + return run_no_group_hstu_forward(arg_parser, is_jagged) ? 0 : -2; } };