You probably want PackedTensorAccessors for this.
See e.g. the BatchNorm 5 kernels for details (this currently is called GenericPackedTensorAccessor, if you know your required index size, you can use PackedTensorAccessor32 or 64.
They offer array-like elementwise access in kernels and can be passed (by value(!)) to the kernel instead of pointer + dimensions + strides.
Best regards
Thomas