composable_kernel icon indicating copy to clipboard operation
composable_kernel copied to clipboard

Unify the type specification for scaling values and float constant values at the Device Op API

Open qianfengz opened this issue 3 years ago • 0 comments

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:

  1. At each Device Op API, the arguments of scaling and float constant values are declared to be double
  2. Inside the Device Op codes, the scaling and float constant values are converted to some AccDataType and saved as the AccDataType variables
  3. At each kernel interface, the arguments of scaling and float constant values are declared to be some AccDataType

qianfengz avatar Oct 18 '22 10:10 qianfengz