From 1a6f602c65ce2418bcfb5ed896d3d3a2fd2578a5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= Date: Tue, 30 Sep 2025 05:53:28 +0000 Subject: [PATCH] Remove debug code. --- .../ops/epilogue/cshuffle_epilogue.hpp | 36 ----------- ...ped_convolution_backward_weight_kernel.hpp | 61 ------------------- 2 files changed, 97 deletions(-) diff --git a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index 18049b1757..c7a86ad7d4 100644 --- a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -274,31 +274,6 @@ struct CShuffleEpilogue } } - template - CK_TILE_DEVICE void print_tensor_matrix_format( - const static_distributed_tensor& tensor, - const char* /*name = "tensor_matrix"*/) - { - const auto spans = tensor.get_distributed_spans(); - //static_assert(spans.size() == 2, "This function is for 2D tensors only"); - - const auto dim0_span = spans[number<0>{}]; - const auto dim1_span = spans[number<1>{}]; - - //printf("%s matrix format (tid %u):\n", name, threadIdx.x); - - sweep_tile_span(dim0_span, [&](auto row) { - printf(" "); - sweep_tile_span(dim1_span, [&](auto col) { - constexpr auto distributed_indices = make_tuple(row, col); - const auto value = tensor[distributed_indices]; - printf("tid %u: %.7f\n", threadIdx.x, static_cast(value)); - }); - //printf("\n"); - }); - //printf("\n"); - } - template //template CK_TILE_DEVICE auto operator()(ODramWindow& out_dram_window, @@ -408,13 +383,6 @@ struct CShuffleEpilogue make_tuple(number{}, number{}), make_tuple(number{}, number{})); - if (threadIdx.x == 0) - { - printf("""CShuffleEpilogue::merged_op(): MPerGroup=%d, NPerGroup=%d, Gs=%d, MBlockWidth=%d, NBlockWidth=%d\n", - static_cast(MPerGroup), static_cast(NPerGroup), static_cast(Gs), - static_cast(MBlockWidth), static_cast(NBlockWidth)); - } - // Loop over the groups (diagonal blocks in LDS) static_for<0, Gs, 1>{}([&](auto g) { block_sync_lds(); @@ -440,10 +408,6 @@ struct CShuffleEpilogue auto c_out_tensor = load_tile(lds_window); - // DEBUG: Print out the c_out_tensor contents for debugging - //print_tensor_matrix_format(c_out_tensor, "c_out_tensor"); - //__syncthreads(); - const auto ds_tensor = generate_tuple( [&](auto idx) { return load_tile(d_dram_windows[idx]); }, number{}); diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp index 178a9e3bb7..84d28cc360 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp @@ -804,18 +804,12 @@ struct GroupedConvolutionBackwardWeightKernel // Run Epilogue Pipeline auto& c_block_window = gemm_tile_windows.at(I3); - constexpr index_t Gs = GroupedConvTraitsType_::NumGroupsToMerge; const index_t ConvK = kargs.wei_g_k_c_xs_lengths[number<1>{}]; const index_t ConvC = kargs.wei_g_k_c_xs_lengths[number<2>{}]; const index_t ZYX = kargs.ZYX; const index_t MPerGroup = ConvK; const index_t NPerGroup = ZYX * ConvC; - if (threadIdx.x == 0) - { - printf("MPerGroup: %d, NPerGroup: %d \n", MPerGroup, NPerGroup); - } - // Check that MPerGroup and NPerGroup map to the existing options if (MPerGroup == 1 && NPerGroup == 16) { @@ -857,61 +851,6 @@ struct GroupedConvolutionBackwardWeightKernel EpiloguePipeline{}.template operator()<2, 8, decltype(c_block_window), decltype(c_block_tile)>( c_block_window, c_block_tile, d_block_window, smem_ptr_0); } - - auto print_lds_per_wg = [&]() - { - // Print out LDS contents. - // The LDS corresponds TilePartitioner_::MPerBlock * TilePartitioner_::NPerBlock matrix. - // Print LDS contents as matrix - printf("LDS Contents (%d x %d) for thread block %u:\n", TilePartitioner::MPerBlock, TilePartitioner::NPerBlock, blockIdx.x); - OutDataType* lds_data = reinterpret_cast(smem_ptr_0); - - const int MBlockWidth = TilePartitioner::MPerBlock / Gs; - const int NBlockWidth = TilePartitioner::NPerBlock / Gs; - const int Ncols = Gs; - const int Nrows = Gs; - - for(int c = 0; c < Ncols; ++c) { - printf("Block %d:\n", c); - for(int r = 0; r < Nrows; ++r) { - printf("Row %d:\n", r); - for (int m = 0; m < MBlockWidth; ++m) - { - for(int n = 0; n < NBlockWidth; ++n) - { - int idx = c + Ncols * n + Ncols * NBlockWidth * m + Ncols * MBlockWidth * NBlockWidth * r; - printf("%.7f ", static_cast(lds_data[idx])); - } - printf(" \n"); - } - } - printf("\n\n"); - } - - - // Print out the c_block_window contents for debugging - printf("C Ptr Contents (%d x %d):\n", TilePartitioner::MPerBlock, NPerGroup); - for(int m = 0; m < TilePartitioner::MPerBlock; ++m) { - for(int n = 0; n < NPerGroup; ++n) { - int idx = m * NPerGroup + n; - printf("%.7f ", static_cast(c_ptr[idx])); - if((n + 1) % NPerGroup == 0) printf("\n "); // Line break every NBlockWidth elements for readability - } - printf("\n"); - } - }; - - __syncthreads(); - if (blockIdx.x == 0 && blockIdx.y == 0 && threadIdx.x == 0 && threadIdx.y == 0) - { - print_lds_per_wg(); - } - __syncthreads(); - // if (blockIdx.x == 1 && blockIdx.y == 0 && threadIdx.x == 0 && threadIdx.y == 0) - // { - // print_lds_per_wg(); - // } - // __syncthreads(); } /**