Struct Swizzle 

pub struct Swizzle { /* private fields */ }

Swizzling strategy for a buffer. See the following docs from cutlass:

0bxxxxxxxxxxxxxxxYYYxxxxxxxZZZxxxx ^–^ MBase is the number of least-sig bits to keep constant ^-^ ^-^ BBits is the number of bits in the mask ^———^ SShift is the distance to shift the YYY mask (pos shifts YYY to the right, neg shifts YYY to the left)

Example

Given: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx the result is: 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxAAxxx where AA = ZZ xor YY

Some newer features, as well as cutlass in places, use a different terminology of span and atom. For shared memory swizzle specifically, the parameters map as follows:

• bits = log2(span / atom), or the number of atoms within one span, converted to address bits
• base = log2(atom), the size of the atom, converted to address bits
• shift = log2(all_banks_bytes / atom), or the total number of atoms in all 32 shared memory banks, converted to address bits

For example:

• 32-byte span with a 16-byte atom = [1, 4, 3]
• 128-byte span with a 32-byte atom = [3, 5, 2]

Implementations

impl Swizzle

pub fn new(bits: u32, base: u32, shift: i32) -> Self

Create a new swizzle with comptime parameters

pub fn none() -> Self

Create a new noop swizzle object

pub fn apply(&self, offset: u32, type_size: usize) -> u32

Apply the swizzle to a coordinate with a given item size. This is the size of the full type, including vectorization. offset should be in terms of vectors from the start of the buffer, and the buffer should be aligned to repeats_after. This is to work around the fact we don’t currently support retrieving the actual address of an offset. If you’re using absolute/unvectorized indices, pass E::Scalar::type_size() instead of the full vector size.

pub fn repeats_after(&self) -> u32

After how many elements this pattern repeats. Can be used to align the buffer (i.e. smem) so offsets match addresses.

pub fn __expand_new( scope: &mut Scope, bits: u32, base: u32, shift: i32, ) -> <Self as CubeType>::ExpandType

pub fn __expand_none(scope: &mut Scope) -> <Self as CubeType>::ExpandType

pub fn __expand_apply( scope: &mut Scope, this: <Self as CubeType>::ExpandType, offset: <u32 as CubeType>::ExpandType, type_size: usize, ) -> <u32 as CubeType>::ExpandType

pub fn __expand_repeats_after( scope: &mut Scope, this: <Self as CubeType>::ExpandType, ) -> u32

Trait Implementations

impl Clone for Swizzle

fn clone(&self) -> Swizzle

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl CubeType for Swizzle

type ExpandType = SwizzleExpand

impl LaunchArg for Swizzle

type RuntimeArg<R: Runtime> = SwizzleLaunch<R>

The runtime argument for the kernel.

type CompilationArg = SwizzleCompilationArg

Compilation argument.

fn register<R: Runtime>( arg: Self::RuntimeArg<R>, launcher: &mut KernelLauncher<R>, ) -> Self::CompilationArg

fn expand( arg: &Self::CompilationArg, builder: &mut KernelBuilder, ) -> <Self as CubeType>::ExpandType

Register an input variable during compilation that fill the KernelBuilder.

fn expand_output( arg: &Self::CompilationArg, builder: &mut KernelBuilder, ) -> <Self as CubeType>::ExpandType

Register an output variable during compilation that fill the KernelBuilder.

impl Copy for Swizzle