Struct TileConfig 

pub struct TileConfig {
    pub block_x: u32,
    pub block_y: u32,
    pub stages: u32,
    pub smem_bytes_per_stage: u32,
    pub cp_async: bool,
    pub tma: bool,
}

Tuned launch geometry and software-pipeline depth for a kernel on a given SM.

Replaces the single portable default (block = 256, 1 stage) that crate::ptx_kernels hard-codes. The values follow the per-SM table in the crate TODO.md “Architecture-Specific Deepening” section:

SM	block_x (correlate)	pipeline stages
75 (Turing)	128	1
80 / 86 (Ampere)	256	2
89 (Ada)	256	2
90 (Hopper)	512	3
100 (Blackwell)	512	3

Fields

block_x: u32

Thread-block extent along x (threads per CTA).

block_y: u32

Thread-block extent along y.

stages: u32

Number of software-pipeline stages (1 = no overlap).

smem_bytes_per_stage: u32

Shared-memory bytes the kernel should stage per pipeline buffer.

cp_async: bool

Whether the SM can issue cp.async global→shared copies (Ampere+).

tma: bool

Whether the SM can issue cp.async.bulk TMA copies (Hopper+).

Implementations

impl TileConfig

pub const fn portable_default() -> Self

The portable default used historically by crate::ptx_kernels: 256×1 threads, single stage, no async copy.

pub fn for_sm(sm: u32) -> Self

Tuned configuration for the correlate kernel on the given SM.

correlate reads one column of Φ per thread streaming over the m rows, so the staged buffer holds block_x f32 residual elements per stage.

pub fn grid_x(self, n: usize) -> u32

Grid extent along x to cover n elements with this block size: ceil(n / block_x), at least 1.

pub fn threads_per_block(self) -> u32

Total threads per CTA (block_x * block_y).

pub fn total_smem_bytes(self) -> u32

Total dynamic shared memory the launch should request (stages * smem_bytes_per_stage).

Trait Implementations

impl Clone for TileConfig

fn clone(&self) -> TileConfig

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Copy for TileConfig

impl Debug for TileConfig

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Eq for TileConfig

impl PartialEq for TileConfig

fn eq(&self, other: &TileConfig) -> bool

Equality operator ==.

fn ne(&self, other: &Rhs) -> bool

Inequality operator !=.

impl StructuralPartialEq for TileConfig