/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/warp/smfmac_xdlops_gemm.hpp Source File#
smfmac_xdlops_gemm.hpp
Go to the documentation of this file.
249 MakeCDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(const CDesc_M0_N0_M1_N1_M2_N2& c_desc_m0_n0_m1_n1_m2_n2)
337 __device__ static auto GetLaneId() { return get_thread_local_1d_id() % smfmac_instr.wave_size; }
Definition: ck.hpp:267
__host__ constexpr __device__ auto make_multi_index(Xs &&... xs)
Definition: array_multi_index.hpp:15
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__host__ constexpr __device__ auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:425
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__host__ constexpr __device__ auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:90
@ smfmac_f32_16x16x32bf16
@ smfmac_f32_16x16x32f16
@ smfmac_f32_32x32x16f16
@ smfmac_f32_32x32x16bf16
__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_descriptor.hpp:319
Definition: array.hpp:14
Definition: sequence.hpp:43
Definition: smfmac_xdlops_gemm.hpp:140
static constexpr index_t GetKPerXdlops()
Definition: smfmac_xdlops_gemm.hpp:200
__host__ constexpr __device__ SmfmacSelector()
Definition: smfmac_xdlops_gemm.hpp:174
static constexpr auto GetSmfmac()
static constexpr auto selected_smfmac
Definition: smfmac_xdlops_gemm.hpp:171
static constexpr index_t GetK1PerXdlops()
Definition: smfmac_xdlops_gemm.hpp:206
Definition: smfmac_xdlops_gemm.hpp:215
static __device__ auto GetLaneId()
Definition: smfmac_xdlops_gemm.hpp:337
static constexpr __device__ index_t GetRegSizePerXdlops()
Definition: smfmac_xdlops_gemm.hpp:317
static constexpr auto K0PerXdlops
Definition: smfmac_xdlops_gemm.hpp:424
static __device__ CIndex4D GetBeginOfThreadBlk4D(index_t, index_t)
Definition: smfmac_xdlops_gemm.hpp:407
__device__ void Run(const FloatA &p_a_wave, const FloatB &p_b_wave, const Idx &idx, FloatC &p_c_thread) const
Definition: smfmac_xdlops_gemm.hpp:326
static constexpr auto K1PerXdlops
Definition: smfmac_xdlops_gemm.hpp:423
static constexpr __device__ index_t GetWaveSize()
Definition: smfmac_xdlops_gemm.hpp:322
static constexpr __device__ index_t GetNumBlks()
Definition: smfmac_xdlops_gemm.hpp:226
static __device__ auto GetBlkIdx()
Definition: smfmac_xdlops_gemm.hpp:339
__host__ static constexpr __device__ auto GetCM0M1M2NThreadBlkLengths()
Definition: smfmac_xdlops_gemm.hpp:426
static __device__ CIndex GetBeginOfThreadBlk(index_t xdlops_i, index_t blk_i)
Definition: smfmac_xdlops_gemm.hpp:394
__host__ static __device__ auto CalculateAThreadOriginDataIndex()
Definition: smfmac_xdlops_gemm.hpp:358
static constexpr auto smfmac_instr
Definition: smfmac_xdlops_gemm.hpp:420
__host__ static constexpr __device__ auto MakeCDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(const CDesc_M0_N0_M1_N1_M2_N2 &c_desc_m0_n0_m1_n1_m2_n2)
Definition: smfmac_xdlops_gemm.hpp:249
static constexpr __device__ index_t GetNumXdlops()
Definition: smfmac_xdlops_gemm.hpp:228
__host__ constexpr __device__ SparseXdlopsGemm()
Definition: smfmac_xdlops_gemm.hpp:234
__host__ static constexpr __device__ auto MakeCDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2(const CDesc_G_M0_N0_M1_N1_M2_N2 &c_desc_g_m0_n0_m1_n1_m2_n2)
Definition: smfmac_xdlops_gemm.hpp:281
static constexpr auto KPerXdlops
Definition: smfmac_xdlops_gemm.hpp:422
__host__ static __device__ auto CalculateBThreadOriginDataIndex()
Definition: smfmac_xdlops_gemm.hpp:376
Definition: integral_constant.hpp:20
Definition: amd_smfmac.hpp:34
Definition: amd_smfmac.hpp:10
Definition: amd_smfmac.hpp:78
Definition: amd_smfmac.hpp:56
Definition: type.hpp:177
__device__ void run(const FloatA &a, const FloatB &b, const index_t &idx, FloatC ®_c) const
Definition: smfmac_xdlops_gemm.hpp:100
__device__ void run(const FloatA &a, const FloatB &b, const index_t &idx, FloatC ®_c) const
Definition: smfmac_xdlops_gemm.hpp:44
__device__ void run(const FloatA &a, const FloatB &b, const index_t &idx, FloatC ®_c) const
Definition: smfmac_xdlops_gemm.hpp:128
__device__ void run(const FloatA &a, const FloatB &b, const index_t &idx, FloatC ®_c) const
Definition: smfmac_xdlops_gemm.hpp:72
Definition: smfmac_xdlops_gemm.hpp:21
Definition: functional2.hpp:33