From cb629a747f105cdef31672786000fdd372739c6c Mon Sep 17 00:00:00 2001 From: Aviral Goel Date: Thu, 11 Dec 2025 20:25:29 +0400 Subject: [PATCH] chore: update copyright header for misc files (#3402) * chore: update copyright header for misc files * fix: typo in kernel resulting in ci failure [ROCm/composable_kernel commit: 4dcc3e59c1c0195dae7ee9da9ab76d18a4cafe9f] --- docs/conceptual/ck_tile/convert_mermaid_to_svg.py | 3 +++ docs/conceptual/ck_tile/convert_raw_html_to_commented.py | 3 +++ docs/conceptual/ck_tile/update_diagrams.py | 3 +++ example/test_old_ck_gpu_reference.cpp | 2 +- experimental/builder/test/test_ckb_conv_builder.cpp | 2 ++ include/ck_tile/ref/conv_common.hpp | 2 +- include/ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp | 2 +- include/ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp | 2 +- include/ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp | 2 +- .../device_grouped_gemm_wmma_splitk_instance.hpp | 2 +- test/ck_tile/gemm_block_scale/test_gemm_quant_aquant.cpp | 2 +- test/ck_tile/gemm_block_scale/test_gemm_quant_bquant.cpp | 2 +- .../test_gemm_quant_bquant_preshuffle.cpp | 2 +- test/ck_tile/gemm_block_scale/test_gemm_quant_rowcol.cpp | 2 +- test/ck_tile/gemm_block_scale/test_gemm_quant_tensor.cpp | 2 +- test/ck_tile/utility/test_fill.cpp | 2 +- test/ck_tile/warp_gemm/CMakeLists.txt | 3 +++ test/ck_tile/warp_gemm/test_f32_16x16x128_fp4.cpp | 2 +- .../practice_gemm_host_pipeline_agmem_bgmem_creg.hpp | 8 ++++---- 19 files changed, 31 insertions(+), 17 deletions(-) diff --git a/docs/conceptual/ck_tile/convert_mermaid_to_svg.py b/docs/conceptual/ck_tile/convert_mermaid_to_svg.py index 1d62405e53..2bfaffdb57 100644 --- a/docs/conceptual/ck_tile/convert_mermaid_to_svg.py +++ b/docs/conceptual/ck_tile/convert_mermaid_to_svg.py @@ -1,4 +1,7 @@ #!/usr/bin/env python3 +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + """ Script to convert all mermaid diagrams in CK Tile docs to SVGs. This script: diff --git a/docs/conceptual/ck_tile/convert_raw_html_to_commented.py b/docs/conceptual/ck_tile/convert_raw_html_to_commented.py index e90bf9def0..8e4a849e7f 100644 --- a/docs/conceptual/ck_tile/convert_raw_html_to_commented.py +++ b/docs/conceptual/ck_tile/convert_raw_html_to_commented.py @@ -1,4 +1,7 @@ #!/usr/bin/env python3 +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + """Convert raw HTML mermaid blocks to commented format for SVG conversion.""" import os diff --git a/docs/conceptual/ck_tile/update_diagrams.py b/docs/conceptual/ck_tile/update_diagrams.py index 2fbe2ef5a9..f78599010e 100644 --- a/docs/conceptual/ck_tile/update_diagrams.py +++ b/docs/conceptual/ck_tile/update_diagrams.py @@ -1,4 +1,7 @@ #!/usr/bin/env python3 +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + """ Helper script to update SVG diagrams from commented mermaid sources in RST files. diff --git a/example/test_old_ck_gpu_reference.cpp b/example/test_old_ck_gpu_reference.cpp index 0bcf43d20b..9f12eaea4d 100644 --- a/example/test_old_ck_gpu_reference.cpp +++ b/example/test_old_ck_gpu_reference.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. // Standalone test program for Old CK GPU references // Tests naive_conv_fwd (existing) and future backward ops diff --git a/experimental/builder/test/test_ckb_conv_builder.cpp b/experimental/builder/test/test_ckb_conv_builder.cpp index e69de29bb2..81e63887c1 100644 --- a/experimental/builder/test/test_ckb_conv_builder.cpp +++ b/experimental/builder/test/test_ckb_conv_builder.cpp @@ -0,0 +1,2 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT diff --git a/include/ck_tile/ref/conv_common.hpp b/include/ck_tile/ref/conv_common.hpp index ed43e87b14..50ae18eb99 100644 --- a/include/ck_tile/ref/conv_common.hpp +++ b/include/ck_tile/ref/conv_common.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/include/ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp b/include/ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp index a5f6a697f2..f75bdda912 100644 --- a/include/ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp +++ b/include/ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/include/ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp b/include/ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp index 2ac9c19892..0839074dd4 100644 --- a/include/ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp +++ b/include/ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/include/ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp b/include/ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp index 720fa40297..f582fcd71a 100644 --- a/include/ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp +++ b/include/ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/library/include/ck/library/tensor_operation_instance/gpu/grouped_gemm/device_grouped_gemm_wmma_splitk_instance.hpp b/library/include/ck/library/tensor_operation_instance/gpu/grouped_gemm/device_grouped_gemm_wmma_splitk_instance.hpp index 6d5da9208b..d0de1c859b 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/grouped_gemm/device_grouped_gemm_wmma_splitk_instance.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/grouped_gemm/device_grouped_gemm_wmma_splitk_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_aquant.cpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_aquant.cpp index 9ba0b9c804..b6e69cd649 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_aquant.cpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_aquant.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host.hpp" #include "ck_tile/ops/gemm.hpp" diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant.cpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant.cpp index ec123364cb..4b1ad068a7 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant.cpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host.hpp" #include "ck_tile/ops/gemm.hpp" diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle.cpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle.cpp index 3a62fc091a..ae01bddf96 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle.cpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_bquant_preshuffle.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host.hpp" #include "ck_tile/ops/gemm.hpp" diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_rowcol.cpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_rowcol.cpp index 5a58ed886a..bb0fa21899 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_rowcol.cpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_rowcol.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host.hpp" #include "ck_tile/ops/gemm.hpp" diff --git a/test/ck_tile/gemm_block_scale/test_gemm_quant_tensor.cpp b/test/ck_tile/gemm_block_scale/test_gemm_quant_tensor.cpp index 0fa4048dab..8b4c90f8b9 100644 --- a/test/ck_tile/gemm_block_scale/test_gemm_quant_tensor.cpp +++ b/test/ck_tile/gemm_block_scale/test_gemm_quant_tensor.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host.hpp" #include "ck_tile/ops/gemm.hpp" diff --git a/test/ck_tile/utility/test_fill.cpp b/test/ck_tile/utility/test_fill.cpp index 18f42c4ad0..3633f8bbff 100644 --- a/test/ck_tile/utility/test_fill.cpp +++ b/test/ck_tile/utility/test_fill.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. #include "ck_tile/host/fill.hpp" #include "ck_tile/host/joinable_thread.hpp" diff --git a/test/ck_tile/warp_gemm/CMakeLists.txt b/test/ck_tile/warp_gemm/CMakeLists.txt index 664ebc003b..5079741e1b 100644 --- a/test/ck_tile/warp_gemm/CMakeLists.txt +++ b/test/ck_tile/warp_gemm/CMakeLists.txt @@ -1,3 +1,6 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + if(GPU_TARGETS MATCHES "gfx95") add_gtest_executable(test_ck_tile_wg_16x16x128_fp4 test_f32_16x16x128_fp4.cpp) endif() diff --git a/test/ck_tile/warp_gemm/test_f32_16x16x128_fp4.cpp b/test/ck_tile/warp_gemm/test_f32_16x16x128_fp4.cpp index 7878fda618..47fa1ff43e 100644 --- a/test/ck_tile/warp_gemm/test_f32_16x16x128_fp4.cpp +++ b/test/ck_tile/warp_gemm/test_f32_16x16x128_fp4.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. #include #include "ck_tile/host.hpp" diff --git a/tutorial/ck_tile/01_naive_gemm/host_level/practice_gemm_host_pipeline_agmem_bgmem_creg.hpp b/tutorial/ck_tile/01_naive_gemm/host_level/practice_gemm_host_pipeline_agmem_bgmem_creg.hpp index 15c1743a86..45f439e8fa 100644 --- a/tutorial/ck_tile/01_naive_gemm/host_level/practice_gemm_host_pipeline_agmem_bgmem_creg.hpp +++ b/tutorial/ck_tile/01_naive_gemm/host_level/practice_gemm_host_pipeline_agmem_bgmem_creg.hpp @@ -28,9 +28,9 @@ struct PracticeGemmHostPipeline { // Size of the entire problem - const auto M = a_dram.get_tensor_descriptor().get_length(number<0>{}); // M x K - const auto N = c_dram_ref.get_tensor_descriptor().get_length(number<1>{}); // M x N - const auto K = a_dram.get_tensor_descriptor().get_length(number<1>{}); // M x K + const auto M = a_dram.get_tensor_descriptor().get_length(number<0>{}); // M x K + const auto N = c_dram.get_tensor_descriptor().get_length(number<1>{}); // M x N + const auto K = a_dram.get_tensor_descriptor().get_length(number<1>{}); // M x K // Size of the block tile const auto MPerBlock = BlockTile::at(number<0>{}); @@ -83,7 +83,7 @@ struct PracticeGemmHostPipeline __shared__ char p_smem_char[block_gemm_pipeline.GetStaticLDSSize()]; const auto c_block_tile = block_gemm_pipeline(a_block_window, b_block_window, num_loops_k, p_smem_char); - auto c_window = make_tile_window(c_dram_ref, + auto c_window = make_tile_window(c_dram, make_tuple(number{}, number{}), {tile_origin_m, tile_origin_n}); store_tile(c_window, c_block_tile);