composable_kernel
composable_kernel copied to clipboard
Unify the type specification for scaling values and float constant values at the Device Op API
Scaling values (eg. alpha and beta) and float constants (eg. epsilon) are used by our kernels and passed from the user through the Device Op API. The following is a list of how those values are used in C.K
- alpha and beta are used in Reduction and Softmax
- epsilon is used in BatchNorm and LayerNorm
- averageFactor is used in BatchNorm
And currently, different component might use different methods to define the types of those values at the Device Op API layer. A unified specification to define these values is necessary to make our Device Op API stable and brings less potential issues.
The following is my suggested method to specify the types of scaling values and float constants:
- At each Device Op API, the arguments of scaling and float constant values are declared to be
double - Inside the Device Op codes, the scaling and float constant values are converted to some
AccDataTypeand saved as theAccDataTypevariables - At each kernel interface, the arguments of scaling and float constant values are declared to be some
AccDataType