ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-07 08:15:04 +00:00
Code Issues Packages Projects Releases Wiki Activity

Review fixes

Browse Source
This commit is contained in:
Enrico Degregori
2025-09-03 11:05:08 +00:00
parent 80f72391c5
commit 85570f98a0
1 changed files with 11 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files

34
include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_wmma_cshuffle_v3.hpp
View File

@@ -328,23 +328,15 @@ struct DeviceBatchedGemmMultiD_Wmma_CShuffleV3
}
}();
if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4)
{
// TODO
}
else
{
hip_check_error(hipOccupancyMaxActiveBlocksPerMultiprocessor(
&max_occupancy,
kernel_batched_gemm_multi_d_wmma_cshuffle_v3<
GridwiseGemm,
ComputePtrOffsetOfStridedBatch,
true,
InMemoryDataOperationEnum::AtomicAdd,
minimum_occupancy>,
BlockSize,
dynamic_smem_size));
}
hip_check_error(hipOccupancyMaxActiveBlocksPerMultiprocessor(
&max_occupancy,
kernel_batched_gemm_multi_d_wmma_cshuffle_v3<GridwiseGemm,
ComputePtrOffsetOfStridedBatch,
true,
InMemoryDataOperationEnum::AtomicAdd,
minimum_occupancy>,
BlockSize,
dynamic_smem_size));
max_occupancy_ = std::max(1, max_occupancy);
}
@@ -428,7 +420,7 @@ struct DeviceBatchedGemmMultiD_Wmma_CShuffleV3
rotating_mem.Next();
// clear c mem
if(arg_.KBatch > 1)
hipGetErrorString(
HIP_CHECK_ERROR(
hipMemsetAsync(arg_.p_e_grid,
0,
arg.Batch * arg_.M * arg_.N * sizeof(EDataType),
@@ -449,7 +441,7 @@ struct DeviceBatchedGemmMultiD_Wmma_CShuffleV3
{
const auto clear_workspace = [&]() {
if(arg.KBatch > 1)
hipGetErrorString(
HIP_CHECK_ERROR(
hipMemsetAsync(arg.p_e_grid,
0,
arg.Batch * arg.M * arg.N * sizeof(EDataType),
@@ -510,10 +502,6 @@ struct DeviceBatchedGemmMultiD_Wmma_CShuffleV3
Run(kernel);
}
}
else
{
// TODO: Implement
}
}
else
{