Skip to content

GitLab

Commit 0cd0835c authored by Po-Yen, Chen's avatar Po-Yen, Chen
Browse files

Re-format files

parent 8d5e79fb
Hide whitespace changes
Inline Side-by-side
Showing with 14 additions and 10 deletions
include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp
View file @ 0cd0835c
...@@ -67,11 +67,13 @@ __global__ void ...@@ -67,11 +67,13 @@ __global__ void
__shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()]; __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
const auto a_grid_desc_k0_m_k1 = amd_wave_read_first_lane(GridwiseGemm::MakeAGridDescriptor_K0_M_K1( const auto a_grid_desc_k0_m_k1 =
karg.M, karg.MPadded, karg.K, karg.K0, karg.StrideA)); amd_wave_read_first_lane(GridwiseGemm::MakeAGridDescriptor_K0_M_K1(
const auto b_grid_desc_k0_n_k1 = amd_wave_read_first_lane(GridwiseGemm::MakeBGridDescriptor_K0_N_K1( karg.M, karg.MPadded, karg.K, karg.K0, karg.StrideA));
karg.K, karg.N, karg.NPadded, karg.K0, karg.StrideB)); const auto b_grid_desc_k0_n_k1 =
const auto c_grid_desc_m_n = amd_wave_read_first_lane(GridwiseGemm::MakeCGridDescriptor_M_N( amd_wave_read_first_lane(GridwiseGemm::MakeBGridDescriptor_K0_N_K1(
karg.K, karg.N, karg.NPadded, karg.K0, karg.StrideB));
const auto c_grid_desc_m_n = amd_wave_read_first_lane(GridwiseGemm::MakeCGridDescriptor_M_N(
karg.M, karg.MPadded, karg.N, karg.NPadded, karg.StrideC)); karg.M, karg.MPadded, karg.N, karg.NPadded, karg.StrideC));
GridwiseGemm::template Run<HasMainKBlockLoop>(karg.p_a_grid + a_batch_offset, GridwiseGemm::template Run<HasMainKBlockLoop>(karg.p_a_grid + a_batch_offset,
......
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp
View file @ 0cd0835c
...@@ -67,11 +67,13 @@ __global__ void ...@@ -67,11 +67,13 @@ __global__ void
defined(__gfx940__)) defined(__gfx940__))
__shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()]; __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
const auto a_grid_desc_k0_m_k1 = amd_wave_read_first_lane(GridwiseGemm::MakeAGridDescriptor_K0_M_K1( const auto a_grid_desc_k0_m_k1 =
karg.M, karg.MPadded, karg.K, karg.K0, karg.StrideA)); amd_wave_read_first_lane(GridwiseGemm::MakeAGridDescriptor_K0_M_K1(
const auto b_grid_desc_k0_n_k1 = amd_wave_read_first_lane(GridwiseGemm::MakeBGridDescriptor_K0_N_K1( karg.M, karg.MPadded, karg.K, karg.K0, karg.StrideA));
karg.K, karg.N, karg.NPadded, karg.K0, karg.StrideB)); const auto b_grid_desc_k0_n_k1 =
const auto c_grid_desc_m_n = amd_wave_read_first_lane(GridwiseGemm::MakeCGridDescriptor_M_N( amd_wave_read_first_lane(GridwiseGemm::MakeBGridDescriptor_K0_N_K1(
karg.K, karg.N, karg.NPadded, karg.K0, karg.StrideB));
const auto c_grid_desc_m_n = amd_wave_read_first_lane(GridwiseGemm::MakeCGridDescriptor_M_N(
karg.M, karg.MPadded, karg.N, karg.NPadded, karg.StrideC)); karg.M, karg.MPadded, karg.N, karg.NPadded, karg.StrideC));
GridwiseGemm::template Run<HasMainKBlockLoop>(karg.p_a_grid, GridwiseGemm::template Run<HasMainKBlockLoop>(karg.p_a_grid,
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment