ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-18 09:38:17 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove debug code.

Browse Source
This commit is contained in:
Ville Pietilä
2025-09-30 05:53:28 +00:00
parent 193907fd85
commit 1a6f602c65
2 changed files with 0 additions and 97 deletions
Download Patch File Download Diff File Expand all files Collapse all files

36
include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp
View File

@@ -274,31 +274,6 @@ struct CShuffleEpilogue
}
}
template <typename DataType, typename StaticTileDistribution>
CK_TILE_DEVICE void print_tensor_matrix_format(
const static_distributed_tensor<DataType, StaticTileDistribution>& 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<float>(value));
});
//printf("\n");
});
//printf("\n");
}
template <index_t MPerGroup, index_t NPerGroup, typename ODramWindow, typename OAccTile, typename DsDramWindows>
//template <typename ODramWindow, typename OAccTile, typename DsDramWindows>
CK_TILE_DEVICE auto operator()(ODramWindow& out_dram_window,
@@ -408,13 +383,6 @@ struct CShuffleEpilogue
make_tuple(number<MBlockWidth>{}, number<NBlockWidth>{}),
make_tuple(number<Gs * NBlockWidth>{}, number<Gs>{}));
if (threadIdx.x == 0)
{
printf("""CShuffleEpilogue::merged_op(): MPerGroup=%d, NPerGroup=%d, Gs=%d, MBlockWidth=%d, NBlockWidth=%d\n",
static_cast<int>(MPerGroup), static_cast<int>(NPerGroup), static_cast<int>(Gs),
static_cast<int>(MBlockWidth), static_cast<int>(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<NumDTensor>{});

61
include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp
View File

@@ -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<OutDataType*>(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<float>(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<float>(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();
}
/**