From 1bec1dd0912ed8db19e435c03c7b9c3bd44a48b3 Mon Sep 17 00:00:00 2001 From: Aviral Goel Date: Sat, 22 Nov 2025 19:38:27 -0500 Subject: [PATCH] chore(copyright): update copyright header for test directory (#3265) [ROCm/composable_kernel commit: f6c999bddb9e0ae468c7b45bc68cc1410472dcf5] --- .../smoothquant_fp16_n3072_instance.cpp | 4 +- .../smoothquant_fp16_n4096_instance.cpp | 4 +- .../smoothquant_fp16_n4096_tp_instance.cpp | 4 +- .../smoothquant_fp16_n512_instance.cpp | 4 +- .../smoothquant_fp16_n64_n128_instance.cpp | 4 +- .../smoothquant_fp16_n768_instance.cpp | 4 +- .../instances/smoothquant_fwd_api.cpp | 4 +- .../instances/smoothquant_instance_common.hpp | 4 +- test/ck_tile/smoothquant/smoothquant.hpp | 4 +- test/ck_tile/smoothquant/test_smoothquant.cpp | 4 +- .../smoothquant/test_smoothquant_cases.inc | 4 +- .../smoothquant/test_smoothquant_types.hpp | 4 +- .../smoothquant/test_smoothquant_util.hpp | 4 +- .../topk_softmax/test_topk_softmax.hpp | 2 +- .../topk_softmax/test_topk_softmax_api.cpp | 2 +- .../topk_softmax/test_topk_softmax_api.hpp | 2 +- .../topk_softmax/test_topk_softmax_bf16.cpp | 2 +- .../topk_softmax/test_topk_softmax_fp16.cpp | 2 +- .../utility/print/test_print_array.cpp | 2 +- .../utility/print/test_print_basic_types.cpp | 2 +- .../utility/print/test_print_buffer_view.cpp | 2 +- .../utility/print/test_print_common.hpp | 2 +- .../print/test_print_coordinate_transform.cpp | 2 +- .../utility/print/test_print_sequence.cpp | 2 +- .../test_print_static_encoding_pattern.cpp | 2 +- .../utility/print/test_print_tile_window.cpp | 2 +- .../utility/print/test_print_tuple.cpp | 2 +- .../test_contraction_interface_xdl.cpp | 2 +- test/contraction/test_contraction_xdl.cpp | 2 +- .../test_conv_tensor_rearrange.cpp | 2 +- .../test_conv_tensor_rearrange_interface.cpp | 2 +- test/conv_util/conv_util.cpp | 2 +- test/convnd_bwd_data/convnd_bwd_data_xdl.cpp | 2 +- test/convnd_fwd/convnd_fwd_xdl.cpp | 2 +- test/data_type/test_bf6.cpp | 2 +- test/data_type/test_bf8_fnuz.cpp | 2 +- test/data_type/test_bf8_ocp.cpp | 2 +- test/data_type/test_bhalf.cpp | 2 +- test/data_type/test_custom_type.cpp | 2 +- test/data_type/test_e8m0.cpp | 3 + test/data_type/test_fp4.cpp | 2 +- test/data_type/test_fp6.cpp | 2 +- test/data_type/test_fp8_fnuz.cpp | 2 +- test/data_type/test_fp8_ocp.cpp | 2 +- test/data_type/test_int4.cpp | 2 +- test/data_type/test_mx_bf8.cpp | 2 +- test/data_type/test_mx_fp4.cpp | 2 +- test/data_type/test_mx_fp8.cpp | 2 +- test/data_type/test_pk_i4.cpp | 2 +- test/data_type/type_convert_const.cpp | 2 +- .../test_elementwise_layernorm_fp16.cpp | 2 +- test/gemm/gemm_bf16.cpp | 2 +- test/gemm/gemm_fp16.cpp | 2 +- test/gemm/gemm_fp32.cpp | 2 +- test/gemm/gemm_fp64.cpp | 2 +- test/gemm/gemm_int8.cpp | 2 +- test/gemm/gemm_standalone_xdl_fp16.cpp | 2 +- test/gemm/gemm_util.hpp | 2 +- test/gemm/instance/gemm_f16_nn_instance.cpp | 2 +- test/gemm/instance/gemm_f16_nn_instance.hpp | 2 +- test/gemm/instance/gemm_f16_nt_instance.cpp | 2 +- test/gemm/instance/gemm_f16_nt_instance.hpp | 2 +- test/gemm/instance/gemm_f16_tn_instance.cpp | 2 +- test/gemm/instance/gemm_f16_tn_instance.hpp | 2 +- test/gemm/instance/gemm_f16_tt_instance.cpp | 2 +- test/gemm/instance/gemm_f16_tt_instance.hpp | 2 +- .../instance/gemm_wavelet_f16_tn_instance.cpp | 2 +- .../instance/gemm_wavelet_f16_tn_instance.hpp | 2 +- test/gemm/run_gemm_test.inc | 2 +- .../test_gemm_add_add_fastgelu_wmma.cpp | 2 +- test/gemm_add/test_gemm_add_fastgelu_wmma.cpp | 2 +- test/gemm_add/test_gemm_add_fastgelu_xdl.cpp | 2 +- test/gemm_add/test_gemm_add_multiply_wmma.cpp | 2 +- test/gemm_add/test_gemm_add_relu_wmma.cpp | 2 +- test/gemm_add/test_gemm_add_relu_xdl.cpp | 2 +- test/gemm_add/test_gemm_add_silu_wmma.cpp | 3 +- test/gemm_add/test_gemm_add_silu_xdl.cpp | 2 +- test/gemm_add/test_gemm_add_wmma.cpp | 2 +- test/gemm_add/test_gemm_add_xdl.cpp | 2 +- test/gemm_add/test_gemm_bilinear_wmma.cpp | 2 +- test/gemm_add/test_gemm_common.hpp | 2 +- test/gemm_add/test_gemm_fastgelu_wmma.cpp | 2 +- test/gemm_add/test_gemm_multiply_add_wmma.cpp | 2 +- .../test_gemm_multiply_multiply_wmma.cpp | 2 +- .../test_gemm_b_scale_ut_cases.inc | 3 + test/gemm_b_scale/test_gemm_b_scale_util.hpp | 2 +- test/gemm_b_scale/test_gemm_b_scale_wmma.cpp | 2 +- test/gemm_b_scale/test_gemm_b_scale_xdl.cpp | 2 +- .../test_gemm_blockscale_wp_xdl_fp8.cpp | 2 +- test/gemm_blockscale_wp/test_gemm_common.hpp | 2 +- .../test_gemm_add_relu_add_layernorm_fp16.cpp | 2 +- test/gemm_multi_abd/test_gemm_common.hpp | 2 +- .../test_gemm_multi_abd_wmma.cpp | 2 +- .../test_gemm_multi_abd_xdl.cpp | 2 +- .../test_gemm_common.hpp | 2 +- ...test_gemm_multiply_multiply_wp_xdl_fp8.cpp | 2 +- test/gemm_mx/test_gemm_mx.cpp | 2 +- test/gemm_mx/test_gemm_mx_util.hpp | 2 +- test/gemm_reduce/gemm_reduce_fp16.cpp | 2 +- .../test_gemm_splitk_ut_cases.inc | 3 + test/gemm_split_k/test_gemm_splitk_util.hpp | 2 +- test/gemm_split_k/test_gemm_splitk_xdl.cpp | 2 +- .../test_gemm_universal_ut_cases_bf16.inc | 4 +- .../test_gemm_universal_ut_cases_fp16.inc | 4 +- .../test_gemm_universal_ut_cases_fp8.inc | 4 +- .../test_gemm_universal_util.hpp | 2 +- .../test_gemm_universal_wmma_bf16.cpp | 2 +- .../test_gemm_universal_wmma_fp16.cpp | 2 +- .../test_gemm_universal_wmma_fp8.cpp | 2 +- .../test_gemm_universal_xdl_bf16.cpp | 2 +- .../test_gemm_universal_xdl_fp16.cpp | 2 +- .../test_gemm_universal_xdl_fp8.cpp | 2 +- .../test_gemm_common.hpp | 2 +- ...test_gemm_universal_preshuffle_xdl_fp8.cpp | 2 +- ...st_gemm_universal_reduce_bf16A_i8_wmma.cpp | 2 +- .../test_gemm_universal_reduce_bf16_wmma.cpp | 2 +- .../test_gemm_universal_reduce_fp16_wmma.cpp | 2 +- ...t_gemm_universal_streamk_ut_cases_bf16.inc | 4 +- ...t_gemm_universal_streamk_ut_cases_fp16.inc | 4 +- ...st_gemm_universal_streamk_ut_cases_fp8.inc | 4 +- .../test_gemm_universal_streamk_util.hpp | 2 +- .../test_gemm_universal_streamk_xdl_bf16.cpp | 2 +- .../test_gemm_universal_streamk_xdl_fp16.cpp | 2 +- .../test_gemm_universal_streamk_xdl_fp8.cpp | 2 +- ...grouped_convnd_bwd_data_interface_wmma.cpp | 2 +- ..._grouped_convnd_bwd_data_interface_xdl.cpp | 2 +- .../test_grouped_convnd_bwd_data_wmma.cpp | 2 +- .../test_grouped_convnd_bwd_data_xdl.cpp | 2 +- ...rouped_convnd_bwd_data_xdl_large_cases.cpp | 2 +- ...t_grouped_conv_bwd_weight_xdl_bilinear.cpp | 2 +- .../test_grouped_convnd_bwd_weight.cpp | 2 +- ...ouped_convnd_bwd_weight_interface_wmma.cpp | 2 +- ...rouped_convnd_bwd_weight_interface_xdl.cpp | 2 +- ...ped_convnd_bwd_weight_v3_interface_xdl.cpp | 2 +- .../test_grouped_convnd_fwd.cpp | 2 +- .../test_grouped_convnd_fwd_dataset_xdl.cpp | 2 +- ...est_grouped_convnd_fwd_large_cases_xdl.cpp | 2 +- ..._grouped_convnd_fwd_multi_ab_interface.cpp | 2 +- ...lti_d_interface_compatibility_xdl_wmma.cpp | 2 +- ...st_grouped_convnd_fwd_bias_bnorm_clamp.cpp | 2 +- .../test_grouped_convnd_fwd_bias_clamp.cpp | 2 +- ...uped_convnd_fwd_bias_clamp_large_cases.cpp | 2 +- .../test_grouped_convnd_fwd_clamp.cpp | 2 +- ...grouped_convnd_fwd_gk_bias_bnorm_clamp.cpp | 2 +- .../test_grouped_convnd_fwd_gk_bias_clamp.cpp | 2 +- .../test_grouped_gemm_interface_xdl.cpp | 2 +- .../test_grouped_gemm_splitk_xdl.cpp | 2 +- ...d_gemm_two_stage_multiple_d_splitk_xdl.cpp | 2 +- .../test_grouped_gemm_two_stage_ut_cases.inc | 3 + .../test_grouped_gemm_ut_cases.inc | 3 + test/grouped_gemm/test_grouped_gemm_util.hpp | 2 +- .../magic_number_division.cpp | 2 +- test/mx_mfma_op/mx_mfma_op.cpp | 2 +- test/mx_mfma_op/mx_mfma_op.hpp | 2 +- .../test_groupnorm_bwd_data_fp32.cpp | 2 +- .../test_layernorm2d_bwd_data_fp32.cpp | 2 +- .../test_groupnorm_bwd_gamma_beta_fp32.cpp | 2 +- .../test_layernorm2d_bwd_gamma_beta_fp32.cpp | 2 +- .../test_groupnorm_fwd_fp16.cpp | 2 +- .../test_groupnorm_fwd_fp32.cpp | 2 +- .../test_layernorm2d_fwd_fp16.cpp | 2 +- .../test_layernorm2d_fwd_fp32.cpp | 2 +- .../test_layernorm4d_fwd_fp16.cpp | 2 +- test/permute_scale/test_permute_scale.cpp | 200 +++---- test/pool/test_avg_pool2d_bwd.cpp | 2 +- test/pool/test_avg_pool2d_fwd.cpp | 2 +- test/pool/test_avg_pool3d_bwd.cpp | 2 +- test/pool/test_avg_pool3d_fwd.cpp | 2 +- test/pool/test_max_pool2d_bwd.cpp | 2 +- test/pool/test_max_pool2d_fwd.cpp | 2 +- test/pool/test_max_pool3d_bwd.cpp | 2 +- test/pool/test_max_pool3d_fwd.cpp | 2 +- test/pool/test_pool_fwd_common.hpp | 2 +- .../position_embedding/position_embedding.cpp | 2 +- .../gemm/test_gemm_quantization.cpp | 2 +- .../gemm/test_gemm_quantization_ut_cases.inc | 3 + .../gemm/test_gemm_quantization_util.hpp | 2 +- test/reduce/reduce_no_index.cpp | 2 +- test/reduce/reduce_with_index.cpp | 2 +- .../reference_conv_fwd/reference_conv_fwd.cpp | 2 +- test/s_prefetch_op/s_prefetch_op.cpp | 132 ++--- test/s_prefetch_op/s_prefetch_op_util.hpp | 498 +++++++++--------- test/scatter_gather/scatter_gather.cpp | 2 +- test/smfmac_op/smfmac_op.cpp | 2 +- test/smfmac_op/smfmac_op_util.hpp | 2 +- test/smfmac_op/smfmac_op_xdl.cpp | 2 +- test/softmax/test_softmax_interface.cpp | 2 +- test/softmax/test_softmax_rank3.cpp | 2 +- test/softmax/test_softmax_rank4.cpp | 2 +- test/softmax/test_softmax_ut_cases.inc | 3 + test/softmax/test_softmax_util.hpp | 2 +- .../space_filling_curve.cpp | 2 +- test/transpose/test_transpose_xdl.cpp | 71 +-- test/wmma_op/wmma_op.cpp | 2 +- test/wmma_op/wmma_op_util.hpp | 2 +- test/wrapper/test_wrapper_copy.cpp | 2 +- test/wrapper/test_wrapper_gemm_xdl.cpp | 2 +- test/wrapper/test_wrapper_layout.cpp | 2 +- test/wrapper/test_wrapper_partition.cpp | 2 +- test/wrapper/test_wrapper_tensor.cpp | 2 +- 200 files changed, 680 insertions(+), 659 deletions(-) diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n3072_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n3072_instance.cpp index 99257bc322..1d245cec68 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n3072_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n3072_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_instance.cpp index dec70cefb2..a6c9621cbf 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_tp_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_tp_instance.cpp index b85e864523..062e51efa6 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_tp_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n4096_tp_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n512_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n512_instance.cpp index 8d64ae043f..dd01ee6556 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n512_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n512_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n64_n128_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n64_n128_instance.cpp index 4675a31c25..20c2524fb8 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n64_n128_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n64_n128_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n768_instance.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n768_instance.cpp index f0f71fa717..857e8dbc26 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fp16_n768_instance.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fp16_n768_instance.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "smoothquant_instance_common.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_fwd_api.cpp b/test/ck_tile/smoothquant/instances/smoothquant_fwd_api.cpp index 04e6732a7e..e7880397b2 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_fwd_api.cpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_fwd_api.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include #include "smoothquant.hpp" diff --git a/test/ck_tile/smoothquant/instances/smoothquant_instance_common.hpp b/test/ck_tile/smoothquant/instances/smoothquant_instance_common.hpp index 138afcffaf..cadd8d401b 100644 --- a/test/ck_tile/smoothquant/instances/smoothquant_instance_common.hpp +++ b/test/ck_tile/smoothquant/instances/smoothquant_instance_common.hpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include #include "smoothquant.hpp" diff --git a/test/ck_tile/smoothquant/smoothquant.hpp b/test/ck_tile/smoothquant/smoothquant.hpp index ef0c36aa98..82af8b82ac 100644 --- a/test/ck_tile/smoothquant/smoothquant.hpp +++ b/test/ck_tile/smoothquant/smoothquant.hpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/ck_tile/smoothquant/test_smoothquant.cpp b/test/ck_tile/smoothquant/test_smoothquant.cpp index 6cce425e1b..d3789fcb79 100644 --- a/test/ck_tile/smoothquant/test_smoothquant.cpp +++ b/test/ck_tile/smoothquant/test_smoothquant.cpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "test_smoothquant_types.hpp" #include "test_smoothquant_util.hpp" diff --git a/test/ck_tile/smoothquant/test_smoothquant_cases.inc b/test/ck_tile/smoothquant/test_smoothquant_cases.inc index 27a7ea4676..be83250a0b 100644 --- a/test/ck_tile/smoothquant/test_smoothquant_cases.inc +++ b/test/ck_tile/smoothquant/test_smoothquant_cases.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/ck_tile/smoothquant/test_smoothquant_types.hpp b/test/ck_tile/smoothquant/test_smoothquant_types.hpp index 7f79ce3ff9..ae0dda6a88 100644 --- a/test/ck_tile/smoothquant/test_smoothquant_types.hpp +++ b/test/ck_tile/smoothquant/test_smoothquant_types.hpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include #include "ck_tile/host.hpp" diff --git a/test/ck_tile/smoothquant/test_smoothquant_util.hpp b/test/ck_tile/smoothquant/test_smoothquant_util.hpp index 5c1b733e03..7a7e292f7f 100644 --- a/test/ck_tile/smoothquant/test_smoothquant_util.hpp +++ b/test/ck_tile/smoothquant/test_smoothquant_util.hpp @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include "ck_tile/host.hpp" #include "smoothquant.hpp" diff --git a/test/ck_tile/topk_softmax/test_topk_softmax.hpp b/test/ck_tile/topk_softmax/test_topk_softmax.hpp index 73f0703534..2a47fe6291 100644 --- a/test/ck_tile/topk_softmax/test_topk_softmax.hpp +++ b/test/ck_tile/topk_softmax/test_topk_softmax.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/ck_tile/topk_softmax/test_topk_softmax_api.cpp b/test/ck_tile/topk_softmax/test_topk_softmax_api.cpp index 7e9ed6301a..2e216162fc 100644 --- a/test/ck_tile/topk_softmax/test_topk_softmax_api.cpp +++ b/test/ck_tile/topk_softmax/test_topk_softmax_api.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "test_topk_softmax_api.hpp" diff --git a/test/ck_tile/topk_softmax/test_topk_softmax_api.hpp b/test/ck_tile/topk_softmax/test_topk_softmax_api.hpp index c98a887736..ebefa43b78 100644 --- a/test/ck_tile/topk_softmax/test_topk_softmax_api.hpp +++ b/test/ck_tile/topk_softmax/test_topk_softmax_api.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once #include "ck_tile/core.hpp" diff --git a/test/ck_tile/topk_softmax/test_topk_softmax_bf16.cpp b/test/ck_tile/topk_softmax/test_topk_softmax_bf16.cpp index c541f6d9a4..5733b68e6c 100644 --- a/test/ck_tile/topk_softmax/test_topk_softmax_bf16.cpp +++ b/test/ck_tile/topk_softmax/test_topk_softmax_bf16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "test_topk_softmax.hpp" diff --git a/test/ck_tile/topk_softmax/test_topk_softmax_fp16.cpp b/test/ck_tile/topk_softmax/test_topk_softmax_fp16.cpp index 401b3c0013..3368d02daf 100644 --- a/test/ck_tile/topk_softmax/test_topk_softmax_fp16.cpp +++ b/test/ck_tile/topk_softmax/test_topk_softmax_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "test_topk_softmax.hpp" diff --git a/test/ck_tile/utility/print/test_print_array.cpp b/test/ck_tile/utility/print/test_print_array.cpp index 2fe9bc2a0c..8cc0916dd2 100644 --- a/test/ck_tile/utility/print/test_print_array.cpp +++ b/test/ck_tile/utility/print/test_print_array.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 "test_print_common.hpp" #include "ck_tile/core/container/array.hpp" diff --git a/test/ck_tile/utility/print/test_print_basic_types.cpp b/test/ck_tile/utility/print/test_print_basic_types.cpp index 7a26b6371a..8048087187 100644 --- a/test/ck_tile/utility/print/test_print_basic_types.cpp +++ b/test/ck_tile/utility/print/test_print_basic_types.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 "test_print_common.hpp" #include "ck_tile/core/utility/print.hpp" diff --git a/test/ck_tile/utility/print/test_print_buffer_view.cpp b/test/ck_tile/utility/print/test_print_buffer_view.cpp index 66668a2103..d0609417dc 100644 --- a/test/ck_tile/utility/print/test_print_buffer_view.cpp +++ b/test/ck_tile/utility/print/test_print_buffer_view.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 "test_print_common.hpp" #include "ck_tile/core/tensor/buffer_view.hpp" diff --git a/test/ck_tile/utility/print/test_print_common.hpp b/test/ck_tile/utility/print/test_print_common.hpp index 3ba2270802..25fd284315 100644 --- a/test/ck_tile/utility/print/test_print_common.hpp +++ b/test/ck_tile/utility/print/test_print_common.hpp @@ -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. #pragma once diff --git a/test/ck_tile/utility/print/test_print_coordinate_transform.cpp b/test/ck_tile/utility/print/test_print_coordinate_transform.cpp index 639b113eb7..94a4f5cf66 100644 --- a/test/ck_tile/utility/print/test_print_coordinate_transform.cpp +++ b/test/ck_tile/utility/print/test_print_coordinate_transform.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 "test_print_common.hpp" #include "ck_tile/core/algorithm/coordinate_transform.hpp" diff --git a/test/ck_tile/utility/print/test_print_sequence.cpp b/test/ck_tile/utility/print/test_print_sequence.cpp index e73a9f7e33..bd725e3c47 100644 --- a/test/ck_tile/utility/print/test_print_sequence.cpp +++ b/test/ck_tile/utility/print/test_print_sequence.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 "test_print_common.hpp" #include "ck_tile/core/utility/print.hpp" diff --git a/test/ck_tile/utility/print/test_print_static_encoding_pattern.cpp b/test/ck_tile/utility/print/test_print_static_encoding_pattern.cpp index 3b1b6ffb6d..ad41f9504b 100644 --- a/test/ck_tile/utility/print/test_print_static_encoding_pattern.cpp +++ b/test/ck_tile/utility/print/test_print_static_encoding_pattern.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 "test_print_common.hpp" #include "ck_tile/core/algorithm/static_encoding_pattern.hpp" diff --git a/test/ck_tile/utility/print/test_print_tile_window.cpp b/test/ck_tile/utility/print/test_print_tile_window.cpp index 0921525734..6dc99f4ed1 100644 --- a/test/ck_tile/utility/print/test_print_tile_window.cpp +++ b/test/ck_tile/utility/print/test_print_tile_window.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 "test_print_common.hpp" #include "ck_tile/core.hpp" diff --git a/test/ck_tile/utility/print/test_print_tuple.cpp b/test/ck_tile/utility/print/test_print_tuple.cpp index 79aaf1b3af..58d12f399d 100644 --- a/test/ck_tile/utility/print/test_print_tuple.cpp +++ b/test/ck_tile/utility/print/test_print_tuple.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 "test_print_common.hpp" #include "ck_tile/core/container/tuple.hpp" diff --git a/test/contraction/test_contraction_interface_xdl.cpp b/test/contraction/test_contraction_interface_xdl.cpp index 16812ce809..b6b220d093 100644 --- a/test/contraction/test_contraction_interface_xdl.cpp +++ b/test/contraction/test_contraction_interface_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/contraction/test_contraction_xdl.cpp b/test/contraction/test_contraction_xdl.cpp index 3a65b57b0e..373aaa2597 100644 --- a/test/contraction/test_contraction_xdl.cpp +++ b/test/contraction/test_contraction_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp b/test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp index 8904b58d8d..c8a53e418d 100644 --- a/test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp +++ b/test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp b/test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp index 5de0cc13a7..0082aacab9 100644 --- a/test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp +++ b/test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/conv_util/conv_util.cpp b/test/conv_util/conv_util.cpp index 124efb6b8a..10d2d2a3b3 100644 --- a/test/conv_util/conv_util.cpp +++ b/test/conv_util/conv_util.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/convnd_bwd_data/convnd_bwd_data_xdl.cpp b/test/convnd_bwd_data/convnd_bwd_data_xdl.cpp index 5ad4f63d30..98f466a2b3 100644 --- a/test/convnd_bwd_data/convnd_bwd_data_xdl.cpp +++ b/test/convnd_bwd_data/convnd_bwd_data_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/convnd_fwd/convnd_fwd_xdl.cpp b/test/convnd_fwd/convnd_fwd_xdl.cpp index 6d507211ce..a2fdcaf870 100644 --- a/test/convnd_fwd/convnd_fwd_xdl.cpp +++ b/test/convnd_fwd/convnd_fwd_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/data_type/test_bf6.cpp b/test/data_type/test_bf6.cpp index 51361ddbb6..58df189b15 100644 --- a/test/data_type/test_bf6.cpp +++ b/test/data_type/test_bf6.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 "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_bf8_fnuz.cpp b/test/data_type/test_bf8_fnuz.cpp index f028c0da73..d1dfe23a24 100644 --- a/test/data_type/test_bf8_fnuz.cpp +++ b/test/data_type/test_bf8_fnuz.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_bf8_ocp.cpp b/test/data_type/test_bf8_ocp.cpp index 9b9b6e3b53..306cf20663 100644 --- a/test/data_type/test_bf8_ocp.cpp +++ b/test/data_type/test_bf8_ocp.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/library/utility/device_memory.hpp" diff --git a/test/data_type/test_bhalf.cpp b/test/data_type/test_bhalf.cpp index 6d3d9e0f0b..84f840a5e7 100644 --- a/test/data_type/test_bhalf.cpp +++ b/test/data_type/test_bhalf.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 "gtest/gtest.h" diff --git a/test/data_type/test_custom_type.cpp b/test/data_type/test_custom_type.cpp index b8c0d402a2..b081798c70 100644 --- a/test/data_type/test_custom_type.cpp +++ b/test/data_type/test_custom_type.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_e8m0.cpp b/test/data_type/test_e8m0.cpp index 83bb0e38b8..1f1c47f8f6 100644 --- a/test/data_type/test_e8m0.cpp +++ b/test/data_type/test_e8m0.cpp @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #include #include "ck/utility/e8m0.hpp" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_fp4.cpp b/test/data_type/test_fp4.cpp index 3fc74a2ef3..9fe7210838 100644 --- a/test/data_type/test_fp4.cpp +++ b/test/data_type/test_fp4.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 "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_fp6.cpp b/test/data_type/test_fp6.cpp index 7a51ac514f..7256e5031f 100644 --- a/test/data_type/test_fp6.cpp +++ b/test/data_type/test_fp6.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 "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_fp8_fnuz.cpp b/test/data_type/test_fp8_fnuz.cpp index 0cf775f947..63c0cd19e9 100644 --- a/test/data_type/test_fp8_fnuz.cpp +++ b/test/data_type/test_fp8_fnuz.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/data_type/test_fp8_ocp.cpp b/test/data_type/test_fp8_ocp.cpp index 552b4ca871..7a22f6239d 100644 --- a/test/data_type/test_fp8_ocp.cpp +++ b/test/data_type/test_fp8_ocp.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/library/utility/device_memory.hpp" diff --git a/test/data_type/test_int4.cpp b/test/data_type/test_int4.cpp index b0ec1a6646..52364e3060 100644 --- a/test/data_type/test_int4.cpp +++ b/test/data_type/test_int4.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/data_type/test_mx_bf8.cpp b/test/data_type/test_mx_bf8.cpp index cc0cbb2b3b..139cc03f69 100644 --- a/test/data_type/test_mx_bf8.cpp +++ b/test/data_type/test_mx_bf8.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/library/utility/device_memory.hpp" diff --git a/test/data_type/test_mx_fp4.cpp b/test/data_type/test_mx_fp4.cpp index 1a31e10fac..6435be39bc 100644 --- a/test/data_type/test_mx_fp4.cpp +++ b/test/data_type/test_mx_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 "gtest/gtest.h" #include "ck/library/utility/device_memory.hpp" diff --git a/test/data_type/test_mx_fp8.cpp b/test/data_type/test_mx_fp8.cpp index 7d6c660bc8..7ecb20b807 100644 --- a/test/data_type/test_mx_fp8.cpp +++ b/test/data_type/test_mx_fp8.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/library/utility/device_memory.hpp" diff --git a/test/data_type/test_pk_i4.cpp b/test/data_type/test_pk_i4.cpp index 101c9f926f..aa82a9446e 100644 --- a/test/data_type/test_pk_i4.cpp +++ b/test/data_type/test_pk_i4.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 diff --git a/test/data_type/type_convert_const.cpp b/test/data_type/type_convert_const.cpp index 8b9c34861a..6e29e58023 100644 --- a/test/data_type/type_convert_const.cpp +++ b/test/data_type/type_convert_const.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/utility/data_type.hpp" diff --git a/test/elementwise_normalization/test_elementwise_layernorm_fp16.cpp b/test/elementwise_normalization/test_elementwise_layernorm_fp16.cpp index 43192ed139..ea8148e1e4 100644 --- a/test/elementwise_normalization/test_elementwise_layernorm_fp16.cpp +++ b/test/elementwise_normalization/test_elementwise_layernorm_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_elementwise_layernorm_impl.hpp" diff --git a/test/gemm/gemm_bf16.cpp b/test/gemm/gemm_bf16.cpp index d06735a097..f8eff8507b 100644 --- a/test/gemm/gemm_bf16.cpp +++ b/test/gemm/gemm_bf16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/gemm_fp16.cpp b/test/gemm/gemm_fp16.cpp index 185412ab65..0380a9a420 100644 --- a/test/gemm/gemm_fp16.cpp +++ b/test/gemm/gemm_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/gemm_fp32.cpp b/test/gemm/gemm_fp32.cpp index cf2d0bd01d..b8f6cd27d0 100644 --- a/test/gemm/gemm_fp32.cpp +++ b/test/gemm/gemm_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/gemm_fp64.cpp b/test/gemm/gemm_fp64.cpp index 7bf89d9c20..5c80de530c 100644 --- a/test/gemm/gemm_fp64.cpp +++ b/test/gemm/gemm_fp64.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/gemm_int8.cpp b/test/gemm/gemm_int8.cpp index f1a19dd61a..ae6403b129 100644 --- a/test/gemm/gemm_int8.cpp +++ b/test/gemm/gemm_int8.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/gemm_standalone_xdl_fp16.cpp b/test/gemm/gemm_standalone_xdl_fp16.cpp index bee2d1ec80..90a5a325b8 100644 --- a/test/gemm/gemm_standalone_xdl_fp16.cpp +++ b/test/gemm/gemm_standalone_xdl_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gemm_util.hpp" diff --git a/test/gemm/gemm_util.hpp b/test/gemm/gemm_util.hpp index 043eca0e83..df22531bc3 100644 --- a/test/gemm/gemm_util.hpp +++ b/test/gemm/gemm_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/gemm/instance/gemm_f16_nn_instance.cpp b/test/gemm/instance/gemm_f16_nn_instance.cpp index 9016257f13..701a281a9b 100644 --- a/test/gemm/instance/gemm_f16_nn_instance.cpp +++ b/test/gemm/instance/gemm_f16_nn_instance.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_nn_instance.hpp b/test/gemm/instance/gemm_f16_nn_instance.hpp index e174b99a1d..db6942ff0e 100644 --- a/test/gemm/instance/gemm_f16_nn_instance.hpp +++ b/test/gemm/instance/gemm_f16_nn_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_nt_instance.cpp b/test/gemm/instance/gemm_f16_nt_instance.cpp index 27103b88d4..460b35992f 100644 --- a/test/gemm/instance/gemm_f16_nt_instance.cpp +++ b/test/gemm/instance/gemm_f16_nt_instance.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_nt_instance.hpp b/test/gemm/instance/gemm_f16_nt_instance.hpp index c624425e69..234bcb6fce 100644 --- a/test/gemm/instance/gemm_f16_nt_instance.hpp +++ b/test/gemm/instance/gemm_f16_nt_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_tn_instance.cpp b/test/gemm/instance/gemm_f16_tn_instance.cpp index 5b11f4dad9..06fd0ea109 100644 --- a/test/gemm/instance/gemm_f16_tn_instance.cpp +++ b/test/gemm/instance/gemm_f16_tn_instance.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_tn_instance.hpp b/test/gemm/instance/gemm_f16_tn_instance.hpp index 563e10600a..cddcfd1105 100644 --- a/test/gemm/instance/gemm_f16_tn_instance.hpp +++ b/test/gemm/instance/gemm_f16_tn_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_tt_instance.cpp b/test/gemm/instance/gemm_f16_tt_instance.cpp index 9032150f0c..18167fec88 100644 --- a/test/gemm/instance/gemm_f16_tt_instance.cpp +++ b/test/gemm/instance/gemm_f16_tt_instance.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_f16_tt_instance.hpp b/test/gemm/instance/gemm_f16_tt_instance.hpp index 62914d7ac2..100d497de7 100644 --- a/test/gemm/instance/gemm_f16_tt_instance.hpp +++ b/test/gemm/instance/gemm_f16_tt_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_wavelet_f16_tn_instance.cpp b/test/gemm/instance/gemm_wavelet_f16_tn_instance.cpp index 23ff89da60..5536cf7935 100644 --- a/test/gemm/instance/gemm_wavelet_f16_tn_instance.cpp +++ b/test/gemm/instance/gemm_wavelet_f16_tn_instance.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm/instance/gemm_wavelet_f16_tn_instance.hpp b/test/gemm/instance/gemm_wavelet_f16_tn_instance.hpp index ef269d78ee..d93ef2508f 100644 --- a/test/gemm/instance/gemm_wavelet_f16_tn_instance.hpp +++ b/test/gemm/instance/gemm_wavelet_f16_tn_instance.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/gemm/run_gemm_test.inc b/test/gemm/run_gemm_test.inc index 0ab2a63367..0165472bb9 100644 --- a/test/gemm/run_gemm_test.inc +++ b/test/gemm/run_gemm_test.inc @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. int run_gemm_test(int argc, char* argv[]) { diff --git a/test/gemm_add/test_gemm_add_add_fastgelu_wmma.cpp b/test/gemm_add/test_gemm_add_add_fastgelu_wmma.cpp index 25da138a04..d9ffee68b0 100644 --- a/test/gemm_add/test_gemm_add_add_fastgelu_wmma.cpp +++ b/test/gemm_add/test_gemm_add_add_fastgelu_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_fastgelu_wmma.cpp b/test/gemm_add/test_gemm_add_fastgelu_wmma.cpp index df70a0cc99..e72b1c3761 100644 --- a/test/gemm_add/test_gemm_add_fastgelu_wmma.cpp +++ b/test/gemm_add/test_gemm_add_fastgelu_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_fastgelu_xdl.cpp b/test/gemm_add/test_gemm_add_fastgelu_xdl.cpp index 0e034f46b5..21c5b47f88 100644 --- a/test/gemm_add/test_gemm_add_fastgelu_xdl.cpp +++ b/test/gemm_add/test_gemm_add_fastgelu_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_multiply_wmma.cpp b/test/gemm_add/test_gemm_add_multiply_wmma.cpp index be4a99d69f..4d3532d207 100644 --- a/test/gemm_add/test_gemm_add_multiply_wmma.cpp +++ b/test/gemm_add/test_gemm_add_multiply_wmma.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_relu_wmma.cpp b/test/gemm_add/test_gemm_add_relu_wmma.cpp index 76c66a11b1..1d099f8bcf 100644 --- a/test/gemm_add/test_gemm_add_relu_wmma.cpp +++ b/test/gemm_add/test_gemm_add_relu_wmma.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_relu_xdl.cpp b/test/gemm_add/test_gemm_add_relu_xdl.cpp index 4b445e8e41..d87ac74188 100644 --- a/test/gemm_add/test_gemm_add_relu_xdl.cpp +++ b/test/gemm_add/test_gemm_add_relu_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_silu_wmma.cpp b/test/gemm_add/test_gemm_add_silu_wmma.cpp index 7afa68dfe9..f68f67a36f 100644 --- a/test/gemm_add/test_gemm_add_silu_wmma.cpp +++ b/test/gemm_add/test_gemm_add_silu_wmma.cpp @@ -1,6 +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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_silu_xdl.cpp b/test/gemm_add/test_gemm_add_silu_xdl.cpp index 6bd0ee422d..3af279c286 100644 --- a/test/gemm_add/test_gemm_add_silu_xdl.cpp +++ b/test/gemm_add/test_gemm_add_silu_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_wmma.cpp b/test/gemm_add/test_gemm_add_wmma.cpp index ae08d50fcc..bc440a9ae8 100644 --- a/test/gemm_add/test_gemm_add_wmma.cpp +++ b/test/gemm_add/test_gemm_add_wmma.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_add_xdl.cpp b/test/gemm_add/test_gemm_add_xdl.cpp index 6696c1ccf6..873e87edd4 100644 --- a/test/gemm_add/test_gemm_add_xdl.cpp +++ b/test/gemm_add/test_gemm_add_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_bilinear_wmma.cpp b/test/gemm_add/test_gemm_bilinear_wmma.cpp index dfa8ac7121..30949cc555 100644 --- a/test/gemm_add/test_gemm_bilinear_wmma.cpp +++ b/test/gemm_add/test_gemm_bilinear_wmma.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_common.hpp b/test/gemm_add/test_gemm_common.hpp index 9ab6c335e9..30cd42eeb0 100644 --- a/test/gemm_add/test_gemm_common.hpp +++ b/test/gemm_add/test_gemm_common.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_fastgelu_wmma.cpp b/test/gemm_add/test_gemm_fastgelu_wmma.cpp index d8dd218ec6..138fa1c245 100644 --- a/test/gemm_add/test_gemm_fastgelu_wmma.cpp +++ b/test/gemm_add/test_gemm_fastgelu_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_add/test_gemm_multiply_add_wmma.cpp b/test/gemm_add/test_gemm_multiply_add_wmma.cpp index 2531464f72..c9ba5c1e27 100644 --- a/test/gemm_add/test_gemm_multiply_add_wmma.cpp +++ b/test/gemm_add/test_gemm_multiply_add_wmma.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 "gtest/gtest.h" #include "test_gemm_common.hpp" diff --git a/test/gemm_add/test_gemm_multiply_multiply_wmma.cpp b/test/gemm_add/test_gemm_multiply_multiply_wmma.cpp index 74e900c43f..6999c2a1df 100644 --- a/test/gemm_add/test_gemm_multiply_multiply_wmma.cpp +++ b/test/gemm_add/test_gemm_multiply_multiply_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_b_scale/test_gemm_b_scale_ut_cases.inc b/test/gemm_b_scale/test_gemm_b_scale_ut_cases.inc index b9b4ea7b9d..9fa2c9d471 100644 --- a/test/gemm_b_scale/test_gemm_b_scale_ut_cases.inc +++ b/test/gemm_b_scale/test_gemm_b_scale_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TYPED_TEST(TestGemmBScale_MK_NK, SmallM) diff --git a/test/gemm_b_scale/test_gemm_b_scale_util.hpp b/test/gemm_b_scale/test_gemm_b_scale_util.hpp index ec47470b84..5bec27d5ca 100644 --- a/test/gemm_b_scale/test_gemm_b_scale_util.hpp +++ b/test/gemm_b_scale/test_gemm_b_scale_util.hpp @@ -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. #pragma once diff --git a/test/gemm_b_scale/test_gemm_b_scale_wmma.cpp b/test/gemm_b_scale/test_gemm_b_scale_wmma.cpp index 38a3540925..93eb128bb0 100644 --- a/test/gemm_b_scale/test_gemm_b_scale_wmma.cpp +++ b/test/gemm_b_scale/test_gemm_b_scale_wmma.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 diff --git a/test/gemm_b_scale/test_gemm_b_scale_xdl.cpp b/test/gemm_b_scale/test_gemm_b_scale_xdl.cpp index 38a3540925..93eb128bb0 100644 --- a/test/gemm_b_scale/test_gemm_b_scale_xdl.cpp +++ b/test/gemm_b_scale/test_gemm_b_scale_xdl.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 diff --git a/test/gemm_blockscale_wp/test_gemm_blockscale_wp_xdl_fp8.cpp b/test/gemm_blockscale_wp/test_gemm_blockscale_wp_xdl_fp8.cpp index 5d88e04690..7aab5b13c1 100644 --- a/test/gemm_blockscale_wp/test_gemm_blockscale_wp_xdl_fp8.cpp +++ b/test/gemm_blockscale_wp/test_gemm_blockscale_wp_xdl_fp8.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 diff --git a/test/gemm_blockscale_wp/test_gemm_common.hpp b/test/gemm_blockscale_wp/test_gemm_common.hpp index 25ed67a737..63f102a9b3 100644 --- a/test/gemm_blockscale_wp/test_gemm_common.hpp +++ b/test/gemm_blockscale_wp/test_gemm_common.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp b/test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp index 93142155d5..125ad102a0 100644 --- a/test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.cpp +++ b/test/gemm_layernorm/test_gemm_add_relu_add_layernorm_fp16.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_multi_abd/test_gemm_common.hpp b/test/gemm_multi_abd/test_gemm_common.hpp index 030fbcac77..12c76f2a45 100644 --- a/test/gemm_multi_abd/test_gemm_common.hpp +++ b/test/gemm_multi_abd/test_gemm_common.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_multi_abd/test_gemm_multi_abd_wmma.cpp b/test/gemm_multi_abd/test_gemm_multi_abd_wmma.cpp index 42584ecc02..ed3fbbf087 100644 --- a/test/gemm_multi_abd/test_gemm_multi_abd_wmma.cpp +++ b/test/gemm_multi_abd/test_gemm_multi_abd_wmma.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 diff --git a/test/gemm_multi_abd/test_gemm_multi_abd_xdl.cpp b/test/gemm_multi_abd/test_gemm_multi_abd_xdl.cpp index 42584ecc02..ed3fbbf087 100644 --- a/test/gemm_multi_abd/test_gemm_multi_abd_xdl.cpp +++ b/test/gemm_multi_abd/test_gemm_multi_abd_xdl.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 diff --git a/test/gemm_multiply_multiply_wp/test_gemm_common.hpp b/test/gemm_multiply_multiply_wp/test_gemm_common.hpp index 37e2b353e6..8018861022 100644 --- a/test/gemm_multiply_multiply_wp/test_gemm_common.hpp +++ b/test/gemm_multiply_multiply_wp/test_gemm_common.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_multiply_multiply_wp/test_gemm_multiply_multiply_wp_xdl_fp8.cpp b/test/gemm_multiply_multiply_wp/test_gemm_multiply_multiply_wp_xdl_fp8.cpp index bf9b909628..e2c69f0583 100644 --- a/test/gemm_multiply_multiply_wp/test_gemm_multiply_multiply_wp_xdl_fp8.cpp +++ b/test/gemm_multiply_multiply_wp/test_gemm_multiply_multiply_wp_xdl_fp8.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 diff --git a/test/gemm_mx/test_gemm_mx.cpp b/test/gemm_mx/test_gemm_mx.cpp index b63fd880c1..e307ffe6dc 100644 --- a/test/gemm_mx/test_gemm_mx.cpp +++ b/test/gemm_mx/test_gemm_mx.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 diff --git a/test/gemm_mx/test_gemm_mx_util.hpp b/test/gemm_mx/test_gemm_mx_util.hpp index c2b56bb01f..1a52d00853 100644 --- a/test/gemm_mx/test_gemm_mx_util.hpp +++ b/test/gemm_mx/test_gemm_mx_util.hpp @@ -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. #pragma once diff --git a/test/gemm_reduce/gemm_reduce_fp16.cpp b/test/gemm_reduce/gemm_reduce_fp16.cpp index 30657c87c5..aafbc00a50 100644 --- a/test/gemm_reduce/gemm_reduce_fp16.cpp +++ b/test/gemm_reduce/gemm_reduce_fp16.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 diff --git a/test/gemm_split_k/test_gemm_splitk_ut_cases.inc b/test/gemm_split_k/test_gemm_splitk_ut_cases.inc index d583a3e377..2f5c7d362a 100644 --- a/test/gemm_split_k/test_gemm_splitk_ut_cases.inc +++ b/test/gemm_split_k/test_gemm_splitk_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TYPED_TEST(TestGemmSplitK_MK_KN, SmallM) diff --git a/test/gemm_split_k/test_gemm_splitk_util.hpp b/test/gemm_split_k/test_gemm_splitk_util.hpp index f994f146c7..65fee427af 100644 --- a/test/gemm_split_k/test_gemm_splitk_util.hpp +++ b/test/gemm_split_k/test_gemm_splitk_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/gemm_split_k/test_gemm_splitk_xdl.cpp b/test/gemm_split_k/test_gemm_splitk_xdl.cpp index 3ff32977fa..84599a0949 100644 --- a/test/gemm_split_k/test_gemm_splitk_xdl.cpp +++ b/test/gemm_split_k/test_gemm_splitk_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal/test_gemm_universal_ut_cases_bf16.inc b/test/gemm_universal/test_gemm_universal_ut_cases_bf16.inc index c344d10434..a45e74e7c8 100644 --- a/test/gemm_universal/test_gemm_universal_ut_cases_bf16.inc +++ b/test/gemm_universal/test_gemm_universal_ut_cases_bf16.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal/test_gemm_universal_ut_cases_fp16.inc b/test/gemm_universal/test_gemm_universal_ut_cases_fp16.inc index 309b212249..25d95cda3d 100644 --- a/test/gemm_universal/test_gemm_universal_ut_cases_fp16.inc +++ b/test/gemm_universal/test_gemm_universal_ut_cases_fp16.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal/test_gemm_universal_ut_cases_fp8.inc b/test/gemm_universal/test_gemm_universal_ut_cases_fp8.inc index 770107a2df..7bbbc4088e 100644 --- a/test/gemm_universal/test_gemm_universal_ut_cases_fp8.inc +++ b/test/gemm_universal/test_gemm_universal_ut_cases_fp8.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal/test_gemm_universal_util.hpp b/test/gemm_universal/test_gemm_universal_util.hpp index 12835805b3..a6621b2aa2 100644 --- a/test/gemm_universal/test_gemm_universal_util.hpp +++ b/test/gemm_universal/test_gemm_universal_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/gemm_universal/test_gemm_universal_wmma_bf16.cpp b/test/gemm_universal/test_gemm_universal_wmma_bf16.cpp index 5e7aa7ddc7..e9f25df162 100644 --- a/test/gemm_universal/test_gemm_universal_wmma_bf16.cpp +++ b/test/gemm_universal/test_gemm_universal_wmma_bf16.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 diff --git a/test/gemm_universal/test_gemm_universal_wmma_fp16.cpp b/test/gemm_universal/test_gemm_universal_wmma_fp16.cpp index e530e5bbc2..e8b4e3317f 100644 --- a/test/gemm_universal/test_gemm_universal_wmma_fp16.cpp +++ b/test/gemm_universal/test_gemm_universal_wmma_fp16.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 diff --git a/test/gemm_universal/test_gemm_universal_wmma_fp8.cpp b/test/gemm_universal/test_gemm_universal_wmma_fp8.cpp index 81695258f6..5d54144747 100644 --- a/test/gemm_universal/test_gemm_universal_wmma_fp8.cpp +++ b/test/gemm_universal/test_gemm_universal_wmma_fp8.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 diff --git a/test/gemm_universal/test_gemm_universal_xdl_bf16.cpp b/test/gemm_universal/test_gemm_universal_xdl_bf16.cpp index 9e643df7b8..18031cd762 100644 --- a/test/gemm_universal/test_gemm_universal_xdl_bf16.cpp +++ b/test/gemm_universal/test_gemm_universal_xdl_bf16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal/test_gemm_universal_xdl_fp16.cpp b/test/gemm_universal/test_gemm_universal_xdl_fp16.cpp index cabf6fb38d..9e99b45e80 100644 --- a/test/gemm_universal/test_gemm_universal_xdl_fp16.cpp +++ b/test/gemm_universal/test_gemm_universal_xdl_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2025, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal/test_gemm_universal_xdl_fp8.cpp b/test/gemm_universal/test_gemm_universal_xdl_fp8.cpp index d99f25eb12..49a0670528 100644 --- a/test/gemm_universal/test_gemm_universal_xdl_fp8.cpp +++ b/test/gemm_universal/test_gemm_universal_xdl_fp8.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal_preshuffle/test_gemm_common.hpp b/test/gemm_universal_preshuffle/test_gemm_common.hpp index 367c1a9c7e..5f64444bf3 100644 --- a/test/gemm_universal_preshuffle/test_gemm_common.hpp +++ b/test/gemm_universal_preshuffle/test_gemm_common.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/gemm_universal_preshuffle/test_gemm_universal_preshuffle_xdl_fp8.cpp b/test/gemm_universal_preshuffle/test_gemm_universal_preshuffle_xdl_fp8.cpp index 06dca026ee..18f6c386a1 100644 --- a/test/gemm_universal_preshuffle/test_gemm_universal_preshuffle_xdl_fp8.cpp +++ b/test/gemm_universal_preshuffle/test_gemm_universal_preshuffle_xdl_fp8.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 diff --git a/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16A_i8_wmma.cpp b/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16A_i8_wmma.cpp index ec4c0dc784..425855bf05 100644 --- a/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16A_i8_wmma.cpp +++ b/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16A_i8_wmma.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 diff --git a/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16_wmma.cpp b/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16_wmma.cpp index cbc7860fd9..ab143695b1 100644 --- a/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16_wmma.cpp +++ b/test/gemm_universal_reduce/test_gemm_universal_reduce_bf16_wmma.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 diff --git a/test/gemm_universal_reduce/test_gemm_universal_reduce_fp16_wmma.cpp b/test/gemm_universal_reduce/test_gemm_universal_reduce_fp16_wmma.cpp index 731bee89ed..483a7d198b 100644 --- a/test/gemm_universal_reduce/test_gemm_universal_reduce_fp16_wmma.cpp +++ b/test/gemm_universal_reduce/test_gemm_universal_reduce_fp16_wmma.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 diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_bf16.inc b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_bf16.inc index 5cefd911a7..36b3adcf0d 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_bf16.inc +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_bf16.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp16.inc b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp16.inc index 6deb867cd3..d8a7f2210e 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp16.inc +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp16.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp8.inc b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp8.inc index 43140e0ef4..f73e660386 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp8.inc +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_ut_cases_fp8.inc @@ -1,5 +1,5 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_util.hpp b/test/gemm_universal_streamk/test_gemm_universal_streamk_util.hpp index 6fdf8aa71f..680c43d7e0 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_util.hpp +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_bf16.cpp b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_bf16.cpp index 5675413862..2010e4082e 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_bf16.cpp +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_bf16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp16.cpp b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp16.cpp index b6262c95c9..a86e6cfffd 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp16.cpp +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp8.cpp b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp8.cpp index a9ea93bfa6..159c84548e 100644 --- a/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp8.cpp +++ b/test/gemm_universal_streamk/test_gemm_universal_streamk_xdl_fp8.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_wmma.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_wmma.cpp index fbb6ffc6f5..871c41e706 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_wmma.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_xdl.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_xdl.cpp index 7903c17b22..969960275f 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_xdl.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_interface_xdl.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 #include diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_wmma.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_wmma.cpp index f335183a52..9becb48be2 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_wmma.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl.cpp index 17839887bb..14ed1b8939 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl.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 #include diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.cpp index 73d793cc5f..ffea6516db 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_xdl_large_cases.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 diff --git a/test/grouped_convnd_bwd_weight/test_grouped_conv_bwd_weight_xdl_bilinear.cpp b/test/grouped_convnd_bwd_weight/test_grouped_conv_bwd_weight_xdl_bilinear.cpp index 2c4b847a5d..fe71ba86c0 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_conv_bwd_weight_xdl_bilinear.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_conv_bwd_weight_xdl_bilinear.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp index f0b3b28020..4d4fcb300d 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.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 #include diff --git a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_wmma.cpp b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_wmma.cpp index 2e2f5332ae..9f1f43e780 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_wmma.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_xdl.cpp b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_xdl.cpp index 354d1fc23b..bce6da4b68 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_xdl.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_xdl.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 #include diff --git a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_v3_interface_xdl.cpp b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_v3_interface_xdl.cpp index bfd55a7c55..b2fa8e26b1 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_v3_interface_xdl.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_v3_interface_xdl.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 diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd.cpp index ca78cf4af3..b5c5248df5 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd.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 #include diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp index a1ffdaa441..0928256817 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include // Standard C library (exit codes, malloc) #include // C++ I/O streams (cout, cerr) diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp index d017a40bce..c549945d82 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_large_cases_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp index 9f4c33b34e..e483b4983c 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_ab_interface.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility_xdl_wmma.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility_xdl_wmma.cpp index f153479685..dfb8850922 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility_xdl_wmma.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_multi_d_interface_compatibility_xdl_wmma.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_bnorm_clamp.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_bnorm_clamp.cpp index 614f88d44e..c54b218739 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_bnorm_clamp.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_bnorm_clamp.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 #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp.cpp index af4f8b67f3..955ff72413 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp.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 #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp_large_cases.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp_large_cases.cpp index dcc3ec1cae..7c45d585ce 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp_large_cases.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_bias_clamp_large_cases.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 #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_clamp.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_clamp.cpp index 71fc017f1e..dc4b0f4503 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_clamp.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_clamp.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 #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_bnorm_clamp.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_bnorm_clamp.cpp index 23ab359648..8d0024354b 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_bnorm_clamp.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_bnorm_clamp.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 #include diff --git a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_clamp.cpp b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_clamp.cpp index 33273568e1..726ce30a1b 100644 --- a/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_clamp.cpp +++ b/test/grouped_convnd_fwd_activation/test_grouped_convnd_fwd_gk_bias_clamp.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 #include diff --git a/test/grouped_gemm/test_grouped_gemm_interface_xdl.cpp b/test/grouped_gemm/test_grouped_gemm_interface_xdl.cpp index ef07e2c348..1683e16323 100644 --- a/test/grouped_gemm/test_grouped_gemm_interface_xdl.cpp +++ b/test/grouped_gemm/test_grouped_gemm_interface_xdl.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 #include diff --git a/test/grouped_gemm/test_grouped_gemm_splitk_xdl.cpp b/test/grouped_gemm/test_grouped_gemm_splitk_xdl.cpp index a44214be96..c237fd562e 100644 --- a/test/grouped_gemm/test_grouped_gemm_splitk_xdl.cpp +++ b/test/grouped_gemm/test_grouped_gemm_splitk_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_gemm/test_grouped_gemm_two_stage_multiple_d_splitk_xdl.cpp b/test/grouped_gemm/test_grouped_gemm_two_stage_multiple_d_splitk_xdl.cpp index 67ecbaea30..8e1eef2304 100644 --- a/test/grouped_gemm/test_grouped_gemm_two_stage_multiple_d_splitk_xdl.cpp +++ b/test/grouped_gemm/test_grouped_gemm_two_stage_multiple_d_splitk_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/grouped_gemm/test_grouped_gemm_two_stage_ut_cases.inc b/test/grouped_gemm/test_grouped_gemm_two_stage_ut_cases.inc index 40d48f4ec0..a4f5b08fc5 100644 --- a/test/grouped_gemm/test_grouped_gemm_two_stage_ut_cases.inc +++ b/test/grouped_gemm/test_grouped_gemm_two_stage_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TEST_P(RRR_BF16_BF16_BF16, MNKPadded) diff --git a/test/grouped_gemm/test_grouped_gemm_ut_cases.inc b/test/grouped_gemm/test_grouped_gemm_ut_cases.inc index 3a42638e30..16c4ad5909 100644 --- a/test/grouped_gemm/test_grouped_gemm_ut_cases.inc +++ b/test/grouped_gemm/test_grouped_gemm_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TYPED_TEST(TestGroupedGemm, TinyCases) diff --git a/test/grouped_gemm/test_grouped_gemm_util.hpp b/test/grouped_gemm/test_grouped_gemm_util.hpp index e6a9981671..912066ee80 100644 --- a/test/grouped_gemm/test_grouped_gemm_util.hpp +++ b/test/grouped_gemm/test_grouped_gemm_util.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/magic_number_division/magic_number_division.cpp b/test/magic_number_division/magic_number_division.cpp index 1311f40d32..ceb76ba334 100644 --- a/test/magic_number_division/magic_number_division.cpp +++ b/test/magic_number_division/magic_number_division.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/mx_mfma_op/mx_mfma_op.cpp b/test/mx_mfma_op/mx_mfma_op.cpp index 9decfe14ac..0d2f58a40a 100644 --- a/test/mx_mfma_op/mx_mfma_op.cpp +++ b/test/mx_mfma_op/mx_mfma_op.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 "gtest/gtest.h" diff --git a/test/mx_mfma_op/mx_mfma_op.hpp b/test/mx_mfma_op/mx_mfma_op.hpp index 47b4419379..35a97e78e7 100644 --- a/test/mx_mfma_op/mx_mfma_op.hpp +++ b/test/mx_mfma_op/mx_mfma_op.hpp @@ -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. #pragma once diff --git a/test/normalization_bwd_data/test_groupnorm_bwd_data_fp32.cpp b/test/normalization_bwd_data/test_groupnorm_bwd_data_fp32.cpp index e6b5c918ba..c0f10b967b 100644 --- a/test/normalization_bwd_data/test_groupnorm_bwd_data_fp32.cpp +++ b/test/normalization_bwd_data/test_groupnorm_bwd_data_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_groupnorm_bwd_data_impl.hpp" diff --git a/test/normalization_bwd_data/test_layernorm2d_bwd_data_fp32.cpp b/test/normalization_bwd_data/test_layernorm2d_bwd_data_fp32.cpp index 6786c83938..a763d89961 100644 --- a/test/normalization_bwd_data/test_layernorm2d_bwd_data_fp32.cpp +++ b/test/normalization_bwd_data/test_layernorm2d_bwd_data_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_layernorm_bwd_data_impl.hpp" diff --git a/test/normalization_bwd_gamma_beta/test_groupnorm_bwd_gamma_beta_fp32.cpp b/test/normalization_bwd_gamma_beta/test_groupnorm_bwd_gamma_beta_fp32.cpp index ab9cb29891..84e063ee9d 100644 --- a/test/normalization_bwd_gamma_beta/test_groupnorm_bwd_gamma_beta_fp32.cpp +++ b/test/normalization_bwd_gamma_beta/test_groupnorm_bwd_gamma_beta_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_groupnorm_bwd_gamma_beta_impl.hpp" diff --git a/test/normalization_bwd_gamma_beta/test_layernorm2d_bwd_gamma_beta_fp32.cpp b/test/normalization_bwd_gamma_beta/test_layernorm2d_bwd_gamma_beta_fp32.cpp index 6123efbe8d..037e8fae0e 100644 --- a/test/normalization_bwd_gamma_beta/test_layernorm2d_bwd_gamma_beta_fp32.cpp +++ b/test/normalization_bwd_gamma_beta/test_layernorm2d_bwd_gamma_beta_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_layernorm_bwd_gamma_beta_impl.hpp" diff --git a/test/normalization_fwd/test_groupnorm_fwd_fp16.cpp b/test/normalization_fwd/test_groupnorm_fwd_fp16.cpp index e835668bd6..7fda26a2d2 100644 --- a/test/normalization_fwd/test_groupnorm_fwd_fp16.cpp +++ b/test/normalization_fwd/test_groupnorm_fwd_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_groupnorm_fwd_impl.hpp" diff --git a/test/normalization_fwd/test_groupnorm_fwd_fp32.cpp b/test/normalization_fwd/test_groupnorm_fwd_fp32.cpp index fcb9102fac..825ee7fc0b 100644 --- a/test/normalization_fwd/test_groupnorm_fwd_fp32.cpp +++ b/test/normalization_fwd/test_groupnorm_fwd_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_groupnorm_fwd_impl.hpp" diff --git a/test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp b/test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp index 1d8bd560b7..4ef4afc2c4 100644 --- a/test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp +++ b/test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_layernorm_fwd_impl.hpp" diff --git a/test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp b/test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp index 10ffeb762e..d12fcac35c 100644 --- a/test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp +++ b/test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_layernorm_fwd_impl.hpp" diff --git a/test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp b/test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp index b7355de96b..f04feb25cf 100644 --- a/test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp +++ b/test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_layernorm_fwd_impl.hpp" diff --git a/test/permute_scale/test_permute_scale.cpp b/test/permute_scale/test_permute_scale.cpp index e40d4861cf..b7dfb1c223 100644 --- a/test/permute_scale/test_permute_scale.cpp +++ b/test/permute_scale/test_permute_scale.cpp @@ -1,100 +1,100 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "gtest/gtest.h" -#include "profiler/profile_permute_scale_impl.hpp" - -using F16 = ck::half_t; -using F32 = float; -using ck::index_t; - -template -class TestPermute : public ::testing::Test -{ - protected: - using ADataType = std::tuple_element_t<0, Tuple>; - using BDataType = std::tuple_element_t<1, Tuple>; - - constexpr bool skip_case() - { -#ifndef CK_ENABLE_FP16 - if constexpr(ck::is_same_v || ck::is_same_v) - { - return true; - } -#endif -#ifndef CK_ENABLE_FP32 - if constexpr(ck::is_same_v || ck::is_same_v) - { - return true; - } -#endif - return false; - } - - template - void Run(std::vector lengths, - std::vector input_strides, - std::vector output_strides) - { - if(!skip_case()) - { - bool success = ck::profiler::profile_permute_scale_impl( - true, 2, false, false, lengths, input_strides, output_strides); - EXPECT_TRUE(success); - } - } -}; - -using KernelTypes = ::testing::Types, std::tuple>; - -TYPED_TEST_SUITE(TestPermute, KernelTypes); -TYPED_TEST(TestPermute, Test1D) -{ - constexpr ck::index_t NumDims = 1; - this->template Run({16}, {1}, {1}); - this->template Run({16}, {1}, {2}); - this->template Run({1}, {1}, {1}); -} - -TYPED_TEST(TestPermute, Test2D) -{ - constexpr ck::index_t NumDims = 2; - this->template Run({8, 16}, {16, 1}, {1, 8}); - this->template Run({8, 16}, {1, 8}, {16, 1}); - this->template Run({1, 1}, {1, 1}, {1, 1}); -} - -TYPED_TEST(TestPermute, Test3D) -{ - constexpr ck::index_t NumDims = 3; - this->template Run({8, 2, 8}, {16, 8, 1}, {1, 8, 16}); - this->template Run({8, 2, 8}, {1, 8, 16}, {16, 8, 1}); - this->template Run({1, 1, 1}, {1, 1, 1}, {1, 1, 1}); -} - -TYPED_TEST(TestPermute, Test4D) -{ - constexpr ck::index_t NumDims = 4; - this->template Run({8, 2, 3, 8}, {48, 24, 8, 1}, {1, 8, 16, 48}); - this->template Run({8, 2, 3, 8}, {1, 8, 16, 48}, {48, 24, 8, 1}); - this->template Run({1, 1, 1, 1}, {1, 1, 1, 1}, {1, 1, 1, 1}); -} - -TYPED_TEST(TestPermute, Test5D) -{ - constexpr ck::index_t NumDims = 5; - this->template Run({8, 2, 3, 4, 8}, {192, 96, 32, 8, 1}, {1, 8, 16, 48, 192}); - this->template Run({8, 2, 3, 4, 8}, {1, 8, 16, 48, 192}, {192, 96, 32, 8, 1}); - this->template Run({1, 1, 1, 1, 1}, {1, 1, 1, 1, 1}, {1, 1, 1, 1, 1}); -} - -TYPED_TEST(TestPermute, Test6D) -{ - constexpr ck::index_t NumDims = 6; - this->template Run( - {8, 2, 3, 4, 5, 8}, {960, 480, 160, 40, 8, 1}, {1, 8, 16, 48, 192, 960}); - this->template Run( - {8, 2, 3, 4, 5, 8}, {1, 8, 16, 48, 192, 960}, {960, 480, 160, 40, 8, 1}); - this->template Run({1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1}); -} +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "gtest/gtest.h" +#include "profiler/profile_permute_scale_impl.hpp" + +using F16 = ck::half_t; +using F32 = float; +using ck::index_t; + +template +class TestPermute : public ::testing::Test +{ + protected: + using ADataType = std::tuple_element_t<0, Tuple>; + using BDataType = std::tuple_element_t<1, Tuple>; + + constexpr bool skip_case() + { +#ifndef CK_ENABLE_FP16 + if constexpr(ck::is_same_v || ck::is_same_v) + { + return true; + } +#endif +#ifndef CK_ENABLE_FP32 + if constexpr(ck::is_same_v || ck::is_same_v) + { + return true; + } +#endif + return false; + } + + template + void Run(std::vector lengths, + std::vector input_strides, + std::vector output_strides) + { + if(!skip_case()) + { + bool success = ck::profiler::profile_permute_scale_impl( + true, 2, false, false, lengths, input_strides, output_strides); + EXPECT_TRUE(success); + } + } +}; + +using KernelTypes = ::testing::Types, std::tuple>; + +TYPED_TEST_SUITE(TestPermute, KernelTypes); +TYPED_TEST(TestPermute, Test1D) +{ + constexpr ck::index_t NumDims = 1; + this->template Run({16}, {1}, {1}); + this->template Run({16}, {1}, {2}); + this->template Run({1}, {1}, {1}); +} + +TYPED_TEST(TestPermute, Test2D) +{ + constexpr ck::index_t NumDims = 2; + this->template Run({8, 16}, {16, 1}, {1, 8}); + this->template Run({8, 16}, {1, 8}, {16, 1}); + this->template Run({1, 1}, {1, 1}, {1, 1}); +} + +TYPED_TEST(TestPermute, Test3D) +{ + constexpr ck::index_t NumDims = 3; + this->template Run({8, 2, 8}, {16, 8, 1}, {1, 8, 16}); + this->template Run({8, 2, 8}, {1, 8, 16}, {16, 8, 1}); + this->template Run({1, 1, 1}, {1, 1, 1}, {1, 1, 1}); +} + +TYPED_TEST(TestPermute, Test4D) +{ + constexpr ck::index_t NumDims = 4; + this->template Run({8, 2, 3, 8}, {48, 24, 8, 1}, {1, 8, 16, 48}); + this->template Run({8, 2, 3, 8}, {1, 8, 16, 48}, {48, 24, 8, 1}); + this->template Run({1, 1, 1, 1}, {1, 1, 1, 1}, {1, 1, 1, 1}); +} + +TYPED_TEST(TestPermute, Test5D) +{ + constexpr ck::index_t NumDims = 5; + this->template Run({8, 2, 3, 4, 8}, {192, 96, 32, 8, 1}, {1, 8, 16, 48, 192}); + this->template Run({8, 2, 3, 4, 8}, {1, 8, 16, 48, 192}, {192, 96, 32, 8, 1}); + this->template Run({1, 1, 1, 1, 1}, {1, 1, 1, 1, 1}, {1, 1, 1, 1, 1}); +} + +TYPED_TEST(TestPermute, Test6D) +{ + constexpr ck::index_t NumDims = 6; + this->template Run( + {8, 2, 3, 4, 5, 8}, {960, 480, 160, 40, 8, 1}, {1, 8, 16, 48, 192, 960}); + this->template Run( + {8, 2, 3, 4, 5, 8}, {1, 8, 16, 48, 192, 960}, {960, 480, 160, 40, 8, 1}); + this->template Run({1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1}); +} diff --git a/test/pool/test_avg_pool2d_bwd.cpp b/test/pool/test_avg_pool2d_bwd.cpp index 54c75a5553..518ed1b614 100644 --- a/test/pool/test_avg_pool2d_bwd.cpp +++ b/test/pool/test_avg_pool2d_bwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_avg_pool2d_bwd_impl.hpp" diff --git a/test/pool/test_avg_pool2d_fwd.cpp b/test/pool/test_avg_pool2d_fwd.cpp index ba78973042..727b3c826a 100644 --- a/test/pool/test_avg_pool2d_fwd.cpp +++ b/test/pool/test_avg_pool2d_fwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_pool2d_fwd_impl.hpp" diff --git a/test/pool/test_avg_pool3d_bwd.cpp b/test/pool/test_avg_pool3d_bwd.cpp index 7fa1c4907a..cfc36a0f50 100644 --- a/test/pool/test_avg_pool3d_bwd.cpp +++ b/test/pool/test_avg_pool3d_bwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_avg_pool3d_bwd_impl.hpp" diff --git a/test/pool/test_avg_pool3d_fwd.cpp b/test/pool/test_avg_pool3d_fwd.cpp index 12e83f8e5f..b20fd9639d 100644 --- a/test/pool/test_avg_pool3d_fwd.cpp +++ b/test/pool/test_avg_pool3d_fwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_pool3d_fwd_impl.hpp" diff --git a/test/pool/test_max_pool2d_bwd.cpp b/test/pool/test_max_pool2d_bwd.cpp index e6a53d0d64..d67e81f518 100644 --- a/test/pool/test_max_pool2d_bwd.cpp +++ b/test/pool/test_max_pool2d_bwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_max_pool2d_bwd_impl.hpp" diff --git a/test/pool/test_max_pool2d_fwd.cpp b/test/pool/test_max_pool2d_fwd.cpp index 4bf2a1cf8d..90f24782aa 100644 --- a/test/pool/test_max_pool2d_fwd.cpp +++ b/test/pool/test_max_pool2d_fwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_pool2d_fwd_impl.hpp" diff --git a/test/pool/test_max_pool3d_bwd.cpp b/test/pool/test_max_pool3d_bwd.cpp index 1ae2270272..5897c893c9 100644 --- a/test/pool/test_max_pool3d_bwd.cpp +++ b/test/pool/test_max_pool3d_bwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_max_pool3d_bwd_impl.hpp" diff --git a/test/pool/test_max_pool3d_fwd.cpp b/test/pool/test_max_pool3d_fwd.cpp index e7f5614d12..eb82d94d55 100644 --- a/test/pool/test_max_pool3d_fwd.cpp +++ b/test/pool/test_max_pool3d_fwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "profiler/profile_pool3d_fwd_impl.hpp" diff --git a/test/pool/test_pool_fwd_common.hpp b/test/pool/test_pool_fwd_common.hpp index b510b2f214..be6d442061 100644 --- a/test/pool/test_pool_fwd_common.hpp +++ b/test/pool/test_pool_fwd_common.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/position_embedding/position_embedding.cpp b/test/position_embedding/position_embedding.cpp index 4e13225dd3..134d2e5f37 100644 --- a/test/position_embedding/position_embedding.cpp +++ b/test/position_embedding/position_embedding.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/quantization/gemm/test_gemm_quantization.cpp b/test/quantization/gemm/test_gemm_quantization.cpp index 9981ae8a41..21c4b4c05b 100644 --- a/test/quantization/gemm/test_gemm_quantization.cpp +++ b/test/quantization/gemm/test_gemm_quantization.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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/quantization/gemm/test_gemm_quantization_ut_cases.inc b/test/quantization/gemm/test_gemm_quantization_ut_cases.inc index 83a13e4a85..2d347276f2 100644 --- a/test/quantization/gemm/test_gemm_quantization_ut_cases.inc +++ b/test/quantization/gemm/test_gemm_quantization_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TYPED_TEST(TestGemmQuantization, SmallM) diff --git a/test/quantization/gemm/test_gemm_quantization_util.hpp b/test/quantization/gemm/test_gemm_quantization_util.hpp index e1ca0de2db..de79971b48 100644 --- a/test/quantization/gemm/test_gemm_quantization_util.hpp +++ b/test/quantization/gemm/test_gemm_quantization_util.hpp @@ -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 "gtest/gtest.h" #include "ck/ck.hpp" diff --git a/test/reduce/reduce_no_index.cpp b/test/reduce/reduce_no_index.cpp index 20f7ee3d57..655593228a 100644 --- a/test/reduce/reduce_no_index.cpp +++ b/test/reduce/reduce_no_index.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/reduce/reduce_with_index.cpp b/test/reduce/reduce_with_index.cpp index fa539b4026..f46bfc4f30 100644 --- a/test/reduce/reduce_with_index.cpp +++ b/test/reduce/reduce_with_index.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include diff --git a/test/reference_conv_fwd/reference_conv_fwd.cpp b/test/reference_conv_fwd/reference_conv_fwd.cpp index 298cb96af4..53f7a200d6 100644 --- a/test/reference_conv_fwd/reference_conv_fwd.cpp +++ b/test/reference_conv_fwd/reference_conv_fwd.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/s_prefetch_op/s_prefetch_op.cpp b/test/s_prefetch_op/s_prefetch_op.cpp index fc0ae84132..ffcde5525f 100644 --- a/test/s_prefetch_op/s_prefetch_op.cpp +++ b/test/s_prefetch_op/s_prefetch_op.cpp @@ -1,66 +1,66 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. - -#include "ck/ck.hpp" -#include "ck/host_utility/device_prop.hpp" - -#include "s_prefetch_op_util.hpp" - -template -bool run_test(bool time_kernels) -{ - bool pass = true; - - const auto s_prefetch_kernel = - ck::s_prefetch_op_util::kernel_with_prefetch>; - const auto s_buffer_prefetch_kernel = ck::s_prefetch_op_util::kernel_with_prefetch< - T, - NUM_THREADS, - NUM_SCALARS, - ck::s_prefetch_op_util::SBufferPrefetchDataOp>; - - const auto prefetch_kernel_container = - std::make_tuple(s_prefetch_kernel, s_buffer_prefetch_kernel); - - ck::static_for<0, 2, 1>{}([&](auto i) { - std::string kernel_name = (i == 1 ? "s_buffer_prefetch" : "s_prefetch"); - - auto kernel = std::get{}>(prefetch_kernel_container); - - pass &= ck::s_prefetch_op_util:: - test_prefetch_impl( - time_kernels, kernel, kernel_name); - }); - - return pass; -} - -int main(int argc, char* argv[]) -{ - if(!ck::is_gfx12_supported()) - { - std::cout << "This feature is not supported by current HW, skipping tests." << std::endl; - return 0; - } - - bool time_kernels = false; - - if(argc == 2) - { - time_kernels = std::stoi(argv[1]); - } - - bool pass = true; - - std::cout << "=== Testing Constant Cache Prefetch ===" << std::endl; - - // Test different data types - pass &= run_test(time_kernels); - pass &= run_test(time_kernels); - - std::cout << "TestConstantPrefetch ..... " << (pass ? "SUCCESS" : "FAILURE") << std::endl; - return pass ? 0 : 1; -} +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "ck/ck.hpp" +#include "ck/host_utility/device_prop.hpp" + +#include "s_prefetch_op_util.hpp" + +template +bool run_test(bool time_kernels) +{ + bool pass = true; + + const auto s_prefetch_kernel = + ck::s_prefetch_op_util::kernel_with_prefetch>; + const auto s_buffer_prefetch_kernel = ck::s_prefetch_op_util::kernel_with_prefetch< + T, + NUM_THREADS, + NUM_SCALARS, + ck::s_prefetch_op_util::SBufferPrefetchDataOp>; + + const auto prefetch_kernel_container = + std::make_tuple(s_prefetch_kernel, s_buffer_prefetch_kernel); + + ck::static_for<0, 2, 1>{}([&](auto i) { + std::string kernel_name = (i == 1 ? "s_buffer_prefetch" : "s_prefetch"); + + auto kernel = std::get{}>(prefetch_kernel_container); + + pass &= ck::s_prefetch_op_util:: + test_prefetch_impl( + time_kernels, kernel, kernel_name); + }); + + return pass; +} + +int main(int argc, char* argv[]) +{ + if(!ck::is_gfx12_supported()) + { + std::cout << "This feature is not supported by current HW, skipping tests." << std::endl; + return 0; + } + + bool time_kernels = false; + + if(argc == 2) + { + time_kernels = std::stoi(argv[1]); + } + + bool pass = true; + + std::cout << "=== Testing Constant Cache Prefetch ===" << std::endl; + + // Test different data types + pass &= run_test(time_kernels); + pass &= run_test(time_kernels); + + std::cout << "TestConstantPrefetch ..... " << (pass ? "SUCCESS" : "FAILURE") << std::endl; + return pass ? 0 : 1; +} diff --git a/test/s_prefetch_op/s_prefetch_op_util.hpp b/test/s_prefetch_op/s_prefetch_op_util.hpp index 077b876b1a..26ffffab1d 100644 --- a/test/s_prefetch_op/s_prefetch_op_util.hpp +++ b/test/s_prefetch_op/s_prefetch_op_util.hpp @@ -1,249 +1,249 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. - -#include -#include -#include -#include -#include -#include - -#include "ck/ck.hpp" -#include "ck/library/utility/device_memory.hpp" -#include "ck/library/utility/host_tensor.hpp" -#include "ck/library/utility/check_err.hpp" -#include "ck/host_utility/hip_check_error.hpp" - -#include - -namespace ck { -namespace s_prefetch_op_util { - -// Enable scalar prefetch in hardware (required on gfx12 before using s_prefetch) -__device__ __forceinline__ void enable_scalar_prefetch() -{ -#if defined(__gfx12__) - // SCALAR_PREFETCH_EN is bit 24 in MODE register (hwreg 1) - // Set 1 bit at offset 24 to value 1 - __builtin_amdgcn_s_setreg(1 | (24 << 6), 1); // Set bit to 1 -#endif -} - -template -struct SPrefetchDataOp -{ - // Prefetch to constant cache using AMD builtin with cachelines to prefetch(1..32) - __device__ __forceinline__ void operator()(const T CK_CONSTANT_ADDRESS_SPACE* addr, - unsigned int num_cachelines) const - { -#if defined(__gfx12__) - assert(num_cachelines > 0 && num_cachelines <= 32); - __builtin_amdgcn_s_prefetch_data(addr, num_cachelines - 1); // we need to pass 0..31 -#else - // ignore - not supported - (void)addr; - (void)num_cachelines; -#endif - } -}; - -template -struct SBufferPrefetchDataOp -{ - // Prefetch to constant cache using AMD builtin with cachelines to prefetch(1..32) - __device__ __forceinline__ void operator()(const T CK_CONSTANT_ADDRESS_SPACE* addr, - unsigned int num_cachelines) const - { -#if defined(__gfx12__) - __amdgpu_buffer_rsrc_t buf_res = make_wave_buffer_resource_new(addr, NUM_SCALARS); - assert(num_cachelines > 0 && num_cachelines <= 32); - __builtin_amdgcn_s_buffer_prefetch_data(buf_res, 0, num_cachelines - 1); -#else - // ignore - not supported - (void)addr; - (void)num_cachelines; -#endif - } -}; - -template -__global__ void kernel_with_prefetch(const T* src, - T* dst, - const T CK_CONSTANT_ADDRESS_SPACE* scalar_data, - bool enable_prefetch) -{ - uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x; - - // Calculate number of 128B cachelines needed to cover num_scalars elements - constexpr index_t cachelineSize = 128; - constexpr index_t elements_per_cachelineSize = cachelineSize / sizeof(T); - constexpr unsigned int cachelinesNeeded = - (NUM_SCALARS + elements_per_cachelineSize - 1) / elements_per_cachelineSize; - - // Prefetch all scalar data at once - if(threadIdx.x == 0) - { - if(enable_prefetch) - { - enable_scalar_prefetch(); - } - PrefetchOp{}(scalar_data, cachelinesNeeded); - } - - T sum = 0; - if(tid < NUM_THREADS) - { - sum = src[tid]; // load from global mem to give time for prefetch to finish or be close to - // finishs - } - __syncthreads(); // waits on loads from global mem - if(tid < NUM_THREADS) - { - // Access prefetched scalar data - for(uint32_t i = 0; i < NUM_SCALARS; i++) - { - sum += scalar_data[i]; // should be fast due to scalars being preloaded - } - - dst[tid] = sum; - } -} - -template -bool test_prefetch_impl(bool time_kernels, - const PrefetchKernel& prefetch_kernel, - const std::string& kernel_name) -{ - // TODO: maybe add more prefetch instructions inside kernel to support more values - assert(NUM_SCALARS / sizeof(T) < (128 * 32)); - constexpr index_t num_elements = NUM_THREADS; - constexpr index_t num_scalars = NUM_SCALARS; - constexpr index_t block_size = 256; - constexpr index_t grid_size = (num_elements + block_size - 1) / block_size; - - std::cout << "Testing " << kernel_name << " to constant cache for type: " << typeid(T).name() - << std::endl; - std::cout << "Elements: " << num_elements << ", Scalars: " << num_scalars << std::endl; - - // Host data - std::vector h_src(num_elements); - std::vector h_scalar(num_scalars); - std::vector h_dst_with_prefetch_chunks(num_elements); - std::vector h_expected(num_elements); - - // Initialize data - for(index_t i = 0; i < num_elements; i++) - { - h_src[i] = static_cast(i % 100); - } - - T scalar_sum = 0; - for(index_t i = 0; i < num_scalars; i++) - { - h_scalar[i] = static_cast(i + 1); - scalar_sum += h_scalar[i]; - } - - // Expected results - for(index_t i = 0; i < num_elements; i++) - { - h_expected[i] = h_src[i] + scalar_sum; - } - - // Device memory - DeviceMem d_src(sizeof(T) * num_elements); - DeviceMem d_scalar(sizeof(T) * num_scalars); - DeviceMem d_dst_with_prefetch_chunks(sizeof(T) * num_elements); - - d_src.ToDevice(h_src.data()); - d_scalar.ToDevice(h_scalar.data()); - - hipStream_t stream; - hip_check_error(hipStreamCreate(&stream)); - - if(time_kernels) - { - ck::static_for<0, 2, 1>{}([&](auto static_i) { - constexpr bool prefetch_enabled = static_i == 0; - std::cout << "PREFETCH " << (prefetch_enabled ? "ENABLED!" : "DISABLED!") << std::endl; - - constexpr int num_warmup = 1; - constexpr int num_iterations = 10; - - // Warmup runs - for(int i = 0; i < num_warmup; i++) - { - prefetch_kernel<<>>( - static_cast(d_src.GetDeviceBuffer()), - static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), - cast_pointer_to_constant_address_space( - static_cast(d_scalar.GetDeviceBuffer())), - prefetch_enabled); - } - hip_check_error(hipStreamSynchronize(stream)); - - // Performance measurement - hipEvent_t start, stop; - hip_check_error(hipEventCreate(&start)); - hip_check_error(hipEventCreate(&stop)); - - hip_check_error(hipEventRecord(start, stream)); - for(int i = 0; i < num_iterations; i++) - { - prefetch_kernel<<>>( - static_cast(d_src.GetDeviceBuffer()), - static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), - cast_pointer_to_constant_address_space( - static_cast(d_scalar.GetDeviceBuffer())), - prefetch_enabled); - } - hip_check_error(hipEventRecord(stop, stream)); - - hip_check_error(hipStreamSynchronize(stream)); - - float elapsed_ms = 0; - hip_check_error(hipEventElapsedTime(&elapsed_ms, start, stop)); - - float avg_time_us = (elapsed_ms * 1000.0f) / num_iterations; - float total_bytes = (num_elements * sizeof(T) + num_scalars * sizeof(T)); // read - float bandwidth_gb_s = (total_bytes / (avg_time_us * 1e-6)) / 1e9; - float ops_per_iteration = num_elements * num_scalars; // adds - float gflops = (ops_per_iteration / (avg_time_us * 1e-6)) / 1e9; - - std::cout << " Performance: " << std::endl; - std::cout << " Average kernel time: " << avg_time_us << " us" << std::endl; - std::cout << " Effective bandwidth: " << bandwidth_gb_s << " GB/s" << std::endl; - std::cout << " Compute throughput: " << gflops << " GFLOPS" << std::endl; - - hip_check_error(hipEventDestroy(start)); - hip_check_error(hipEventDestroy(stop)); - }); - } - else - { - prefetch_kernel<<>>( - static_cast(d_src.GetDeviceBuffer()), - static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), - cast_pointer_to_constant_address_space( - static_cast(d_scalar.GetDeviceBuffer())), - true); - - hip_check_error(hipStreamSynchronize(stream)); - } - - // Copy results back - d_dst_with_prefetch_chunks.FromDevice(h_dst_with_prefetch_chunks.data()); - - // Verify results - bool pass = ck::utils::check_err(h_dst_with_prefetch_chunks, h_expected); - - std::cout << " Correctness: " << (pass ? "PASS" : "FAIL") << std::endl; - std::cout << std::endl; - - hip_check_error(hipStreamDestroy(stream)); - - return pass; -} - -} // namespace s_prefetch_op_util -} // namespace ck +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include +#include +#include + +#include "ck/ck.hpp" +#include "ck/library/utility/device_memory.hpp" +#include "ck/library/utility/host_tensor.hpp" +#include "ck/library/utility/check_err.hpp" +#include "ck/host_utility/hip_check_error.hpp" + +#include + +namespace ck { +namespace s_prefetch_op_util { + +// Enable scalar prefetch in hardware (required on gfx12 before using s_prefetch) +__device__ __forceinline__ void enable_scalar_prefetch() +{ +#if defined(__gfx12__) + // SCALAR_PREFETCH_EN is bit 24 in MODE register (hwreg 1) + // Set 1 bit at offset 24 to value 1 + __builtin_amdgcn_s_setreg(1 | (24 << 6), 1); // Set bit to 1 +#endif +} + +template +struct SPrefetchDataOp +{ + // Prefetch to constant cache using AMD builtin with cachelines to prefetch(1..32) + __device__ __forceinline__ void operator()(const T CK_CONSTANT_ADDRESS_SPACE* addr, + unsigned int num_cachelines) const + { +#if defined(__gfx12__) + assert(num_cachelines > 0 && num_cachelines <= 32); + __builtin_amdgcn_s_prefetch_data(addr, num_cachelines - 1); // we need to pass 0..31 +#else + // ignore - not supported + (void)addr; + (void)num_cachelines; +#endif + } +}; + +template +struct SBufferPrefetchDataOp +{ + // Prefetch to constant cache using AMD builtin with cachelines to prefetch(1..32) + __device__ __forceinline__ void operator()(const T CK_CONSTANT_ADDRESS_SPACE* addr, + unsigned int num_cachelines) const + { +#if defined(__gfx12__) + __amdgpu_buffer_rsrc_t buf_res = make_wave_buffer_resource_new(addr, NUM_SCALARS); + assert(num_cachelines > 0 && num_cachelines <= 32); + __builtin_amdgcn_s_buffer_prefetch_data(buf_res, 0, num_cachelines - 1); +#else + // ignore - not supported + (void)addr; + (void)num_cachelines; +#endif + } +}; + +template +__global__ void kernel_with_prefetch(const T* src, + T* dst, + const T CK_CONSTANT_ADDRESS_SPACE* scalar_data, + bool enable_prefetch) +{ + uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x; + + // Calculate number of 128B cachelines needed to cover num_scalars elements + constexpr index_t cachelineSize = 128; + constexpr index_t elements_per_cachelineSize = cachelineSize / sizeof(T); + constexpr unsigned int cachelinesNeeded = + (NUM_SCALARS + elements_per_cachelineSize - 1) / elements_per_cachelineSize; + + // Prefetch all scalar data at once + if(threadIdx.x == 0) + { + if(enable_prefetch) + { + enable_scalar_prefetch(); + } + PrefetchOp{}(scalar_data, cachelinesNeeded); + } + + T sum = 0; + if(tid < NUM_THREADS) + { + sum = src[tid]; // load from global mem to give time for prefetch to finish or be close to + // finishs + } + __syncthreads(); // waits on loads from global mem + if(tid < NUM_THREADS) + { + // Access prefetched scalar data + for(uint32_t i = 0; i < NUM_SCALARS; i++) + { + sum += scalar_data[i]; // should be fast due to scalars being preloaded + } + + dst[tid] = sum; + } +} + +template +bool test_prefetch_impl(bool time_kernels, + const PrefetchKernel& prefetch_kernel, + const std::string& kernel_name) +{ + // TODO: maybe add more prefetch instructions inside kernel to support more values + assert(NUM_SCALARS / sizeof(T) < (128 * 32)); + constexpr index_t num_elements = NUM_THREADS; + constexpr index_t num_scalars = NUM_SCALARS; + constexpr index_t block_size = 256; + constexpr index_t grid_size = (num_elements + block_size - 1) / block_size; + + std::cout << "Testing " << kernel_name << " to constant cache for type: " << typeid(T).name() + << std::endl; + std::cout << "Elements: " << num_elements << ", Scalars: " << num_scalars << std::endl; + + // Host data + std::vector h_src(num_elements); + std::vector h_scalar(num_scalars); + std::vector h_dst_with_prefetch_chunks(num_elements); + std::vector h_expected(num_elements); + + // Initialize data + for(index_t i = 0; i < num_elements; i++) + { + h_src[i] = static_cast(i % 100); + } + + T scalar_sum = 0; + for(index_t i = 0; i < num_scalars; i++) + { + h_scalar[i] = static_cast(i + 1); + scalar_sum += h_scalar[i]; + } + + // Expected results + for(index_t i = 0; i < num_elements; i++) + { + h_expected[i] = h_src[i] + scalar_sum; + } + + // Device memory + DeviceMem d_src(sizeof(T) * num_elements); + DeviceMem d_scalar(sizeof(T) * num_scalars); + DeviceMem d_dst_with_prefetch_chunks(sizeof(T) * num_elements); + + d_src.ToDevice(h_src.data()); + d_scalar.ToDevice(h_scalar.data()); + + hipStream_t stream; + hip_check_error(hipStreamCreate(&stream)); + + if(time_kernels) + { + ck::static_for<0, 2, 1>{}([&](auto static_i) { + constexpr bool prefetch_enabled = static_i == 0; + std::cout << "PREFETCH " << (prefetch_enabled ? "ENABLED!" : "DISABLED!") << std::endl; + + constexpr int num_warmup = 1; + constexpr int num_iterations = 10; + + // Warmup runs + for(int i = 0; i < num_warmup; i++) + { + prefetch_kernel<<>>( + static_cast(d_src.GetDeviceBuffer()), + static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), + cast_pointer_to_constant_address_space( + static_cast(d_scalar.GetDeviceBuffer())), + prefetch_enabled); + } + hip_check_error(hipStreamSynchronize(stream)); + + // Performance measurement + hipEvent_t start, stop; + hip_check_error(hipEventCreate(&start)); + hip_check_error(hipEventCreate(&stop)); + + hip_check_error(hipEventRecord(start, stream)); + for(int i = 0; i < num_iterations; i++) + { + prefetch_kernel<<>>( + static_cast(d_src.GetDeviceBuffer()), + static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), + cast_pointer_to_constant_address_space( + static_cast(d_scalar.GetDeviceBuffer())), + prefetch_enabled); + } + hip_check_error(hipEventRecord(stop, stream)); + + hip_check_error(hipStreamSynchronize(stream)); + + float elapsed_ms = 0; + hip_check_error(hipEventElapsedTime(&elapsed_ms, start, stop)); + + float avg_time_us = (elapsed_ms * 1000.0f) / num_iterations; + float total_bytes = (num_elements * sizeof(T) + num_scalars * sizeof(T)); // read + float bandwidth_gb_s = (total_bytes / (avg_time_us * 1e-6)) / 1e9; + float ops_per_iteration = num_elements * num_scalars; // adds + float gflops = (ops_per_iteration / (avg_time_us * 1e-6)) / 1e9; + + std::cout << " Performance: " << std::endl; + std::cout << " Average kernel time: " << avg_time_us << " us" << std::endl; + std::cout << " Effective bandwidth: " << bandwidth_gb_s << " GB/s" << std::endl; + std::cout << " Compute throughput: " << gflops << " GFLOPS" << std::endl; + + hip_check_error(hipEventDestroy(start)); + hip_check_error(hipEventDestroy(stop)); + }); + } + else + { + prefetch_kernel<<>>( + static_cast(d_src.GetDeviceBuffer()), + static_cast(d_dst_with_prefetch_chunks.GetDeviceBuffer()), + cast_pointer_to_constant_address_space( + static_cast(d_scalar.GetDeviceBuffer())), + true); + + hip_check_error(hipStreamSynchronize(stream)); + } + + // Copy results back + d_dst_with_prefetch_chunks.FromDevice(h_dst_with_prefetch_chunks.data()); + + // Verify results + bool pass = ck::utils::check_err(h_dst_with_prefetch_chunks, h_expected); + + std::cout << " Correctness: " << (pass ? "PASS" : "FAIL") << std::endl; + std::cout << std::endl; + + hip_check_error(hipStreamDestroy(stream)); + + return pass; +} + +} // namespace s_prefetch_op_util +} // namespace ck diff --git a/test/scatter_gather/scatter_gather.cpp b/test/scatter_gather/scatter_gather.cpp index 874c4d86c0..31465d3d9d 100644 --- a/test/scatter_gather/scatter_gather.cpp +++ b/test/scatter_gather/scatter_gather.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/smfmac_op/smfmac_op.cpp b/test/smfmac_op/smfmac_op.cpp index de4f9414af..5016cc320d 100644 --- a/test/smfmac_op/smfmac_op.cpp +++ b/test/smfmac_op/smfmac_op.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/smfmac_op/smfmac_op_util.hpp b/test/smfmac_op/smfmac_op_util.hpp index 44122c551d..3b1ff27775 100644 --- a/test/smfmac_op/smfmac_op_util.hpp +++ b/test/smfmac_op/smfmac_op_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/smfmac_op/smfmac_op_xdl.cpp b/test/smfmac_op/smfmac_op_xdl.cpp index 3cf2fb73b9..144c67f310 100644 --- a/test/smfmac_op/smfmac_op_xdl.cpp +++ b/test/smfmac_op/smfmac_op_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/softmax/test_softmax_interface.cpp b/test/softmax/test_softmax_interface.cpp index 25f666f0ea..1822ca2ce1 100644 --- a/test/softmax/test_softmax_interface.cpp +++ b/test/softmax/test_softmax_interface.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/softmax/test_softmax_rank3.cpp b/test/softmax/test_softmax_rank3.cpp index a8b950ce6b..e328f5dcd8 100644 --- a/test/softmax/test_softmax_rank3.cpp +++ b/test/softmax/test_softmax_rank3.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/softmax/test_softmax_rank4.cpp b/test/softmax/test_softmax_rank4.cpp index cbd790ac9b..9891d16101 100644 --- a/test/softmax/test_softmax_rank4.cpp +++ b/test/softmax/test_softmax_rank4.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/softmax/test_softmax_ut_cases.inc b/test/softmax/test_softmax_ut_cases.inc index 46154eb445..68a97834ae 100644 --- a/test/softmax/test_softmax_ut_cases.inc +++ b/test/softmax/test_softmax_ut_cases.inc @@ -1,3 +1,6 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + #pragma once TYPED_TEST(TestSoftmax, ReduceOutermostDim) diff --git a/test/softmax/test_softmax_util.hpp b/test/softmax/test_softmax_util.hpp index 96c8fe588d..2c7b9e6ffe 100644 --- a/test/softmax/test_softmax_util.hpp +++ b/test/softmax/test_softmax_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/space_filling_curve/space_filling_curve.cpp b/test/space_filling_curve/space_filling_curve.cpp index a192ecb28f..5cb2d3c289 100644 --- a/test/space_filling_curve/space_filling_curve.cpp +++ b/test/space_filling_curve/space_filling_curve.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/transpose/test_transpose_xdl.cpp b/test/transpose/test_transpose_xdl.cpp index ead622b4d7..e6ecf13521 100644 --- a/test/transpose/test_transpose_xdl.cpp +++ b/test/transpose/test_transpose_xdl.cpp @@ -1,35 +1,36 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved. -#include "gtest/gtest.h" -#include "profiler/profile_transpose_impl.hpp" - -using F16 = ck::half_t; -using F32 = float; -using ck::index_t; - -template -class TestTranspose : public ::testing::Test -{ - protected: - using ADataType = std::tuple_element_t<0, Tuple>; - using BDataType = std::tuple_element_t<1, Tuple>; - - void Run() - { - std::vector> lengths = { - {4, 16, 16, 32, 5}, {8, 16, 16, 32, 8} /**{32, 16, 16, 32, 8},**/}; - - for(auto length : lengths) - { - bool success = ck::profiler::profile_transpose_impl( - true, 2, false, false, length); - EXPECT_TRUE(success); - } - } -}; - -using KernelTypes = ::testing::Types, std::tuple>; - -TYPED_TEST_SUITE(TestTranspose, KernelTypes); -TYPED_TEST(TestTranspose, Test_FP16) { this->Run(); } -TYPED_TEST(TestTranspose, Test_FP32) { this->Run(); } +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "gtest/gtest.h" +#include "profiler/profile_transpose_impl.hpp" + +using F16 = ck::half_t; +using F32 = float; +using ck::index_t; + +template +class TestTranspose : public ::testing::Test +{ + protected: + using ADataType = std::tuple_element_t<0, Tuple>; + using BDataType = std::tuple_element_t<1, Tuple>; + + void Run() + { + std::vector> lengths = { + {4, 16, 16, 32, 5}, {8, 16, 16, 32, 8} /**{32, 16, 16, 32, 8},**/}; + + for(auto length : lengths) + { + bool success = ck::profiler::profile_transpose_impl( + true, 2, false, false, length); + EXPECT_TRUE(success); + } + } +}; + +using KernelTypes = ::testing::Types, std::tuple>; + +TYPED_TEST_SUITE(TestTranspose, KernelTypes); +TYPED_TEST(TestTranspose, Test_FP16) { this->Run(); } +TYPED_TEST(TestTranspose, Test_FP32) { this->Run(); } diff --git a/test/wmma_op/wmma_op.cpp b/test/wmma_op/wmma_op.cpp index 47d8c7ed6f..c9d1598e91 100644 --- a/test/wmma_op/wmma_op.cpp +++ b/test/wmma_op/wmma_op.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/wmma_op/wmma_op_util.hpp b/test/wmma_op/wmma_op_util.hpp index 3e511ab5bf..64692b31b1 100644 --- a/test/wmma_op/wmma_op_util.hpp +++ b/test/wmma_op/wmma_op_util.hpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #pragma once diff --git a/test/wrapper/test_wrapper_copy.cpp b/test/wrapper/test_wrapper_copy.cpp index b25bf8f1f0..ad71b9eb38 100644 --- a/test/wrapper/test_wrapper_copy.cpp +++ b/test/wrapper/test_wrapper_copy.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/wrapper/test_wrapper_gemm_xdl.cpp b/test/wrapper/test_wrapper_gemm_xdl.cpp index a11b3c00ce..b9d4bc3e57 100644 --- a/test/wrapper/test_wrapper_gemm_xdl.cpp +++ b/test/wrapper/test_wrapper_gemm_xdl.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/wrapper/test_wrapper_layout.cpp b/test/wrapper/test_wrapper_layout.cpp index 0b07303299..41bcb2086e 100644 --- a/test/wrapper/test_wrapper_layout.cpp +++ b/test/wrapper/test_wrapper_layout.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/wrapper/test_wrapper_partition.cpp b/test/wrapper/test_wrapper_partition.cpp index 08d196c4ca..edfb923fe5 100644 --- a/test/wrapper/test_wrapper_partition.cpp +++ b/test/wrapper/test_wrapper_partition.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include diff --git a/test/wrapper/test_wrapper_tensor.cpp b/test/wrapper/test_wrapper_tensor.cpp index 25b8626cac..d0595a82bd 100644 --- a/test/wrapper/test_wrapper_tensor.cpp +++ b/test/wrapper/test_wrapper_tensor.cpp @@ -1,5 +1,5 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. #include #include