From 104572c1f3c884688bee20fcbfd3baee7b0df0c1 Mon Sep 17 00:00:00 2001 From: BoboFang Date: Tue, 22 Apr 2025 15:24:55 +0000 Subject: [PATCH] Fix error after clang-format --- .../99_toy_example/01_add/CMakeLists.txt | 33 ++- .../99_toy_example/02_gemm/CMakeLists.txt | 41 ++- .../03_flash_attention_fwd/CMakeLists.txt | 31 +- .../CMakeLists.txt | 85 +++--- .../generate.py | 272 +++++++++--------- 5 files changed, 234 insertions(+), 228 deletions(-) mode change 100644 => 100755 example/ck_tile/99_toy_example/03_flash_attention_fwd/CMakeLists.txt mode change 100644 => 100755 example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/CMakeLists.txt diff --git a/example/ck_tile/99_toy_example/01_add/CMakeLists.txt b/example/ck_tile/99_toy_example/01_add/CMakeLists.txt index ad51b75885..f241173fe7 100644 --- a/example/ck_tile/99_toy_example/01_add/CMakeLists.txt +++ b/example/ck_tile/99_toy_example/01_add/CMakeLists.txt @@ -1,23 +1,22 @@ set(EXAMPLE_REDUCE "add") -#not using add_example_executable() to add this target, since we don't want this to have -#to be included in "make all/install/check" - message("adding example ${EXAMPLE_REDUCE}") +# not using add_example_executable() to add this target, since we don't want this to have +# to be included in "make all/install/check" +message("adding example ${EXAMPLE_REDUCE}") - add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL add.cpp) target_include_directories(${ - EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) set(EXAMPLE_REDUCE_COMPILE_OPTIONS) +add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL add.cpp) +target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) +set(EXAMPLE_REDUCE_COMPILE_OPTIONS) -#generate assembly -#list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - v-- save - temps - Wno - gnu - line - marker) +# generate assembly +# list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS -v --save-temps -Wno-gnu-line-marker) -#NOTE : we turn off undefined - func - \ - template to let source compile without explicit declare function specializations - list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - Wno - undefined - func - template - Wno - - float - equal) +# NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations +list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) - target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${EXAMPLE_REDUCE_COMPILE_OPTIONS}) +target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${EXAMPLE_REDUCE_COMPILE_OPTIONS}) -#TODO : we have to turn off this global prop, otherwise the progress bar generated -#by cmake will print too many files, execvp : / bin / sh : Argument list too long -#however, this property may affect global -#TODO : consider codegen a makefile by us - set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) +# TODO: we have to turn off this global prop, otherwise the progress bar generated +# by cmake will print too many files, execvp: /bin/sh: Argument list too long +# however, this property may affect global +# TODO: consider codegen a makefile by us +set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) diff --git a/example/ck_tile/99_toy_example/02_gemm/CMakeLists.txt b/example/ck_tile/99_toy_example/02_gemm/CMakeLists.txt index afb32d49a9..02863398fe 100644 --- a/example/ck_tile/99_toy_example/02_gemm/CMakeLists.txt +++ b/example/ck_tile/99_toy_example/02_gemm/CMakeLists.txt @@ -1,28 +1,27 @@ set(EXAMPLE_REDUCE "basic_gemm") -#not using add_example_executable() to add this target, since we don't want this to have -#to be included in "make all/install/check" - message("adding example ${EXAMPLE_REDUCE}") +# not using add_example_executable() to add this target, since we don't want this to have +# to be included in "make all/install/check" +message("adding example ${EXAMPLE_REDUCE}") - add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL gemm.cpp) target_include_directories(${ - EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) set(EXAMPLE_REDUCE_COMPILE_OPTIONS) +add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL gemm.cpp) +target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) +set(EXAMPLE_REDUCE_COMPILE_OPTIONS) -#generate assembly -#list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - v-- save - temps - Wno - gnu - line - marker) +# generate assembly +# list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS -v --save-temps -Wno-gnu-line-marker) -#NOTE : we turn off undefined - func - \ - template to let source compile without explicit declare function specializations - list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - Wno - undefined - func - template - Wno - - float - equal) +# NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations +list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) - if(DEFINED kernel) message("Compiling with Kernel: ${kernel}") - target_compile_definitions(${EXAMPLE_REDUCE} PRIVATE KERNEL_${kernel} = 1) - endif() +if(DEFINED kernel) + message("Compiling with Kernel: ${kernel}") + target_compile_definitions(${EXAMPLE_REDUCE} PRIVATE KERNEL_${kernel}=1) +endif() - target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${ - EXAMPLE_REDUCE_COMPILE_OPTIONS}) +target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${EXAMPLE_REDUCE_COMPILE_OPTIONS}) -#TODO : we have to turn off this global prop, otherwise the progress bar generated -#by cmake will print too many files, execvp : / bin / sh : Argument list too long -#however, this property may affect global -#TODO : consider codegen a makefile by us - set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) +# TODO: we have to turn off this global prop, otherwise the progress bar generated +# by cmake will print too many files, execvp: /bin/sh: Argument list too long +# however, this property may affect global +# TODO: consider codegen a makefile by us +set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) diff --git a/example/ck_tile/99_toy_example/03_flash_attention_fwd/CMakeLists.txt b/example/ck_tile/99_toy_example/03_flash_attention_fwd/CMakeLists.txt old mode 100644 new mode 100755 index 5ab8e1e314..44dfac099c --- a/example/ck_tile/99_toy_example/03_flash_attention_fwd/CMakeLists.txt +++ b/example/ck_tile/99_toy_example/03_flash_attention_fwd/CMakeLists.txt @@ -1,22 +1,19 @@ set(EXAMPLE_REDUCE "basic_flash_attention_fwd") -#not using add_example_executable() to add this target, since we don't want this to have -#to be included in "make all/install/check" - message("adding example ${EXAMPLE_REDUCE}") +# not using add_example_executable() to add this target, since we don't want this to have +# to be included in "make all/install/check" +message("adding example ${EXAMPLE_REDUCE}") - add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL flash_attention_fwd.cpp) - target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) - set(EXAMPLE_REDUCE_COMPILE_OPTIONS) +add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL flash_attention_fwd.cpp) +target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) +set(EXAMPLE_REDUCE_COMPILE_OPTIONS) -#NOTE : we turn off undefined - func - \ - template to let source compile without explicit declare function specializations - list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - Wno - undefined - func - template - - Wno - float - equal) +# NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations +list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) - target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${ - EXAMPLE_REDUCE_COMPILE_OPTIONS}) +target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${EXAMPLE_REDUCE_COMPILE_OPTIONS}) -#TODO : we have to turn off this global prop, otherwise the progress bar generated -#by cmake will print too many files, execvp : / bin / sh : Argument list too long -#however, this property may affect global -#TODO : consider codegen a makefile by us - set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) +# TODO: we have to turn off this global prop, otherwise the progress bar generated +# by cmake will print too many files, execvp: /bin/sh: Argument list too long +# however, this property may affect global +# TODO: consider codegen a makefile by us +set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/CMakeLists.txt b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/CMakeLists.txt old mode 100644 new mode 100755 index 87ec6628cd..e9a697fde3 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/CMakeLists.txt +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/CMakeLists.txt @@ -1,48 +1,59 @@ -set(FLASH_ATTENTION_FWD_KNOWN_APIS - "fwd") set(FLASH_ATTENTION_FWD_ENABLE_APIS - "fwd" CACHE STRING - "semicolon-separated list of APIs to generate (${FLASH_ATTENTION_FWD_KNOWN_APIS}) & " - "link, or \"all\".") if(FLASH_ATTENTION_FWD_ENABLE_APIS STREQUAL - "all") set(FLASH_ATTENTION_FWD_ENABLE_APIS ${ - FLASH_ATTENTION_FWD_KNOWN_APIS}) endif() +set(FLASH_ATTENTION_FWD_KNOWN_APIS "fwd") +set(FLASH_ATTENTION_FWD_ENABLE_APIS "fwd" CACHE STRING + "semicolon-separated list of APIs to generate (${FLASH_ATTENTION_FWD_KNOWN_APIS}) & link, or \"all\".") +if(FLASH_ATTENTION_FWD_ENABLE_APIS STREQUAL "all") + set(FLASH_ATTENTION_FWD_ENABLE_APIS ${FLASH_ATTENTION_FWD_KNOWN_APIS}) +endif() - execute_process( - COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR} / - generate.py-- api $ { FLASH_ATTENTION_FWD_ENABLE_APIS } --working_path $ { - CMAKE_CURRENT_BINARY_DIR - } --list_blobs RESULT_VARIABLE ret) if(ret AND NOT ret EQUAL 0) - message(FATAL_ERROR "Failed to list Flash Attention kernels via Python. ${ret}") endif() +execute_process( + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/generate.py + --api ${FLASH_ATTENTION_FWD_ENABLE_APIS} + --working_path ${CMAKE_CURRENT_BINARY_DIR} + --list_blobs + RESULT_VARIABLE ret +) +if(ret AND NOT ret EQUAL 0) + message(FATAL_ERROR "Failed to list Flash Attention kernels via Python. ${ret}") +endif() - file(STRINGS ${CMAKE_CURRENT_BINARY_DIR} / - flash_attention_fwd_blobs.txt FLASH_ATTENTION_FWD_GEN_BLOBS) +file(STRINGS ${CMAKE_CURRENT_BINARY_DIR}/flash_attention_fwd_blobs.txt FLASH_ATTENTION_FWD_GEN_BLOBS) - add_custom_command( - OUTPUT ${FLASH_ATTENTION_FWD_GEN_BLOBS} COMMAND ${Python3_EXECUTABLE} ${ - CMAKE_CURRENT_LIST_DIR} / - generate.py-- api $ { FLASH_ATTENTION_FWD_ENABLE_APIS } --working_path $ { - CMAKE_CURRENT_BINARY_DIR - } --gen_blobs) +add_custom_command( + OUTPUT ${FLASH_ATTENTION_FWD_GEN_BLOBS} + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/generate.py + --api ${FLASH_ATTENTION_FWD_ENABLE_APIS} + --working_path ${CMAKE_CURRENT_BINARY_DIR} + --gen_blobs +) - set(EXAMPLE_REDUCE "codegen_basic_flash_attention_fwd") message( - "adding example ${EXAMPLE_REDUCE}") +set(EXAMPLE_REDUCE "codegen_basic_flash_attention_fwd") +message("adding example ${EXAMPLE_REDUCE}") - add_executable(${EXAMPLE_REDUCE} EXCLUDE_FROM_ALL flash_attention_fwd.cpp) +add_executable(${EXAMPLE_REDUCE} + EXCLUDE_FROM_ALL + flash_attention_fwd.cpp +) - target_include_directories(${EXAMPLE_REDUCE} PRIVATE ${ - CMAKE_CURRENT_LIST_DIR}) +target_include_directories(${EXAMPLE_REDUCE} + PRIVATE + ${CMAKE_CURRENT_LIST_DIR} +) - target_sources(${EXAMPLE_REDUCE} PRIVATE ${ - FLASH_ATTENTION_FWD_GEN_BLOBS}) +target_sources(${EXAMPLE_REDUCE} PRIVATE ${FLASH_ATTENTION_FWD_GEN_BLOBS}) - message("FLASH_ATTENTION_FWD_GEN_BLOBS = " - "${FLASH_ATTENTION_FWD_GEN_BLOBS}") +message("FLASH_ATTENTION_FWD_GEN_BLOBS = ${FLASH_ATTENTION_FWD_GEN_BLOBS}") - set(EXAMPLE_REDUCE_COMPILE_OPTIONS) - list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS - Wno - - undefined - func - template - Wno - float - - equal-- offload - compress) - target_compile_options(${EXAMPLE_REDUCE} PRIVATE ${ - EXAMPLE_REDUCE_COMPILE_OPTIONS}) +set(EXAMPLE_REDUCE_COMPILE_OPTIONS) +list(APPEND EXAMPLE_REDUCE_COMPILE_OPTIONS + -Wno-undefined-func-template + -Wno-float-equal + --offload-compress +) - set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) \ No newline at end of file +target_compile_options(${EXAMPLE_REDUCE} + PRIVATE + ${EXAMPLE_REDUCE_COMPILE_OPTIONS} +) + +set_property(GLOBAL PROPERTY RULE_MESSAGES OFF) \ No newline at end of file diff --git a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/generate.py b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/generate.py index ac6c16c797..040133f8e9 100644 --- a/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/generate.py +++ b/example/ck_tile/99_toy_example/04_codegen_flash_attention_fwd/generate.py @@ -1,15 +1,23 @@ -#SPDX - License - Identifier : MIT -#Copyright(c) 2025, Advanced Micro Devices, Inc.All rights reserved. +# SPDX-License-Identifier: MIT +# Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. -import argparse from enum import IntEnum from pathlib import Path import sys from typing import List, Optional, Any import functools import itertools import copy from dataclasses import dataclass +import argparse +from enum import IntEnum +from pathlib import Path +import sys +from typing import List, Optional, Any +import functools +import itertools +import copy +from dataclasses import dataclass -#def get_if_str(idx, total, last_else = True): -#if idx == 0: -#return 'if' -#elif idx < total - 1: -#return 'else if' -#else: -#return 'else' if last_else else 'else if' +# def get_if_str(idx, total, last_else=True): +# if idx == 0: +# return 'if' +# elif idx < total - 1: +# return 'else if' +# else: +# return 'else' if last_else else 'else if' def get_if_str(size_, total, last_else=True): if size_ == "small": @@ -26,7 +34,7 @@ def BOOL_MAP(b_) -> str: class FlashAttentionFwdCodegen: API_TRAITS_DEFINE = """ - + template ; """ -#API_COMMON_HEADER = "" \ - " -#// SPDX-License-Identifier: MIT -#// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. +# API_COMMON_HEADER = """ +# // SPDX-License-Identifier: MIT +# // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. -# #include +# #include # #include "flash_attention_fwd.hpp" -# #include +# #include # #pragma once -#using S = ck_tile::stream_config; -#using A = FlashAttnArgs; +# using S = ck_tile::stream_config; +# using A = FlashAttnArgs; -#{F_traits_define } +# {F_traits_define} -#template < typename QDataType, -#typename KDataType, -#typename VDataType, -#typename ODataType, -#typename Traits_> -#float flash_attention_fwd_(const FlashAttnArgs& a, -#const ck_tile::stream_config& stream_config){{ -#using SaccDataType = typename Traits_::SaccDataType; -#using SMPLComputeDataType = typename Traits_::SMPLComputeDataType; -#using PDataType = typename Traits_::PDataType; -#using OaccDataType = typename Traits_::OaccDataType; +# template +# float flash_attention_fwd_(const FlashAttnArgs& a, +# const ck_tile::stream_config& stream_config) {{ +# using SaccDataType = typename Traits_::SaccDataType; +# using SMPLComputeDataType = typename Traits_::SMPLComputeDataType; +# using PDataType = typename Traits_::PDataType; +# using OaccDataType = typename Traits_::OaccDataType; -#index_t kGridSize = a.Batch * (a.M0 / Traits_::kM0PerBlock) * (a.N1 / Traits_::kN1PerBlock); +# index_t kGridSize = a.Batch * (a.M0 / Traits_::kM0PerBlock) * (a.N1 / Traits_::kN1PerBlock); -#if (stream_config.log_level_ > 0) -#std::cout << ", " << "FlashAttentionFwd<" << Traits_::kBlockSize << "," << Traits_::kHeadDim \ - << ">" << std::flush; +# if(stream_config.log_level_ > 0) +# std::cout << ", " << "FlashAttentionFwd<" << Traits_::kBlockSize << "," << Traits_::kHeadDim << ">" << std::flush; -#return ck_tile::launch_kernel(stream_config, -#ck_tile::make_kernel < Traits_::kBlockSize, Traits_::kBlockPerCu>( -#ck_tile::FlashAttentionFwd < QDataType, -#KDataType, -#VDataType, -#SaccDataType, -#SMPLComputeDataType, -#PDataType, -#OaccDataType, -#ODataType, -#Traits_::kBlockSize, -#Traits_::kHeadDim, -#Traits_::kM0PerBlock, -#Traits_::kN0PerBlock, -#Traits_::kK0PerBlock, -#Traits_::kN1PerBlock, -#Traits_::kK1PerBlock>{{} }, -#kGridSize, -#Traits_::kBlockSize, -# 0, -#a.q_ptr, -#a.k_ptr, -#a.v_ptr, -#a.o_ptr, -#a.M0, -#a.N0, -#a.K0, -#a.N1, -#a.Batch, -#a.strideQ, // StrideQ -#a.strideK, // StrideK -#a.strideV, // StrideV -#a.strideO, // StrideO -#a.batchStrideQ, // BatchStrideQ -#a.batchStrideK, // BatchStrideK -#a.batchStrideV, // BatchStrideV -#a.batchStrideO)); // BatchStrideO -#} } -#"" \ - " +# return ck_tile::launch_kernel(stream_config, +# ck_tile::make_kernel( +# ck_tile::FlashAttentionFwd{{}}, +# kGridSize, +# Traits_::kBlockSize, +# 0, +# a.q_ptr, +# a.k_ptr, +# a.v_ptr, +# a.o_ptr, +# a.M0, +# a.N0, +# a.K0, +# a.N1, +# a.Batch, +# a.strideQ, // StrideQ +# a.strideK, // StrideK +# a.strideV, // StrideV +# a.strideO, // StrideO +# a.batchStrideQ, // BatchStrideQ +# a.batchStrideK, // BatchStrideK +# a.batchStrideV, // BatchStrideV +# a.batchStrideO)); // BatchStrideO +# }} +# """ API_BASE = """ // SPDX-License-Identifier: MIT @@ -199,19 +204,14 @@ template float flash_attention_fwd && std::is_same_v && std::is_same_v && std::is_same_v) {{ -#{F_per_size_case } -#} } -#"" \ - " -#API_PER_SIZE_CASE = "" \ - " {F_if} {F_SIZE_COND} {{ -#{F_inner_dispatch } -#} } -#"" \ - " +# API_PER_DTYPE = """ {F_if}(std::is_same_v && std::is_same_v && std::is_same_v && std::is_same_v) {{ +# {F_per_size_case} +# }} +# """ +# API_PER_SIZE_CASE = """ {F_if} {F_SIZE_COND} {{ +# {F_inner_dispatch} +# }} +# """ API_INNER_CASE = """ {F_if} {F_VEC_COND} r = flash_attention_fwd_>(a, stream_config); """ @@ -224,7 +224,7 @@ template float flash_attention_fwd float flash_attention_fwd_(const FlashAttnArgs& a, const ck_tile::stream_config& stream_config) {{ - using SaccDataType = typename Traits_::SaccDataType; - using SMPLComputeDataType = typename Traits_::SMPLComputeDataType; - using PDataType = typename Traits_::PDataType; - using OaccDataType = typename Traits_::OaccDataType; - + using SaccDataType = typename Traits_::SaccDataType; + using SMPLComputeDataType = typename Traits_::SMPLComputeDataType; + using PDataType = typename Traits_::PDataType; + using OaccDataType = typename Traits_::OaccDataType; + index_t kGridSize = a.Batch * (a.M0 / Traits_::kM0PerBlock) * (a.N1 / Traits_::kN1PerBlock); if(stream_config.log_level_ > 0) @@ -430,10 +430,10 @@ float flash_attention_fwd_(const FlashAttnArgs str: -#Sort based on dtype + # Sort based on dtype t_dtype_dict = {} blobs = self.get_blobs(args) - + for blob in blobs: if blob.F_DataTypePair not in t_dtype_dict: t_dtype_dict[blob.F_DataTypePair] = {} @@ -445,16 +445,16 @@ float flash_attention_fwd_(const FlashAttnArgs None: w_p = Path(self.working_path) list_p = w_p / 'flash_attention_fwd_blobs.txt' blobs = self.get_blobs(args) - + with list_p.open('w') as list_f: -#API related files + # API related files list_f.write(str(w_p / (self.name_api + ".cpp")) + "\n") list_f.write(str(w_p / (self.name_common_header + ".hpp")) + "\n") -#Kernel instance files + # Kernel instance files for b in blobs: list_f.write(str(w_p / (b.name + ".cpp")) + "\n") @@ -557,7 +557,7 @@ float flash_attention_fwd_(const FlashAttnArgs