# Kernel writing conventions¶

## Variable suffixes¶

We prefer suffixes to prefixes because:

- avoids looking at the middle of the variable to find the interesting part
- It allows quick comparaison, for example
`q = q_ptr + q_offs`

, everything starts we`q`

and we reduce mistakes - It follows the pronunciation

Suffixes list:

`_idx`

integer representing an index`_size`

integer representing a size`_off`

integer representing an offset`_offs`

vector or matrix of integers to be used as offset for pointers`_ptrs`

vector or matrix of pointers`_ptr`

single pointer`_range_offs`

output of`tl.arange(0, N)`

to be used as offset for pointers`_ptr_mask`

bool matrix or vector to be used at mask for load and store operations`_LOAD_MASK_NEEDED`

boolean, usually used as constexpr to determine if we need to use a mask when loading or storing tensor

## Variable prefixes¶

`is_`

,`has_`

or`should_`

for booleans

## Dimension naming¶

- Dimension is singular
- If dimension follows variable name from a formula. You can use this name. Example MNK for matmul
- Use
`col`

or`row`

singular if you don't have a name for the last two dimensions

## Tensor Stride naming¶

`p_d_stride`

`p`

is the name of the tensor`d`

is the name of the dimension

For example:

`q_batch_stride`

`output_row_stride`

## Tensor Size naming¶

`p_d_size`

`p`

is the name of the tensor`d`

is the name of the dimension

For example:

`q_batch_size`

`output_row_size`