diff --git a/.gitignore b/.gitignore index bcc5888b7f..6641e5bc58 100644 --- a/.gitignore +++ b/.gitignore @@ -36,6 +36,9 @@ tags # Editors .vscode +# Cline +.cline* + # build-in-source directory (see exceptions below) build* diff --git a/Dockerfile.aiter b/Dockerfile.aiter index b61c1e41a5..dab3f9588d 100644 --- a/Dockerfile.aiter +++ b/Dockerfile.aiter @@ -1,9 +1,8 @@ -ARG BASE_DOCKER="rocm/composable_kernel-private:ck_aiter_base" +ARG BASE_DOCKER="rocm/pytorch:latest" FROM $BASE_DOCKER ARG AITER_BRANCH="main" ARG CK_AITER_BRANCH="develop" -RUN groupadd irc && \ - pip install pandas zmq einops && \ +RUN pip install pandas zmq einops ninja && \ pip install numpy==1.26.2 && \ sudo mkdir /home/jenkins && \ sudo mkdir /home/jenkins/workspace && \ @@ -14,6 +13,8 @@ RUN groupadd irc && \ rm -rf 3rdparty/composable_kernel/ && \ git clone -b "$CK_AITER_BRANCH" https://github.com/ROCm/composable_kernel.git 3rdparty/composable_kernel/ && \ python3 setup.py develop && \ - chown -R jenkins:jenkins /home/jenkins/workspace && \ - chmod -R a+rwx /home/jenkins/workspace && \ + groupadd -g 1001 jenkins && \ + useradd -u 1001 -g 1001 -m -s /bin/bash jenkins && \ + chown -R jenkins:jenkins /home/jenkins && \ + chmod -R a+rwx /home/jenkins && \ sudo usermod -aG irc jenkins diff --git a/Jenkinsfile b/Jenkinsfile index 7a8574df05..c642e2d3b1 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -12,6 +12,14 @@ def show_node_info() { """ } +// Error patterns to scan build logs for specific failure types and send detailed notifications. +def failurePatterns = [ + [pattern: /login attempt to .* failed with status: 401 Unauthorized/, description: "Docker registry authentication failed"], + [pattern: /docker login failed/, description: "Docker login failed"], + [pattern: /HTTP request sent .* 404 Not Found/, description: "HTTP request failed with 404"], + [pattern: /cat: .* No such file or directory/, description: "GPU not found"], +] + class Version { int major, minor, patch @Override @@ -1488,7 +1496,7 @@ pipeline { -D GEMM_LAYOUT="rcr;rrr;crr;ccr" \ -D GEMM_MULTI_D_DATATYPE="fp16" \ -D GEMM_MULTI_D_LAYOUT="rcrr;rrrr;crrr;ccrr" \ - -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8" \ + -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8;bf16;bf8" \ -D GEMM_PRESHUFFLE_LAYOUT="rcr" \ -DCMAKE_CXX_FLAGS=" -O3 " .. && \ ninja -j64 benchmark_gemm_all && \ @@ -1528,7 +1536,7 @@ pipeline { -D GEMM_LAYOUT="rcr;rrr;crr;ccr" \ -D GEMM_MULTI_D_DATATYPE="fp16" \ -D GEMM_MULTI_D_LAYOUT="rcrr;rrrr;crrr;ccrr" \ - -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8" \ + -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8;bf16;bf8" \ -D GEMM_PRESHUFFLE_LAYOUT="rcr" \ -DCMAKE_CXX_FLAGS=" -O3 " .. && \ ninja -j64 benchmark_gemm_all && \ @@ -1566,15 +1574,10 @@ pipeline { -D GPU_TARGETS="gfx1201" \ -D GEMM_DATATYPE="fp16" \ -D GEMM_LAYOUT="rcr;rrr;crr;ccr" \ - -DGEMM_CONFIG_FILE=gfx120x_config.json \ -DCMAKE_CXX_FLAGS=" -O3 " .. && \ ninja -j64 benchmark_gemm_all && \ python3 ../tile_engine/ops/gemm/gemm_benchmark.py . --problem-sizes "1024,1024,1024" \ - --warmup 5 --repeat 5 --verbose --json results.json && \ - ninja -j64 benchmark_gemm_fp16_rcr && \ - ninja -j64 benchmark_gemm_fp16_rrr && \ - ninja -j64 benchmark_gemm_fp16_crr && \ - ninja -j64 benchmark_gemm_fp16_ccr """ + --warmup 5 --repeat 5 --verbose --json results.json """ } steps{ buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args) @@ -1853,4 +1856,36 @@ pipeline { } } } + post { + failure { + node(rocmnode("nogpu")) { + script { + // Get the build log. + def buildLog = sh(script: 'wget -q --no-check-certificate -O - ' + BUILD_URL + 'consoleText', returnStdout: true) + // Check for patterns in the log. + def foundPatterns = [] + for (patternMap in failurePatterns) { + def result = checkForPattern(patternMap.pattern, buildLog) + if (result.found) { + foundPatterns.add([ + description: patternMap.description, + matchedLine: result.matchedLine, + context: result.context + ]) + } + } + // Send a notification for each matched failure pattern. + for (patternMap in foundPatterns) { + withCredentials([string(credentialsId: 'ck_ci_errors_webhook_url', variable: 'WEBHOOK_URL')]) { + sh ''' + curl -X POST "${WEBHOOK_URL}" \ + -H 'Content-Type: application/json' \ + -d '{"text": "\\n\\n**Build Failed**\\n\\n**Issues detected:** ''' + patternMap.description + '''\\n\\n**Log context:**\\n```\\n''' + patternMap.context.replace("'", "\\'") + '''\\n```\\n\\n**Job:** ''' + env.JOB_NAME + '''\\n\\n**Build:** #''' + env.BUILD_NUMBER + '''\\n\\n**URL:** ''' + env.RUN_DISPLAY_URL + '''"}' + ''' + } + } + } + } + } + } } diff --git a/codegen/test/rtc/include/rtc/kernel.hpp b/codegen/test/rtc/include/rtc/kernel.hpp index b1ee729f77..96337fe2c1 100644 --- a/codegen/test/rtc/include/rtc/kernel.hpp +++ b/codegen/test/rtc/include/rtc/kernel.hpp @@ -52,7 +52,7 @@ struct kernel template auto launch(hipStream_t stream, std::size_t global, std::size_t local, Ts... zs) const { - return [=](auto&&... xs) { + return [=, this](auto&&... xs) { launch(stream, global, local, std::vector{xs...}, zs...); }; } diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp b/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp index 10d7befc06..57d3f224d8 100644 --- a/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp +++ b/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp @@ -182,6 +182,7 @@ struct GemmConfigPreshuffleDecode : public GemmConfigBase static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Default; static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_PRESHUFFLE_V2; static constexpr bool Preshuffle = true; + static constexpr bool Persistent = true; static constexpr bool DoubleSmemBuffer = true; }; diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm_preshuffle.cpp b/example/ck_tile/17_grouped_gemm/grouped_gemm_preshuffle.cpp index b9d6a4a1bc..52b84737cc 100644 --- a/example/ck_tile/17_grouped_gemm/grouped_gemm_preshuffle.cpp +++ b/example/ck_tile/17_grouped_gemm/grouped_gemm_preshuffle.cpp @@ -167,6 +167,113 @@ float grouped_gemm(const std::vector& gemm_descs, return ave_time = BaseGemmPipeline::TailHandler(RunSplitk, has_hot_loop, tail_num); } +template +float grouped_gemm_tileloop(const ck_tile::stream_config& s, + const ck_tile::index_t num_groups, + void* kargs_ptr, + bool splitk) +{ + using GemmShape = ck_tile::TileGemmShape< + ck_tile::sequence, + ck_tile::sequence, + ck_tile:: + sequence>; + using TilePartitioner = + ck_tile::GemmSpatiallyLocalTilePartitioner; + + using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits; + + float ave_time{0}; + + const auto Run = [&](const auto memory_operation_) { + constexpr auto scheduler = GemmConfig::Scheduler; + constexpr auto memory_operation = memory_operation_.value; + + using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem; + + using GemmPipeline = typename PipelineTypeTraits< + GemmConfig::Pipeline>::template GemmPipeline; + using GemmEpilogue = ck_tile::CShuffleEpilogue, // DsDataType (empty for no D tensors) + AccDataType, + CDataType, + ck_tile::tuple<>, // DsLayout (empty for no D tensors) + CLayout, + ck_tile::element_wise::PassThrough, + TilePartitioner::MPerBlock, + TilePartitioner::NPerBlock, + GemmConfig::M_Warp, + GemmConfig::N_Warp, + GemmConfig::M_Warp_Tile, + GemmConfig::N_Warp_Tile, + GemmConfig::K_Warp_Tile, + UniversalGemmProblem::TransposeC, + memory_operation>>; + using Kernel = ck_tile::GroupedGemmKernel; + const dim3 blocks = Kernel::BlockSize(); + const dim3 grids = Kernel::MaxOccupancyGridSize(s); + + if(s.log_level_ > 0) + { + std::cout << "Launching kernel: " << Kernel::GetName() << " with args:" << " grid: {" + << grids.x << ", " << grids.y << ", " << grids.z << "}" << ", blocks: {" + << blocks.x << ", " << blocks.y << ", " << blocks.z << "}" << std::endl; + } + + ave_time = + ck_tile::launch_kernel(s, + ck_tile::make_kernel( + Kernel{}, + grids, + blocks, + 0, + ck_tile::cast_pointer_to_constant_address_space(kargs_ptr), + num_groups)); + + return ave_time; + }; + + if(splitk) + { + Run(ck_tile::integral_constant{}); + } + else + { + Run(ck_tile::integral_constant{}); + } + + return ave_time; +} + #include "run_grouped_gemm_example.inc" template diff --git a/example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp b/example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp index 3b4258d8b1..1a913fcfc1 100644 --- a/example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp +++ b/example/ck_tile/17_grouped_gemm/quant_grouped_gemm.cpp @@ -29,7 +29,7 @@ template + ck_tile::QuantType QuantMode = ck_tile::QuantType::BQuantGrouped> float grouped_gemm_tileloop(const ck_tile::stream_config& s, const ck_tile::index_t num_groups, void* kargs_ptr) @@ -48,8 +48,8 @@ float grouped_gemm_tileloop(const ck_tile::stream_config& s, using GemmUniversalTraits = ck_tile::TileGemmQuantTraits; + using QuantGemmProblem = typename std::conditional< + QuantMode == ck_tile::QuantType::BQuantGrouped, + ck_tile::GemmBQuantPipelineProblem, // QuantGroupSize + ck_tile::GemmRowColTensorQuantPipelineProblem>::type; - using GemmPipeline = typename PipelineTypeTraits< - GemmConfig::Pipeline>::template GemmPipeline; + using GemmPipeline = + typename std::conditional, + ck_tile::GemmPipelineAgBgCrCompV3>::type; using GemmEpilogue = ck_tile::CShuffleEpilogue< ck_tile::CShuffleEpilogueProblem constexpr ck_tile::index_t get_k_warp_tile() @@ -41,6 +42,14 @@ struct GemmTypeConfig using AccDataType = float; using CDataType = ck_tile::half_t; }; +template <> +struct GemmTypeConfig +{ + using ADataType = ck_tile::bf8_t; + using BDataType = ck_tile::bf8_t; + using AccDataType = float; + using CDataType = ck_tile::half_t; +}; struct GemmConfigBase { @@ -77,24 +86,11 @@ struct GemmConfigComputeV3_2 : public GemmConfigBase static constexpr ck_tile::index_t N_Warp_Tile = 32; static constexpr ck_tile::index_t K_Warp_Tile = get_k_warp_tile(); - static constexpr bool DoubleSmemBuffer = false; - static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_COMPUTE_V3; + static constexpr bool DoubleSmemBuffer = false; static constexpr int kBlockPerCu = 1; }; -template -struct PipelineTypeTraits; - -template <> -struct PipelineTypeTraits -{ - template - using GemmPipeline = ck_tile::GemmPipelineAgBgCrCompV3; - template - using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAgBgCrCompV3; -}; - using grouped_gemm_kargs = ck_tile::QuantGroupedGemmHostArgs; auto create_args(int argc, char* argv[]) @@ -122,8 +118,7 @@ auto create_args(int argc, char* argv[]) .insert("repeat", "100", "number of iterations to benchmark the kernel.") .insert("group_count", "8", "group count.") .insert("kbatch", "1", "kbatch for SplitK") - .insert("quant_mode", "tensor", "Choose tensor (default), or rowcol"); - ; + .insert("quant_mode", "bquant", "Choose bquant (default), tensor, or rowcol"); bool result = arg_parser.parse(argc, argv); return std::make_tuple(result, arg_parser); diff --git a/example/ck_tile/17_grouped_gemm/quant_run_grouped_gemm_example.inc b/example/ck_tile/17_grouped_gemm/quant_run_grouped_gemm_example.inc index 19211ed494..152df38bff 100644 --- a/example/ck_tile/17_grouped_gemm/quant_run_grouped_gemm_example.inc +++ b/example/ck_tile/17_grouped_gemm/quant_run_grouped_gemm_example.inc @@ -43,8 +43,8 @@ template + ck_tile::QuantType QuantMode = ck_tile::QuantType::BQuantGrouped, + typename CDEElementWise = ck_tile::element_wise::PassThrough> float invoke_gemm(int n_warmup, int n_repeat, int group_count, @@ -159,11 +159,12 @@ int run_grouped_gemm_example_with_layouts(int argc, return group_count != 0 && ((args.size() == static_cast(group_count)) && ...); }; - const int group_count = arg_parser.get_int("group_count"); - const int repeat = arg_parser.get_int("repeat"); - const int warmup = arg_parser.get_int("warmup"); - const int kbatch = arg_parser.get_int("kbatch"); - bool validate = arg_parser.get_bool("validate"); + const int group_count = arg_parser.get_int("group_count"); + const int repeat = arg_parser.get_int("repeat"); + const int warmup = arg_parser.get_int("warmup"); + const int kbatch = arg_parser.get_int("kbatch"); + bool validate = arg_parser.get_bool("validate"); + const ck_tile::index_t QuantGroupSize = 128; if(kbatch > 1 && validate && warmup + repeat > 1) { @@ -172,9 +173,11 @@ int run_grouped_gemm_example_with_layouts(int argc, validate = false; } - std::vector Ms = arg_parser.get_int_vec("Ms"); - std::vector Ns = arg_parser.get_int_vec("Ns"); - std::vector Ks = arg_parser.get_int_vec("Ks"); + std::vector Ms = arg_parser.get_int_vec("Ms"); + std::vector Ns = arg_parser.get_int_vec("Ns"); + std::vector Ks = arg_parser.get_int_vec("Ks"); + std::vector AQs; // dimension of AQ tensor is calculated from A tensor + std::vector BQs; // dimension of BQ tensor is calculated from B tensor std::vector stride_As = arg_parser.get_int_vec("stride_As"); std::vector stride_Bs = arg_parser.get_int_vec("stride_Bs"); std::vector stride_Cs = arg_parser.get_int_vec("stride_Cs"); @@ -252,6 +255,15 @@ int run_grouped_gemm_example_with_layouts(int argc, AQK = 1; // Row quantization: tensor shape [M, 1] or [1] BQK = 1; // Column quantization: tensor shape [1, N] or [1] } + else if constexpr(QuantMode == ck_tile::QuantType::BQuantGrouped) + { + AQK = 0; // No A quantization + BQK = K / QuantGroupSize; // Group quantization: BQK = K / GroupSize + if(K % QuantGroupSize != 0) + { + throw std::runtime_error("K must be divisible by 128 for BQuantGrouped mode"); + } + } stride_As[i] = ck_tile::get_default_stride(M, K, stride_As[i], is_row_major(a_layout)); stride_Bs[i] = ck_tile::get_default_stride(K, N, stride_Bs[i], is_row_major(b_layout)); @@ -289,6 +301,13 @@ int run_grouped_gemm_example_with_layouts(int argc, bq_tensors.push_back(ck_tile::HostTensor( ck_tile::host_tensor_descriptor(1, 1, stride_BQs[i], is_row_major(bq_layout)))); } + else if constexpr(QuantMode == ck_tile::QuantType::BQuantGrouped) + { + aq_tensors.push_back(ck_tile::HostTensor( + ck_tile::host_tensor_descriptor(0, AQK, stride_AQs[i], is_row_major(aq_layout)))); + bq_tensors.push_back(ck_tile::HostTensor( + ck_tile::host_tensor_descriptor(BQK, N, stride_BQs[i], is_row_major(bq_layout)))); + } std::cout << "gemm[" << i << "]" << " a_m_k: " << a_m_k_tensors[i].mDesc << " b_k_n: " << b_k_n_tensors[i].mDesc << " c_m_n: " << c_m_n_tensors[i].mDesc @@ -394,6 +413,17 @@ int run_grouped_gemm_example_with_layouts(int argc, bq_tensors[i], c_m_n_host_ref); } + else if constexpr(QuantMode == ck_tile::QuantType::BQuantGrouped) + { + ck_tile::reference_gemm_quant( + a_m_k_tensors[i], bq_tensors[i], b_k_n_tensors[i], c_m_n_host_ref); + } const float max_accumulated_value = *std::max_element(c_m_n_host_ref.mData.begin(), c_m_n_host_ref.mData.end()); @@ -441,42 +471,6 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a QuantMode>( argc, argv, Row{}, Row{}, Col{}, Col{}, Row{}); } - else if(a_layout == "R" && b_layout == "R") - { - return run_grouped_gemm_example_with_layouts( - argc, argv, Row{}, Row{}, Row{}, Col{}, Row{}); - } - else if(a_layout == "C" && b_layout == "R") - { - return run_grouped_gemm_example_with_layouts( - argc, argv, Row{}, Row{}, Col{}, Col{}, Row{}); - } - else if(a_layout == "C" && b_layout == "C") - { - return run_grouped_gemm_example_with_layouts( - argc, argv, Col{}, Col{}, Col{}, Col{}, Row{}); - } else { throw std::runtime_error("Unsupported data layout configuration for A,B and C tensors!"); @@ -513,6 +507,41 @@ int run_grouped_gemm_example(int argc, char* argv[]) ck_tile::QuantType::RowColQuant>( a_layout, b_layout, argc, argv); } + else if(quant_mode == "bquant") + { + return run_gemm_example_prec_type, + ck_tile::fp8_t, + ck_tile::QuantType::BQuantGrouped>( + a_layout, b_layout, argc, argv); + } + else + { + throw std::runtime_error("Unsupported quantization mode!"); + } + } + if(data_type == "bf8") + { + if(quant_mode == "tensor") + { + return run_gemm_example_prec_type, + ck_tile::bf8_t, + ck_tile::QuantType::TensorQuant>( + a_layout, b_layout, argc, argv); + } + else if(quant_mode == "rowcol") + { + return run_gemm_example_prec_type, + ck_tile::bf8_t, + ck_tile::QuantType::RowColQuant>( + a_layout, b_layout, argc, argv); + } + else if(quant_mode == "bquant") + { + return run_gemm_example_prec_type, + ck_tile::bf8_t, + ck_tile::QuantType::BQuantGrouped>( + a_layout, b_layout, argc, argv); + } else { throw std::runtime_error("Unsupported quantization mode!"); diff --git a/example/ck_tile/17_grouped_gemm/run_grouped_gemm_example.inc b/example/ck_tile/17_grouped_gemm/run_grouped_gemm_example.inc index dbdbe80c5d..4eee165d66 100644 --- a/example/ck_tile/17_grouped_gemm/run_grouped_gemm_example.inc +++ b/example/ck_tile/17_grouped_gemm/run_grouped_gemm_example.inc @@ -70,23 +70,13 @@ float invoke_gemm(int n_warmup, } else { - if(GemmConfig::Preshuffle) - { - // not supported yet - throw std::runtime_error( - "Persistent grouped gemm with preshuffle is not supported yet"); - } - - // NOTE: With the persistent TileLoop kernel, we do not necessarily need to haveCollapse - // commentComment on line L74tenpercent commented on Sep 5, 2025 tenpercenton Sep 5, - // 2025ContributorMore actionsdid you intend to remove the comment?Write a replyResolve - // commentCode has comments. Press enter to view. the gemm problems known on the host. - // Instead, we can just pass the pointer to the kernel and let the workgroups figure out - // which tiles to work on. This is useful when the gemm problems are generated dynamically. - // In this example however, we generate the `kargs` using the known gemm_descs, - // and copy the gemm descriptions to the device memory. - // The contents of the memory pointed to by `kargs_ptr` pointer could be - // written by e.g. another kernel from earlier stage. + // NOTE: With the persistent TileLoop kernel, we do not necessarily need to have the gemm + // problems known on the host. Instead, we can just pass the pointer to the kernel and let + // the workgroups figure out which tiles to work on. This is useful when the gemm problems + // are generated dynamically. In this example however, we generate the `kargs` using the + // known gemm_descs, and copy the gemm descriptions to the device memory. The contents of + // the memory pointed to by `kargs_ptr` pointer could be written by e.g. another kernel from + // earlier stage. std::vector> kargs; void* kargs_ptr = gemm_workspace.GetDeviceBuffer(); diff --git a/example/ck_tile/20_grouped_convolution/CMakeLists.txt b/example/ck_tile/20_grouped_convolution/CMakeLists.txt index 10332137e2..e9614061e1 100644 --- a/example/ck_tile/20_grouped_convolution/CMakeLists.txt +++ b/example/ck_tile/20_grouped_convolution/CMakeLists.txt @@ -4,6 +4,9 @@ list(APPEND EXAMPLE_CONV_COMPILE_OPTIONS -mllvm -enable-noalias-to-md-conversion add_executable(tile_example_grouped_conv_fwd EXCLUDE_FROM_ALL grouped_convolution_forward.cpp) target_compile_options(tile_example_grouped_conv_fwd PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS}) +add_executable(tile_example_grouped_conv_fwd_bias_clamp EXCLUDE_FROM_ALL grouped_convolution_forward_bias_clamp.cpp) +target_compile_options(tile_example_grouped_conv_fwd_bias_clamp PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS}) + add_executable(tile_example_grouped_conv_bwd_weight EXCLUDE_FROM_ALL grouped_convolution_backward_weight.cpp) target_compile_options(tile_example_grouped_conv_bwd_weight PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS}) diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp new file mode 100644 index 0000000000..ed215cb178 --- /dev/null +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_forward_bias_clamp.cpp @@ -0,0 +1,58 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include + +#include +#include +#include +#include +#include + +#include "ck_tile/host.hpp" +#include "grouped_convolution_utils.hpp" +#include "grouped_convolution_forward_invoker.hpp" +#include "run_grouped_convolution_fwd_bias_clamp_example.inc" + +template