From 164963bf8323a3622bf4c4dff4ec4bddfd628361 Mon Sep 17 00:00:00 2001 From: macurtis-amd Date: Wed, 2 Oct 2024 15:56:22 -0500 Subject: [PATCH] Fix compilation errors generated by forthcoming Clang changes (#1544) Without this change, the following diagnostic is generated: a template argument list is expected after a name prefixed by the template keyword [-Wmissing-template-arg-list-after-template-kw] See C++17 spec [temp.names] p5. [ROCm/composable_kernel commit: aeb7c91f48a0e8fa1e288d91f719415282c03f03] --- ...ckwise_gemm_pipeline_xdlops_v1_ab_scale.hpp | 9 +++++---- ...ckwise_gemm_pipeline_xdlops_v2_ab_scale.hpp | 18 ++++++++++-------- ...ckwise_gemm_pipeline_xdlops_v3_ab_scale.hpp | 9 +++++---- 3 files changed, 20 insertions(+), 16 deletions(-) diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp index 121593d3cc..821bbb0051 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1_ab_scale.hpp @@ -308,7 +308,7 @@ struct BlockwiseGemmXdlops_pipeline_v1_ab_scale::type; - xdlops_gemm.template Run( + xdlops_gemm.template Run<>( a_thread_vec.template AsType(), b_thread_vec.template AsType(), c_thread_buf_per_scale.GetVectorTypeReference(I0)); @@ -390,9 +390,10 @@ struct BlockwiseGemmXdlops_pipeline_v1_ab_scale::type; - xdlops_gemm.template Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf_per_scale.GetVectorTypeReference(I0)); + xdlops_gemm.template Run<>( + a_thread_vec.template AsType(), + b_thread_vec.template AsType(), + c_thread_buf_per_scale.GetVectorTypeReference(I0)); }); static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) { constexpr index_t c_offset = diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp index cb7cf605be..40fa776484 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v2_ab_scale.hpp @@ -350,7 +350,7 @@ struct BlockwiseGemmXdlops_pipeline_v2_ab_scale::type; - xdlops_gemm.template Run( + xdlops_gemm.template Run<>( a_thread_vec.template AsType(), b_thread_vec.template AsType(), c_thread_buf_per_scale.GetVectorTypeReference(I0)); @@ -443,7 +443,7 @@ struct BlockwiseGemmXdlops_pipeline_v2_ab_scale::type; - xdlops_gemm.template Run( + xdlops_gemm.template Run<>( a_thread_vec.template AsType(), b_thread_vec.template AsType(), c_thread_buf_per_scale.GetVectorTypeReference(I0)); @@ -518,9 +518,10 @@ struct BlockwiseGemmXdlops_pipeline_v2_ab_scale::type; - xdlops_gemm.template Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf_per_scale.GetVectorTypeReference(I0)); + xdlops_gemm.template Run<>( + a_thread_vec.template AsType(), + b_thread_vec.template AsType(), + c_thread_buf_per_scale.GetVectorTypeReference(I0)); }); static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) { constexpr index_t c_offset = @@ -575,9 +576,10 @@ struct BlockwiseGemmXdlops_pipeline_v2_ab_scale::type; - xdlops_gemm.template Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf_per_scale.GetVectorTypeReference(I0)); + xdlops_gemm.template Run<>( + a_thread_vec.template AsType(), + b_thread_vec.template AsType(), + c_thread_buf_per_scale.GetVectorTypeReference(I0)); }); static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) { constexpr index_t c_offset = diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp index 66c9a5c339..de542866a6 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3_ab_scale.hpp @@ -427,7 +427,7 @@ struct BlockwiseGemmXdlops_pipeline_v3_ab_scale::type; - xdlops_gemm.template Run( + xdlops_gemm.template Run<>( a_thread_vec.template AsType(), b_thread_vec.template AsType(), c_thread_buf_per_scale.GetVectorTypeReference(I0)); @@ -504,9 +504,10 @@ struct BlockwiseGemmXdlops_pipeline_v3_ab_scale::type; - xdlops_gemm.template Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf_per_scale.GetVectorTypeReference(I0)); + xdlops_gemm.template Run<>( + a_thread_vec.template AsType(), + b_thread_vec.template AsType(), + c_thread_buf_per_scale.GetVectorTypeReference(I0)); }); static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) { constexpr index_t c_offset =