diff --git a/example/01_gemm/run_gemm_example_streamk_v2.inc b/example/01_gemm/run_gemm_example_streamk_v2.inc index 0a07015c08..6679f95157 100644 --- a/example/01_gemm/run_gemm_example_streamk_v2.inc +++ b/example/01_gemm/run_gemm_example_streamk_v2.inc @@ -116,27 +116,38 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) }; auto f_get_default_stride = - [](std::size_t row, std::size_t col, std::size_t stride, auto layout) { + [](std::size_t row, std::size_t col, ck::index_t stride, auto layout) { if(stride == -1) { - // give a chance if stride is zero, return a default packed stride + // give a chance if stride is -1, return a default packed stride if constexpr(std::is_same_v) { - return col; + return static_cast(col); } else { - return row; + return static_cast(row); } } else - return stride; + return static_cast(stride); }; + auto f_get_default_streamk_policy = [](ck::index_t streamk_sel) { + if(streamk_sel == -1) + { + return static_cast(4); + } + else + return static_cast(streamk_sel); + }; + StrideA = f_get_default_stride(M, K, StrideA, ALayout{}); StrideB = f_get_default_stride(K, N, StrideB, BLayout{}); StrideC = f_get_default_stride(M, N, StrideC, CLayout{}); + Streamk_sel = f_get_default_streamk_policy(Streamk_sel); + Tensor a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout{})); Tensor b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout{})); diff --git a/include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp b/include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp index 828a0d1e9f..56c37b1b72 100644 --- a/include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp +++ b/include/ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp @@ -1517,30 +1517,6 @@ struct BlockToCTileMap_GemmStreamK_v2 equiv_tiles_big = MDiv(upper_big / k_iters_per_tile.get()); equiv_tiles_little = MDiv(upper_little / k_iters_per_tile.get()); } - - if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING))) - { - printf("streamk_sel=%0d,grid_size=%0d, num_tiles:%d, dp_tiles:%d, sk_tiles:%u, " - "sk_num_blocks:%d,dp_num_blocks:%d,sk_num_big_blocks:%d, " - "sk_total_iters:%d, dp_start_block_idx:%d, " - "k_iters_per_tile:%d, k_iters_per_big_block:%d, reduction_start_block_idx:%u, " - " workspace(acc float):%u\n", - streamk_sel, - grid_size, - num_tiles, - dp_tiles, - get_sk_tiles(), - sk_num_blocks, - dp_num_blocks, - sk_num_big_blocks, - sk_total_iters, - dp_start_block_idx, - - k_iters_per_tile.get(), - k_iters_per_big_block, - reduction_start_block_idx, - get_workspace_size(sizeof(float))); - } } __host__ __device__ static constexpr index_t CalculateGridSize(index_t M, index_t N) diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp index c5e7feb6dc..ff10215353 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp @@ -1315,7 +1315,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 num_k_block_main_loop = __builtin_amdgcn_readfirstlane( (a_grid_desc_ak0_m_ak1.GetLength(I0) * a_grid_desc_ak0_m_ak1.GetLength(I2)) / - KPerBlock); :AK0*KPadded/KPerBlock + KPerBlock); blockwise_gemm_pipeline.template Run( a_grid_desc_ak0_m_ak1,