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 fa9a09d423..a1e5da7f02 100755 --- 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 @@ -1224,11 +1224,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 const auto b_grid_buf = make_dynamic_buffer(p_b_grid, b_grid_desc_bk0_n_bk1.GetElementSpaceSize()); // Create mapping from block IDs to C matrix tiles using StreamK scheduling - Block2CTileMap_streamk block_2_ctile_map_streamk(problem.M, - problem.N, - AK0Number * problem.KPadded, - problem.Grid_size, - problem.Streamk_sel); + Block2CTileMap_streamk block_2_ctile_map_streamk(problem.M, problem.N, AK0Number * problem.KPadded, problem.Grid_size, problem.Streamk_sel); // Variables for iterating over K dimension uint32_t iter_start, iter_end; @@ -1246,9 +1242,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 auto c_grid_buf = make_dynamic_buffer(p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize()); // Semaphores for synchronization between blocks - uint32_t* p_semaphore = reinterpret_cast( - reinterpret_cast(p_workspace) + - block_2_ctile_map_streamk.get_workspace_size_for_acc(sizeof(AccDataType))); + uint32_t* p_semaphore = reinterpret_cast( reinterpret_cast(p_workspace) + block_2_ctile_map_streamk.get_workspace_size_for_acc(sizeof(AccDataType))); for(auto block_idx = get_block_1d_id(); @@ -1292,19 +1286,13 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 const auto thread_n_cluster_id = reduce_thread_cluster_idx[I1]; // Calculate iterations needed for reduction - constexpr auto MReduceIters = math::integer_divide_ceil( - Number{}, cluster_length_reduce.At(I0)); - constexpr auto NReduceIters = math::integer_divide_ceil( - Number{}, - cluster_length_reduce.At(I1) * - Number{}); + constexpr auto MReduceIters = math::integer_divide_ceil(Number{}, cluster_length_reduce.At(I0)); + constexpr auto NReduceIters = math::integer_divide_ceil(Number{}, cluster_length_reduce.At(I1) *Number{}); // Descriptors for thread buffer transfers - constexpr auto acc_thread_buf_load_desc = make_naive_tensor_descriptor_packed( - make_tuple(I1, Number{})); - constexpr auto acc_thread_buf_store_desc = - make_naive_tensor_descriptor_packed(make_tuple( - I1, I1, I1, Number{})); + constexpr auto acc_thread_buf_load_desc = make_naive_tensor_descriptor_packed(make_tuple(I1, Number{})); + + constexpr auto acc_thread_buf_store_desc = make_naive_tensor_descriptor_packed(make_tuple( I1, I1, I1, Number{})); constexpr auto c_partial_acc_block_m_n = GetPartialAccBlockDescriptor(); @@ -1341,6 +1329,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 CShuffleBlockTransferScalarPerVector_NPerBlock, true> parcial_acc_buf; + StaticBuffer; @@ -1474,7 +1465,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 // printf(" gridwise_gemm_xdl_cshuffle line %d , Block %d , reduction_idx %d, i_m %d, i_n_reduce %d, thread_m_cluster_id %d, thread_n_cluster_id %d\n", // __LINE__, blockIdx.x, reduction_idx, i_m, i_n_reduce, thread_m_cluster_id, thread_n_cluster_id); - printf(" gridwise_gemm_xdl_cshuffle line %d , Block %d , reduction_idx %d, i_m %d, i_n_reduce %d, thread_m_cluster_id %d, thread_n_cluster_id %d\n", + printf("Gridwise_gemm_xdl_cshuffle line %d , Block %d , reduction_idx %d, i_m %d, i_n_reduce %d, thread_m_cluster_id %d, thread_n_cluster_id %d\n", __LINE__, blockIdx.x, reduction_idx, i_m, static_cast(i_n_reduce), thread_m_cluster_id, thread_n_cluster_id); // Print values from acc_buf (up to 8 values to avoid excessive output) @@ -1482,14 +1473,14 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 { switch(i) { - case 0: printf("%.4f ", static_cast(acc_buf[Number<0>{}])); break; - case 1: printf("%.4f ", static_cast(acc_buf[Number<1>{}])); break; - case 2: printf("%.4f ", static_cast(acc_buf[Number<2>{}])); break; - case 3: printf("%.4f ", static_cast(acc_buf[Number<3>{}])); break; - case 4: printf("%.4f ", static_cast(acc_buf[Number<4>{}])); break; - case 5: printf("%.4f ", static_cast(acc_buf[Number<5>{}])); break; - case 6: printf("%.4f ", static_cast(acc_buf[Number<6>{}])); break; - case 7: printf("%.4f ", static_cast(acc_buf[Number<7>{}])); break; + case 0: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<0>{}])); break; + case 1: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<1>{}])); break; + case 2: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<2>{}])); break; + case 3: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<3>{}])); break; + case 4: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<4>{}])); break; + case 5: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<5>{}])); break; + case 6: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<6>{}])); break; + case 7: printf("acc_buf[0] = %.4f \n", static_cast(acc_buf[Number<7>{}])); break; // Add more cases if CShuffleBlockTransferScalarPerVector_NPerBlock is larger than 8 } } @@ -1497,7 +1488,7 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 // Print matrix coordinates auto global_m = spatial_idx[I0] * MPerBlock + thread_m_cluster_id; auto global_n = spatial_idx[I1] * NPerBlock + thread_n_cluster_id * CShuffleBlockTransferScalarPerVector_NPerBlock; - printf(" at C[%d,%d]\n", global_m, global_n); + printf(" at C[%d,%d] \n", global_m, global_n); } #endif diff --git a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp index 3a096df4c8..c75040a1b8 100755 --- a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp +++ b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp @@ -224,8 +224,9 @@ struct ThreadwiseTensorSliceTransfer_v6r1r2 auto dst_fp16_value = dst_value[0]; float dst_fp32_value = static_cast(dst_fp16_value); - printf("BlockId %d - Dst Vector Data being copied to dst_buf at idx %d: %f\n", + printf("BlockId %d - Line %d: (After dst_buf Update)- Dst Vector Data being copied to dst_buf at idx %d: %f\n", static_cast(blockIdx.x), + __LINE__, static_cast(idx_1d.value), dst_fp32_value);