Mojo module

swizzle

Defines swizzle layouts for optimizing memory access patterns, particularly in shared memory for GPU kernels, to mitigate bank conflicts.

This module provides tools to create and apply swizzle transformations to memory indices. Swizzling rearranges the memory access order to distribute accesses across different memory banks, thereby reducing bank contention and improving memory access efficiency.

The module includes:

• Swizzle struct: Represents a swizzle transformation with configurable parameters like bits, base, and shift.
• Helper functions: make_ldmatrix_swizzle, make_swizzle to create predefined swizzle patterns optimized for specific scenarios like ldmatrix instructions and general 2D memory access.
• ComposedLayout struct: Allows combining a base layout with a swizzle layout for more complex memory access optimizations.

The primary use case is in GPU kernel development where shared memory bank conflicts can significantly degrade performance. By applying appropriate swizzle layouts, memory access patterns can be optimized to achieve higher throughput.

Structs

• ​ComposedLayout: A layout composed of two layouts applied in sequence.
• ​Swizzle: A swizzle functor for memory access pattern optimization.

Functions

• ​eval_composed: Evaluate a composed layout with swizzle.
• ​make_ldmatrix_swizzle: Make a swizzle to avoid bank conflict for ldmatrix.
• ​make_swizzle: Create a 2D swizzle to avoid bank conflicts.
• ​shiftl: Shift left or right based on sign of shift amount.
• ​shiftr: Shift right or left based on sign of shift amount.

Was this page helpful?

Thank you! We'll create more content like this.

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit