device_normalization_fwd_impl.hpp Source File#
device_normalization_fwd_impl.hpp
Go to the documentation of this file.
float launch_and_time_kernel(const StreamConfig &stream_config, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition host_utility/kernel_launch.hpp:14
__host__ __device__ constexpr auto integer_least_multiple(X x, Y y)
Definition utility/math.hpp:78
__host__ __device__ constexpr auto integer_divide_ceil(X x, Y y)
Definition utility/math.hpp:72
Definition convolution_backward_data_specialization.hpp:8
auto make_tuple_from_array(const std::vector< index_t > &lengths, Number< arraySize >)
Definition device_reduce_common.hpp:65
std::pair< long_index_t, long_index_t > get_2d_lengths(const std::vector< index_t > &inLengths)
Definition device_reduce_common.hpp:20
std::vector< index_t > shuffle_tensor_dimensions(const std::vector< index_t > &origLengthsStrides, const std::vector< int > &reduceDims)
Definition device_reduce_common.hpp:75
auto make_tuple_from_array_and_index_seq(const std::vector< index_t > &lengths, Sequence< Ns... >)
Definition device_reduce_common.hpp:59
Definition convolution_backward_data_specialization.hpp:7
Definition ck.hpp:268
__host__ __device__ constexpr auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition tensor_descriptor_helper.hpp:49
__host__ __device__ constexpr auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:37
__host__ __device__ constexpr auto make_merge_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:55
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211
__host__ __device__ constexpr auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition tensor_description/tensor_descriptor.hpp:319
__host__ __device__ constexpr auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:90
auto NormalizationKernelSelector(bool isSweepOnce)
Definition gridwise_normalization_selector.hpp:78
Definition ck/stream_config.hpp:10
Definition utility/sequence.hpp:43
typename conditional< kHasContent, type0, type1 >::type type
Definition utility/sequence.hpp:271
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_normalization_fwd.hpp:23
Definition device_normalization_fwd_impl.hpp:177
const GammaDataType * p_gamma_
Definition device_normalization_fwd_impl.hpp:239
const XDataType * p_x_
Definition device_normalization_fwd_impl.hpp:238
YDataType * p_y_
Definition device_normalization_fwd_impl.hpp:241
GridDesc_M save_mean_grid_desc_m_
Definition device_normalization_fwd_impl.hpp:262
Argument(const std::vector< index_t > lengths, const std::vector< index_t > xStrides, const std::vector< index_t > gammaStrides, const std::vector< index_t > betaStrides, const std::vector< index_t > yStrides, const std::vector< index_t > saveMeanStrides, const std::vector< index_t > saveInvStdStrides, const std::vector< index_t > reduceDims, YElementwiseOperation y_elementwise_op, double epsilon, const XDataType *p_x, const GammaDataType *p_gamma, const BetaDataType *p_beta, YDataType *p_y, SaveMeanInvStdDataType *p_saveMean, SaveMeanInvStdDataType *p_saveInvStd)
Definition device_normalization_fwd_impl.hpp:178
index_t MRaw_
Definition device_normalization_fwd_impl.hpp:266
ComputeDataType epsilon_
Definition device_normalization_fwd_impl.hpp:236
std::vector< index_t > saveInvStdStrides_
Definition device_normalization_fwd_impl.hpp:251
GridDesc_M_K x_grid_desc_m_k_
Definition device_normalization_fwd_impl.hpp:258
SaveMeanInvStdDataType * p_saveInvStd_
Definition device_normalization_fwd_impl.hpp:243
std::vector< index_t > Lengths_
Definition device_normalization_fwd_impl.hpp:245
SaveMeanInvStdDataType * p_saveMean_
Definition device_normalization_fwd_impl.hpp:242
index_t invariant_lowest_length_
Definition device_normalization_fwd_impl.hpp:269
std::vector< index_t > betaStrides_
Definition device_normalization_fwd_impl.hpp:248
std::vector< index_t > saveMeanStrides_
Definition device_normalization_fwd_impl.hpp:250
bool isSweeponce_
Definition device_normalization_fwd_impl.hpp:264
std::vector< index_t > xStrides_
Definition device_normalization_fwd_impl.hpp:246
index_t KRaw_
Definition device_normalization_fwd_impl.hpp:267
std::vector< index_t > gammaStrides_
Definition device_normalization_fwd_impl.hpp:247
GridDesc_M_K y_grid_desc_m_k_
Definition device_normalization_fwd_impl.hpp:261
GridDesc_M_K gamma_grid_desc_m_k_
Definition device_normalization_fwd_impl.hpp:259
GridDesc_M save_inv_std_grid_desc_m_
Definition device_normalization_fwd_impl.hpp:263
YElementwiseOperation y_elementwise_op_
Definition device_normalization_fwd_impl.hpp:253
GridDesc_M_K beta_grid_desc_m_k_
Definition device_normalization_fwd_impl.hpp:260
int numBlockTileIteration_
Definition device_normalization_fwd_impl.hpp:255
std::vector< index_t > yStrides_
Definition device_normalization_fwd_impl.hpp:249
const BetaDataType * p_beta_
Definition device_normalization_fwd_impl.hpp:240
size_t gridSize_
Definition device_normalization_fwd_impl.hpp:256
Definition device_normalization_fwd_impl.hpp:273
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_normalization_fwd_impl.hpp:326
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_normalization_fwd_impl.hpp:274
Definition device_normalization_fwd_impl.hpp:57
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_normalization_fwd_impl.hpp:452
static constexpr index_t M_BlockTileSize
Definition device_normalization_fwd_impl.hpp:76
decltype(MakeSaveMeanInvStdDescriptor_M({1}, {1})) GridDesc_M
Definition device_normalization_fwd_impl.hpp:174
tensor_operation::element_wise::PassThrough PassThrough
Definition device_normalization_fwd_impl.hpp:73
static auto MakeSrc2dDescriptor(const std::vector< index_t > &inLengths, const std::vector< index_t > &inStrides, int numBlockTileIteration)
Definition device_normalization_fwd_impl.hpp:82
static constexpr index_t NumInvariantDim
Definition device_normalization_fwd_impl.hpp:75
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_normalization_fwd_impl.hpp:333
std::unique_ptr< BaseArgument > MakeArgumentPointer(const std::vector< index_t > lengths, const std::vector< index_t > xStrides, const std::vector< index_t > gammaStrides, const std::vector< index_t > betaStrides, const std::vector< index_t > yStrides, const std::vector< index_t > saveMeanStrides, const std::vector< index_t > saveInvStdStrides, const std::vector< index_t > reduceDims, double epsilon, const void *p_x, const void *p_gamma, const void *p_beta, void *p_y, void *p_saveMean, void *p_saveInvStd, YElementwiseOperation y_elementwise_op) override
Definition device_normalization_fwd_impl.hpp:412
static auto MakeSaveMeanInvStdDescriptor_M(const std::vector< index_t > &lengths, const std::vector< index_t > &strides)
Definition device_normalization_fwd_impl.hpp:144
static constexpr index_t K_BlockTileSize
Definition device_normalization_fwd_impl.hpp:77
decltype(MakeSrc2dDescriptor({1}, {1}, 1)) GridDesc_M_K
Definition device_normalization_fwd_impl.hpp:173
std::string GetTypeString() const override
Definition device_normalization_fwd_impl.hpp:457
static constexpr bool reduceAllDim
Definition device_normalization_fwd_impl.hpp:79
Definition tensor_operation/gpu/element/unary_element_wise_operation.hpp:340