diff --git a/example/ck_tile/02_layernorm2d/CMakeLists.txt b/example/ck_tile/02_layernorm2d/CMakeLists.txt index 1bf74bc055..fa69ac0f7a 100644 --- a/example/ck_tile/02_layernorm2d/CMakeLists.txt +++ b/example/ck_tile/02_layernorm2d/CMakeLists.txt @@ -33,7 +33,7 @@ target_sources(${EXAMPLE_LAYERNORM2D_FWD} PRIVATE ${LAYERNORM2D_FWD_GEN_BLOBS}) set(EXAMPLE_LAYERNORM2D_FWD_COMPILE_OPTIONS) # NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations -list(APPEND EXAMPLE_LAYERNORM2D_FWD_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) +list(APPEND EXAMPLE_LAYERNORM2D_FWD_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal --offload-compress) target_compile_options(${EXAMPLE_LAYERNORM2D_FWD} PRIVATE ${EXAMPLE_LAYERNORM2D_FWD_COMPILE_OPTIONS}) diff --git a/example/ck_tile/02_layernorm2d/generate.py b/example/ck_tile/02_layernorm2d/generate.py index 6db52c93e1..700b007fad 100644 --- a/example/ck_tile/02_layernorm2d/generate.py +++ b/example/ck_tile/02_layernorm2d/generate.py @@ -39,7 +39,8 @@ def get_if_str(idx, total, lase_else = True): DATA_TYPE_MAP = {'fp32' : 'float', 'fp16' : 'ck_tile::fp16_t', 'bf16' : 'ck_tile::bf16_t', - 'int8' : 'ck_tile::int8_t'} + 'int8' : 'ck_tile::int8_t', + 'fp8' : 'ck_tile::fp8_t'} def BOOL_MAP(b_) -> str: if b_: @@ -504,12 +505,13 @@ def get_blobs(self, args): h_traits = layernorm_fwd_codegen.h_traits h_instance = layernorm_fwd_codegen.h_instance - dynamic_quant_out_dtype = ['int8'] + dynamic_quant_out_dtype = ['int8', 'fp8'] # some predefined support range # (prec_i,prec_o) for simplicity this string will be used as key for dict scale_list = [('fp32,fp32')] dtype_list = [('fp16,fp16'), ('bf16,bf16'), - ('fp16,int8'), ('bf16,int8')] # NOTE: only fused-dynamic-quant use int8 out + ('fp16,int8'), ('bf16,int8'), + ('fp16,fp8'), ('bf16,fp8')] # NOTE: only fused-dynamic-quant use int8 or fp8 out types_8bit = ('int8', 'fp8') types_16bit = ('int16', 'fp16', 'bf16') #fused_add_list = [0, 1, 2] diff --git a/example/ck_tile/02_layernorm2d/layernorm2d_fwd.cpp b/example/ck_tile/02_layernorm2d/layernorm2d_fwd.cpp index a5ec55299b..b72485222e 100644 --- a/example/ck_tile/02_layernorm2d/layernorm2d_fwd.cpp +++ b/example/ck_tile/02_layernorm2d/layernorm2d_fwd.cpp @@ -20,6 +20,14 @@ auto get_elimit() return ck_tile::make_tuple(rtol, atol); } +template <> +auto get_elimit() +{ + double rtol = 1e-2; + double atol = 1.0; + return ck_tile::make_tuple(rtol, atol); +} + auto create_args(int argc, char* argv[]) { ck_tile::ArgParser arg_parser; @@ -97,9 +105,11 @@ bool run(const ck_tile::ArgParser& arg_parser) int xbias = arg_parser.get_int("xbias"); int fused_add = arg_parser.get_int("fadd"); int fused_quant = arg_parser.get_int("fquant"); - if(fused_quant == 1 && prec_o != "int8") + if(fused_quant == 1 && prec_o != "int8" && prec_o != "fp8") { - std::cout << "if fused_quant is 1, only support \"-prec_o=int8\" case" << std::endl; + std::cout + << "if fused_quant is 1 or 2, only support \"-prec_o=int8\" or \"-prec_o=fp8\" cases." + << std::endl; return false; } @@ -291,7 +301,11 @@ bool run(const ck_tile::ArgParser& arg_parser) absmax = a > absmax ? a : absmax; } // printf("cpu:absmax:%f\n", absmax); - ComputeDataType y_scale = absmax / static_cast(127.0); + constexpr ComputeDataType kMaxY = + std::is_same::value ? 240.0 + : std::is_same::value ? 127.0 + : 0.0; + ComputeDataType y_scale = absmax / kMaxY; y_scale_host_ref(m_) = ck_tile::type_convert(y_scale); for(int n_ = 0; n_ < N_; n_++) { @@ -334,7 +348,7 @@ bool run(const ck_tile::ArgParser& arg_parser) y_residual_buf.FromDevice(y_residual_host_dev.data()); } - auto [rtol, atol] = get_elimit(); + auto [rtol, atol] = get_elimit(); if(x_stride == n) { @@ -452,6 +466,16 @@ int main(int argc, char* argv[]) { return run(arg_parser) ? 0 : -2; } + else if(prec_i == "fp16" && prec_o == "fp8" && prec_sm == "fp32" && prec_sy == "fp32" && + !save_mv) + { + return run(arg_parser) ? 0 : -2; + } + else if(prec_i == "bf16" && prec_o == "fp8" && prec_sm == "fp32" && prec_sy == "fp32" && + !save_mv) + { + return run(arg_parser) ? 0 : -2; + } return -3; } diff --git a/example/ck_tile/02_layernorm2d/script/smoke_test.sh b/example/ck_tile/02_layernorm2d/script/smoke_test.sh index 3f5c3eb134..ceaf262bd9 100755 --- a/example/ck_tile/02_layernorm2d/script/smoke_test.sh +++ b/example/ck_tile/02_layernorm2d/script/smoke_test.sh @@ -1,7 +1,7 @@ #!/bin/sh EXE="$(find . -name tile_example_layernorm2d_fwd -type f | head -n 1)" -for fquant in "" "-fquant=1 -prec_o=int8"; do +for fquant in "" "-fquant=1 -prec_o=int8" "-fquant=1 -prec_o=fp8"; do for pr_i in "fp16" "bf16" ; do for fadd in "0" "1"; do $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=99 -n=13 diff --git a/example/ck_tile/10_rmsnorm2d/CMakeLists.txt b/example/ck_tile/10_rmsnorm2d/CMakeLists.txt index df7b5ff11a..5684c9b2e0 100644 --- a/example/ck_tile/10_rmsnorm2d/CMakeLists.txt +++ b/example/ck_tile/10_rmsnorm2d/CMakeLists.txt @@ -33,7 +33,7 @@ target_sources(${TILE_RMSNORM2D_FWD} PRIVATE ${RMSNORM2D_FWD_GEN_BLOBS}) set(TILE_RMSNORM2D_FWD_COMPILE_OPTIONS) # NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations -list(APPEND TILE_RMSNORM2D_FWD_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) +list(APPEND TILE_RMSNORM2D_FWD_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal --offload-compress) target_compile_options(${TILE_RMSNORM2D_FWD} PRIVATE ${TILE_RMSNORM2D_FWD_COMPILE_OPTIONS}) diff --git a/example/ck_tile/10_rmsnorm2d/generate.py b/example/ck_tile/10_rmsnorm2d/generate.py index ce1243c52d..dadb2268b2 100644 --- a/example/ck_tile/10_rmsnorm2d/generate.py +++ b/example/ck_tile/10_rmsnorm2d/generate.py @@ -37,7 +37,8 @@ def get_if_str(idx, total, lase_else = True): DATA_TYPE_MAP = {'fp32' : 'float', 'fp16' : 'ck_tile::fp16_t', 'bf16' : 'ck_tile::bf16_t', - 'int8' : 'ck_tile::int8_t'} + 'int8' : 'ck_tile::int8_t', + 'fp8' : 'ck_tile::fp8_t'} def BOOL_MAP(b_) -> str: if b_: @@ -477,12 +478,13 @@ def get_blobs(self): h_traits = rmsnorm_fwd_codegen.h_traits h_instance = rmsnorm_fwd_codegen.h_instance - dynamic_quant_out_dtype = ['int8'] + dynamic_quant_out_dtype = ['int8', 'fp8'] # some predefined support range # (prec_i,prec_o) for simplicity this string will be used as key for dict scale_list = [('fp32,fp32')] dtype_list = [('fp16,fp16'), ('bf16,bf16'), - ('fp16,int8'), ('bf16,int8')] # NOTE: only fused-dynamic-quant use int8 out + ('fp16,int8'), ('bf16,int8'), + ('fp16,fp8'), ('bf16,fp8')] # NOTE: only fused-dynamic-quant use int8 out #fused_add_list = [0, 1, 2] #fused_sweep_list = [0, 1, 2] # NOTE: only single pass can use fused (smooth) dynamic quant fused_add_list = [0, 1] diff --git a/example/ck_tile/10_rmsnorm2d/rmsnorm2d_fwd.cpp b/example/ck_tile/10_rmsnorm2d/rmsnorm2d_fwd.cpp index 63890a8a47..cdee6dfb80 100644 --- a/example/ck_tile/10_rmsnorm2d/rmsnorm2d_fwd.cpp +++ b/example/ck_tile/10_rmsnorm2d/rmsnorm2d_fwd.cpp @@ -105,9 +105,11 @@ bool run(const ck_tile::ArgParser& arg_parser) prec_sy = "fp32"; } - if((fused_quant == 1 || fused_quant == 2) && prec_o != "int8") + if((fused_quant == 1 || fused_quant == 2) && prec_o != "int8" && prec_o != "fp8") { - std::cout << "if fused_quant is 1, only support \"-prec_o=int8\" case" << std::endl; + std::cout + << "if fused_quant is 1 or 2, only support \"-prec_o=int8\" or \"-prec_o=fp8\" cases." + << std::endl; return false; } @@ -248,7 +250,11 @@ bool run(const ck_tile::ArgParser& arg_parser) absmax = a > absmax ? a : absmax; } // printf("cpu:absmax:%f\n", absmax); - ComputeDataType y_scale = absmax / static_cast(127.0); + constexpr ComputeDataType kMaxY = + std::is_same::value ? 240.0 + : std::is_same::value ? 127.0 + : 0.0; + ComputeDataType y_scale = absmax / kMaxY; y_scale_host_ref(m_) = ck_tile::type_convert(y_scale); for(int n_ = 0; n_ < N_; n_++) { @@ -400,6 +406,16 @@ int main(int argc, char* argv[]) { return run(arg_parser) ? 0 : -2; } + else if(prec_i == "fp16" && prec_o == "fp8" && prec_sm == "fp32" && prec_sy == "fp32" && + !save_rms) + { + return run(arg_parser) ? 0 : -2; + } + else if(prec_i == "bf16" && prec_o == "fp8" && prec_sm == "fp32" && prec_sy == "fp32" && + !save_rms) + { + return run(arg_parser) ? 0 : -2; + } return -3; } diff --git a/example/ck_tile/10_rmsnorm2d/script/smoke_test.sh b/example/ck_tile/10_rmsnorm2d/script/smoke_test.sh index e5e2bd1184..ab890738b3 100755 --- a/example/ck_tile/10_rmsnorm2d/script/smoke_test.sh +++ b/example/ck_tile/10_rmsnorm2d/script/smoke_test.sh @@ -1,7 +1,7 @@ #!/bin/sh EXE="$(find . -name tile_rmsnorm2d_fwd -type f | head -n 1)" -for fquant in "" "-fquant=1 -prec_o=int8" "-fquant=2 -prec_o=int8"; do +for fquant in "" "-fquant=1 -prec_o=int8" "-fquant=2 -prec_o=int8" "-fquant=1 -prec_o=fp8" "-fquant=2 -prec_o=fp8"; do for pr_i in "fp16" "bf16" ; do for fadd in "0" "1"; do $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=99 -n=13 @@ -27,7 +27,7 @@ $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=7 -n=2734 $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=1 -n=3182 $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=9 -n=4096 $EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=3 -n=8192 -#$EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=1 -n=10547 +$EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=1 -n=10547 #$EXE -prec_i=$pr_i -fadd=$fadd $fquant -m=3 -n=17134 done done diff --git a/include/ck_tile/host/check_err.hpp b/include/ck_tile/host/check_err.hpp index 529bfdff25..08b8e39635 100644 --- a/include/ck_tile/host/check_err.hpp +++ b/include/ck_tile/host/check_err.hpp @@ -1,5 +1,5 @@ // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once @@ -337,7 +337,11 @@ std::enable_if_t<(std::is_same_v, ranges::range_val } if(!res) { - std::cerr << std::setw(12) << std::setprecision(7) << "max err: " << max_err << std::endl; + const float error_percent = + static_cast(err_count) / static_cast(out.size()) * 100.f; + std::cerr << "max err: " << max_err; + std::cerr << ", number of errors: " << err_count; + std::cerr << ", " << error_percent << "% wrong values" << std::endl; } return res; }