Updates based on feedback
156
docs/conceptual/ck_tile/MERMAID_DIAGRAMS.md
Normal file
@@ -0,0 +1,156 @@
|
||||
# Mermaid Diagram Management
|
||||
|
||||
This document explains how to manage mermaid diagrams in the CK Tile documentation.
|
||||
|
||||
## Overview
|
||||
|
||||
All mermaid diagrams in the CK Tile documentation have been converted to SVG files for better rendering compatibility. The original mermaid source code is preserved as commented blocks in the RST files, allowing easy updates when needed.
|
||||
|
||||
## Directory Structure
|
||||
|
||||
- `docs/conceptual/ck_tile/diagrams/` - Contains all SVG diagram files
|
||||
- `docs/conceptual/ck_tile/convert_mermaid_to_svg.py` - Initial conversion script (one-time use)
|
||||
- `docs/conceptual/ck_tile/update_diagrams.py` - Helper script to regenerate diagrams from comments
|
||||
|
||||
## Diagram Format in RST Files
|
||||
|
||||
Each diagram follows this format:
|
||||
|
||||
```rst
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
A --> B
|
||||
B --> C
|
||||
|
||||
.. image:: diagrams/diagram_name.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
```
|
||||
|
||||
The commented mermaid block won't appear in the rendered documentation but serves as the source for regenerating the SVG.
|
||||
|
||||
## Updating Diagrams
|
||||
|
||||
### When to Update
|
||||
|
||||
You need to regenerate SVG files when:
|
||||
- Modifying the mermaid source in a commented block
|
||||
- Adding new diagrams
|
||||
- Updating diagram styling
|
||||
|
||||
### How to Update
|
||||
|
||||
1. **Edit the commented mermaid source** in the RST file
|
||||
2. **Run the update script**:
|
||||
```bash
|
||||
# Update all diagrams
|
||||
python docs/conceptual/ck_tile/update_diagrams.py
|
||||
|
||||
# Update diagrams in a specific file
|
||||
python docs/conceptual/ck_tile/update_diagrams.py transforms.rst
|
||||
|
||||
# Force regenerate all diagrams (even if SVGs exist)
|
||||
python docs/conceptual/ck_tile/update_diagrams.py --force
|
||||
```
|
||||
|
||||
### Prerequisites
|
||||
|
||||
The update script requires [mermaid-cli](https://github.com/mermaid-js/mermaid-cli):
|
||||
|
||||
```bash
|
||||
npm install -g @mermaid-js/mermaid-cli
|
||||
```
|
||||
|
||||
## Adding New Diagrams
|
||||
|
||||
To add a new mermaid diagram:
|
||||
|
||||
1. **Create the commented block** in your RST file:
|
||||
```rst
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
A --> B
|
||||
```
|
||||
|
||||
2. **Add the image reference** immediately after:
|
||||
```rst
|
||||
.. image:: diagrams/my_new_diagram.svg
|
||||
:alt: My New Diagram
|
||||
:align: center
|
||||
```
|
||||
|
||||
3. **Generate the SVG**:
|
||||
```bash
|
||||
python docs/conceptual/ck_tile/update_diagrams.py your_file.rst
|
||||
```
|
||||
|
||||
## Current Diagrams
|
||||
|
||||
The following RST files contain mermaid diagrams (40 total):
|
||||
|
||||
- `adaptors.rst` (2 diagrams)
|
||||
- `convolution_example.rst` (1 diagram)
|
||||
- `coordinate_movement.rst` (1 diagram)
|
||||
- `descriptors.rst` (2 diagrams)
|
||||
- `encoding_internals.rst` (2 diagrams)
|
||||
- `lds_index_swapping.rst` (3 diagrams)
|
||||
- `load_store_traits.rst` (2 diagrams)
|
||||
- `space_filling_curve.rst` (1 diagram)
|
||||
- `static_distributed_tensor.rst` (1 diagram)
|
||||
- `sweep_tile.rst` (4 diagrams)
|
||||
- `tensor_coordinates.rst` (2 diagrams)
|
||||
- `thread_mapping.rst` (2 diagrams)
|
||||
- `tile_window.rst` (5 diagrams)
|
||||
- `transforms.rst` (12 diagrams)
|
||||
|
||||
## Troubleshooting
|
||||
|
||||
### SVG not generated
|
||||
|
||||
- Check that mermaid-cli is installed: `mmdc --version`
|
||||
- Verify the mermaid syntax is valid
|
||||
- Look for error messages in the script output
|
||||
|
||||
### Diagram not updating
|
||||
|
||||
- Use `--force` flag to regenerate: `python docs/update_diagrams.py --force`
|
||||
- Check that the image reference matches the generated filename
|
||||
|
||||
### Pattern not matching
|
||||
|
||||
If the update script can't find your commented diagram:
|
||||
- Ensure proper indentation (3 spaces for comment block content)
|
||||
- Verify the `.. mermaid::` directive is commented
|
||||
- Check that the image reference immediately follows the comment block
|
||||
|
||||
## Script Details
|
||||
|
||||
### update_diagrams.py
|
||||
|
||||
This script:
|
||||
1. Scans RST files for commented mermaid blocks
|
||||
2. Extracts the mermaid source code
|
||||
3. Converts to SVG using `mmdc`
|
||||
4. Saves to the diagrams directory
|
||||
|
||||
**Usage:**
|
||||
- `python docs/conceptual/ck_tile/update_diagrams.py` - Check all files, update missing SVGs
|
||||
- `python docs/conceptual/ck_tile/update_diagrams.py --force` - Regenerate all SVGs
|
||||
- `python docs/conceptual/ck_tile/update_diagrams.py <file.rst>` - Update specific file
|
||||
|
||||
### convert_mermaid_to_svg.py
|
||||
|
||||
This was the initial conversion script. It:
|
||||
1. Found all active `.. mermaid::` directives
|
||||
2. Converted them to SVGs
|
||||
3. Replaced directives with commented source + image references
|
||||
|
||||
This script was used once for the initial conversion and typically doesn't need to be run again.
|
||||
@@ -6,7 +6,7 @@ Tensor Adaptors - Chaining Transformations
|
||||
Overview
|
||||
--------
|
||||
|
||||
While individual :ref:`transforms <ck_tile_transforms>` are powerful, TensorAdaptors enable the chaining of multiple transforms together to create complex coordinate transformations. Think of adaptors as transformation pipelines that can reshape, reorder, and restructure tensors in sophisticated ways.
|
||||
While individual :ref:`transforms <ck_tile_transforms>` are effective, TensorAdaptors enable the chaining of multiple transforms together to create complex coordinate transformations. Think of adaptors as transformation pipelines that can reshape, reorder, and restructure tensors in advanced ways.
|
||||
|
||||
TensorAdaptors serve as the bridge between individual transforms and the high-level tensor operations used in real applications. They provide a composable abstraction that allows developers to build complex data access patterns from simple building blocks.
|
||||
|
||||
@@ -15,33 +15,41 @@ TensorAdaptor Basics
|
||||
|
||||
A TensorAdaptor encapsulates a sequence of :ref:`coordinate transformations <ck_tile_coordinate_systems>`, managing the flow of coordinates through multiple transform stages:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Adaptor Composition"
|
||||
subgraph "Single Transform"
|
||||
direction TB
|
||||
I1["Input Coords<br/>[0,1,2]"]
|
||||
T1["Transform<br/>(e.g., Transpose)"]
|
||||
O1["Output Coords<br/>[2,0,1]"]
|
||||
I1 --> T1 --> O1
|
||||
end
|
||||
|
||||
subgraph "Chained Transforms"
|
||||
direction TB
|
||||
I2["Input<br/>2D"]
|
||||
T2A["Transform A<br/>(e.g., Merge)"]
|
||||
M2["Intermediate<br/>1D"]
|
||||
T2B["Transform B<br/>(e.g., Pad)"]
|
||||
O2["Output<br/>1D Padded"]
|
||||
I2 --> T2A --> M2 --> T2B --> O2
|
||||
end
|
||||
end
|
||||
|
||||
style T1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T2A fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style T2B fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Adaptor Composition"
|
||||
subgraph "Single Transform"
|
||||
direction TB
|
||||
I1["Input Coords<br/>[0,1,2]"]
|
||||
T1["Transform<br/>(e.g., Transpose)"]
|
||||
O1["Output Coords<br/>[2,0,1]"]
|
||||
I1 --> T1 --> O1
|
||||
end
|
||||
|
||||
subgraph "Chained Transforms"
|
||||
direction TB
|
||||
I2["Input<br/>2D"]
|
||||
T2A["Transform A<br/>(e.g., Merge)"]
|
||||
M2["Intermediate<br/>1D"]
|
||||
T2B["Transform B<br/>(e.g., Pad)"]
|
||||
O2["Output<br/>1D Padded"]
|
||||
I2 --> T2A --> M2 --> T2B --> O2
|
||||
end
|
||||
end
|
||||
|
||||
style T1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T2A fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style T2B fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/adaptors_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Core Components
|
||||
~~~~~~~~~~~~~~~
|
||||
|
||||
@@ -105,44 +113,52 @@ Custom adaptors can be created by specifying exactly which transforms to use and
|
||||
Chaining Adaptors: Building Complex Transformations
|
||||
---------------------------------------------------
|
||||
|
||||
The real power of adaptors comes from chaining multiple transformations together to create sophisticated data access patterns:
|
||||
The real power of adaptors comes from chaining multiple transformations together to create advanced data access patterns:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Adaptor Chaining Flow"
|
||||
subgraph "Adaptor 1"
|
||||
A1I["Bottom Dims<br/>[0,1]"]
|
||||
A1T["Transform:<br/>Merge[2,3]"]
|
||||
A1O["Top Dims<br/>[0]"]
|
||||
end
|
||||
|
||||
subgraph "Adaptor 2"
|
||||
A2I["Bottom Dims<br/>[0]"]
|
||||
A2T["Transform:<br/>Unmerge[2,3]"]
|
||||
A2O["Top Dims<br/>[0,1]"]
|
||||
end
|
||||
|
||||
subgraph "Chained Result"
|
||||
CI["Input 2D<br/>Bottom[0,1]"]
|
||||
CO["Output 2D<br/>Top[0,1]"]
|
||||
end
|
||||
end
|
||||
|
||||
A1I --> A1T
|
||||
A1T --> A1O
|
||||
A1O --> A2I
|
||||
A2I --> A2T
|
||||
A2T --> A2O
|
||||
|
||||
CI --> A1I
|
||||
A2O --> CO
|
||||
|
||||
style A1T fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style A2T fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style CI fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style CO fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Adaptor Chaining Flow"
|
||||
subgraph "Adaptor 1"
|
||||
A1I["Bottom Dims<br/>[0,1]"]
|
||||
A1T["Transform:<br/>Merge[2,3]"]
|
||||
A1O["Top Dims<br/>[0]"]
|
||||
end
|
||||
|
||||
subgraph "Adaptor 2"
|
||||
A2I["Bottom Dims<br/>[0]"]
|
||||
A2T["Transform:<br/>Unmerge[2,3]"]
|
||||
A2O["Top Dims<br/>[0,1]"]
|
||||
end
|
||||
|
||||
subgraph "Chained Result"
|
||||
CI["Input 2D<br/>Bottom[0,1]"]
|
||||
CO["Output 2D<br/>Top[0,1]"]
|
||||
end
|
||||
end
|
||||
|
||||
A1I --> A1T
|
||||
A1T --> A1O
|
||||
A1O --> A2I
|
||||
A2I --> A2T
|
||||
A2T --> A2O
|
||||
|
||||
CI --> A1I
|
||||
A2O --> CO
|
||||
|
||||
style A1T fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style A2T fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style CI fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style CO fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/adaptors_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
.. code-block:: cpp
|
||||
|
||||
// Start with a 2D descriptor
|
||||
@@ -200,7 +216,7 @@ Advanced Patterns
|
||||
Complex Nested Transforms
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
CK Tile supports complex nested transform patterns that enable sophisticated data layouts:
|
||||
CK Tile supports complex nested transform patterns that enable advanced data layouts:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -343,7 +359,7 @@ Key C++ Patterns in Composable Kernel
|
||||
3. **Type Safety**: Template metaprogramming ensures coordinate transformations are type-safe
|
||||
4. **GPU Optimization**: Transform chains are designed for efficient GPU memory access patterns (see :ref:`ck_tile_lds_bank_conflicts` for LDS optimization)
|
||||
|
||||
TensorAdaptors bridge the gap between low-level transforms and high-level tensor operations, providing the flexibility to create sophisticated data layouts and access patterns that are essential for efficient GPU computing. They build upon the foundation of :ref:`BufferViews <ck_tile_buffer_views>` and :ref:`TensorViews <ck_tile_tensor_views>` to provide complex transformation capabilities.
|
||||
TensorAdaptors bridge the gap between low-level transforms and high-level tensor operations, providing the flexibility to create advanced data layouts and access patterns that are essential for efficient GPU computing. They build upon the foundation of :ref:`BufferViews <ck_tile_buffer_views>` and :ref:`TensorViews <ck_tile_tensor_views>` to provide complex transformation capabilities.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
@@ -6,13 +6,13 @@ Buffer Views - Raw Memory Access
|
||||
Overview
|
||||
--------
|
||||
|
||||
At the foundation of the CK Tile system lies BufferView, a sophisticated abstraction that provides structured access to raw memory regions within GPU kernels. This fundamental building block serves as the bridge between the hardware's physical memory model and the higher-level abstractions that enable efficient GPU programming. BufferView encapsulates the complexity of GPU memory hierarchies while exposing a unified interface that works seamlessly across different memory address spaces—whether accessing global memory shared across the entire device, local data share (LDS) memory shared within a workgroup, or the ultra-fast register files private to each thread.
|
||||
At the foundation of the CK Tile system lies BufferView, a compile-time abstraction that provides structured access to raw memory regions within GPU kernels. This fundamental building block serves as the bridge between the hardware's physical memory model and the higher-level abstractions that enable efficient GPU programming. BufferView encapsulates the complexity of GPU memory hierarchies while exposing a unified interface that works seamlessly across different memory address spaces—whether accessing global memory shared across the entire device, local data share (LDS) memory shared within a workgroup, or the ultra-fast register files private to each thread.
|
||||
|
||||
BufferView serves as the foundation for :ref:`ck_tile_tensor_views`, which add multi-dimensional structure on top of raw memory access. Understanding BufferView is essential before moving on to more complex abstractions like :ref:`ck_tile_distribution` and :ref:`ck_tile_window`.
|
||||
|
||||
The design of BufferView reflects a deep understanding of GPU architecture and the performance characteristics that distinguish efficient from inefficient memory access patterns. By providing compile-time knowledge of buffer properties through template metaprogramming, BufferView enables the compiler to generate optimal machine code for each specific use case. This zero-overhead abstraction ensures that the convenience of a high-level interface comes with no runtime performance penalty.
|
||||
|
||||
One of BufferView's most critical features is its sophisticated handling of out-of-bounds memory access. Unlike CPU programming where such accesses typically result in segmentation faults or undefined behavior, GPU programming must gracefully handle cases where threads attempt to access memory beyond allocated boundaries. BufferView provides configurable strategies for these scenarios, allowing developers to choose between returning numerical zero values or custom sentinel values for invalid accesses. This flexibility proves essential for algorithms that naturally extend beyond data boundaries, such as convolutions with padding or matrix operations with non-aligned dimensions.
|
||||
One of BufferView's most critical features is its advanced handling of out-of-bounds memory access. Unlike CPU programming where such accesses typically result in segmentation faults or undefined behavior, GPU programming must gracefully handle cases where threads attempt to access memory beyond allocated boundaries. BufferView provides configurable strategies for these scenarios, allowing developers to choose between returning numerical zero values or custom sentinel values for invalid accesses. This flexibility proves essential for algorithms that naturally extend beyond data boundaries, such as convolutions with padding or matrix operations with non-aligned dimensions.
|
||||
|
||||
The abstraction extends beyond simple memory access to encompass both scalar and vector data types. Modern GPUs achieve their highest efficiency when loading or storing multiple data elements in a single instruction. BufferView seamlessly supports these vectorized operations, automatically selecting the appropriate hardware instructions based on the data type and access pattern. This capability transforms what would be multiple memory transactions into single, efficient operations that fully utilize the available memory bandwidth.
|
||||
|
||||
@@ -117,7 +117,7 @@ The address space template parameter represents another crucial design decision.
|
||||
Out-of-Bounds Handling
|
||||
~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
The sophisticated out-of-bounds handling in BufferView represents a critical innovation for robust GPU programming. Traditional approaches to bounds checking often involve conditional branches that can severely impact performance on GPU architectures, where divergent execution paths within a warp lead to serialization. BufferView's approach elegantly sidesteps this problem through two carefully designed modes that maintain performance while providing predictable behavior.
|
||||
The advanced out-of-bounds handling in BufferView represents a critical innovation for robust GPU programming. Traditional approaches to bounds checking often involve conditional branches that can severely impact performance on GPU architectures, where divergent execution paths within a warp lead to serialization. BufferView's approach cleanly sidesteps this problem through two carefully designed modes that maintain performance while providing predictable behavior.
|
||||
|
||||
The Zero Value Mode leverages the mathematical property that zero often serves as a neutral element in computations. When an access falls outside the valid buffer range, this mode returns numerical zero without branching. This approach proves particularly effective for algorithms like convolution, where out-of-bounds accesses naturally correspond to zero-padding. The branchless implementation ensures that all threads in a warp follow the same execution path, maintaining the SIMD efficiency that is crucial for GPU performance.
|
||||
|
||||
@@ -152,11 +152,11 @@ Get Operations
|
||||
Scalar Access
|
||||
~~~~~~~~~~~~~
|
||||
|
||||
The get operations in BufferView form the cornerstone of memory access patterns in CK Tile. These operations embody a sophisticated understanding of GPU memory systems and the patterns that lead to optimal performance. The scalar access interface, while appearing simple, incorporates multiple layers of optimization and safety mechanisms that work together to provide both performance and correctness.
|
||||
The get operations in BufferView form the cornerstone of memory access patterns in CK Tile. These operations embody a advanced understanding of GPU memory systems and the patterns that lead to optimal performance. The scalar access interface, while appearing simple, incorporates multiple layers of optimization and safety mechanisms that work together to provide both performance and correctness.
|
||||
|
||||
The parameter structure of scalar access operations reflects careful design choices aimed at maximizing flexibility while maintaining efficiency. The base index parameter ``i`` represents the primary offset into the buffer, expressed in terms of elements of type T rather than raw bytes. This type-aware indexing prevents common errors related to pointer arithmetic and ensures that vector types are handled correctly. The additional ``linear_offset`` parameter provides fine-grained control over the final access location, enabling complex access patterns without requiring expensive index calculations in the kernel code.
|
||||
|
||||
The ``is_valid_element`` parameter represents a particularly elegant solution to conditional memory access. Rather than using traditional if-statements that would cause warp divergence, this boolean parameter enables predicated execution where the memory access occurs unconditionally but the result is conditionally used. This approach maintains uniform control flow across all threads in a warp, preserving the SIMD execution model that is fundamental to GPU performance.
|
||||
The ``is_valid_element`` parameter represents a particularly clean solution to conditional memory access. Rather than using traditional if-statements that would cause warp divergence, this boolean parameter enables predicated execution where the memory access occurs unconditionally but the result is conditionally used. This approach maintains uniform control flow across all threads in a warp, preserving the SIMD execution model that is fundamental to GPU performance.
|
||||
|
||||
The invalid value modes provide a crucial mechanism for handling the boundary conditions that naturally arise in parallel algorithms. When ``InvalidElementUseNumericalZeroValue`` is set to true, the system returns zero for any invalid access, whether due to the ``is_valid_element`` flag or out-of-bounds indexing. This mode proves invaluable for algorithms where zero serves as a natural extension value, such as in image processing with zero-padding or sparse matrix operations where missing elements are implicitly zero.
|
||||
|
||||
@@ -167,7 +167,7 @@ Out-of-bounds handling leverages AMD GPU hardware capabilities to provide safety
|
||||
Vector Access
|
||||
~~~~~~~~~~~~~
|
||||
|
||||
Vector memory operations represent one of the most powerful optimizations available in modern GPU programming, and BufferView's vector access interface exposes this capability with elegant simplicity. By using template parameters to specify vector types through constructs like ``ext_vector_t<float, N>``, the interface enables compile-time selection of optimal load and store instructions that can transfer multiple data elements in a single memory transaction. This vectorization is crucial for :ref:`ck_tile_load_store_traits`, which automatically selects optimal access patterns.
|
||||
Vector memory operations represent one of the most critical optimizations available in modern GPU programming, and BufferView's vector access interface exposes this capability with clean simplicity. By using template parameters to specify vector types through constructs like ``ext_vector_t<float, N>``, the interface enables compile-time selection of optimal load and store instructions that can transfer multiple data elements in a single memory transaction. This vectorization is crucial for :ref:`ck_tile_load_store_traits`, which automatically selects optimal access patterns.
|
||||
|
||||
The significance of vector operations extends beyond simple bandwidth improvements. Modern GPUs are designed with wide memory buses that can transfer 128, 256, or even 512 bits per transaction. When scalar operations access only 32 bits at a time, they utilize only a fraction of this available bandwidth. Vector operations align naturally with these wide buses, enabling full bandwidth utilization and reducing the total number of memory transactions required.
|
||||
|
||||
@@ -395,7 +395,7 @@ C++ Atomic Operations
|
||||
Summary
|
||||
-------
|
||||
|
||||
The BufferView abstraction represents a foundational achievement in GPU programming, providing a sophisticated yet accessible interface to the complex memory hierarchies of modern GPUs. Through its careful design and implementation, BufferView demonstrates that high-level abstractions need not come at the cost of performance—indeed, they can enable optimizations that would be impractical or impossible with lower-level approaches.
|
||||
The BufferView abstraction represents a foundational achievement in GPU programming, providing a advanced yet accessible interface to the complex memory hierarchies of modern GPUs. Through its careful design and implementation, BufferView demonstrates that high-level abstractions need not come at the cost of performance—indeed, they can enable optimizations that would be impractical or impossible with lower-level approaches.
|
||||
|
||||
The unified interface that BufferView provides across all memory spaces represents a significant simplification for GPU developers. Whether accessing global memory with its large capacity and high latency, shared memory with its limited size but superior bandwidth, or register files with their extreme speed but strict limitations, developers use the same consistent API. This uniformity reduces cognitive load, prevents errors, and enables code reuse across different memory hierarchies. The abstraction ensures that algorithms can be written once and adapted to different memory configurations through simple template parameter changes.
|
||||
|
||||
@@ -405,7 +405,7 @@ The performance characteristics of BufferView demonstrate the power of thoughtfu
|
||||
|
||||
Flexibility in handling edge cases and boundary conditions proves crucial for real-world algorithms. BufferView's support for configurable invalid value handling, runtime bounds checking when needed, and conditional access patterns accommodates the full spectrum of GPU algorithms. From simple, regular computations to complex, data-dependent access patterns, BufferView provides the tools needed while maintaining performance. This flexibility extends to the atomic operations that are essential for parallel algorithms, enabling thread-safe updates without sacrificing the benefits of the abstraction.
|
||||
|
||||
The true achievement of BufferView lies not in any single feature but in the harmonious integration of all these capabilities. By hiding the complexity of different memory spaces while exposing precisely the operations needed for high-performance GPU computing, BufferView establishes a pattern that the rest of CK Tile follows: sophisticated abstractions that enhance rather than compromise performance. As we build upon this foundation with :ref:`ck_tile_tensor_views` and :ref:`ck_tile_distribution`, the wisdom of this approach becomes increasingly apparent—each layer adds capability while maintaining the efficiency established at the base. For hardware-specific details about memory hierarchies, see :ref:`ck_tile_gpu_basics`.
|
||||
The true achievement of BufferView lies not in any single feature but in the harmonious integration of all these capabilities. By hiding the complexity of different memory spaces while exposing precisely the operations needed for high-performance GPU computing, BufferView establishes a pattern that the rest of CK Tile follows: compile-time abstractions that enhance rather than compromise performance. As we build upon this foundation with :ref:`ck_tile_tensor_views` and :ref:`ck_tile_distribution`, the wisdom of this approach becomes increasingly apparent—each layer adds capability while maintaining the efficiency established at the base. For hardware-specific details about memory hierarchies, see :ref:`ck_tile_gpu_basics`.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
206
docs/conceptual/ck_tile/convert_mermaid_to_svg.py
Normal file
@@ -0,0 +1,206 @@
|
||||
#!/usr/bin/env python3
|
||||
"""
|
||||
Script to convert all mermaid diagrams in CK Tile docs to SVGs.
|
||||
This script:
|
||||
1. Finds all mermaid blocks in RST files
|
||||
2. Converts them to SVG using mmdc
|
||||
3. Updates RST files to use SVG images with commented mermaid source
|
||||
"""
|
||||
|
||||
import os
|
||||
import re
|
||||
import subprocess
|
||||
import tempfile
|
||||
from pathlib import Path
|
||||
|
||||
# Configuration
|
||||
DOCS_DIR = Path(__file__).parent
|
||||
DIAGRAMS_DIR = DOCS_DIR / 'diagrams'
|
||||
RST_FILES = [
|
||||
'convolution_example.rst',
|
||||
'encoding_internals.rst',
|
||||
'lds_index_swapping.rst',
|
||||
'space_filling_curve.rst',
|
||||
'sweep_tile.rst',
|
||||
'tensor_coordinates.rst',
|
||||
'thread_mapping.rst',
|
||||
'static_distributed_tensor.rst',
|
||||
'load_store_traits.rst',
|
||||
'tile_window.rst',
|
||||
'transforms.rst',
|
||||
'descriptors.rst',
|
||||
'coordinate_movement.rst',
|
||||
'adaptors.rst',
|
||||
]
|
||||
|
||||
# Pattern to find mermaid blocks
|
||||
MERMAID_PATTERN = re.compile(
|
||||
r'^\.\. mermaid::\s*\n((?:(?:\n| .*))*)',
|
||||
re.MULTILINE
|
||||
)
|
||||
|
||||
|
||||
def extract_mermaid_content(block):
|
||||
"""Extract the actual mermaid code from the block, removing RST indentation."""
|
||||
lines = block.split('\n')
|
||||
# Remove the leading spaces (RST indentation)
|
||||
content_lines = []
|
||||
for line in lines:
|
||||
if line.startswith(' '):
|
||||
content_lines.append(line[3:]) # Remove 3 spaces
|
||||
elif line.strip() == '':
|
||||
content_lines.append('')
|
||||
return '\n'.join(content_lines).strip()
|
||||
|
||||
|
||||
def generate_diagram_name(file_path, diagram_index, total_in_file):
|
||||
"""Generate a descriptive name for the diagram."""
|
||||
base_name = file_path.stem
|
||||
if total_in_file == 1:
|
||||
return f"{base_name}.svg"
|
||||
else:
|
||||
return f"{base_name}_{diagram_index + 1}.svg"
|
||||
|
||||
|
||||
def convert_mermaid_to_svg(mermaid_code, output_path):
|
||||
"""Convert mermaid code to SVG using mmdc."""
|
||||
# Create a temporary file for the mermaid code
|
||||
with tempfile.NamedTemporaryFile(mode='w', suffix='.mmd', delete=False, encoding='utf-8') as tmp:
|
||||
tmp.write(mermaid_code)
|
||||
tmp_path = tmp.name
|
||||
|
||||
try:
|
||||
# Run mmdc to convert to SVG (use shell=True on Windows for .cmd files)
|
||||
result = subprocess.run(
|
||||
['mmdc', '-i', tmp_path, '-o', str(output_path), '-t', 'neutral', '-b', 'transparent'],
|
||||
capture_output=True,
|
||||
text=True,
|
||||
check=True,
|
||||
shell=True # Required for Windows .cmd files
|
||||
)
|
||||
print(f" ✓ Generated: {output_path.name}")
|
||||
return True
|
||||
except subprocess.CalledProcessError as e:
|
||||
print(f" ✗ Error converting diagram: {e.stderr}")
|
||||
return False
|
||||
finally:
|
||||
# Clean up temp file
|
||||
os.unlink(tmp_path)
|
||||
|
||||
|
||||
def update_rst_file(file_path, diagrams_info):
|
||||
"""Update RST file to replace mermaid blocks with commented source + image reference."""
|
||||
with open(file_path, 'r', encoding='utf-8') as f:
|
||||
content = f.read()
|
||||
|
||||
# Sort diagrams by position (reverse order to maintain positions)
|
||||
diagrams_info.sort(key=lambda x: x['position'], reverse=True)
|
||||
|
||||
for info in diagrams_info:
|
||||
# Find the mermaid block
|
||||
match = info['match']
|
||||
start_pos = match.start()
|
||||
end_pos = match.end()
|
||||
|
||||
# Create the replacement text
|
||||
mermaid_block = match.group(0)
|
||||
mermaid_content = match.group(1)
|
||||
|
||||
# Create commented mermaid block
|
||||
commented_lines = ['.. ', ' Original mermaid diagram (edit here, then run update_diagrams.py)', ' ']
|
||||
for line in mermaid_block.split('\n'):
|
||||
commented_lines.append(f' {line}')
|
||||
|
||||
# Add image reference
|
||||
svg_rel_path = f"diagrams/{info['svg_name']}"
|
||||
image_block = [
|
||||
'',
|
||||
f'.. image:: {svg_rel_path}',
|
||||
' :alt: Diagram',
|
||||
' :align: center',
|
||||
''
|
||||
]
|
||||
|
||||
replacement = '\n'.join(commented_lines + image_block)
|
||||
|
||||
# Replace in content
|
||||
content = content[:start_pos] + replacement + content[end_pos:]
|
||||
|
||||
# Write back
|
||||
with open(file_path, 'w', encoding='utf-8') as f:
|
||||
f.write(content)
|
||||
|
||||
print(f" ✓ Updated: {file_path.name}")
|
||||
|
||||
|
||||
def process_file(file_path):
|
||||
"""Process a single RST file."""
|
||||
print(f"\nProcessing {file_path.name}...")
|
||||
|
||||
with open(file_path, 'r', encoding='utf-8') as f:
|
||||
content = f.read()
|
||||
|
||||
# Find all mermaid blocks
|
||||
matches = list(MERMAID_PATTERN.finditer(content))
|
||||
|
||||
if not matches:
|
||||
print(f" No mermaid diagrams found.")
|
||||
return
|
||||
|
||||
print(f" Found {len(matches)} diagram(s)")
|
||||
|
||||
diagrams_info = []
|
||||
|
||||
# Process each mermaid block
|
||||
for idx, match in enumerate(matches):
|
||||
mermaid_content = extract_mermaid_content(match.group(1))
|
||||
svg_name = generate_diagram_name(file_path, idx, len(matches))
|
||||
svg_path = DIAGRAMS_DIR / svg_name
|
||||
|
||||
# Convert to SVG
|
||||
if convert_mermaid_to_svg(mermaid_content, svg_path):
|
||||
diagrams_info.append({
|
||||
'match': match,
|
||||
'svg_name': svg_name,
|
||||
'position': match.start()
|
||||
})
|
||||
|
||||
# Update the RST file
|
||||
if diagrams_info:
|
||||
update_rst_file(file_path, diagrams_info)
|
||||
|
||||
|
||||
def main():
|
||||
"""Main function."""
|
||||
print("CK Tile Mermaid to SVG Converter")
|
||||
print("=" * 50)
|
||||
|
||||
# Verify mmdc is available
|
||||
try:
|
||||
subprocess.run(['mmdc', '--version'], capture_output=True, check=True, shell=True)
|
||||
except (subprocess.CalledProcessError, FileNotFoundError):
|
||||
print("Error: mermaid-cli (mmdc) not found. Please install it:")
|
||||
print(" npm install -g @mermaid-js/mermaid-cli")
|
||||
return 1
|
||||
|
||||
# Ensure diagrams directory exists
|
||||
DIAGRAMS_DIR.mkdir(parents=True, exist_ok=True)
|
||||
|
||||
# Process each file
|
||||
total_diagrams = 0
|
||||
for rst_file in RST_FILES:
|
||||
file_path = DOCS_DIR / rst_file
|
||||
if file_path.exists():
|
||||
process_file(file_path)
|
||||
else:
|
||||
print(f"\n⚠ Warning: {rst_file} not found")
|
||||
|
||||
print("\n" + "=" * 50)
|
||||
print("✓ Conversion complete!")
|
||||
print(f"SVG files saved to: {DIAGRAMS_DIR}")
|
||||
|
||||
return 0
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
exit(main())
|
||||
@@ -15,40 +15,48 @@ This chapter demonstrates how CK Tile's :ref:`tensor descriptor <ck_tile_descrip
|
||||
|
||||
The key insight is that convolution can be transformed from a complex nested loop operation into a highly parallel matrix multiplication through the im2col (image to column) transformation. CK Tile's tensor descriptors provide the perfect abstraction for implementing this transformation efficiently without data duplication.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Convolution Process"
|
||||
I["Input Image<br/>6×6"]
|
||||
K["Kernel<br/>3×3"]
|
||||
SW["Sliding Window<br/>Extract 3×3 patches"]
|
||||
DP["Dot Product<br/>Element-wise multiply & sum"]
|
||||
O["Output<br/>4×4"]
|
||||
end
|
||||
|
||||
subgraph "Im2col Optimization"
|
||||
W["Windows Matrix<br/>16×9<br/>(all patches)"]
|
||||
KF["Kernel Flattened<br/>9×1"]
|
||||
MM["Matrix Multiply<br/>W @ K"]
|
||||
OF["Output Flattened<br/>16×1"]
|
||||
end
|
||||
|
||||
I --> SW
|
||||
K --> DP
|
||||
SW --> DP
|
||||
DP --> O
|
||||
|
||||
SW --> W
|
||||
K --> KF
|
||||
W --> MM
|
||||
KF --> MM
|
||||
MM --> OF
|
||||
OF --> O
|
||||
|
||||
style I fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style O fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style MM fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Convolution Process"
|
||||
I["Input Image<br/>6×6"]
|
||||
K["Kernel<br/>3×3"]
|
||||
SW["Sliding Window<br/>Extract 3×3 patches"]
|
||||
DP["Dot Product<br/>Element-wise multiply & sum"]
|
||||
O["Output<br/>4×4"]
|
||||
end
|
||||
|
||||
subgraph "Im2col Optimization"
|
||||
W["Windows Matrix<br/>16×9<br/>(all patches)"]
|
||||
KF["Kernel Flattened<br/>9×1"]
|
||||
MM["Matrix Multiply<br/>W @ K"]
|
||||
OF["Output Flattened<br/>16×1"]
|
||||
end
|
||||
|
||||
I --> SW
|
||||
K --> DP
|
||||
SW --> DP
|
||||
DP --> O
|
||||
|
||||
SW --> W
|
||||
K --> KF
|
||||
W --> MM
|
||||
KF --> MM
|
||||
MM --> OF
|
||||
OF --> O
|
||||
|
||||
style I fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style O fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style MM fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/convolution_example.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Understanding Sliding Windows
|
||||
=============================
|
||||
|
||||
@@ -164,7 +172,7 @@ This implementation directly follows the mathematical definition but has poor me
|
||||
Window Extraction with Tensor Descriptors
|
||||
=========================================
|
||||
|
||||
CK Tile's tensor descriptors provide an elegant way to extract convolution windows:
|
||||
CK Tile's tensor descriptors provide an clean way to extract convolution windows:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -364,7 +372,7 @@ Combining all components into an optimized convolution implementation:
|
||||
Multi-Channel Convolution
|
||||
=========================
|
||||
|
||||
Real-world convolutions involve multiple input and output channels. CK Tile handles this elegantly:
|
||||
Real-world convolutions involve multiple input and output channels. CK Tile handles this cleanly:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
|
||||
@@ -13,38 +13,46 @@ Overview
|
||||
|
||||
Advanced coordinate operations form the bridge between mathematical transformations and practical tensor manipulation in CK Tile. These operations enable efficient navigation through complex tensor layouts without recalculating entire transformation chains. Understanding coordinate movement is essential for implementing high-performance GPU kernels that traverse multi-dimensional data structures.
|
||||
|
||||
The coordinate movement system provides two key abstractions: TensorCoordinate for descriptor-aware navigation and TensorAdaptorCoordinate for tracking positions through transformation chains. Together with movement functions, they enable sophisticated access patterns while maintaining optimal performance through incremental updates rather than full recalculation.
|
||||
The coordinate movement system provides two key abstractions: TensorCoordinate for descriptor-aware navigation and TensorAdaptorCoordinate for tracking positions through transformation chains. Together with movement functions, they enable advanced access patterns while maintaining optimal performance through incremental updates rather than full recalculation.
|
||||
|
||||
For the mathematical foundations of coordinate systems, see :ref:`ck_tile_coordinate_systems`. For simpler coordinate concepts, see :ref:`ck_tile_tensor_coordinates`.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Coordinate Movement System"
|
||||
TC["TensorCoordinate<br/>Position + Descriptor Context"]
|
||||
TAC["TensorAdaptorCoordinate<br/>Position + Transform Context"]
|
||||
MC["move_coordinate()<br/>Efficient Navigation"]
|
||||
end
|
||||
|
||||
subgraph "Movement Example"
|
||||
S["Start: [1,1]<br/>Offset: 5"]
|
||||
M1["Move [0,1]<br/>→ [1,2]<br/>Offset: 6"]
|
||||
M2["Move [1,0]<br/>→ [2,2]<br/>Offset: 10"]
|
||||
M3["Move [1,1]<br/>→ [3,3]<br/>Offset: 15"]
|
||||
end
|
||||
|
||||
TC --> MC
|
||||
TAC --> MC
|
||||
|
||||
S --> M1
|
||||
M1 --> M2
|
||||
M2 --> M3
|
||||
|
||||
style TC fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style TAC fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style MC fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Coordinate Movement System"
|
||||
TC["TensorCoordinate<br/>Position + Descriptor Context"]
|
||||
TAC["TensorAdaptorCoordinate<br/>Position + Transform Context"]
|
||||
MC["move_coordinate()<br/>Efficient Navigation"]
|
||||
end
|
||||
|
||||
subgraph "Movement Example"
|
||||
S["Start: [1,1]<br/>Offset: 5"]
|
||||
M1["Move [0,1]<br/>→ [1,2]<br/>Offset: 6"]
|
||||
M2["Move [1,0]<br/>→ [2,2]<br/>Offset: 10"]
|
||||
M3["Move [1,1]<br/>→ [3,3]<br/>Offset: 15"]
|
||||
end
|
||||
|
||||
TC --> MC
|
||||
TAC --> MC
|
||||
|
||||
S --> M1
|
||||
M1 --> M2
|
||||
M2 --> M3
|
||||
|
||||
style TC fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style TAC fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style MC fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/coordinate_movement.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
TensorCoordinate: Descriptor-Aware Navigation
|
||||
=============================================
|
||||
|
||||
@@ -299,7 +307,7 @@ Movement Through Adaptors
|
||||
Advanced Movement Patterns
|
||||
==========================
|
||||
|
||||
Real-world applications use sophisticated movement patterns for optimal memory access. These patterns often relate to :ref:`ck_tile_tile_window` operations and :ref:`ck_tile_tile_distribution` concepts:
|
||||
Real-world applications use advanced movement patterns for optimal memory access. These patterns often relate to :ref:`ck_tile_tile_window` operations and :ref:`ck_tile_tile_distribution` concepts:
|
||||
|
||||
Tiled Access Pattern
|
||||
--------------------
|
||||
@@ -519,7 +527,7 @@ Advanced coordinate operations provide the foundation for efficient tensor navig
|
||||
- **TensorCoordinate**: Combines position with descriptor context for validated navigation
|
||||
- **TensorAdaptorCoordinate**: Tracks coordinates through transformation chains
|
||||
- **move_tensor_coordinate**: Enables efficient incremental updates without recalculation
|
||||
- **Movement Patterns**: Support sophisticated access patterns like tiling and space-filling curves
|
||||
- **Movement Patterns**: Support advanced access patterns like tiling and space-filling curves
|
||||
- **Performance**: Incremental updates are orders of magnitude faster than coordinate recreation
|
||||
- **Integration**: Seamlessly works with tile windows, distributions, and other CK Tile components
|
||||
|
||||
|
||||
@@ -6,7 +6,7 @@ Coordinate Systems - The Mathematical Foundation
|
||||
Overview
|
||||
--------
|
||||
|
||||
At the heart of the Composable Kernel framework lies a sophisticated mathematical foundation based on coordinate transformations. This foundation enables the automatic generation of optimal memory access patterns while maintaining a clear separation between algorithmic intent and hardware implementation details. The coordinate system framework represents one of CK's most powerful innovations, transforming the complex task of GPU work distribution into a series of well-defined mathematical transformations.
|
||||
At the heart of the Composable Kernel framework lies a advanced mathematical foundation based on coordinate transformations. This foundation enables the automatic generation of optimal memory access patterns while maintaining a clear separation between algorithmic intent and hardware implementation details. The coordinate system framework represents one of CK's most key innovations, transforming the complex task of GPU work distribution into a series of well-defined mathematical transformations.
|
||||
|
||||
Understanding these coordinate systems is essential for mastering :ref:`tile distribution <ck_tile_distribution>`. They provide the mathematical machinery that maps from abstract thread identities to concrete memory addresses, ensuring that every memory access is optimized for the underlying hardware. This systematic approach eliminates the error-prone manual calculations that plague traditional GPU programming while enabling optimizations that would be impractical to implement by hand.
|
||||
|
||||
@@ -314,7 +314,7 @@ This transformation is highly configurable through the distribution encoding, en
|
||||
R-Space: Replication and Cooperation
|
||||
------------------------------------
|
||||
|
||||
R-space (Replication Space) introduces a sophisticated mechanism for expressing data sharing and cooperation patterns between threads. Unlike the other coordinate spaces which map to unique data elements, R-space enables multiple processing elements to work on the same data, facilitating efficient communication and reduction operations.
|
||||
R-space (Replication Space) introduces a advanced mechanism for expressing data sharing and cooperation patterns between threads. Unlike the other coordinate spaces which map to unique data elements, R-space enables multiple processing elements to work on the same data, facilitating efficient communication and reduction operations.
|
||||
|
||||
Replication Patterns
|
||||
~~~~~~~~~~~~~~~~~~~~
|
||||
@@ -351,7 +351,7 @@ Replication Patterns
|
||||
}
|
||||
}
|
||||
|
||||
R-space enables sophisticated cooperation patterns that would be difficult to express otherwise. By providing a systematic way to identify which threads share data, it enables automatic generation of efficient communication patterns.
|
||||
R-space enables advanced cooperation patterns that would be difficult to express otherwise. By providing a systematic way to identify which threads share data, it enables automatic generation of efficient communication patterns.
|
||||
|
||||
D-Space: Memory Linearization
|
||||
-----------------------------
|
||||
@@ -514,7 +514,7 @@ The coordinate system framework enables several critical optimizations:
|
||||
Summary
|
||||
-------
|
||||
|
||||
The coordinate system framework represents the mathematical foundation that enables CK's remarkable performance and productivity benefits. Through the systematic transformation from thread identity (P-space) through logical work organization (Y-space) to physical tensor coordinates (X-space) and finally to linear memory addresses (D-space), this framework solves the fundamental challenges of GPU programming.
|
||||
The coordinate system framework represents the mathematical foundation that enables CK's high performance and productivity benefits. Through the systematic transformation from thread identity (P-space) through logical work organization (Y-space) to physical tensor coordinates (X-space) and finally to linear memory addresses (D-space), this framework solves the fundamental challenges of GPU programming.
|
||||
|
||||
Key insights from the coordinate system framework:
|
||||
|
||||
|
||||
@@ -96,33 +96,41 @@ The Pipeline Concept
|
||||
|
||||
Every TensorDescriptor in CK Tile can be thought of as a **transformation pipeline**. The functions above create the *first stage* of this pipeline: they define the initial :ref:`transformation <ck_tile_transforms>` that takes a simple, one-dimensional block of memory and presents it as a logical, multi-dimensional tensor view.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Pipeline Stages"
|
||||
S1["Stage 1<br/>Base Layout<br/>[M, N]"]
|
||||
S2["Stage 2<br/>Transform<br/>Unmerge"]
|
||||
S3["Stage 3<br/>New View<br/>[M1, M2, N]"]
|
||||
S4["Stage N<br/>Final View<br/>[...]"]
|
||||
end
|
||||
|
||||
subgraph "Same Data"
|
||||
D["Physical Memory<br/>No data movement"]
|
||||
end
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
S1 -.-> D
|
||||
S2 -.-> D
|
||||
S3 -.-> D
|
||||
S4 -.-> D
|
||||
|
||||
style D fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
style S1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style S3 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Pipeline Stages"
|
||||
S1["Stage 1<br/>Base Layout<br/>[M, N]"]
|
||||
S2["Stage 2<br/>Transform<br/>Unmerge"]
|
||||
S3["Stage 3<br/>New View<br/>[M1, M2, N]"]
|
||||
S4["Stage N<br/>Final View<br/>[...]"]
|
||||
end
|
||||
|
||||
subgraph "Same Data"
|
||||
D["Physical Memory<br/>No data movement"]
|
||||
end
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
S1 -.-> D
|
||||
S2 -.-> D
|
||||
S3 -.-> D
|
||||
S4 -.-> D
|
||||
|
||||
style D fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
style S1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style S3 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/descriptors_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
The Initial Pipeline Stage
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
@@ -186,38 +194,46 @@ To get from [2, 6] to [2, 2, 3], we need:
|
||||
Analysis of the Final Pipeline
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Transform Pipeline"
|
||||
T0["Transform 0<br/>Base Unmerge<br/>Input: [0]<br/>Output: [1,2]"]
|
||||
T1["Transform 1<br/>PassThrough<br/>Input: [1]<br/>Output: [3]"]
|
||||
T2["Transform 2<br/>Unmerge<br/>Input: [2]<br/>Output: [4,5]"]
|
||||
end
|
||||
|
||||
subgraph "Hidden Dimensions"
|
||||
H0["Hidden ID 0<br/>Raw Buffer"]
|
||||
H1["Hidden ID 1<br/>Dim 0 (size 2)"]
|
||||
H2["Hidden ID 2<br/>Dim 1 (size 6)"]
|
||||
H3["Hidden ID 3<br/>Final Dim 0"]
|
||||
H4["Hidden ID 4<br/>Final Dim 1"]
|
||||
H5["Hidden ID 5<br/>Final Dim 2"]
|
||||
end
|
||||
|
||||
H0 --> T0
|
||||
T0 --> H1
|
||||
T0 --> H2
|
||||
H1 --> T1
|
||||
H2 --> T2
|
||||
T1 --> H3
|
||||
T2 --> H4
|
||||
T2 --> H5
|
||||
|
||||
style H0 fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
style H3 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style H4 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style H5 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Transform Pipeline"
|
||||
T0["Transform 0<br/>Base Unmerge<br/>Input: [0]<br/>Output: [1,2]"]
|
||||
T1["Transform 1<br/>PassThrough<br/>Input: [1]<br/>Output: [3]"]
|
||||
T2["Transform 2<br/>Unmerge<br/>Input: [2]<br/>Output: [4,5]"]
|
||||
end
|
||||
|
||||
subgraph "Hidden Dimensions"
|
||||
H0["Hidden ID 0<br/>Raw Buffer"]
|
||||
H1["Hidden ID 1<br/>Dim 0 (size 2)"]
|
||||
H2["Hidden ID 2<br/>Dim 1 (size 6)"]
|
||||
H3["Hidden ID 3<br/>Final Dim 0"]
|
||||
H4["Hidden ID 4<br/>Final Dim 1"]
|
||||
H5["Hidden ID 5<br/>Final Dim 2"]
|
||||
end
|
||||
|
||||
H0 --> T0
|
||||
T0 --> H1
|
||||
T0 --> H2
|
||||
H1 --> T1
|
||||
H2 --> T2
|
||||
T1 --> H3
|
||||
T2 --> H4
|
||||
T2 --> H5
|
||||
|
||||
style H0 fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
style H3 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style H4 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style H5 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/descriptors_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
The pipeline now has three stages:
|
||||
|
||||
1. **Base UnmergeTransform**: Converts raw buffer to [2, 6] layout
|
||||
@@ -320,7 +336,7 @@ Tensor Slicing
|
||||
Key Concepts Summary
|
||||
--------------------
|
||||
|
||||
TensorDescriptors provide a powerful abstraction for tensor manipulation:
|
||||
TensorDescriptors provide a key abstraction for tensor manipulation:
|
||||
|
||||
- **Pipeline Architecture**: Each descriptor is a transformation pipeline
|
||||
- **Zero-Copy Views**: All transformations are logical, no data movement
|
||||
|
||||
1
docs/conceptual/ck_tile/diagrams/adaptors_1.svg
Normal file
|
After Width: | Height: | Size: 17 KiB |
1
docs/conceptual/ck_tile/diagrams/adaptors_2.svg
Normal file
|
After Width: | Height: | Size: 18 KiB |
1
docs/conceptual/ck_tile/diagrams/convolution_example.svg
Normal file
|
After Width: | Height: | Size: 20 KiB |
1
docs/conceptual/ck_tile/diagrams/coordinate_movement.svg
Normal file
|
After Width: | Height: | Size: 15 KiB |
1
docs/conceptual/ck_tile/diagrams/descriptors_1.svg
Normal file
|
After Width: | Height: | Size: 16 KiB |
1
docs/conceptual/ck_tile/diagrams/descriptors_2.svg
Normal file
|
After Width: | Height: | Size: 18 KiB |
|
After Width: | Height: | Size: 24 KiB |
|
After Width: | Height: | Size: 16 KiB |
|
After Width: | Height: | Size: 16 KiB |
|
After Width: | Height: | Size: 17 KiB |
|
After Width: | Height: | Size: 18 KiB |
1
docs/conceptual/ck_tile/diagrams/load_store_traits_1.svg
Normal file
|
After Width: | Height: | Size: 22 KiB |
1
docs/conceptual/ck_tile/diagrams/load_store_traits_2.svg
Normal file
|
After Width: | Height: | Size: 16 KiB |
1
docs/conceptual/ck_tile/diagrams/space_filling_curve.svg
Normal file
|
After Width: | Height: | Size: 16 KiB |
|
After Width: | Height: | Size: 17 KiB |
1
docs/conceptual/ck_tile/diagrams/sweep_tile_1.svg
Normal file
|
After Width: | Height: | Size: 15 KiB |
1
docs/conceptual/ck_tile/diagrams/sweep_tile_2.svg
Normal file
|
After Width: | Height: | Size: 20 KiB |
1
docs/conceptual/ck_tile/diagrams/sweep_tile_3.svg
Normal file
|
After Width: | Height: | Size: 19 KiB |
1
docs/conceptual/ck_tile/diagrams/sweep_tile_4.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
|
After Width: | Height: | Size: 17 KiB |
|
After Width: | Height: | Size: 16 KiB |
1
docs/conceptual/ck_tile/diagrams/thread_mapping_1.svg
Normal file
|
After Width: | Height: | Size: 21 KiB |
1
docs/conceptual/ck_tile/diagrams/thread_mapping_2.svg
Normal file
|
After Width: | Height: | Size: 18 KiB |
1
docs/conceptual/ck_tile/diagrams/tile_window_1.svg
Normal file
|
After Width: | Height: | Size: 25 KiB |
1
docs/conceptual/ck_tile/diagrams/tile_window_2.svg
Normal file
|
After Width: | Height: | Size: 16 KiB |
1
docs/conceptual/ck_tile/diagrams/tile_window_3.svg
Normal file
|
After Width: | Height: | Size: 16 KiB |
1
docs/conceptual/ck_tile/diagrams/tile_window_4.svg
Normal file
|
After Width: | Height: | Size: 18 KiB |
1
docs/conceptual/ck_tile/diagrams/tile_window_5.svg
Normal file
|
After Width: | Height: | Size: 14 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_1.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_10.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_11.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_12.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_2.svg
Normal file
|
After Width: | Height: | Size: 19 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_3.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_4.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_5.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_6.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_7.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_8.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
1
docs/conceptual/ck_tile/diagrams/transforms_9.svg
Normal file
|
After Width: | Height: | Size: 12 KiB |
@@ -8,60 +8,65 @@
|
||||
Encoding Internals
|
||||
******************
|
||||
|
||||
|
||||
The tile distribution encoding system represents the core mathematical framework that transforms high-level tensor distribution specifications into concrete, optimized GPU kernel implementations. This sophisticated compile-time machinery bridges the gap between abstract mathematical descriptions and executable coordinate transformations, enabling the Composable Kernel framework to generate highly efficient code for complex tensor operations.
|
||||
Overview
|
||||
|
||||
The tile distribution encoding system represents the core mathematical framework that transforms high-level tensor distribution specifications into concrete, optimized GPU kernel implementations. This sophisticated compile-time machinery bridges the gap between abstract mathematical descriptions and executable coordinate transformations, enabling the Composable Kernel framework to generate highly efficient code for complex tensor operations.
|
||||
========
|
||||
|
||||
The tile distribution encoding system represents the core mathematical framework that transforms high-level tensor distribution specifications into concrete, optimized GPU kernel implementations. This sophisticated compile-time machinery bridges the gap between abstract mathematical descriptions and executable coordinate transformations, enabling the Composable Kernel framework to generate highly efficient code for complex tensor operations.
|
||||
The tile distribution encoding system represents the core mathematical framework that transforms high-level tensor distribution specifications into concrete, optimized GPU kernel implementations. This advanced compile-time machinery bridges the gap between abstract mathematical descriptions and executable coordinate transformations, enabling the Composable Kernel framework to generate highly efficient code for complex tensor operations.
|
||||
|
||||
At its heart, the encoding system defines how multi-dimensional tensor data is distributed across GPU processing elements through a hierarchical decomposition scheme. By specifying relationships between different coordinate spaces - replication (R), hierarchical (H), partition (P), and yield (Y) dimensions (see :ref:`ck_tile_coordinate_systems` for detailed explanation) - the encoding provides a complete blueprint for data layout and access patterns that can be resolved entirely at compile time. This is the internal mechanism behind :ref:`ck_tile_tile_distribution`.
|
||||
|
||||
.. mermaid::
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Encoding Components"
|
||||
RS["R-space Lengths<br/>Replication dimensions"]
|
||||
HS["H-space Lengths<br/>Hierarchical decomposition<br/>[[2,2],[2,2]]"]
|
||||
P2RH["P→RH Mappings<br/>Thread to hierarchy<br/>Major/Minor"]
|
||||
Y2RH["Y→RH Mappings<br/>Element to hierarchy<br/>Major/Minor"]
|
||||
end
|
||||
|
||||
subgraph "Generated Components"
|
||||
ADAPTOR["ps_ys_to_xs_adaptor<br/>Coordinate transformer"]
|
||||
DESC["ys_to_d_descriptor<br/>Memory linearizer"]
|
||||
ENC["Encoding<br/>Original specification"]
|
||||
end
|
||||
|
||||
subgraph "Transformation Chain"
|
||||
T1["Replicate<br/>Transform"]
|
||||
T2["Unmerge<br/>Transform"]
|
||||
T3["Merge<br/>Transform"]
|
||||
end
|
||||
|
||||
RS --> T1
|
||||
HS --> T2
|
||||
P2RH --> ADAPTOR
|
||||
Y2RH --> ADAPTOR
|
||||
|
||||
T1 --> T2
|
||||
T2 --> T3
|
||||
T3 --> ADAPTOR
|
||||
|
||||
HS --> DESC
|
||||
Y2RH --> DESC
|
||||
|
||||
style RS fill:#fce4ec,stroke:#c2185b,stroke-width:2px
|
||||
style HS fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style ADAPTOR fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style DESC fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
|
||||
|
||||
|
||||
graph TB
|
||||
subgraph "Encoding Components"
|
||||
RS["R-space Lengths<br/>Replication dimensions"]
|
||||
HS["H-space Lengths<br/>Hierarchical decomposition<br/>[[2,2],[2,2]]"]
|
||||
P2RH["P→RH Mappings<br/>Thread to hierarchy<br/>Major/Minor"]
|
||||
Y2RH["Y→RH Mappings<br/>Element to hierarchy<br/>Major/Minor"]
|
||||
end
|
||||
|
||||
subgraph "Generated Components"
|
||||
ADAPTOR["ps_ys_to_xs_adaptor<br/>Coordinate transformer"]
|
||||
DESC["ys_to_d_descriptor<br/>Memory linearizer"]
|
||||
ENC["Encoding<br/>Original specification"]
|
||||
end
|
||||
|
||||
subgraph "Transformation Chain"
|
||||
T1["Replicate<br/>Transform"]
|
||||
T2["Unmerge<br/>Transform"]
|
||||
T3["Merge<br/>Transform"]
|
||||
end
|
||||
|
||||
RS --> T1
|
||||
HS --> T2
|
||||
P2RH --> ADAPTOR
|
||||
Y2RH --> ADAPTOR
|
||||
|
||||
T1 --> T2
|
||||
T2 --> T3
|
||||
T3 --> ADAPTOR
|
||||
|
||||
HS --> DESC
|
||||
Y2RH --> DESC
|
||||
|
||||
style RS fill:#fce4ec,stroke:#c2185b,stroke-width:2px
|
||||
style HS fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style ADAPTOR fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style DESC fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
.. image:: diagrams/encoding_internals_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
|
||||
Encoding Structure
|
||||
==================
|
||||
|
||||
The tile distribution encoding employs a sophisticated type system that captures the complete specification of tensor distribution patterns at compile time:
|
||||
The tile distribution encoding employs a template-based type system that captures the complete specification of tensor distribution patterns at compile time:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -93,7 +98,7 @@ The tile distribution encoding employs a sophisticated type system that captures
|
||||
// Precomputed mappings and transformations
|
||||
static constexpr auto get_h_dim_lengths_prefix_sum();
|
||||
static constexpr auto get_uniformed_idx_y_to_h();
|
||||
// ... extensive compile-time computation ...
|
||||
// ... compile-time computation ...
|
||||
};
|
||||
};
|
||||
|
||||
@@ -193,44 +198,49 @@ The ``Ys2RHsMajor`` and ``Ys2RHsMinor`` define the user-facing interface:
|
||||
// Y[1,0] → H1[0], H2[1]
|
||||
// Y[1,1] → H1[1], H2[1]
|
||||
|
||||
|
||||
The encoding generates a transformation pipeline that converts coordinates:
|
||||
Transformation Pipeline
|
||||
|
||||
The encoding generates a transformation pipeline that converts coordinates (using the concepts from :ref:`ck_tile_transforms` and :ref:`ck_tile_adaptors`):
|
||||
=======================
|
||||
|
||||
The encoding generates a transformation pipeline that converts coordinates:
|
||||
The encoding generates a transformation pipeline that converts coordinates (using the concepts from :ref:`ck_tile_transforms` and :ref:`ck_tile_adaptors`):
|
||||
|
||||
.. mermaid::
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart LR
|
||||
subgraph "Input Coordinates"
|
||||
P["P-coordinates<br/>[warp_id, lane_id]"]
|
||||
Y["Y-coordinates<br/>[y0, y1, y2, y3]"]
|
||||
end
|
||||
|
||||
subgraph "Transformation Pipeline"
|
||||
C1["Combine P+Y"]
|
||||
T1["Replicate<br/>Transform<br/>(if R-dims exist)"]
|
||||
T2["Unmerge<br/>Transform<br/>(break into H-dims)"]
|
||||
T3["Merge<br/>Transform<br/>(combine to X-dims)"]
|
||||
end
|
||||
|
||||
subgraph "Output"
|
||||
X["X-coordinates<br/>[x0, x1]<br/>Tensor position"]
|
||||
end
|
||||
|
||||
P --> C1
|
||||
Y --> C1
|
||||
C1 --> T1
|
||||
T1 --> T2
|
||||
T2 --> T3
|
||||
T3 --> X
|
||||
|
||||
style P fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style Y fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style X fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
flowchart LR
|
||||
subgraph "Input Coordinates"
|
||||
P["P-coordinates<br/>[warp_id, lane_id]"]
|
||||
Y["Y-coordinates<br/>[y0, y1, y2, y3]"]
|
||||
end
|
||||
|
||||
subgraph "Transformation Pipeline"
|
||||
C1["Combine P+Y"]
|
||||
T1["Replicate<br/>Transform<br/>(if R-dims exist)"]
|
||||
T2["Unmerge<br/>Transform<br/>(break into H-dims)"]
|
||||
T3["Merge<br/>Transform<br/>(combine to X-dims)"]
|
||||
end
|
||||
|
||||
subgraph "Output"
|
||||
X["X-coordinates<br/>[x0, x1]<br/>Tensor position"]
|
||||
end
|
||||
|
||||
P --> C1
|
||||
Y --> C1
|
||||
C1 --> T1
|
||||
T1 --> T2
|
||||
T2 --> T3
|
||||
T3 --> X
|
||||
|
||||
style P fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style Y fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style X fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
.. image:: diagrams/encoding_internals_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
|
||||
Building the Transformation Chain
|
||||
---------------------------------
|
||||
@@ -290,14 +300,10 @@ Transform Implementation Example
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
The Y→D descriptor handles memory layout within each thread:
|
||||
Y to D Linearization
|
||||
|
||||
The Y→D descriptor handles memory layout within each thread (building on :ref:`ck_tile_descriptors` concepts):
|
||||
====================
|
||||
|
||||
The Y→D descriptor handles memory layout within each thread:
|
||||
The Y→D descriptor handles memory layout within each thread (building on :ref:`ck_tile_descriptors` concepts):
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -422,14 +428,10 @@ Example 2: GEMM Distribution
|
||||
Sequence<0, 1, 0, 1> // Y components
|
||||
>;
|
||||
|
||||
|
||||
The encoding system is designed for maximum GPU performance:
|
||||
Performance Implications
|
||||
|
||||
The encoding system is designed for maximum GPU performance (see :ref:`ck_tile_gpu_basics` for hardware fundamentals):
|
||||
========================
|
||||
|
||||
The encoding system is designed for maximum GPU performance:
|
||||
The encoding system is designed for maximum GPU performance (see :ref:`ck_tile_gpu_basics` for hardware fundamentals):
|
||||
|
||||
Memory Access Patterns
|
||||
----------------------
|
||||
@@ -474,7 +476,7 @@ Compile-Time Optimization
|
||||
Summary
|
||||
=======
|
||||
|
||||
The tile distribution encoding system demonstrates sophisticated compile-time computation:
|
||||
The tile distribution encoding system demonstrates compile-time computation:
|
||||
|
||||
- **Mathematical Foundation**: Complete specification through dimensional relationships
|
||||
- **Zero Overhead**: All computations resolve at compile time
|
||||
|
||||
@@ -42,14 +42,10 @@ Our input will be uniformly distributed random data on the interval [-1, 1]:
|
||||
initializeMatrix(A.data(), M, K, -1.0, 1.0);
|
||||
initializeMatrix(B.data(), N, K, -1.0, 1.0);
|
||||
|
||||
|
||||
On the AMD **MI300** GPU architecture, each Compute Unit (CU) contains **four SIMD units**. Each SIMD unit can execute a single **wavefront** of 64 threads in parallel. Since there are four wavefronts per CU, a CU can therefore sustain the execution of up to **256 concurrent threads**.
|
||||
Simple Matmul
|
||||
|
||||
On the AMD **MI300** GPU architecture (see :ref:`ck_tile_gpu_basics` for details), each Compute Unit (CU) contains **four SIMD units**. Each SIMD unit can execute a single **wavefront** of 64 threads in parallel. Since there are four wavefronts per CU, a CU can therefore sustain the execution of up to **256 concurrent threads**.
|
||||
=============
|
||||
|
||||
On the AMD **MI300** GPU architecture, each Compute Unit (CU) contains **four SIMD units**. Each SIMD unit can execute a single **wavefront** of 64 threads in parallel. Since there are four wavefronts per CU, a CU can therefore sustain the execution of up to **256 concurrent threads**.
|
||||
On the AMD **MI300** GPU series (see :ref:`ck_tile_gpu_basics` for details), each Compute Unit (CU) contains **four SIMD units**. Each SIMD unit can execute a single **wavefront** of 64 threads in parallel. Since there are four wavefronts per CU, a CU can therefore sustain the execution of up to **256 concurrent threads**.
|
||||
|
||||
These 256 threads then can be logically grouped into a **thread block**, which is responsible for computing a **sub-block (tile)** of the output matrix ``C``. A block of 256 threads can be arranged as a **16×16 thread block**, where each thread computes one element of a **16×16 tile** of the result matrix ``C``. Multiple thread blocks are then organized into a **grid**, such that the collection of blocks covers the entire output matrix.
|
||||
|
||||
@@ -230,7 +226,7 @@ To maximize performance, the flow for this kernel uses a **pipeline** or **doubl
|
||||
|
||||
* **Stage 4: Computation with MFMA:** The final stage is the core computation. The **MFMA** (Matrix-FMA) intrinsic uses the data from the VGPRs to perform the actual matrix multiplication and accumulation.
|
||||
|
||||
By using this pipelined approach, the different stages of data movement and computation happen in parallel. While the current VGPRs are being consumed by the MFMA operation, the next set of data is already being moved from LDS to another set of VGPRs, and the next tile of data is being loaded from global memory into a third set of VGPRs. This overlapping of operations is key to keeping the GPU's powerful compute units fully utilized.
|
||||
By using this pipelined approach, the different stages of data movement and computation happen in parallel. While the current VGPRs are being consumed by the MFMA operation, the next set of data is already being moved from LDS to another set of VGPRs, and the next tile of data is being loaded from global memory into a third set of VGPRs. This overlapping of operations is key to keeping the GPU's compute units fully utilized.
|
||||
|
||||
CK Tile Implementation
|
||||
======================
|
||||
|
||||
@@ -13,7 +13,7 @@ The AMD CDNA architecture is a specialized GPU design for high-performance compu
|
||||
XCD (eXtreme Chiplet Design)
|
||||
=============================
|
||||
|
||||
A fundamental element of CDNA is the **eXtreme Chiplet Design (XCD)**. This design breaks the GPU into smaller, modular chiplets. Each XCD contains a portion of the GPU's compute resources and a dedicated slice of the L2 cache. This modular approach allows for greater manufacturing flexibility, higher yields, and improved scalability, as multiple XCDs can be connected together to form a single, powerful GPU.
|
||||
A fundamental element of CDNA is the **eXtreme Chiplet Design (XCD)**. This design breaks the GPU into smaller, modular chiplets. Each XCD contains a portion of the GPU's compute resources and a dedicated slice of the L2 cache. This modular approach allows for greater manufacturing flexibility, higher yields, and improved scalability, as multiple XCDs can be connected together to form a single, high-performance GPU.
|
||||
|
||||
MI300 incorporates 8 accelerator complex dies (XCD) with 40 compute units (CUs) per XCD, however 2 of them stay disabled which brings total to 304 CUs.
|
||||
|
||||
@@ -86,6 +86,23 @@ Understanding the CDNA architecture is crucial for effective use of CK Tile:
|
||||
|
||||
By understanding these architectural features, developers can better appreciate how CK Tile's abstractions map to hardware capabilities and why certain design decisions were made in the framework.
|
||||
|
||||
Further Reading
|
||||
|
||||
For comprehensive documentation on AMD GPU architecture and programming:
|
||||
|
||||
**AMD GPU Architecture Programming Documentation**
|
||||
https://gpuopen.com/amd-gpu-architecture-programming-documentation/
|
||||
|
||||
This resource provides in-depth guides covering:
|
||||
|
||||
- **CDNA Architecture Whitepapers**: Detailed technical specifications for MI200 and MI300 series
|
||||
- **ISA (Instruction Set Architecture) Documentation**: Complete instruction reference for AMD GPUs
|
||||
- **Optimization Guides**: Best practices for achieving peak performance
|
||||
- **Programming Guides**: ROCm and HIP programming models
|
||||
- **Performance Tools**: Profiling and analysis tool documentation
|
||||
|
||||
These guides complement the CK Tile documentation by providing the low-level hardware details that inform the design of CK's high-level abstractions. Understanding the underlying architecture enables developers to make informed decisions about tile sizes, distribution patterns, and optimization strategies.
|
||||
|
||||
Related Topics
|
||||
|
||||
- :ref:`ck_tile_thread_mapping` - How threads are organized and mapped to hardware
|
||||
|
||||
@@ -3,7 +3,7 @@
|
||||
CK Tile Conceptual Documentation
|
||||
================================
|
||||
|
||||
Welcome to the conceptual documentation for CK Tile, the core abstraction layer of Composable Kernel that enables efficient GPU programming through sophisticated coordinate transformations and tile-based data distribution.
|
||||
Welcome to the conceptual documentation for CK Tile, the core abstraction layer of Composable Kernel that enables efficient GPU programming through compile-time coordinate transformations and tile-based data distribution.
|
||||
|
||||
Overview
|
||||
--------
|
||||
|
||||
@@ -8,11 +8,11 @@ Overview
|
||||
|
||||
The evolution of GPU computing has brought unprecedented computational power to modern applications, yet harnessing this power efficiently remains one of the most challenging aspects of high-performance computing. At the heart of this challenge lies a fundamental mismatch between how developers conceptualize algorithms and how GPU hardware executes them. While developers think in terms of mathematical operations on multi-dimensional data structures, GPUs operate through thousands of threads accessing memory in complex patterns that must satisfy stringent hardware constraints.
|
||||
|
||||
This conceptual gap manifests most acutely in memory access patterns. Modern GPUs achieve their remarkable performance through massive parallelism, with thousands of threads executing simultaneously. However, this parallelism comes with a critical constraint: memory bandwidth. Despite continuous improvements in computational throughput, memory bandwidth has not scaled proportionally, creating what is often called the "memory wall." The efficiency with which threads access memory determines whether a GPU kernel achieves a few percent or near 100% of the hardware's theoretical performance.
|
||||
This conceptual gap manifests most acutely in memory access patterns. Modern GPUs achieve their high performance through massive parallelism, with thousands of threads executing simultaneously. However, this parallelism comes with a critical constraint: memory bandwidth. Despite continuous improvements in computational throughput, memory bandwidth has not scaled proportionally, creating what is often called the "memory wall." The efficiency with which threads access memory determines whether a GPU kernel achieves a few percent or near 100% of the hardware's theoretical performance.
|
||||
|
||||
The Composable Kernel (CK) framework addresses this challenge through its tile distribution system, a sophisticated abstraction that automatically generates optimal memory access patterns while preserving the natural expression of algorithms. This documentation explores the mathematical foundations and practical implementation of tile distribution, demonstrating how it bridges the gap between algorithmic intent and hardware reality.
|
||||
The Composable Kernel (CK) framework addresses this challenge through its tile distribution system, a compile-time abstraction that automatically generates optimal memory access patterns while preserving the natural expression of algorithms. This documentation explores the mathematical foundations and practical implementation of tile distribution, demonstrating how it bridges the gap between algorithmic intent and hardware reality.
|
||||
|
||||
In this introduction, we establish the fundamental problems that tile distribution solves, explore why these problems are critical for GPU performance, and provide the conceptual framework necessary to understand the sophisticated coordinate transformation system that powers CK's approach to efficient GPU computation.
|
||||
In this introduction, we establish the fundamental problems that tile distribution solves, explore why these problems are critical for GPU performance, and provide the conceptual framework necessary to understand the compile-time coordinate transformation system that powers CK's approach to efficient GPU computation.
|
||||
|
||||
The GPU Memory Problem
|
||||
----------------------
|
||||
@@ -93,14 +93,14 @@ The architecture of modern GPUs represents a study in trade-offs. While these de
|
||||
|
||||
GPU memory systems are designed around the assumption of regular, predictable access patterns. The memory controller can service requests from 32 threads (a warp on AMD GPUs) in a single transaction when these threads access consecutive memory locations. This optimization, known as memory coalescing, can improve effective memory bandwidth by up to 32x compared to random access patterns. However, when threads within a warp access memory locations that are scattered throughout the address space, each access requires a separate memory transaction, reducing the effective bandwidth to a fraction of the theoretical maximum.
|
||||
|
||||
The impact extends beyond raw bandwidth. Modern GPUs employ sophisticated cache hierarchies to reduce memory latency, but these caches are effective only when access patterns exhibit spatial or temporal locality. Random access patterns defeat these optimizations, causing frequent cache misses that expose the full latency of global memory access, which can be hundreds of cycles. During these stalls, the computational units sit idle, unable to hide the latency even with the GPU's massive thread count.
|
||||
The impact extends beyond raw bandwidth. Modern GPUs employ advanced cache hierarchies to reduce memory latency, but these caches are effective only when access patterns exhibit spatial or temporal locality. Random access patterns defeat these optimizations, causing frequent cache misses that expose the full latency of global memory access, which can be hundreds of cycles. During these stalls, the computational units sit idle, unable to hide the latency even with the GPU's massive thread count.
|
||||
|
||||
Furthermore, the GPU's SIMT (Single Instruction, Multiple Thread) execution model means that all threads in a warp must execute the same instruction at the same time. When threads access memory in unpredictable patterns, the memory controller cannot optimize the requests, leading to serialization of what should be parallel operations. This serialization effect compounds with each level of the memory hierarchy, from L1 cache through L2 cache to global memory, multiplying the performance impact.
|
||||
|
||||
The Thread Cooperation Challenge
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
The challenge of efficient thread cooperation becomes particularly evident when examining a fundamental operation like matrix multiplication. Consider a scenario where 256 threads must cooperate to multiply two matrices. The naive approach, where each thread computes one element of the output matrix, illustrates precisely why GPU programming requires sophisticated abstractions.
|
||||
The challenge of efficient thread cooperation becomes particularly evident when examining a fundamental operation like matrix multiplication. Consider a scenario where 256 threads must cooperate to multiply two matrices. The naive approach, where each thread computes one element of the output matrix, illustrates precisely why GPU programming requires compile-time abstractions.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -142,7 +142,7 @@ The lack of coordination between threads exacerbates the problem. While all thre
|
||||
|
||||
Cache utilization suffers dramatically under this access pattern. Each thread traces a unique path through memory, with no overlap between threads' working sets. The L1 and L2 caches, designed to capture and exploit locality, instead thrash continuously as each thread's accesses evict data needed by others. The effective cache capacity approaches zero, exposing every memory access to the full latency of global memory.
|
||||
|
||||
Perhaps most critically, this approach fails to utilize the available memory bandwidth efficiently. Modern GPUs can achieve memory bandwidths exceeding 1 TB/s, but only when accesses are properly structured. The random access pattern of the naive implementation might achieve less than 10% of this theoretical maximum, effectively reducing a powerful GPU to the performance level of a much simpler processor.
|
||||
Perhaps most critically, this approach fails to utilize the available memory bandwidth efficiently. Modern GPUs can achieve memory bandwidths exceeding 1 TB/s, but only when accesses are properly structured. The random access pattern of the naive implementation might achieve less than 10% of this theoretical maximum, effectively reducing a high-performance GPU to the performance level of a much simpler processor.
|
||||
|
||||
The Tile Distribution Solution
|
||||
------------------------------
|
||||
@@ -192,7 +192,7 @@ The essence of tile distribution is the recognition that efficient GPU computati
|
||||
});
|
||||
}
|
||||
|
||||
The transformation from inefficient to efficient memory access is profound. Where the naive implementation scattered memory requests across the address space, tile distribution ensures that adjacent threads access adjacent memory locations. This transformation happens through a sophisticated encoding system that captures the hierarchical nature of both the computation and the hardware.
|
||||
The transformation from inefficient to efficient memory access is profound. Where the naive implementation scattered memory requests across the address space, tile distribution ensures that adjacent threads access adjacent memory locations. This transformation happens through a advanced encoding system that captures the hierarchical nature of both the computation and the hardware.
|
||||
|
||||
The encoding shown above demonstrates the multi-level hierarchy that tile distribution employs. The sequence<4, 2, 8, 4> represents a four-level decomposition: four repetitions per thread, two warps per block, eight threads per warp, and four elements per vector operation. This hierarchical structure maps directly to the GPU's hardware organization, ensuring that each level of the hierarchy operates at maximum efficiency.
|
||||
|
||||
@@ -200,9 +200,9 @@ Memory access patterns become predictable and regular under tile distribution. T
|
||||
|
||||
Thread cooperation emerges naturally from the tile distribution structure. Threads within a warp work on adjacent data, enabling efficient data sharing through register shuffle operations. Warps within a block coordinate through shared memory, with access patterns that avoid bank conflicts. This cooperation transforms what was a collection of independent computations into a unified, efficient operation.
|
||||
|
||||
Cache utilization improves dramatically as well. The structured access patterns ensure that data loaded into cache by one thread is likely to be used by neighboring threads. Temporal locality emerges from the tile-based processing, where all operations on a tile complete before moving to the next tile. This locality transforms the cache from a liability into a powerful performance accelerator.
|
||||
Cache utilization improves dramatically as well. The structured access patterns ensure that data loaded into cache by one thread is likely to be used by neighboring threads. Temporal locality emerges from the tile-based processing, where all operations on a tile complete before moving to the next tile. This locality transforms the cache from a liability into a high performance accelerator.
|
||||
|
||||
The scalability of tile distribution across different GPU architectures represents one of its most powerful features. The same high-level code can achieve near-optimal performance on GPUs with different numbers of compute units, different cache sizes, and different memory bandwidths. The compile-time nature of the encoding allows the compiler to generate architecture-specific optimizations while maintaining portable source code.
|
||||
The scalability of tile distribution across different GPU architectures represents one of its most key features. The same high-level code can achieve near-optimal performance on GPUs with different numbers of compute units, different cache sizes, and different memory bandwidths. The compile-time nature of the encoding allows the compiler to generate architecture-specific optimizations while maintaining portable source code.
|
||||
|
||||
The Coordinate Mapping Insight
|
||||
------------------------------
|
||||
@@ -251,12 +251,12 @@ The elegance of this approach emerges from its separation of concerns. Each coor
|
||||
|
||||
The transformative power of tile distribution emerges from the composition of these mappings. The **P + Y → X** transformation combines a thread's position with its local data coordinates to determine global data positions. This transformation encodes the distribution strategy, determining how work is partitioned across threads. The subsequent **X → D** transformation converts these logical positions into physical memory addresses, incorporating layout optimizations that ensure efficient memory access patterns.
|
||||
|
||||
The mathematical rigor of this framework enables powerful optimizations. Because each transformation is well-defined and composable, the compiler can analyze the complete transformation chain and generate optimal code. The framework can automatically ensure memory coalescing by structuring the P + Y → X transformation appropriately. It can minimize bank conflicts in shared memory by carefully designing the X → D mapping. Most importantly, it can adapt these optimizations to different hardware architectures by adjusting the transformation parameters while keeping the high-level algorithm description unchanged.
|
||||
The mathematical rigor of this framework enables critical optimizations. Because each transformation is well-defined and composable, the compiler can analyze the complete transformation chain and generate optimal code. The framework can automatically ensure memory coalescing by structuring the P + Y → X transformation appropriately. It can minimize bank conflicts in shared memory by carefully designing the X → D mapping. Most importantly, it can adapt these optimizations to different hardware architectures by adjusting the transformation parameters while keeping the high-level algorithm description unchanged.
|
||||
|
||||
What's Coming Next
|
||||
------------------
|
||||
|
||||
Having established the fundamental motivation for tile distribution and its coordinate mapping framework, we now embark on a systematic journey through the complete CK Tile system. This journey is carefully structured to build understanding layer by layer, starting from the most basic abstractions and progressing to sophisticated optimization techniques.
|
||||
Having established the fundamental motivation for tile distribution and its coordinate mapping framework, we now embark on a systematic journey through the complete CK Tile system. This journey is carefully structured to build understanding layer by layer, starting from the most basic abstractions and progressing to advanced optimization techniques.
|
||||
|
||||
The foundation of our exploration begins with raw memory access through :ref:`ck_tile_buffer_views`, the fundamental abstraction that provides type-safe, address-space-aware access to GPU memory. Understanding BufferView is crucial because it establishes the patterns and principles that permeate the entire CK Tile system. From there, we progress to :ref:`ck_tile_tensor_views`, which adds multi-dimensional structure to raw memory, enabling natural expression of algorithms while maintaining the efficiency of the underlying buffer operations.
|
||||
|
||||
@@ -266,11 +266,11 @@ The high-level :ref:`ck_tile_distribution` APIs represent the culmination of the
|
||||
|
||||
Our exploration of coordinate systems goes beyond the basic P, Y, X, D framework to encompass advanced topics such as multi-level tiling, replication strategies, and specialized coordinate systems for specific algorithm classes. The :ref:`ck_tile_encoding_internals` reveals the mathematical foundations, while :ref:`ck_tile_thread_mapping` shows how these abstractions map to hardware. This comprehensive treatment ensures that developers can handle not just common cases but also novel algorithms that require custom distribution strategies.
|
||||
|
||||
The implementation details reveal the sophisticated template metaprogramming techniques that enable CK Tile's zero-overhead abstractions. Topics like :ref:`ck_tile_descriptors`, :ref:`ck_tile_load_store_traits`, and :ref:`ck_tile_static_distributed_tensor` show how these abstractions achieve zero overhead. By understanding these implementation strategies, advanced developers can extend the framework, contribute optimizations, and debug performance issues at the deepest level.
|
||||
The implementation details reveal the template metaprogramming techniques that enable CK Tile's zero-overhead abstractions. Topics like :ref:`ck_tile_descriptors`, :ref:`ck_tile_load_store_traits`, and :ref:`ck_tile_static_distributed_tensor` show how these abstractions achieve zero overhead. By understanding these implementation strategies, advanced developers can extend the framework, contribute optimizations, and debug performance issues at the deepest level.
|
||||
|
||||
The connection between abstract coordinate transformations and concrete hardware thread mapping represents a critical piece of the puzzle. We'll examine how logical thread organizations map to physical GPU resources, how to avoid common pitfalls like bank conflicts (see :ref:`ck_tile_lds_bank_conflicts` and :ref:`ck_tile_lds_index_swapping`) and divergent execution, and how to structure computations for maximum hardware utilization. The :ref:`ck_tile_hardware` section provides deep dives into architecture-specific optimizations.
|
||||
|
||||
Finally, our advanced topics section explores cutting-edge optimization techniques, including :ref:`ck_tile_space_filling_curve` for optimal memory traversal, :ref:`ck_tile_sweep_tile` for elegant iteration patterns, and practical examples like :ref:`ck_tile_convolution_example` and :ref:`ck_tile_gemm_optimization`. These topics prepare developers to push the boundaries of GPU performance and contribute to the ongoing evolution of high-performance computing.
|
||||
Finally, our advanced topics section explores cutting-edge optimization techniques, including :ref:`ck_tile_space_filling_curve` for optimal memory traversal, :ref:`ck_tile_sweep_tile` for clean iteration patterns, and practical examples like :ref:`ck_tile_convolution_example` and :ref:`ck_tile_gemm_optimization`. These topics prepare developers to push the boundaries of GPU performance and contribute to the ongoing evolution of high-performance computing.
|
||||
|
||||
Summary
|
||||
-------
|
||||
@@ -285,7 +285,7 @@ Perhaps most importantly, tile distribution provides a predictable optimization
|
||||
|
||||
The coordinate transformation insight—the systematic mapping through P-space, Y-space, X-space, and D-space—provides a mental model that clarifies the entire GPU computation process. This model enables developers to reason about their code at multiple levels of abstraction simultaneously, understanding both the high-level algorithmic behavior and the low-level hardware execution patterns.
|
||||
|
||||
As we prepare to dive deeper into the implementation details, starting with the foundational BufferView abstraction, remember that each component we'll explore serves the larger purpose of enabling efficient, scalable GPU computation. The journey from raw memory to sophisticated tile distributions mirrors the evolution of GPU programming itself—from low-level, hardware-specific optimizations to high-level, portable abstractions that preserve efficiency.
|
||||
As we prepare to dive deeper into the implementation details, starting with the foundational BufferView abstraction, remember that each component we'll explore serves the larger purpose of enabling efficient, scalable GPU computation. The journey from raw memory to advanced tile distributions mirrors the evolution of GPU programming itself—from low-level, hardware-specific optimizations to high-level, portable abstractions that preserve efficiency.
|
||||
|
||||
The promise of tile distribution is nothing less than democratizing high-performance GPU computing. By providing a systematic framework for achieving optimal memory access patterns, it enables a new generation of developers to harness the full power of modern GPUs without becoming experts in every detail of the hardware architecture. Ready to see how this promise becomes reality? Let's begin our technical journey with the foundation of it all: raw memory access through BufferView.
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ LDS Index Swapping
|
||||
Overview
|
||||
========
|
||||
|
||||
LDS (Local Data Share) index swapping, also known as XOR preshuffle, is a critical optimization technique in CK Tile for resolving bank conflicts in shared memory. Bank conflicts occur when multiple threads in a warp attempt to access different addresses within the same memory bank simultaneously, causing serialization and performance degradation. CK Tile generalizes the XOR preshuffle technique through a sophisticated coordinate transformation system that automatically handles complex access patterns.
|
||||
LDS (Local Data Share) index swapping, also known as XOR preshuffle, is a critical optimization technique in CK Tile for resolving bank conflicts in shared memory. Bank conflicts occur when multiple threads in a warp attempt to access different addresses within the same memory bank simultaneously, causing serialization and performance degradation. CK Tile generalizes the XOR preshuffle technique through a compile-time coordinate transformation system that automatically handles complex access patterns.
|
||||
|
||||
The key insight is that by transforming the logical 2D coordinates used to access LDS into a different 2D coordinate space, we can ensure that threads accessing data simultaneously hit different banks. This transformation is implemented through CK Tile's composable transform system (see :ref:`ck_tile_transforms` and :ref:`ck_tile_coordinate_systems`), making it both flexible and efficient.
|
||||
|
||||
@@ -25,40 +25,48 @@ Step 1: XOR Transform
|
||||
|
||||
The original K coordinate is split into K0 and K1, where K1 represents the thread vector size along the K dimension (KPack), and K0 is KPerBlock/KPack.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "3D LDS coordinate [K0, M, K1]"
|
||||
K0["KPerBlock/KPack * MLdsLayer<br/>K0"]
|
||||
M["MPerBlock/MLdsLayer<br/>M"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "XOR Transform"
|
||||
XT["make_xor_transform"]
|
||||
end
|
||||
|
||||
subgraph "Update K0 with XOR transformation"
|
||||
K01["KPerBlock/KPack * MLdsLayer<br/>K0'"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K11["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
K0 --> XT
|
||||
M --> XT
|
||||
K1 --> K11
|
||||
|
||||
XT --> K01
|
||||
XT --> M1
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style M fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style M1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style K1 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style K11 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "3D LDS coordinate [K0, M, K1]"
|
||||
K0["KPerBlock/KPack * MLdsLayer<br/>K0"]
|
||||
M["MPerBlock/MLdsLayer<br/>M"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "XOR Transform"
|
||||
XT["make_xor_transform"]
|
||||
end
|
||||
|
||||
subgraph "Update K0 with XOR transformation"
|
||||
K01["KPerBlock/KPack * MLdsLayer<br/>K0'"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K11["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
K0 --> XT
|
||||
M --> XT
|
||||
K1 --> K11
|
||||
|
||||
XT --> K01
|
||||
XT --> M1
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style M fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style M1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style K1 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style K11 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/lds_index_swapping_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
The XOR transformation updates the K0 coordinate using the formula:
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -72,43 +80,51 @@ Step 2: Unmerge Transform
|
||||
|
||||
The transformed K0' is split into L and K0'' components, creating an intermediate 4D coordinate space. This is necessary when MLdsLayer > 1, allowing multiple rows to share the same set of memory banks for better utilization with smaller tile sizes.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "3D LDS coordinate [K0', M, K1]"
|
||||
K0["KPerBlock/KPack * MLdsLayer<br/>K0'"]
|
||||
M["MPerBlock/MLdsLayer<br/>M"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "Unmerge into 2 components"
|
||||
UM["make_unmerge_transform"]
|
||||
end
|
||||
|
||||
subgraph "4D intermediate transformation space"
|
||||
L["MLdsLayer<br/>L"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K01["KPerBlock/KPack<br/>K0''"]
|
||||
K11["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
K0 --> UM
|
||||
M --> M1
|
||||
K1 --> K11
|
||||
|
||||
UM --> L
|
||||
UM --> K01
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style L fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style M fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style M1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
style K1 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style K11 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "3D LDS coordinate [K0', M, K1]"
|
||||
K0["KPerBlock/KPack * MLdsLayer<br/>K0'"]
|
||||
M["MPerBlock/MLdsLayer<br/>M"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "Unmerge into 2 components"
|
||||
UM["make_unmerge_transform"]
|
||||
end
|
||||
|
||||
subgraph "4D intermediate transformation space"
|
||||
L["MLdsLayer<br/>L"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K01["KPerBlock/KPack<br/>K0''"]
|
||||
K11["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
K0 --> UM
|
||||
M --> M1
|
||||
K1 --> K11
|
||||
|
||||
UM --> L
|
||||
UM --> K01
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style L fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style M fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style M1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
style K1 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style K11 fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/lds_index_swapping_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
The unmerge operation:
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -123,48 +139,56 @@ Step 3: Merge Transform
|
||||
|
||||
The final step merges the 4D coordinates back into 2D transformed coordinates (M', K').
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "4D LDS Coordinates [L, M, K0'', K1]"
|
||||
L["MLdsLayer<br/>L"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K0["KPerBlock/KPack<br/>K0''"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "Merge into 1 component"
|
||||
ME0["make_merge_transform"]
|
||||
end
|
||||
|
||||
subgraph "Merge into 1 component"
|
||||
ME1["make_merge_transform"]
|
||||
end
|
||||
|
||||
subgraph "Transformed 2D coordinates [M', K']"
|
||||
M11["MPerBlock<br/>M'"]
|
||||
K01["KPerBlock<br/>K'"]
|
||||
end
|
||||
|
||||
L --> ME0
|
||||
M1 --> ME0
|
||||
|
||||
K0 --> ME1
|
||||
K1 --> ME1
|
||||
|
||||
ME0 --> M11
|
||||
ME1 --> K01
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style M1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style L fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style M11 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "4D LDS Coordinates [L, M, K0'', K1]"
|
||||
L["MLdsLayer<br/>L"]
|
||||
M1["MPerBlock/MLdsLayer<br/>M"]
|
||||
K0["KPerBlock/KPack<br/>K0''"]
|
||||
K1["KPack<br/>K1"]
|
||||
end
|
||||
|
||||
subgraph "Merge into 1 component"
|
||||
ME0["make_merge_transform"]
|
||||
end
|
||||
|
||||
subgraph "Merge into 1 component"
|
||||
ME1["make_merge_transform"]
|
||||
end
|
||||
|
||||
subgraph "Transformed 2D coordinates [M', K']"
|
||||
M11["MPerBlock<br/>M'"]
|
||||
K01["KPerBlock<br/>K'"]
|
||||
end
|
||||
|
||||
L --> ME0
|
||||
M1 --> ME0
|
||||
|
||||
K0 --> ME1
|
||||
K1 --> ME1
|
||||
|
||||
ME0 --> M11
|
||||
ME1 --> K01
|
||||
|
||||
style K0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style K01 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
|
||||
style M1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style L fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style M11 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/lds_index_swapping_3.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Here's how the complete transformation chain is implemented in CK Tile:
|
||||
C++ Implementation
|
||||
|
||||
@@ -427,7 +451,7 @@ LDS index swapping works seamlessly with CK Tile's distribution system:
|
||||
Summary
|
||||
=======
|
||||
|
||||
LDS index swapping in CK Tile provides a powerful and flexible solution to the bank conflict problem:
|
||||
LDS index swapping in CK Tile provides a effective and flexible solution to the bank conflict problem:
|
||||
|
||||
- **Generalized XOR Preshuffle**: Extends the basic XOR technique through composable transforms
|
||||
- **Multi-Step Pipeline**: Coordinates flow through XOR → Unmerge → Merge transformations
|
||||
|
||||
@@ -8,7 +8,7 @@ Overview
|
||||
|
||||
LoadStoreTraits is a critical optimization component that analyzes :ref:`tile distributions <ck_tile_tile_distribution>` to determine the most efficient memory access patterns. It serves as the intelligence behind :ref:`TileWindow's <ck_tile_tile_window>` high-performance data movement, automatically identifying the best dimension for vectorization and creating optimized access sequences using :ref:`space-filling curves <ck_tile_space_filling_curve>`.
|
||||
|
||||
At compile time, LoadStoreTraits performs sophisticated analysis of the distribution pattern to extract key information about memory access opportunities. This analysis determines how many elements can be loaded or stored in a single instruction, which dimension provides the best vectorization opportunity, and what traversal order maximizes cache utilization. The result is a set of compile-time constants and methods that guide the runtime execution of load and store operations.
|
||||
At compile time, LoadStoreTraits performs compile-time analysis of the distribution pattern to extract key information about memory access opportunities. This analysis determines how many elements can be loaded or stored in a single instruction, which dimension provides the best vectorization opportunity, and what traversal order maximizes cache utilization. The result is a set of compile-time constants and methods that guide the runtime execution of load and store operations.
|
||||
|
||||
Key Concepts
|
||||
------------
|
||||
@@ -101,27 +101,35 @@ The LoadStoreTraits class analyzes distribution patterns at compile time:
|
||||
Vectorization Selection Algorithm
|
||||
---------------------------------
|
||||
|
||||
LoadStoreTraits employs a sophisticated algorithm to select the best dimension for vectorization:
|
||||
LoadStoreTraits employs a advanced algorithm to select the best dimension for vectorization:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TD
|
||||
A[Analyze Distribution] --> B{Check Each Dimension}
|
||||
B --> C[Calculate Stride]
|
||||
C --> D{Stride == 1?}
|
||||
D -->|Yes| E[Candidate for Vectorization]
|
||||
D -->|No| F[Skip Dimension]
|
||||
E --> G[Check Alignment]
|
||||
G --> H[Check Vector Size]
|
||||
H --> I[Score Dimension]
|
||||
F --> B
|
||||
I --> J[Select Best Dimension]
|
||||
J --> K[Configure Vector Access]
|
||||
|
||||
style A fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style J fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style K fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TD
|
||||
A[Analyze Distribution] --> B{Check Each Dimension}
|
||||
B --> C[Calculate Stride]
|
||||
C --> D{Stride == 1?}
|
||||
D -->|Yes| E[Candidate for Vectorization]
|
||||
D -->|No| F[Skip Dimension]
|
||||
E --> G[Check Alignment]
|
||||
G --> H[Check Vector Size]
|
||||
H --> I[Score Dimension]
|
||||
F --> B
|
||||
I --> J[Select Best Dimension]
|
||||
J --> K[Configure Vector Access]
|
||||
|
||||
style A fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style J fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style K fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/load_store_traits_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**Example: Comparing Different Memory Layouts**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -163,34 +171,42 @@ Memory Access Patterns
|
||||
|
||||
LoadStoreTraits creates efficient access patterns using space-filling curves:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Traversal"
|
||||
L1["0→1→2→3"]
|
||||
L2["4→5→6→7"]
|
||||
L3["Cache miss"]
|
||||
L4["8→9→10→11"]
|
||||
end
|
||||
|
||||
subgraph "Snake Pattern"
|
||||
S1["0→1→2→3"]
|
||||
S2["7←6←5←4"]
|
||||
S3["Cache hit!"]
|
||||
S4["8→9→10→11"]
|
||||
end
|
||||
|
||||
L1 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
style L3 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style S3 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Traversal"
|
||||
L1["0→1→2→3"]
|
||||
L2["4→5→6→7"]
|
||||
L3["Cache miss"]
|
||||
L4["8→9→10→11"]
|
||||
end
|
||||
|
||||
subgraph "Snake Pattern"
|
||||
S1["0→1→2→3"]
|
||||
S2["7←6←5←4"]
|
||||
S3["Cache hit!"]
|
||||
S4["8→9→10→11"]
|
||||
end
|
||||
|
||||
L1 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
style L3 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style S3 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/load_store_traits_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Access Pattern Example:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -449,7 +465,7 @@ LoadStoreTraits provides:
|
||||
- **Hardware adaptation**: Adjusts to different data types and :ref:`architectures <ck_tile_gpu_basics>`
|
||||
- **Performance transparency**: Clear metrics for memory efficiency
|
||||
|
||||
The sophisticated analysis performed by LoadStoreTraits ensures that every memory operation in CK Tile achieves near-optimal performance, making it a critical component in the high-performance computing stack.
|
||||
The compile-time analysis performed by LoadStoreTraits ensures that every memory operation in CK Tile achieves near-optimal performance, making it a critical component in the high-performance computing stack.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
@@ -185,34 +185,42 @@ Snake Pattern for Cache Optimization
|
||||
|
||||
The snake pattern reverses traversal direction on alternate rows, minimizing the distance between consecutive accesses:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Pattern"
|
||||
L1["Row 0: →"]
|
||||
L2["Row 1: →"]
|
||||
L3["Jump back"]
|
||||
L4["Row 2: →"]
|
||||
end
|
||||
|
||||
subgraph "Snake Pattern"
|
||||
S1["Row 0: →"]
|
||||
S2["Row 1: ←"]
|
||||
S3["Continue"]
|
||||
S4["Row 2: →"]
|
||||
end
|
||||
|
||||
L1 --> L3
|
||||
L3 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S4
|
||||
|
||||
style L3 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style S3 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Pattern"
|
||||
L1["Row 0: →"]
|
||||
L2["Row 1: →"]
|
||||
L3["Jump back"]
|
||||
L4["Row 2: →"]
|
||||
end
|
||||
|
||||
subgraph "Snake Pattern"
|
||||
S1["Row 0: →"]
|
||||
S2["Row 1: ←"]
|
||||
S3["Continue"]
|
||||
S4["Row 2: →"]
|
||||
end
|
||||
|
||||
L1 --> L3
|
||||
L3 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S4
|
||||
|
||||
style L3 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style S3 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/space_filling_curve.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
.. code-block:: cpp
|
||||
|
||||
using SnakeCurve = space_filling_curve<
|
||||
@@ -483,7 +491,7 @@ Space-filling curves provide:
|
||||
- **Flexibility**: Adaptable to different :ref:`tensor shapes <ck_tile_descriptors>` and access patterns
|
||||
- **Performance**: Compile-time optimization with zero runtime overhead
|
||||
|
||||
The sophisticated traversal patterns enabled by space-filling curves are fundamental to achieving high performance in GPU kernels, ensuring that memory access patterns align with :ref:`hardware capabilities <ck_tile_gpu_basics>`.
|
||||
The advanced traversal patterns enabled by space-filling curves are fundamental to achieving high performance in GPU kernels, ensuring that memory access patterns align with :ref:`hardware capabilities <ck_tile_gpu_basics>`.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
@@ -20,7 +20,7 @@ This design enables several critical optimizations. First, it maximizes register
|
||||
Thread-Local Storage Model
|
||||
==========================
|
||||
|
||||
The static distributed tensor implements a sophisticated storage model that maps multi-dimensional tensor data to thread-local arrays:
|
||||
The static distributed tensor implements a advanced storage model that maps multi-dimensional tensor data to thread-local arrays:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -85,19 +85,27 @@ Understanding how static distributed tensors organize memory is crucial for perf
|
||||
|
||||
The memory layout follows a hierarchical pattern:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TD
|
||||
A[Global Tensor 64x64] --> B[Thread Block 16x16]
|
||||
B --> C[Thread 0,0<br/>Elements 0:3,0:3]
|
||||
B --> D[Thread 0,1<br/>Elements 0:3,4:7]
|
||||
B --> E[Thread 1,0<br/>Elements 4:7,0:3]
|
||||
B --> F[...]
|
||||
|
||||
C --> G[Local Array<br/>16 elements]
|
||||
D --> H[Local Array<br/>16 elements]
|
||||
E --> I[Local Array<br/>16 elements]
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TD
|
||||
A[Global Tensor 64x64] --> B[Thread Block 16x16]
|
||||
B --> C[Thread 0,0<br/>Elements 0:3,0:3]
|
||||
B --> D[Thread 0,1<br/>Elements 0:3,4:7]
|
||||
B --> E[Thread 1,0<br/>Elements 4:7,0:3]
|
||||
B --> F[...]
|
||||
|
||||
C --> G[Local Array<br/>16 elements]
|
||||
D --> H[Local Array<br/>16 elements]
|
||||
E --> I[Local Array<br/>16 elements]
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/static_distributed_tensor.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Element Access and Indexing
|
||||
===========================
|
||||
|
||||
|
||||
@@ -11,37 +11,45 @@ Sweep Tile
|
||||
Overview
|
||||
========
|
||||
|
||||
Sweep operations are the elegant way to iterate over distributed data in CK Tile. They complete the tile distribution workflow by providing clean, efficient iteration patterns that automatically handle all the complex indexing details. Sweep operations are like forEach() for distributed tensors - give them a function, and they'll call it for every element in the optimal order.
|
||||
Sweep operations are the clean way to iterate over distributed data in CK Tile. They complete the tile distribution workflow by providing clean, efficient iteration patterns that automatically handle all the complex indexing details. Sweep operations are like forEach() for distributed tensors - give them a function, and they'll call it for every element in the optimal order.
|
||||
|
||||
The key insight is the "load once, use many times" pattern. Load X data once into registers, then sweep through Y positions while keeping X in fast memory. This maximizes data reuse and minimizes memory bandwidth requirements.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart LR
|
||||
subgraph "X-Tile (Reused)"
|
||||
XT["X data loaded once<br/>Stays in registers"]
|
||||
end
|
||||
|
||||
subgraph "Y-Sweep"
|
||||
Y1["Y position 0"]
|
||||
Y2["Y position 1"]
|
||||
Y3["Y position 2"]
|
||||
YN["Y position N"]
|
||||
end
|
||||
|
||||
subgraph "Computation"
|
||||
C["Process(X, Y)"]
|
||||
end
|
||||
|
||||
XT --> C
|
||||
Y1 --> C
|
||||
Y2 --> C
|
||||
Y3 --> C
|
||||
YN --> C
|
||||
|
||||
style XT fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style C fill:#e0e7ff,stroke:#4338ca,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart LR
|
||||
subgraph "X-Tile (Reused)"
|
||||
XT["X data loaded once<br/>Stays in registers"]
|
||||
end
|
||||
|
||||
subgraph "Y-Sweep"
|
||||
Y1["Y position 0"]
|
||||
Y2["Y position 1"]
|
||||
Y3["Y position 2"]
|
||||
YN["Y position N"]
|
||||
end
|
||||
|
||||
subgraph "Computation"
|
||||
C["Process(X, Y)"]
|
||||
end
|
||||
|
||||
XT --> C
|
||||
Y1 --> C
|
||||
Y2 --> C
|
||||
Y3 --> C
|
||||
YN --> C
|
||||
|
||||
style XT fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style C fill:#e0e7ff,stroke:#4338ca,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/sweep_tile_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
The Complete GPU Workflow
|
||||
=========================
|
||||
|
||||
@@ -52,7 +60,7 @@ Sweep operations are the final piece of the distributed computing puzzle:
|
||||
3. **Sweep Operations**: "Here's how to process every element"
|
||||
4. **Your code**: "Thanks! *does computation*"
|
||||
|
||||
Without sweep operations, you need manual nested loops, complex index calculations, and risk missing elements or double-processing. With sweep operations, you get elegant lambda-based iteration with automatic handling of all elements.
|
||||
Without sweep operations, you need manual nested loops, complex index calculations, and risk missing elements or double-processing. With sweep operations, you get lambda-based iteration with automatic handling of all elements.
|
||||
|
||||
For understanding the underlying coordinate systems, see :ref:`ck_tile_coordinate_systems`.
|
||||
|
||||
@@ -111,28 +119,36 @@ The sweep pattern provides significant memory efficiency benefits. This is parti
|
||||
|
||||
The sweep pattern provides significant memory efficiency benefits:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Traditional Approach"
|
||||
T1["Load X[0]"] --> P1["Process"]
|
||||
T2["Load Y[0]"] --> P1
|
||||
T3["Load X[0]"] --> P2["Process"]
|
||||
T4["Load Y[1]"] --> P2
|
||||
T5["Load X[0]"] --> P3["Process"]
|
||||
T6["Load Y[2]"] --> P3
|
||||
Note1["X loaded 3 times!"]
|
||||
end
|
||||
|
||||
subgraph "Sweep Approach"
|
||||
S1["Load X[0]"] --> SP["Process with<br/>Y[0], Y[1], Y[2]"]
|
||||
S2["Load Y[0,1,2]"] --> SP
|
||||
Note2["X loaded once!"]
|
||||
end
|
||||
|
||||
style Note1 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style Note2 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Traditional Approach"
|
||||
T1["Load X[0]"] --> P1["Process"]
|
||||
T2["Load Y[0]"] --> P1
|
||||
T3["Load X[0]"] --> P2["Process"]
|
||||
T4["Load Y[1]"] --> P2
|
||||
T5["Load X[0]"] --> P3["Process"]
|
||||
T6["Load Y[2]"] --> P3
|
||||
Note1["X loaded 3 times!"]
|
||||
end
|
||||
|
||||
subgraph "Sweep Approach"
|
||||
S1["Load X[0]"] --> SP["Process with<br/>Y[0], Y[1], Y[2]"]
|
||||
S2["Load Y[0,1,2]"] --> SP
|
||||
Note2["X loaded once!"]
|
||||
end
|
||||
|
||||
style Note1 fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
style Note2 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/sweep_tile_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Practical Sweep Patterns
|
||||
========================
|
||||
|
||||
@@ -362,35 +378,43 @@ Performance Characteristics
|
||||
|
||||
Sweep operations provide several performance benefits:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Sweep Performance Benefits"
|
||||
B1["Zero runtime overhead<br/>Compile-time unrolling"]
|
||||
B2["Perfect memory coalescing<br/>Sequential access patterns"]
|
||||
B3["Automatic vectorization<br/>Compiler optimizations"]
|
||||
B4["Register reuse<br/>X data stays in VGPR"]
|
||||
end
|
||||
|
||||
subgraph "Use Cases"
|
||||
U1["Matrix Multiplication<br/>Reuse A columns"]
|
||||
U2["Convolution<br/>Reuse filter weights"]
|
||||
U3["Reduction<br/>Accumulate over Y"]
|
||||
U4["Broadcast<br/>Apply X to all Y"]
|
||||
end
|
||||
|
||||
B1 --> Performance["High Performance"]
|
||||
B2 --> Performance
|
||||
B3 --> Performance
|
||||
B4 --> Performance
|
||||
|
||||
Performance --> U1
|
||||
Performance --> U2
|
||||
Performance --> U3
|
||||
Performance --> U4
|
||||
|
||||
style Performance fill:#d1fae5,stroke:#10b981,stroke-width:3px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Sweep Performance Benefits"
|
||||
B1["Zero runtime overhead<br/>Compile-time unrolling"]
|
||||
B2["Perfect memory coalescing<br/>Sequential access patterns"]
|
||||
B3["Automatic vectorization<br/>Compiler optimizations"]
|
||||
B4["Register reuse<br/>X data stays in VGPR"]
|
||||
end
|
||||
|
||||
subgraph "Use Cases"
|
||||
U1["Matrix Multiplication<br/>Reuse A columns"]
|
||||
U2["Convolution<br/>Reuse filter weights"]
|
||||
U3["Reduction<br/>Accumulate over Y"]
|
||||
U4["Broadcast<br/>Apply X to all Y"]
|
||||
end
|
||||
|
||||
B1 --> Performance["High Performance"]
|
||||
B2 --> Performance
|
||||
B3 --> Performance
|
||||
B4 --> Performance
|
||||
|
||||
Performance --> U1
|
||||
Performance --> U2
|
||||
Performance --> U3
|
||||
Performance --> U4
|
||||
|
||||
style Performance fill:#d1fae5,stroke:#10b981,stroke-width:3px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/sweep_tile_3.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Compiler Optimizations
|
||||
----------------------
|
||||
|
||||
@@ -426,26 +450,34 @@ Integration with CK Tile Components
|
||||
|
||||
Complete workflow example:
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart TB
|
||||
subgraph "Complete Workflow"
|
||||
TD["TileDistribution<br/>Define data layout"]
|
||||
TW["TileWindow<br/>Create view"]
|
||||
DT["DistributedTensor<br/>Load X data"]
|
||||
ST["SweepTile<br/>Iterate Y positions"]
|
||||
R["Results<br/>Store outputs"]
|
||||
end
|
||||
|
||||
TD --> TW
|
||||
TW --> DT
|
||||
DT --> ST
|
||||
ST --> R
|
||||
|
||||
style TD fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style ST fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style R fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart TB
|
||||
subgraph "Complete Workflow"
|
||||
TD["TileDistribution<br/>Define data layout"]
|
||||
TW["TileWindow<br/>Create view"]
|
||||
DT["DistributedTensor<br/>Load X data"]
|
||||
ST["SweepTile<br/>Iterate Y positions"]
|
||||
R["Results<br/>Store outputs"]
|
||||
end
|
||||
|
||||
TD --> TW
|
||||
TW --> DT
|
||||
DT --> ST
|
||||
ST --> R
|
||||
|
||||
style TD fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style ST fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style R fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/sweep_tile_4.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
.. code-block:: cpp
|
||||
|
||||
template<typename DataType>
|
||||
@@ -495,13 +527,13 @@ Complete workflow example:
|
||||
}
|
||||
|
||||
|
||||
SweepTile provides elegant and efficient iteration over distributed data:
|
||||
SweepTile provides clean and efficient iteration over distributed data:
|
||||
Summary
|
||||
|
||||
SweepTile provides elegant and efficient iteration over distributed data:
|
||||
SweepTile provides clean and efficient iteration over distributed data:
|
||||
=======
|
||||
|
||||
SweepTile provides elegant and efficient iteration over distributed data:
|
||||
SweepTile provides clean and efficient iteration over distributed data:
|
||||
|
||||
- **Efficiency**: Load once, use many times pattern
|
||||
- **Simplicity**: Clean lambda-based iteration abstraction
|
||||
@@ -515,4 +547,4 @@ Key benefits:
|
||||
3. **Code Clarity**: Express algorithms naturally
|
||||
4. **Compiler Optimization**: Enable aggressive optimizations
|
||||
|
||||
The sweep pattern is fundamental to high-performance GPU kernels, turning complex iteration patterns into simple, efficient operations. Combined with TileDistribution and TileWindow, sweep operations complete the toolkit for elegant and performant GPU computing.
|
||||
The sweep pattern is fundamental to high-performance GPU kernels, turning complex iteration patterns into simple, efficient operations. Combined with TileDistribution and TileWindow, sweep operations complete the toolkit for clean and performant GPU computing.
|
||||
|
||||
@@ -8,14 +8,10 @@
|
||||
Memory Swizzling with Morton Ordering
|
||||
**************************************
|
||||
|
||||
|
||||
This chapter demonstrates a practical application of tensor descriptors for implementing memory swizzling patterns, specifically Morton ordering (Z-order curve) within tiles. Memory swizzling is crucial for optimizing GPU memory access patterns and reducing bank conflicts. Morton ordering provides a space-filling curve that maintains spatial locality while enabling efficient parallel access.
|
||||
Overview
|
||||
|
||||
This chapter demonstrates a practical application of tensor descriptors for implementing memory swizzling patterns, specifically Morton ordering (Z-order curve) within tiles. Memory swizzling is crucial for optimizing GPU memory access patterns and reducing bank conflicts (see :ref:`ck_tile_lds_bank_conflicts`). Morton ordering provides a space-filling curve that maintains spatial locality while enabling efficient parallel access (see :ref:`ck_tile_space_filling_curve` for theoretical background).
|
||||
========
|
||||
|
||||
This chapter demonstrates a practical application of tensor descriptors for implementing memory swizzling patterns, specifically Morton ordering (Z-order curve) within tiles. Memory swizzling is crucial for optimizing GPU memory access patterns and reducing bank conflicts. Morton ordering provides a space-filling curve that maintains spatial locality while enabling efficient parallel access.
|
||||
This chapter demonstrates a practical application of tensor descriptors for implementing memory swizzling patterns, specifically Morton ordering (Z-order curve) within tiles. Memory swizzling is crucial for optimizing GPU memory access patterns and reducing bank conflicts (see :ref:`ck_tile_lds_bank_conflicts`). Morton ordering provides a space-filling curve that maintains spatial locality while enabling efficient parallel access (see :ref:`ck_tile_space_filling_curve` for theoretical background).
|
||||
|
||||
Morton ordering is widely used in:
|
||||
|
||||
@@ -90,14 +86,10 @@ Bit pattern breakdown:
|
||||
(1,2) = (01, 10) → 6 = 0110
|
||||
(1,3) = (01, 11) → 7 = 0111
|
||||
|
||||
|
||||
First, we split our texture into tiles using tensor descriptors. This creates a hierarchical structure: (Y_blk, y_in, X_blk, x_in).
|
||||
Stage 1: Tiling with UnmergeTransform
|
||||
|
||||
First, we split our texture into tiles using tensor descriptors (see :ref:`ck_tile_descriptors` and :ref:`ck_tile_transforms`). This creates a hierarchical structure: (Y_blk, y_in, X_blk, x_in).
|
||||
======================================
|
||||
|
||||
First, we split our texture into tiles using tensor descriptors. This creates a hierarchical structure: (Y_blk, y_in, X_blk, x_in).
|
||||
First, we split our texture into tiles using tensor descriptors (see :ref:`ck_tile_descriptors` and :ref:`ck_tile_transforms`). This creates a hierarchical structure: (Y_blk, y_in, X_blk, x_in).
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -357,14 +349,10 @@ Here's a complete GPU kernel using Morton ordering for optimized memory access:
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
Morton ordering is particularly effective for reducing shared memory bank conflicts:
|
||||
Bank Conflict Reduction
|
||||
|
||||
Morton ordering is particularly effective for reducing shared memory bank conflicts (complementing the XOR preshuffle technique described in :ref:`ck_tile_lds_index_swapping`):
|
||||
=======================
|
||||
|
||||
Morton ordering is particularly effective for reducing shared memory bank conflicts:
|
||||
Morton ordering is particularly effective for reducing shared memory bank conflicts (complementing the XOR preshuffle technique described in :ref:`ck_tile_lds_index_swapping`):
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -487,12 +475,12 @@ For complete GEMM optimization techniques, see :ref:`ck_tile_gemm_optimization`.
|
||||
Summary
|
||||
=======
|
||||
|
||||
Morton ordering with CK Tile provides powerful memory optimization capabilities:
|
||||
Morton ordering with CK Tile provides memory optimization capabilities:
|
||||
|
||||
- **Spatial Locality**: Z-order curve maintains 2D locality in 1D memory layout
|
||||
- **Bank Conflict Reduction**: Distributed access patterns across memory banks
|
||||
- **Cache Efficiency**: Better utilization of cache lines for 2D access patterns
|
||||
- **Mathematical Framework**: Tensor descriptors express swizzling elegantly
|
||||
- **Mathematical Framework**: Tensor descriptors express swizzling cleanly
|
||||
- **Practical Implementation**: Bit manipulation provides reliable results
|
||||
|
||||
Key implementation insights:
|
||||
|
||||
@@ -8,51 +8,56 @@
|
||||
Tensor Coordinates
|
||||
*******************
|
||||
|
||||
|
||||
Before diving into transforms and adaptors, it's essential to understand the basic coordinate system in CK Tile. MultiIndex is the fundamental building block used throughout the system - a container that extends the C++ array with additional operations for multi-dimensional indexing. Think of MultiIndex as GPS coordinates for tensors, providing a universal way to specify positions in N-dimensional space.
|
||||
Overview
|
||||
|
||||
Before diving into transforms and adaptors (see :ref:`ck_tile_transforms` and :ref:`ck_tile_adaptors`), it's essential to understand the basic coordinate system in CK Tile. MultiIndex is the fundamental building block used throughout the system - a container that extends the C++ array with additional operations for multi-dimensional indexing. Think of MultiIndex as GPS coordinates for tensors, providing a universal way to specify positions in N-dimensional space.
|
||||
========
|
||||
|
||||
Before diving into transforms and adaptors, it's essential to understand the basic coordinate system in CK Tile. MultiIndex is the fundamental building block used throughout the system - a container that extends the C++ array with additional operations for multi-dimensional indexing. Think of MultiIndex as GPS coordinates for tensors, providing a universal way to specify positions in N-dimensional space.
|
||||
Before diving into transforms and adaptors (see :ref:`ck_tile_transforms` and :ref:`ck_tile_adaptors`), it's essential to understand the basic coordinate system in CK Tile. MultiIndex is the fundamental building block used throughout the system - a container that extends the C++ array with additional operations for multi-dimensional indexing. Think of MultiIndex as GPS coordinates for tensors, providing a universal way to specify positions in N-dimensional space.
|
||||
|
||||
MultiIndex serves as the common currency between different coordinate spaces (see :ref:`ck_tile_coordinate_systems`), enabling seamless transformation and navigation through complex tensor layouts. Every transform, adaptor, and descriptor in CK Tile operates on these coordinate containers.
|
||||
|
||||
.. mermaid::
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "MultiIndex Structure"
|
||||
MI["MultiIndex<br/>Container for N integers"]
|
||||
D0["Dimension 0"]
|
||||
D1["Dimension 1"]
|
||||
D2["Dimension 2"]
|
||||
DN["Dimension N-1"]
|
||||
end
|
||||
|
||||
subgraph "Usage Context"
|
||||
T["Transforms<br/>"]
|
||||
A["Adaptors<br/>"]
|
||||
TV["Tensors<br/>"]
|
||||
end
|
||||
|
||||
MI --> D0
|
||||
MI --> D1
|
||||
MI --> D2
|
||||
MI --> DN
|
||||
|
||||
T --> MI
|
||||
A --> MI
|
||||
TV --> MI
|
||||
|
||||
style MI fill:#f3e5f5,stroke:#7b1fa2,stroke-width:3px
|
||||
style D0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D2 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style DN fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style A fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style TV fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
graph TB
|
||||
subgraph "MultiIndex Structure"
|
||||
MI["MultiIndex<br/>Container for N integers"]
|
||||
D0["Dimension 0"]
|
||||
D1["Dimension 1"]
|
||||
D2["Dimension 2"]
|
||||
DN["Dimension N-1"]
|
||||
end
|
||||
|
||||
subgraph "Usage Context"
|
||||
T["Transforms<br/>"]
|
||||
A["Adaptors<br/>"]
|
||||
TV["Tensors<br/>"]
|
||||
end
|
||||
|
||||
MI --> D0
|
||||
MI --> D1
|
||||
MI --> D2
|
||||
MI --> DN
|
||||
|
||||
T --> MI
|
||||
A --> MI
|
||||
TV --> MI
|
||||
|
||||
style MI fill:#f3e5f5,stroke:#7b1fa2,stroke-width:3px
|
||||
style D0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D1 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D2 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style DN fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style A fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style TV fill:#ffebee,stroke:#d32f2f,stroke-width:2px
|
||||
.. image:: diagrams/tensor_coordinates_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
|
||||
MultiIndex Implementation
|
||||
=========================
|
||||
@@ -171,27 +176,36 @@ MultiIndex in Coordinate Flow
|
||||
|
||||
MultiIndex serves as the interface between user code and the transformation pipeline:
|
||||
|
||||
.. mermaid::
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart TB
|
||||
subgraph CF ["Coordinate Flow"]
|
||||
direction LR
|
||||
UI["User Input<br/>[1, 2, 3]"] --> MI["MultiIndex<br/>Storage"]
|
||||
MI --> TR["Transform<br/>Processing"]
|
||||
TR --> MO["MultiIndex<br/>Output"]
|
||||
MO --> TA["Tensor Access<br/>element(coord)"]
|
||||
end
|
||||
|
||||
subgraph EX ["Example: 3D Tensor Access"]
|
||||
direction LR
|
||||
T3D["3D Tensor<br/>shape=[4,5,6]"] --> COORD["MultiIndex(3, [1,2,3])"]
|
||||
COORD --> ELEM["Element at<br/>position [1,2,3]"]
|
||||
end
|
||||
|
||||
style UI fill:#e0e7ff,stroke:#4338ca,stroke-width:2px
|
||||
style MI fill:#f3e5f5,stroke:#7b1fa2,stroke-width:2px
|
||||
style MO fill:#f3e5f5,stroke:#7b1fa2,stroke-width:2px
|
||||
style COORD fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
flowchart TB
|
||||
subgraph CF ["Coordinate Flow"]
|
||||
direction LR
|
||||
UI["User Input<br/>[1, 2, 3]"] --> MI["MultiIndex<br/>Storage"]
|
||||
MI --> TR["Transform<br/>Processing"]
|
||||
TR --> MO["MultiIndex<br/>Output"]
|
||||
MO --> TA["Tensor Access<br/>element(coord)"]
|
||||
end
|
||||
|
||||
subgraph EX ["Example: 3D Tensor Access"]
|
||||
direction LR
|
||||
T3D["3D Tensor<br/>shape=[4,5,6]"] --> COORD["MultiIndex(3, [1,2,3])"]
|
||||
COORD --> ELEM["Element at<br/>position [1,2,3]"]
|
||||
end
|
||||
|
||||
style UI fill:#e0e7ff,stroke:#4338ca,stroke-width:2px
|
||||
style MI fill:#f3e5f5,stroke:#7b1fa2,stroke-width:2px
|
||||
style MO fill:#f3e5f5,stroke:#7b1fa2,stroke-width:2px
|
||||
style COORD fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
.. image:: diagrams/tensor_coordinates_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
|
||||
Common Usage Patterns
|
||||
=====================
|
||||
@@ -349,14 +363,10 @@ Specialized Coordinates
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
MultiIndex is the foundation for all tensor operations in CK Tile:
|
||||
Integration with Tensor Operations
|
||||
|
||||
MultiIndex is the foundation for all tensor operations in CK Tile (see :ref:`ck_tile_tensor_views` and :ref:`ck_tile_buffer_views` for tensor abstractions):
|
||||
==================================
|
||||
|
||||
MultiIndex is the foundation for all tensor operations in CK Tile:
|
||||
MultiIndex is the foundation for all tensor operations in CK Tile (see :ref:`ck_tile_tensor_views` and :ref:`ck_tile_buffer_views` for tensor abstractions):
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -393,14 +403,10 @@ MultiIndex is the foundation for all tensor operations in CK Tile:
|
||||
});
|
||||
}
|
||||
|
||||
|
||||
MultiIndex is designed for zero-overhead abstraction:
|
||||
Performance Considerations
|
||||
|
||||
MultiIndex is designed for zero-overhead abstraction (see :ref:`ck_tile_gpu_basics` for GPU performance fundamentals):
|
||||
==========================
|
||||
|
||||
MultiIndex is designed for zero-overhead abstraction:
|
||||
MultiIndex is designed for zero-overhead abstraction (see :ref:`ck_tile_gpu_basics` for GPU performance fundamentals):
|
||||
|
||||
1. **Compile-Time Resolution**: When dimensions are known at compile time, all operations are inlined
|
||||
2. **Register Allocation**: Small fixed-size arrays typically stay in registers
|
||||
|
||||
@@ -8,7 +8,7 @@ Overview
|
||||
|
||||
While :ref:`BufferView <ck_tile_buffer_views>` provides the foundation for raw memory access, TensorView elevates this abstraction by adding multi-dimensional structure to flat memory regions. This critical abstraction bridges the gap between how developers conceptualize data—as matrices, tensors, and higher-dimensional structures—and how that data is physically stored in linear memory. TensorView enables coordinate-based access patterns that match the natural structure of algorithms while maintaining the performance characteristics necessary for efficient GPU computation.
|
||||
|
||||
The power of TensorView lies in its ability to present different logical views of the same underlying memory without copying data. A single memory region can be viewed as a row-major matrix, a column-major matrix, or even a transposed matrix, all through different TensorView configurations. This zero-copy abstraction enables powerful transformations and access patterns while maintaining optimal memory bandwidth utilization.
|
||||
The power of TensorView lies in its ability to present different logical views of the same underlying memory without copying data. A single memory region can be viewed as a row-major matrix, a column-major matrix, or even a transposed matrix, all through different TensorView configurations. This zero-copy abstraction enables flexible transformations and access patterns while maintaining optimal memory bandwidth utilization.
|
||||
|
||||
TensorView Architecture
|
||||
-----------------------
|
||||
@@ -52,7 +52,7 @@ The Foundation: BufferView and TensorDescriptor
|
||||
|
||||
TensorView builds upon two fundamental components that work in concert to provide structured access to memory. The :ref:`BufferView <ck_tile_buffer_views>` component handles the low-level memory access, providing type-safe operations with address space awareness. The :ref:`TensorDescriptor <ck_tile_descriptors>` component encodes the multi-dimensional structure, including shape information and stride patterns that determine how coordinates map to memory offsets.
|
||||
|
||||
This separation of concerns enables powerful optimizations. The BufferView can optimize for the specific memory space (global, shared, or register), while the TensorDescriptor can encode complex access patterns without concern for the underlying memory type. Together, they provide a complete abstraction for multi-dimensional data access.
|
||||
This separation of concerns enables critical optimizations. The BufferView can optimize for the specific memory space (global, shared, or register), while the TensorDescriptor can encode complex access patterns without concern for the underlying memory type. Together, they provide a complete abstraction for multi-dimensional data access.
|
||||
|
||||
C++ Implementation
|
||||
------------------
|
||||
@@ -107,7 +107,7 @@ The creation of a TensorView involves combining a BufferView with a TensorDescri
|
||||
Coordinate-Based Access
|
||||
~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
The fundamental operation of TensorView is translating multi-dimensional coordinates into memory accesses. This translation happens through a sophisticated pipeline that maintains efficiency while providing flexibility:
|
||||
The fundamental operation of TensorView is translating multi-dimensional coordinates into memory accesses. This translation happens through a advanced pipeline that maintains efficiency while providing flexibility:
|
||||
|
||||
.. raw:: html
|
||||
|
||||
@@ -142,7 +142,7 @@ The fundamental operation of TensorView is translating multi-dimensional coordin
|
||||
Memory Layouts and Strides
|
||||
--------------------------
|
||||
|
||||
One of the most powerful features of TensorView is its ability to represent different memory layouts through stride manipulation. This capability enables zero-copy transformations that would otherwise require expensive memory operations:
|
||||
One of the most key features of TensorView is its ability to represent different memory layouts through stride manipulation. This capability enables zero-copy transformations that would otherwise require expensive memory operations:
|
||||
|
||||
.. raw:: html
|
||||
|
||||
@@ -216,7 +216,7 @@ Advanced Operations
|
||||
Slicing and Subviews
|
||||
~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
TensorView supports sophisticated slicing operations that create new views of subsets of the data. These operations are essential for algorithms that process data in blocks or tiles (see :ref:`ck_tile_tile_window` for production use):
|
||||
TensorView supports advanced slicing operations that create new views of subsets of the data. These operations are essential for algorithms that process data in blocks or tiles (see :ref:`ck_tile_tile_window` for production use):
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -317,7 +317,7 @@ The efficiency of TensorView operations depends critically on memory access patt
|
||||
Compile-Time Optimization
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
CK's TensorView leverages extensive compile-time optimization to achieve zero-overhead abstraction. When tensor dimensions and strides are known at compile time, the entire coordinate-to-offset calculation can be resolved during compilation:
|
||||
CK's TensorView leverages compile-time optimization to achieve zero-overhead abstraction. When tensor dimensions and strides are known at compile time, the entire coordinate-to-offset calculation can be resolved during compilation:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -397,13 +397,13 @@ TensorView serves as the foundation for :ref:`tile distribution's <ck_tile_tile_
|
||||
Summary
|
||||
-------
|
||||
|
||||
TensorView represents a critical abstraction in the CK framework, bridging the gap between logical multi-dimensional data structures and physical memory layout. Through its sophisticated design, TensorView provides:
|
||||
TensorView represents a critical abstraction in the CK framework, bridging the gap between logical multi-dimensional data structures and physical memory layout. Through its advanced design, TensorView provides:
|
||||
|
||||
**Multi-dimensional Indexing**: Natural coordinate-based access to data, matching how algorithms conceptualize their operations. This abstraction eliminates error-prone manual index calculations while maintaining performance.
|
||||
|
||||
**Flexible Memory Layouts**: Support for row-major, column-major, and custom stride patterns enables algorithms to work with data in its most natural form. Zero-copy transformations like transposition become simple stride manipulations.
|
||||
|
||||
**Zero-Copy Views**: The ability to create different logical views of the same physical memory enables powerful transformations without the overhead of data movement. This capability is essential for efficient GPU programming where memory bandwidth is often the limiting factor.
|
||||
**Zero-Copy Views**: The ability to create different logical views of the same physical memory enables flexible transformations without the overhead of data movement. This capability is essential for efficient GPU programming where memory bandwidth is often the limiting factor.
|
||||
|
||||
**Type Safety**: Dimensions and memory spaces are encoded in the type system, catching errors at compile time rather than runtime. This safety comes without performance overhead thanks to template metaprogramming.
|
||||
|
||||
|
||||
@@ -8,7 +8,7 @@ Overview
|
||||
|
||||
The Composable Kernel framework introduces a rich vocabulary of concepts and abstractions that form the foundation of its approach to high-performance GPU computing. This terminology reference serves as a comprehensive guide to the language of CK, providing detailed explanations of each term along with practical examples of their usage in C++ code. Understanding this vocabulary is not merely an academic exercise—it represents the key to unlocking the full power of the CK framework and writing efficient, maintainable GPU kernels.
|
||||
|
||||
The terminology of CK reflects its layered architecture, with concepts building upon one another in a logical progression. From the fundamental notion of tiles and distributions to the sophisticated coordinate transformation systems, each term represents a carefully designed abstraction that serves a specific purpose in the overall framework. This reference is organized to mirror this conceptual hierarchy, starting with core concepts and progressing through increasingly specialized terminology.
|
||||
The terminology of CK reflects its layered architecture, with concepts building upon one another in a logical progression. From the fundamental notion of tiles and distributions to the compile-time coordinate transformation systems, each term represents a carefully designed abstraction that serves a specific purpose in the overall framework. This reference is organized to mirror this conceptual hierarchy, starting with core concepts and progressing through increasingly specialized terminology.
|
||||
|
||||
As you explore this reference, you'll notice that many terms are interconnected, reflecting the holistic nature of the CK design. A tile is not just a block of data but a fundamental unit of work distribution. A distribution is not merely a pattern but a mathematical framework for optimal resource utilization. These interconnections are intentional and understanding them is crucial for effective use of the framework.
|
||||
|
||||
@@ -23,7 +23,7 @@ The concept of a tile represents the fundamental unit of data organization in th
|
||||
|
||||
Distribution
|
||||
~~~~~~~~~~~~
|
||||
The distribution pattern represents one of the most sophisticated abstractions in the CK framework, defining the precise mapping between logical data elements and the physical processing resources that will operate on them. A distribution is far more than a simple assignment scheme—it embodies a complete strategy for achieving optimal performance on GPU hardware. The distribution determines which threads access which data elements, how those accesses are ordered to maximize memory bandwidth, and how intermediate results are shared between cooperating threads. By encoding these decisions at compile time, distributions enable the generation of highly optimized code that respects hardware constraints while maintaining algorithmic clarity. For a detailed exploration of distribution concepts, see :ref:`ck_tile_distribution`.
|
||||
The distribution pattern represents one of the most compile-time abstractions in the CK framework, defining the precise mapping between logical data elements and the physical processing resources that will operate on them. A distribution is far more than a simple assignment scheme—it embodies a complete strategy for achieving optimal performance on GPU hardware. The distribution determines which threads access which data elements, how those accesses are ordered to maximize memory bandwidth, and how intermediate results are shared between cooperating threads. By encoding these decisions at compile time, distributions enable the generation of highly optimized code that respects hardware constraints while maintaining algorithmic clarity. For a detailed exploration of distribution concepts, see :ref:`ck_tile_distribution`.
|
||||
|
||||
**C++ Type**: ``tile_distribution<...>``
|
||||
|
||||
@@ -42,7 +42,7 @@ P-Space (Partition Space)
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
The Partition Space, or P-space, represents the fundamental abstraction for identifying processing elements within the GPU's execution hierarchy. This coordinate space captures the multi-level organization of GPU computation, from individual threads to warps to thread blocks. P-space typically manifests as either a one-dimensional space containing only lane identifiers for simple distributions, or a two-dimensional space incorporating both warp and lane identifiers for more complex hierarchical distributions. The significance of P-space extends beyond mere thread identification—it forms the foundation for all work distribution decisions, determining which processing elements will collaborate on specific data tiles and how they will coordinate their efforts.
|
||||
|
||||
The dimensions of P-space directly reflect the hardware's execution model. In a one-dimensional P-space, threads are identified solely by their lane ID within a warp, suitable for algorithms where inter-warp coordination is minimal. Two-dimensional P-space adds warp-level coordination, enabling sophisticated tiling strategies that leverage both intra-warp and inter-warp parallelism. The values in P-space are always hardware thread indices, providing a direct mapping to the physical execution resources.
|
||||
The dimensions of P-space directly reflect the hardware's execution model. In a one-dimensional P-space, threads are identified solely by their lane ID within a warp, suitable for algorithms where inter-warp coordination is minimal. Two-dimensional P-space adds warp-level coordination, enabling advanced tiling strategies that leverage both intra-warp and inter-warp parallelism. The values in P-space are always hardware thread indices, providing a direct mapping to the physical execution resources.
|
||||
|
||||
**C++ Example**:
|
||||
|
||||
@@ -53,9 +53,9 @@ The dimensions of P-space directly reflect the hardware's execution model. In a
|
||||
|
||||
Y-Space (Yield Space)
|
||||
~~~~~~~~~~~~~~~~~~~~~
|
||||
The Yield Space, or Y-space, embodies the logical structure of computation within each tile, representing the pattern by which threads traverse their assigned data. Unlike P-space which identifies threads, Y-space defines what each thread does with its assigned work. This abstraction enables the expression of complex access patterns—from simple linear traversals to sophisticated space-filling curves—in a hardware-independent manner. The dimensionality of Y-space varies with the algorithm's requirements, typically ranging from two dimensions for matrix operations to four or more for complex tensor contractions.
|
||||
The Yield Space, or Y-space, embodies the logical structure of computation within each tile, representing the pattern by which threads traverse their assigned data. Unlike P-space which identifies threads, Y-space defines what each thread does with its assigned work. This abstraction enables the expression of complex access patterns—from simple linear traversals to advanced space-filling curves—in a hardware-independent manner. The dimensionality of Y-space varies with the algorithm's requirements, typically ranging from two dimensions for matrix operations to four or more for complex tensor contractions.
|
||||
|
||||
Y-space serves as the primary iteration space for computational kernels. When a thread processes its assigned tile, it iterates through Y-space coordinates, with each coordinate mapping to specific data elements within the tile. This abstraction enables powerful optimizations: the Y-space traversal order can be designed to maximize data reuse, minimize register pressure, or optimize for specific hardware characteristics, all without changing the fundamental algorithm.
|
||||
Y-space serves as the primary iteration space for computational kernels. When a thread processes its assigned tile, it iterates through Y-space coordinates, with each coordinate mapping to specific data elements within the tile. This abstraction enables critical optimizations: the Y-space traversal order can be designed to maximize data reuse, minimize register pressure, or optimize for specific hardware characteristics, all without changing the fundamental algorithm.
|
||||
|
||||
**C++ Example**:
|
||||
|
||||
@@ -79,7 +79,7 @@ The relationship between X-space and physical memory is direct but not necessari
|
||||
|
||||
R-Space (Replication Space)
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
The Replication Space, or R-space, introduces a sophisticated mechanism for expressing redundant computation patterns that enhance performance through data sharing. Unlike the other coordinate spaces which map to unique data elements, R-space enables multiple processing elements to compute the same values, facilitating efficient communication patterns. This replication serves multiple purposes: it can reduce global memory traffic by computing values locally rather than loading them, enable efficient reduction operations by providing private workspace for each thread group, and facilitate complex data exchange patterns that would otherwise require expensive synchronization.
|
||||
The Replication Space, or R-space, introduces a advanced mechanism for expressing redundant computation patterns that enhance performance through data sharing. Unlike the other coordinate spaces which map to unique data elements, R-space enables multiple processing elements to compute the same values, facilitating efficient communication patterns. This replication serves multiple purposes: it can reduce global memory traffic by computing values locally rather than loading them, enable efficient reduction operations by providing private workspace for each thread group, and facilitate complex data exchange patterns that would otherwise require expensive synchronization.
|
||||
|
||||
R-space dimensions are optional and algorithm-specific. A matrix multiplication might use R-space to replicate portions of the input matrices across thread groups, enabling each group to compute partial products independently. The framework automatically manages the complexities of replication, including the allocation of private storage and the coordination of replicated computations.
|
||||
|
||||
@@ -97,7 +97,7 @@ D-Space (Data Space)
|
||||
~~~~~~~~~~~~~~~~~~~~
|
||||
The Data Space, or D-space, represents the final stage of the coordinate transformation pipeline—the linearization of multi-dimensional tile data for efficient storage in thread-local registers. This one-dimensional space serves a critical role in managing the GPU's most precious resource: register files. By transforming the potentially complex Y-space coordinates into a linear D-space index, the framework enables efficient register allocation and access patterns that minimize register bank conflicts and maximize instruction-level parallelism.
|
||||
|
||||
The transformation from Y-space to D-space is more than a simple flattening operation. It incorporates sophisticated strategies for register layout that consider the GPU's register file organization, the kernel's register pressure, and the access patterns of the computation. This transformation ensures that frequently accessed elements are kept in registers, that register bank conflicts are minimized, and that the compiler can generate efficient code for register access.
|
||||
The transformation from Y-space to D-space is more than a simple flattening operation. It incorporates optimized strategies for register layout that consider the GPU's register file organization, the kernel's register pressure, and the access patterns of the computation. This transformation ensures that frequently accessed elements are kept in registers, that register bank conflicts are minimized, and that the compiler can generate efficient code for register access.
|
||||
|
||||
**C++ Example**:
|
||||
|
||||
@@ -111,11 +111,11 @@ Dimension Types
|
||||
|
||||
H-Dimensions (Hierarchical Dimensions)
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
The concept of Hierarchical Dimensions, or H-dimensions, represents one of the most sophisticated aspects of the CK framework's approach to work distribution. These dimensions encode a multi-level decomposition strategy that mirrors the hierarchical nature of GPU hardware, from individual vector operations up through threads, warps, and thread blocks. Each H-dimension group captures how a single tensor dimension is partitioned across these hardware levels, enabling fine-grained control over data access patterns and computational efficiency.
|
||||
The concept of Hierarchical Dimensions, or H-dimensions, represents one of the most key aspects of the CK framework's approach to work distribution. These dimensions encode a multi-level decomposition strategy that mirrors the hierarchical nature of GPU hardware, from individual vector operations up through threads, warps, and thread blocks. Each H-dimension group captures how a single tensor dimension is partitioned across these hardware levels, enabling fine-grained control over data access patterns and computational efficiency.
|
||||
|
||||
The structure of H-dimensions follows a specific pattern that reflects the GPU's execution hierarchy. Each H-dimension is expressed as a sequence of factors, where each factor corresponds to a specific level of the hierarchy. Consider the example ``sequence<4, 2, 8, 4>``. This seemingly simple sequence encodes a sophisticated distribution strategy: the rightmost factor (4) represents vector width, indicating that each memory operation processes 4 elements simultaneously. Moving left, the factor 8 indicates that 8 threads within a warp collaborate on the data. The factor 2 specifies that 2 warps within a block work together. Finally, the leftmost factor 4 indicates that each thread performs 4 iterations, enabling instruction-level parallelism and register reuse.
|
||||
The structure of H-dimensions follows a specific pattern that reflects the GPU's execution hierarchy. Each H-dimension is expressed as a sequence of factors, where each factor corresponds to a specific level of the hierarchy. Consider the example ``sequence<4, 2, 8, 4>``. This seemingly simple sequence encodes a advanced distribution strategy: the rightmost factor (4) represents vector width, indicating that each memory operation processes 4 elements simultaneously. Moving left, the factor 8 indicates that 8 threads within a warp collaborate on the data. The factor 2 specifies that 2 warps within a block work together. Finally, the leftmost factor 4 indicates that each thread performs 4 iterations, enabling instruction-level parallelism and register reuse.
|
||||
|
||||
This hierarchical decomposition enables powerful optimizations. By explicitly encoding the distribution strategy at compile time, the framework can generate code that perfectly matches the hardware's capabilities. The vector width aligns with the GPU's memory transaction size. The thread count per warp matches the hardware's SIMD width. The warp count per block balances parallelism with resource constraints. The repetition factor enables loop unrolling and software pipelining. Together, these factors create a distribution strategy that achieves near-optimal performance.
|
||||
This hierarchical decomposition enables critical optimizations. By explicitly encoding the distribution strategy at compile time, the framework can generate code that perfectly matches the hardware's capabilities. The vector width aligns with the GPU's memory transaction size. The thread count per warp matches the hardware's SIMD width. The warp count per block balances parallelism with resource constraints. The repetition factor enables loop unrolling and software pipelining. Together, these factors create a distribution strategy that achieves near-optimal performance.
|
||||
|
||||
**C++ Example**:
|
||||
|
||||
@@ -139,11 +139,11 @@ Transformations
|
||||
|
||||
Adaptor
|
||||
~~~~~~~
|
||||
An adaptor in the CK framework represents a sophisticated chain of coordinate transformations that bridges different coordinate spaces. Rather than simple one-to-one mappings, adaptors embody complex mathematical transformations that can involve permutations, embeddings, projections, and non-linear mappings. These transformations are composed at compile time, enabling the generation of highly optimized code that performs the complete transformation in a single step without intermediate representations. For detailed information about adaptors and their implementation, see :ref:`ck_tile_adaptors`.
|
||||
An adaptor in the CK framework represents a advanced chain of coordinate transformations that bridges different coordinate spaces. Rather than simple one-to-one mappings, adaptors embody complex mathematical transformations that can involve permutations, embeddings, projections, and non-linear mappings. These transformations are composed at compile time, enabling the generation of highly optimized code that performs the complete transformation in a single step without intermediate representations. For detailed information about adaptors and their implementation, see :ref:`ck_tile_adaptors`.
|
||||
|
||||
The framework provides several specialized adaptor types, each serving a specific role in the coordinate transformation pipeline. The ``ps_ys_to_xs_adaptor`` performs the critical transformation from processing element and yield space coordinates to physical tensor coordinates, implementing the core logic of tile distribution. This adaptor encodes decisions about how threads are assigned to data, how data is traversed within each thread's assignment, and how these patterns map to the global tensor layout. Similarly, the ``ys_to_d_adaptor`` handles the transformation from multi-dimensional yield space to linearized data space, optimizing the layout of data in thread-local registers.
|
||||
|
||||
The power of adaptors lies in their composability. Complex transformations can be built by chaining simpler adaptors, with the framework automatically optimizing the composition. This design enables the expression of sophisticated access patterns—such as transposed access, strided access, or space-filling curves—through the composition of elementary transformations. The compile-time nature of this composition ensures zero runtime overhead while maintaining mathematical clarity.
|
||||
The power of adaptors lies in their composability. Complex transformations can be built by chaining simpler adaptors, with the framework automatically optimizing the composition. This design enables the expression of advanced access patterns—such as transposed access, strided access, or space-filling curves—through the composition of elementary transformations. The compile-time nature of this composition ensures zero runtime overhead while maintaining mathematical clarity.
|
||||
|
||||
**C++ Type**: ``tensor_adaptor<...>``
|
||||
|
||||
@@ -162,15 +162,15 @@ Operations
|
||||
|
||||
Load Tile
|
||||
~~~~~~~~~
|
||||
The load tile operation represents a fundamental building block of GPU kernel design in the CK framework, orchestrating the complex process of transferring data from global memory to thread-local registers. This operation is far more sophisticated than a simple memory copy—it implements the complete distribution strategy encoded in the tile distribution, ensuring that each thread loads exactly the data it needs for its portion of the computation. The load operation automatically handles memory coalescing to maximize bandwidth utilization, coordinates between threads to avoid redundant loads, manages boundary conditions for tiles that extend beyond tensor bounds, and optimizes the access pattern based on the specific distribution strategy.
|
||||
The load tile operation represents a fundamental building block of GPU kernel design in the CK framework, orchestrating the complex process of transferring data from global memory to thread-local registers. This operation is far more advanced than a simple memory copy—it implements the complete distribution strategy encoded in the tile distribution, ensuring that each thread loads exactly the data it needs for its portion of the computation. The load operation automatically handles memory coalescing to maximize bandwidth utilization, coordinates between threads to avoid redundant loads, manages boundary conditions for tiles that extend beyond tensor bounds, and optimizes the access pattern based on the specific distribution strategy.
|
||||
|
||||
The efficiency of the load tile operation stems from its deep integration with the distribution framework. By knowing at compile time exactly which threads will access which data elements, the operation can generate optimal memory access patterns that fully utilize the GPU's memory subsystem. For matrix multiplication, this might mean loading data in a pattern that ensures perfect coalescing. For convolution, it might involve sophisticated patterns that minimize the number of redundant loads while respecting the GPU's cache hierarchy.
|
||||
The efficiency of the load tile operation stems from its deep integration with the distribution framework. By knowing at compile time exactly which threads will access which data elements, the operation can generate optimal memory access patterns that fully utilize the GPU's memory subsystem. For matrix multiplication, this might mean loading data in a pattern that ensures perfect coalescing. For convolution, it might involve complex patterns that minimize the number of redundant loads while respecting the GPU's cache hierarchy.
|
||||
|
||||
**C++ Function**: ``tile_window.load()``
|
||||
|
||||
Store Tile
|
||||
~~~~~~~~~~
|
||||
The store tile operation provides the complementary functionality to load tile, transferring computed results from thread-local registers back to global memory. Like its counterpart, the store operation implements sophisticated strategies that go beyond simple memory writes. It ensures that writes are coalesced for maximum bandwidth efficiency, coordinates between threads to handle overlapping write regions correctly, manages atomic operations when multiple threads write to the same location, and optimizes write patterns to minimize memory traffic.
|
||||
The store tile operation provides the complementary functionality to load tile, transferring computed results from thread-local registers back to global memory. Like its counterpart, the store operation implements optimized strategies that go beyond simple memory writes. It ensures that writes are coalesced for maximum bandwidth efficiency, coordinates between threads to handle overlapping write regions correctly, manages atomic operations when multiple threads write to the same location, and optimizes write patterns to minimize memory traffic.
|
||||
|
||||
The store operation must handle additional complexities compared to loads. While loads can often ignore synchronization issues (reading stale data is usually harmless), stores must ensure correctness when multiple threads write to overlapping regions. The framework provides different store modes for different scenarios: exclusive stores where each element is written by exactly one thread, atomic stores where multiple threads may update the same element, and reduction stores where partial results are accumulated. The choice of store mode is encoded in the distribution strategy and verified at compile time.
|
||||
|
||||
@@ -178,7 +178,7 @@ The store operation must handle additional complexities compared to loads. While
|
||||
|
||||
Sweep Tile
|
||||
~~~~~~~~~~
|
||||
The sweep tile operation embodies a powerful programming paradigm for distributed tensor computation, providing a high-level iteration abstraction over the complex distribution patterns. Rather than requiring manual index calculations and nested loops, sweep tile automatically visits each element in a distributed tensor exactly once, invoking a user-provided function with the appropriate coordinates. This abstraction hides the complexity of the distribution while enabling sophisticated optimizations such as automatic loop unrolling, software pipelining, and register rotation.
|
||||
The sweep tile operation embodies a key programming paradigm for distributed tensor computation, providing a high-level iteration abstraction over the complex distribution patterns. Rather than requiring manual index calculations and nested loops, sweep tile automatically visits each element in a distributed tensor exactly once, invoking a user-provided function with the appropriate coordinates. This abstraction hides the complexity of the distribution while enabling advanced optimizations such as automatic loop unrolling, software pipelining, and register rotation.
|
||||
|
||||
The implementation of sweep tile leverages the compile-time knowledge of the distribution pattern to generate highly optimized iteration code. For simple distributions, this might result in a single unrolled loop. For complex hierarchical distributions, it might generate nested loops with carefully chosen iteration orders that maximize data reuse and minimize register pressure. The beauty of the abstraction is that these optimizations happen transparently—the user simply provides the computation to perform on each element, and the framework handles the rest.
|
||||
|
||||
|
||||
@@ -79,51 +79,59 @@ Composable Kernel abstracts thread identification into partition indices, buildi
|
||||
}
|
||||
};
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "GPU Device"
|
||||
subgraph "Thread Block"
|
||||
subgraph "Warp 0"
|
||||
T0["Thread 0<br/>lane_id=0"]
|
||||
T1["Thread 1<br/>lane_id=1"]
|
||||
T2["..."]
|
||||
T31["Thread 31<br/>lane_id=31"]
|
||||
end
|
||||
|
||||
subgraph "Warp 1"
|
||||
T32["Thread 32<br/>lane_id=0"]
|
||||
T33["Thread 33<br/>lane_id=1"]
|
||||
T34["..."]
|
||||
T63["Thread 63<br/>lane_id=31"]
|
||||
end
|
||||
|
||||
W2["Warp 2"]
|
||||
W3["..."]
|
||||
W7["Warp 7"]
|
||||
end
|
||||
end
|
||||
|
||||
subgraph "Thread Identification"
|
||||
TID["Thread ID = blockIdx.x * blockDim.x + threadIdx.x"]
|
||||
WID["Warp ID = threadIdx.x / 32"]
|
||||
LID["Lane ID = threadIdx.x % 32"]
|
||||
end
|
||||
|
||||
subgraph "P-space Mapping"
|
||||
P["P-coordinates<br/>NDimP=1: [thread_id]<br/>NDimP=2: [warp_id, lane_id]"]
|
||||
end
|
||||
|
||||
T0 --> TID
|
||||
TID --> WID
|
||||
TID --> LID
|
||||
WID --> P
|
||||
LID --> P
|
||||
|
||||
style T0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T32 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style P fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "GPU Device"
|
||||
subgraph "Thread Block"
|
||||
subgraph "Warp 0"
|
||||
T0["Thread 0<br/>lane_id=0"]
|
||||
T1["Thread 1<br/>lane_id=1"]
|
||||
T2["..."]
|
||||
T31["Thread 31<br/>lane_id=31"]
|
||||
end
|
||||
|
||||
subgraph "Warp 1"
|
||||
T32["Thread 32<br/>lane_id=0"]
|
||||
T33["Thread 33<br/>lane_id=1"]
|
||||
T34["..."]
|
||||
T63["Thread 63<br/>lane_id=31"]
|
||||
end
|
||||
|
||||
W2["Warp 2"]
|
||||
W3["..."]
|
||||
W7["Warp 7"]
|
||||
end
|
||||
end
|
||||
|
||||
subgraph "Thread Identification"
|
||||
TID["Thread ID = blockIdx.x * blockDim.x + threadIdx.x"]
|
||||
WID["Warp ID = threadIdx.x / 32"]
|
||||
LID["Lane ID = threadIdx.x % 32"]
|
||||
end
|
||||
|
||||
subgraph "P-space Mapping"
|
||||
P["P-coordinates<br/>NDimP=1: [thread_id]<br/>NDimP=2: [warp_id, lane_id]"]
|
||||
end
|
||||
|
||||
T0 --> TID
|
||||
TID --> WID
|
||||
TID --> LID
|
||||
WID --> P
|
||||
LID --> P
|
||||
|
||||
style T0 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style T32 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style P fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/thread_mapping_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Thread Hierarchy Structure
|
||||
--------------------------
|
||||
|
||||
@@ -168,41 +176,49 @@ Thread-to-Data Mapping
|
||||
|
||||
Once threads know their IDs, they need to map those IDs to specific data elements.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Thread to Data Mapping"
|
||||
subgraph "Thread Grid"
|
||||
T00["Thread[0,0]<br/>Warp 0"]
|
||||
T01["Thread[0,1]<br/>Warp 0"]
|
||||
T10["Thread[1,0]<br/>Warp 1"]
|
||||
T11["Thread[1,1]<br/>Warp 1"]
|
||||
end
|
||||
|
||||
subgraph "Data Tiles"
|
||||
D00["Data[0:4, 0:4]<br/>16 elements"]
|
||||
D01["Data[0:4, 4:8]<br/>16 elements"]
|
||||
D10["Data[4:8, 0:4]<br/>16 elements"]
|
||||
D11["Data[4:8, 4:8]<br/>16 elements"]
|
||||
end
|
||||
|
||||
subgraph "Memory Access"
|
||||
MA["Coalesced Access<br/>Adjacent threads → Adjacent memory"]
|
||||
end
|
||||
end
|
||||
|
||||
T00 --> D00
|
||||
T01 --> D01
|
||||
T10 --> D10
|
||||
T11 --> D11
|
||||
|
||||
D00 --> MA
|
||||
D01 --> MA
|
||||
|
||||
style T00 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D00 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style MA fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Thread to Data Mapping"
|
||||
subgraph "Thread Grid"
|
||||
T00["Thread[0,0]<br/>Warp 0"]
|
||||
T01["Thread[0,1]<br/>Warp 0"]
|
||||
T10["Thread[1,0]<br/>Warp 1"]
|
||||
T11["Thread[1,1]<br/>Warp 1"]
|
||||
end
|
||||
|
||||
subgraph "Data Tiles"
|
||||
D00["Data[0:4, 0:4]<br/>16 elements"]
|
||||
D01["Data[0:4, 4:8]<br/>16 elements"]
|
||||
D10["Data[4:8, 0:4]<br/>16 elements"]
|
||||
D11["Data[4:8, 4:8]<br/>16 elements"]
|
||||
end
|
||||
|
||||
subgraph "Memory Access"
|
||||
MA["Coalesced Access<br/>Adjacent threads → Adjacent memory"]
|
||||
end
|
||||
end
|
||||
|
||||
T00 --> D00
|
||||
T01 --> D01
|
||||
T10 --> D10
|
||||
T11 --> D11
|
||||
|
||||
D00 --> MA
|
||||
D01 --> MA
|
||||
|
||||
style T00 fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style D00 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style MA fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/thread_mapping_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Data Distribution Pattern
|
||||
-------------------------
|
||||
|
||||
|
||||
@@ -6,9 +6,9 @@ Tile Distribution - The Core API
|
||||
Overview
|
||||
--------
|
||||
|
||||
At the heart of Composable Kernels' approach to efficient GPU computation lies TileDistribution, a sophisticated abstraction that fundamentally transforms how developers approach parallel programming on GPUs. Rather than requiring programmers to manually orchestrate the complex choreography of thread coordination, memory access patterns, and data distribution, TileDistribution provides an elegant mathematical framework that automatically maps logical computational coordinates to physical execution resources. This abstraction represents not merely a convenience layer but a fundamental rethinking of how GPU programs should be structured, elevating the level of abstraction while maintaining—and often improving—the performance characteristics of hand-optimized code.
|
||||
At the heart of Composable Kernels' approach to efficient GPU computation lies TileDistribution, a compile-time abstraction that fundamentally transforms how developers approach parallel programming on GPUs. Rather than requiring programmers to manually orchestrate the complex choreography of thread coordination, memory access patterns, and data distribution, TileDistribution provides an mathematical framework that automatically maps logical computational coordinates to physical execution resources. This abstraction represents not merely a convenience layer but a fundamental rethinking of how GPU programs should be structured, elevating the level of abstraction while maintaining—and often improving—the performance characteristics of hand-optimized code.
|
||||
|
||||
The architectural foundation of tile distribution in CK rests upon a sophisticated :ref:`coordinate transformation system <ck_tile_coordinate_systems>` that seamlessly bridges multiple abstract spaces. This system orchestrates the interaction between four primary coordinate dimensions, each serving a distinct purpose in the overall computation model. The X dimensions represent the physical tensor coordinates, capturing the actual layout of data in memory. The Y dimensions encode the tile access patterns, defining how threads traverse their assigned data. The P dimensions map to processing elements, representing the hierarchical organization of threads, warps, and blocks in the :ref:`GPU's execution model <ck_tile_gpu_basics>`. Additionally, the optional R dimensions enable replication strategies for algorithms that benefit from redundant computation to reduce communication overhead.
|
||||
The architectural foundation of tile distribution in CK rests upon a advanced :ref:`coordinate transformation system <ck_tile_coordinate_systems>` that seamlessly bridges multiple abstract spaces. This system orchestrates the interaction between four primary coordinate dimensions, each serving a distinct purpose in the overall computation model. The X dimensions represent the physical tensor coordinates, capturing the actual layout of data in memory. The Y dimensions encode the tile access patterns, defining how threads traverse their assigned data. The P dimensions map to processing elements, representing the hierarchical organization of threads, warps, and blocks in the :ref:`GPU's execution model <ck_tile_gpu_basics>`. Additionally, the optional R dimensions enable replication strategies for algorithms that benefit from redundant computation to reduce communication overhead.
|
||||
|
||||
This multi-dimensional mapping framework enables CK to express arbitrarily complex data access patterns through a mathematically rigorous formalism. The power of this approach becomes evident when considering how traditional GPU programming requires developers to manually calculate memory addresses, ensure coalescing constraints, :ref:`avoid bank conflicts <ck_tile_lds_bank_conflicts>`, and manage the intricate dance of thread cooperation. TileDistribution encapsulates all these concerns within a unified abstraction that can be analyzed, optimized, and verified at compile time.
|
||||
|
||||
@@ -99,9 +99,9 @@ Before diving into code, let's understand the fundamental problem TileDistributi
|
||||
|
||||
The traditional approach without a tile distribution framework requires programmers to manually calculate global memory addresses for each thread, implement complex index arithmetic that accounts for thread hierarchy (threads within warps, warps within blocks), handle edge cases for non-divisible matrix dimensions, and create different implementations for various matrix sizes. This manual approach is not only error-prone but also fails to adapt to different GPU architectures and their specific memory access patterns.
|
||||
|
||||
TileDistribution elegantly solves these challenges through a systematic approach to work distribution. It automatically assigns work to threads based on a hierarchical decomposition of the problem space, generates memory access patterns that respect GPU hardware constraints, provides a uniform interface that works across different tensor sizes and shapes, and ensures optimal thread cooperation by automatically managing data movement to thread-local registers.
|
||||
TileDistribution solves these challenges through a systematic approach to work distribution. It automatically assigns work to threads based on a hierarchical decomposition of the problem space, generates memory access patterns that respect GPU hardware constraints, provides a uniform interface that works across different tensor sizes and shapes, and ensures optimal thread cooperation by automatically managing data movement to thread-local registers.
|
||||
|
||||
The key insight that makes TileDistribution powerful is its ability to abstract the mapping between logical problem coordinates and physical execution resources. Given a thread's position in the GPU's execution hierarchy (specified by warp ID and lane ID within the warp), TileDistribution computes two critical pieces of information: the global memory addresses that this thread should access, and the specific access pattern that ensures efficient memory transactions. This abstraction is implemented in C++ through the following core structure:
|
||||
The key insight that makes TileDistribution effective is its ability to abstract the mapping between logical problem coordinates and physical execution resources. Given a thread's position in the GPU's execution hierarchy (specified by warp ID and lane ID within the warp), TileDistribution computes two critical pieces of information: the global memory addresses that this thread should access, and the specific access pattern that ensures efficient memory transactions. This abstraction is implemented in C++ through the following core structure:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -540,21 +540,21 @@ Performance Comparison
|
||||
Summary
|
||||
-------
|
||||
|
||||
The TileDistribution abstraction represents a paradigm shift in GPU programming methodology, fundamentally altering how developers approach the challenge of parallel computation on massively parallel architectures. Through its sophisticated coordinate transformation framework and compile-time optimization strategies, TileDistribution achieves what has long been considered impossible in systems programming: a high-level abstraction that enhances rather than compromises performance.
|
||||
The TileDistribution abstraction represents a paradigm shift in GPU programming methodology, fundamentally altering how developers approach the challenge of parallel computation on massively parallel architectures. Through its compile-time coordinate transformation framework and compile-time optimization strategies, TileDistribution achieves what has long been considered impossible in systems programming: a high-level abstraction that enhances rather than compromises performance.
|
||||
|
||||
The automatic work distribution capabilities of TileDistribution eliminate one of the most error-prone aspects of GPU programming. Traditional approaches require developers to manually calculate how computational work maps to the physical thread hierarchy, a process fraught with opportunities for subtle bugs that manifest as incorrect results or catastrophic performance degradation. TileDistribution's mathematical framework ensures that every thread knows precisely which data elements it should process, automatically handling the complex index arithmetic that would otherwise consume hundreds of lines of error-prone code.
|
||||
|
||||
Memory access pattern optimization represents perhaps the most significant performance benefit of the TileDistribution approach. Modern GPUs achieve their remarkable computational throughput only when memory accesses follow specific patterns that enable hardware optimizations such as coalescing and broadcast. TileDistribution automatically generates these optimal patterns, ensuring that threads within a warp access contiguous memory locations, that bank conflicts in shared memory are minimized or eliminated, and that the memory subsystem operates at peak efficiency. This optimization happens transparently, freeing developers from the burden of manual memory pattern analysis.
|
||||
Memory access pattern optimization represents perhaps the most significant performance benefit of the TileDistribution approach. Modern GPUs achieve their computational throughput only when memory accesses follow specific patterns that enable hardware optimizations such as coalescing and broadcast. TileDistribution automatically generates these optimal patterns, ensuring that threads within a warp access contiguous memory locations, that bank conflicts in shared memory are minimized or eliminated, and that the memory subsystem operates at peak efficiency. This optimization happens transparently, freeing developers from the burden of manual memory pattern analysis.
|
||||
|
||||
The hierarchical decomposition strategy embedded within TileDistribution reflects a deep understanding of GPU architecture. By encoding the natural hierarchy of threads, warps, and blocks directly into the distribution strategy, the framework ensures that each level of the hierarchy operates optimally. This hierarchical approach enables sophisticated tiling strategies that would be impractical to implement manually, such as multi-level tiling that simultaneously optimizes for L1 cache, L2 cache, and register file usage.
|
||||
The hierarchical decomposition strategy embedded within TileDistribution reflects a deep understanding of GPU architecture. By encoding the natural hierarchy of threads, warps, and blocks directly into the distribution strategy, the framework ensures that each level of the hierarchy operates optimally. This hierarchical approach enables advanced tiling strategies that would be impractical to implement manually, such as multi-level tiling that simultaneously optimizes for L1 cache, L2 cache, and register file usage.
|
||||
|
||||
The zero-overhead nature of TileDistribution, achieved through extensive use of C++ template metaprogramming and compile-time computation, ensures that the abstraction's benefits come without runtime cost. Every aspect of the distribution strategy is resolved at compile time, resulting in machine code that is often more efficient than hand-written implementations. The compiler's ability to see through the abstraction enables optimizations that would be difficult or impossible with runtime-based approaches.
|
||||
The zero-overhead nature of TileDistribution, achieved through use of C++ template metaprogramming and compile-time computation, ensures that the abstraction's benefits come without runtime cost. Every aspect of the distribution strategy is resolved at compile time, resulting in machine code that is often more efficient than hand-written implementations. The compiler's ability to see through the abstraction enables optimizations that would be difficult or impossible with runtime-based approaches.
|
||||
|
||||
Portability across different GPU architectures emerges naturally from TileDistribution's design. The same source code can execute efficiently on GPUs with different warp sizes, different numbers of registers per thread, or different shared memory capacities. This portability extends beyond mere functional correctness to include performance portability—the framework automatically adapts its strategies to match the characteristics of the target architecture.
|
||||
|
||||
The productivity gains from using TileDistribution cannot be overstated. What traditionally requires hundreds of lines of intricate, error-prone code can be expressed in a handful of clear, maintainable lines. This dramatic reduction in code complexity translates directly to faster development cycles, fewer bugs, and easier maintenance. The clear separation between algorithmic logic and distribution strategy also enhances code clarity, making it easier for teams to collaborate and for new developers to understand existing code.
|
||||
|
||||
TileDistribution stands as the foundational technology that enables Composable Kernels to deliver on its promise of combining high performance with high productivity. By solving the fundamental challenge of work distribution in a general, efficient, and elegant manner, it provides the solid foundation upon which the entire CK ecosystem is built. As GPU architectures continue to evolve and grow more complex, the value of this abstraction only increases, providing a stable, high-performance programming model that insulates developers from the growing complexity of the underlying hardware while enabling them to fully exploit its capabilities.
|
||||
TileDistribution stands as the foundational technology that enables Composable Kernels to deliver on its promise of combining high performance with high productivity. By solving the fundamental challenge of work distribution in a general, efficient, and efficient manner, it provides the solid foundation upon which the entire CK ecosystem is built. As GPU architectures continue to evolve and grow more complex, the value of this abstraction only increases, providing a stable, high-performance programming model that insulates developers from the growing complexity of the underlying hardware while enabling them to fully exploit its capabilities.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
@@ -8,59 +8,67 @@ Overview
|
||||
|
||||
The TileWindow abstraction represents the culmination of the CK framework's approach to efficient tensor data access on GPUs. While :ref:`TileDistribution <ck_tile_tile_distribution>` determines the mapping between threads and tensor coordinates, TileWindow provides the actual mechanism for loading and storing data with optimal memory access patterns. This abstraction encapsulates the complexity of coalesced memory accesses, vectorization, and boundary handling into a clean interface that enables developers to focus on algorithmic logic rather than low-level memory management.
|
||||
|
||||
At its core, TileWindow implements a sophisticated windowing mechanism that views a subset of a larger tensor through the lens of a tile distribution. This windowing is not merely a simple sub-tensor extraction but a distribution-aware view that automatically generates the most efficient memory access patterns for the underlying hardware. The system achieves this by combining knowledge of the :ref:`tensor's layout <ck_tile_descriptors>`, the distribution pattern, and the :ref:`GPU's memory subsystem <ck_tile_gpu_basics>` characteristics to generate optimized load and store operations.
|
||||
At its core, TileWindow implements a distribution-aware windowing mechanism that views a subset of a larger tensor through the lens of a tile distribution. This windowing is not merely a simple sub-tensor extraction but a distribution-aware view that automatically generates the most efficient memory access patterns for the underlying hardware. The system achieves this by combining knowledge of the :ref:`tensor's layout <ck_tile_descriptors>`, the distribution pattern, and the :ref:`GPU's memory subsystem <ck_tile_gpu_basics>` characteristics to generate optimized load and store operations.
|
||||
|
||||
TileWindow Architecture
|
||||
-----------------------
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Components"
|
||||
TV["TensorView<br/>Data source"]
|
||||
TD["TileDistribution<br/>Thread mapping"]
|
||||
TW["TileWindow<br/>Access gateway"]
|
||||
LT["LoadStoreTraits<br/>Access optimizer"]
|
||||
DT["DistributedTensor<br/>Register storage"]
|
||||
end
|
||||
|
||||
subgraph "Operations"
|
||||
Load["Load<br/>Global → Registers"]
|
||||
Compute["Compute<br/>In registers"]
|
||||
Store["Store<br/>Registers → Global"]
|
||||
end
|
||||
|
||||
subgraph "Optimizations"
|
||||
Coal["Coalescing<br/>Adjacent access"]
|
||||
Vec["Vectorization<br/>Multi-element ops"]
|
||||
Bank["Bank conflict<br/>avoidance"]
|
||||
SFC["Space-filling<br/>curve traversal"]
|
||||
end
|
||||
|
||||
TV --> TW
|
||||
TD --> TW
|
||||
TW --> LT
|
||||
LT --> DT
|
||||
|
||||
TW --> Load
|
||||
Load --> Compute
|
||||
Compute --> Store
|
||||
|
||||
Load --> Coal
|
||||
Load --> Vec
|
||||
Load --> SFC
|
||||
Store --> Bank
|
||||
|
||||
style TW fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style LT fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style DT fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Components"
|
||||
TV["TensorView<br/>Data source"]
|
||||
TD["TileDistribution<br/>Thread mapping"]
|
||||
TW["TileWindow<br/>Access gateway"]
|
||||
LT["LoadStoreTraits<br/>Access optimizer"]
|
||||
DT["DistributedTensor<br/>Register storage"]
|
||||
end
|
||||
|
||||
subgraph "Operations"
|
||||
Load["Load<br/>Global → Registers"]
|
||||
Compute["Compute<br/>In registers"]
|
||||
Store["Store<br/>Registers → Global"]
|
||||
end
|
||||
|
||||
subgraph "Optimizations"
|
||||
Coal["Coalescing<br/>Adjacent access"]
|
||||
Vec["Vectorization<br/>Multi-element ops"]
|
||||
Bank["Bank conflict<br/>avoidance"]
|
||||
SFC["Space-filling<br/>curve traversal"]
|
||||
end
|
||||
|
||||
TV --> TW
|
||||
TD --> TW
|
||||
TW --> LT
|
||||
LT --> DT
|
||||
|
||||
TW --> Load
|
||||
Load --> Compute
|
||||
Compute --> Store
|
||||
|
||||
Load --> Coal
|
||||
Load --> Vec
|
||||
Load --> SFC
|
||||
Store --> Bank
|
||||
|
||||
style TW fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style LT fill:#fff3e0,stroke:#f57c00,stroke-width:2px
|
||||
style DT fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/tile_window_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
What is a TileWindow?
|
||||
---------------------
|
||||
|
||||
The fundamental challenge in GPU programming lies in the gap between logical tensor operations and the physical realities of memory access. While :ref:`TileDistribution <ck_tile_tile_distribution>` elegantly solves the problem of work assignment by mapping threads to :ref:`tensor coordinates <ck_tile_coordinate_systems>`, it does not address how threads actually access the data at those coordinates. This is where TileWindow enters the picture, serving as the critical bridge between logical work assignment and physical memory operations.
|
||||
The fundamental challenge in GPU programming lies in the gap between logical tensor operations and the physical realities of memory access. While :ref:`TileDistribution <ck_tile_tile_distribution>` solves the problem of work assignment by mapping threads to :ref:`tensor coordinates <ck_tile_coordinate_systems>`, it does not address how threads actually access the data at those coordinates. This is where TileWindow enters the picture, serving as the critical bridge between logical work assignment and physical memory operations.
|
||||
|
||||
TileWindow implements a distribution-aware windowing mechanism that transforms abstract coordinate mappings into concrete memory access patterns. The abstraction understands not just which data elements each thread needs, but also how to access them in a way that maximizes memory bandwidth utilization. This involves sophisticated techniques such as memory coalescing, where adjacent threads access adjacent memory locations, and vectorization, where multiple elements are loaded or stored in a single transaction.
|
||||
TileWindow implements a distribution-aware windowing mechanism that transforms abstract coordinate mappings into concrete memory access patterns. The abstraction understands not just which data elements each thread needs, but also how to access them in a way that maximizes memory bandwidth utilization. This involves optimized techniques such as memory coalescing, where adjacent threads access adjacent memory locations, and vectorization, where multiple elements are loaded or stored in a single transaction.
|
||||
|
||||
**C++ Implementation Overview:**
|
||||
|
||||
@@ -115,7 +123,7 @@ TileWindow implements a distribution-aware windowing mechanism that transforms a
|
||||
LoadStoreTraits - The Access Pattern Engine
|
||||
-------------------------------------------
|
||||
|
||||
Behind every efficient TileWindow operation lies :ref:`LoadStoreTraits <ck_tile_load_store_traits>`, a sophisticated analysis engine that determines the optimal way to access memory. This component bridges the gap between the logical distribution pattern and the physical memory subsystem, analyzing the distribution to find opportunities for vectorization and coalescing.
|
||||
Behind every efficient TileWindow operation lies :ref:`LoadStoreTraits <ck_tile_load_store_traits>`, a compile-time analysis engine that determines the optimal way to access memory. This component bridges the gap between the logical distribution pattern and the physical memory subsystem, analyzing the distribution to find opportunities for vectorization and coalescing.
|
||||
|
||||
LoadStoreTraits performs several critical analyses:
|
||||
|
||||
@@ -159,38 +167,44 @@ LoadStoreTraits performs several critical analyses:
|
||||
Space-Filling Curves for Memory Access
|
||||
--------------------------------------
|
||||
|
||||
One of the most sophisticated aspects of TileWindow is its use of :ref:`space-filling curves <ck_tile_space_filling_curve>` to determine the order in which memory is accessed. This isn't just about iterating through elements - it's about doing so in a way that maximizes cache utilization and minimizes memory latency.
|
||||
One of the most key aspects of TileWindow is its use of :ref:`space-filling curves <ck_tile_space_filling_curve>` to determine the order in which memory is accessed. Rather than simple linear iteration, space-filling curves provide cache-friendly traversal patterns that maximize hardware utilization. The "snake" pattern is particularly effective because it minimizes the distance between consecutive accesses, keeping data in cache longer.
|
||||
|
||||
A space-filling curve is a continuous curve that visits every point in a multi-dimensional space exactly once. In the context of TileWindow, it determines the order in which a thread accesses its assigned elements. The "snake" pattern is particularly effective because it minimizes the distance between consecutive accesses, keeping data in cache longer.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Access Pattern"
|
||||
L1["0→1→2→3"]
|
||||
L2["4→5→6→7"]
|
||||
L3["8→9→10→11"]
|
||||
L4["12→13→14→15"]
|
||||
end
|
||||
|
||||
subgraph "Snake Access Pattern"
|
||||
S1["0→1→2→3"]
|
||||
S2["7←6←5←4"]
|
||||
S3["8→9→10→11"]
|
||||
S4["15←14←13←12"]
|
||||
end
|
||||
|
||||
L1 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
style S1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style S2 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Linear Access Pattern"
|
||||
L1["0→1→2→3"]
|
||||
L2["4→5→6→7"]
|
||||
L3["8→9→10→11"]
|
||||
L4["12→13→14→15"]
|
||||
end
|
||||
|
||||
subgraph "Snake Access Pattern"
|
||||
S1["0→1→2→3"]
|
||||
S2["7←6←5←4"]
|
||||
S3["8→9→10→11"]
|
||||
S4["15←14←13←12"]
|
||||
end
|
||||
|
||||
L1 --> L2
|
||||
L2 --> L3
|
||||
L3 --> L4
|
||||
|
||||
S1 --> S2
|
||||
S2 --> S3
|
||||
S3 --> S4
|
||||
|
||||
style S1 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style S2 fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/tile_window_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Space-Filling Curve Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -221,36 +235,44 @@ A space-filling curve is a continuous curve that visits every point in a multi-d
|
||||
TileWindow Data Flow
|
||||
--------------------
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart LR
|
||||
subgraph "Step 1: Create Window"
|
||||
T["Tensor<br/>[256, 256]"]
|
||||
O["Origin<br/>(64, 64)"]
|
||||
W["Window Size<br/>[32, 32]"]
|
||||
end
|
||||
|
||||
subgraph "Step 2: Apply Distribution"
|
||||
TD["TileDistribution<br/>Thread mapping"]
|
||||
TW["TileWindow<br/>Created"]
|
||||
end
|
||||
|
||||
subgraph "Step 3: Load Data"
|
||||
GM["Global Memory<br/>Window region"]
|
||||
REG["Registers<br/>Distributed tensor"]
|
||||
end
|
||||
|
||||
T --> TW
|
||||
O --> TW
|
||||
W --> TW
|
||||
TD --> TW
|
||||
|
||||
TW --> GM
|
||||
GM -->|"load()"| REG
|
||||
|
||||
style TW fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style REG fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
flowchart LR
|
||||
subgraph "Step 1: Create Window"
|
||||
T["Tensor<br/>[256, 256]"]
|
||||
O["Origin<br/>(64, 64)"]
|
||||
W["Window Size<br/>[32, 32]"]
|
||||
end
|
||||
|
||||
subgraph "Step 2: Apply Distribution"
|
||||
TD["TileDistribution<br/>Thread mapping"]
|
||||
TW["TileWindow<br/>Created"]
|
||||
end
|
||||
|
||||
subgraph "Step 3: Load Data"
|
||||
GM["Global Memory<br/>Window region"]
|
||||
REG["Registers<br/>Distributed tensor"]
|
||||
end
|
||||
|
||||
T --> TW
|
||||
O --> TW
|
||||
W --> TW
|
||||
TD --> TW
|
||||
|
||||
TW --> GM
|
||||
GM -->|"load()"| REG
|
||||
|
||||
style TW fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style REG fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/tile_window_3.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Creating and Using TileWindow
|
||||
-----------------------------
|
||||
|
||||
@@ -298,7 +320,7 @@ Let's explore how to create and use a TileWindow in practice:
|
||||
The Load Operation Deep Dive
|
||||
----------------------------
|
||||
|
||||
The load operation is where all the sophisticated analysis comes together. When you call ``window.load()``, a carefully orchestrated sequence of operations occurs:
|
||||
The load operation is where all the compile-time analysis comes together. When you call ``window.load()``, a carefully orchestrated sequence of operations occurs:
|
||||
|
||||
1. **Distributed tensor creation**: Automatically creates a :ref:`distributed tensor <ck_tile_static_distributed_tensor>` sized for the distribution
|
||||
2. **Coordinate calculation**: Uses precomputed coordinates for efficiency
|
||||
@@ -343,42 +365,50 @@ The load operation is where all the sophisticated analysis comes together. When
|
||||
Load Operation Architecture
|
||||
---------------------------
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Load Analysis"
|
||||
Analyze["Analyze access pattern<br/>Detect coalescing opportunities"]
|
||||
end
|
||||
|
||||
subgraph "Vectorization"
|
||||
V1["Scalar: 4 loads"]
|
||||
V2["Vector2: 2 loads"]
|
||||
V4["Vector4: 1 load"]
|
||||
end
|
||||
|
||||
subgraph "Memory Transaction"
|
||||
Coal["Coalesced access<br/>32 threads → 1 transaction"]
|
||||
NonCoal["Non-coalesced<br/>32 threads → 32 transactions"]
|
||||
end
|
||||
|
||||
subgraph "Result"
|
||||
Reg["Thread registers<br/>Local data"]
|
||||
end
|
||||
|
||||
Analyze --> V1
|
||||
Analyze --> V2
|
||||
Analyze --> V4
|
||||
|
||||
V4 --> Coal
|
||||
V1 --> NonCoal
|
||||
|
||||
Coal --> Reg
|
||||
NonCoal --> Reg
|
||||
|
||||
style V4 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
style Coal fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
style NonCoal fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Load Analysis"
|
||||
Analyze["Analyze access pattern<br/>Detect coalescing opportunities"]
|
||||
end
|
||||
|
||||
subgraph "Vectorization"
|
||||
V1["Scalar: 4 loads"]
|
||||
V2["Vector2: 2 loads"]
|
||||
V4["Vector4: 1 load"]
|
||||
end
|
||||
|
||||
subgraph "Memory Transaction"
|
||||
Coal["Coalesced access<br/>32 threads → 1 transaction"]
|
||||
NonCoal["Non-coalesced<br/>32 threads → 32 transactions"]
|
||||
end
|
||||
|
||||
subgraph "Result"
|
||||
Reg["Thread registers<br/>Local data"]
|
||||
end
|
||||
|
||||
Analyze --> V1
|
||||
Analyze --> V2
|
||||
Analyze --> V4
|
||||
|
||||
V4 --> Coal
|
||||
V1 --> NonCoal
|
||||
|
||||
Coal --> Reg
|
||||
NonCoal --> Reg
|
||||
|
||||
style V4 fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
style Coal fill:#d1fae5,stroke:#10b981,stroke-width:2px
|
||||
style NonCoal fill:#fee2e2,stroke:#ef4444,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/tile_window_4.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Memory Access Patterns
|
||||
----------------------
|
||||
|
||||
@@ -443,7 +473,7 @@ TileWindow supports efficient window movement for sliding window algorithms. The
|
||||
Store Operations with Vectorization
|
||||
-----------------------------------
|
||||
|
||||
Store operations use the same sophisticated analysis as loads. The :ref:`LoadStoreTraits <ck_tile_load_store_traits>` ensures that stores are just as efficient as loads, with the same vectorization and coalescing benefits:
|
||||
Store operations use the same compile-time analysis as loads. The :ref:`LoadStoreTraits <ck_tile_load_store_traits>` ensures that stores are just as efficient as loads, with the same vectorization and coalescing benefits:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
@@ -557,31 +587,39 @@ Here's a complete example showing how all components work together in a :ref:`ma
|
||||
Performance Characteristics
|
||||
---------------------------
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Memory Access Optimization"
|
||||
V["Vectorization<br/>4x fewer transactions"]
|
||||
C["Coalescing<br/>32x bandwidth efficiency"]
|
||||
P["Precomputation<br/>Zero overhead addressing"]
|
||||
S["Space-filling<br/>Optimal cache usage"]
|
||||
end
|
||||
|
||||
subgraph "Hardware Utilization"
|
||||
BW["Memory Bandwidth<br/>Near 100% utilization"]
|
||||
L["Latency Hiding<br/>Overlapped operations"]
|
||||
R["Register Reuse<br/>Minimal spills"]
|
||||
end
|
||||
|
||||
V --> BW
|
||||
C --> BW
|
||||
P --> L
|
||||
S --> R
|
||||
|
||||
style V fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style C fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style BW fill:#d1fae5,stroke:#10b981,stroke-width:3px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph LR
|
||||
subgraph "Memory Access Optimization"
|
||||
V["Vectorization<br/>4x fewer transactions"]
|
||||
C["Coalescing<br/>32x bandwidth efficiency"]
|
||||
P["Precomputation<br/>Zero overhead addressing"]
|
||||
S["Space-filling<br/>Optimal cache usage"]
|
||||
end
|
||||
|
||||
subgraph "Hardware Utilization"
|
||||
BW["Memory Bandwidth<br/>Near 100% utilization"]
|
||||
L["Latency Hiding<br/>Overlapped operations"]
|
||||
R["Register Reuse<br/>Minimal spills"]
|
||||
end
|
||||
|
||||
V --> BW
|
||||
C --> BW
|
||||
P --> L
|
||||
S --> R
|
||||
|
||||
style V fill:#e3f2fd,stroke:#1976d2,stroke-width:2px
|
||||
style C fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
style BW fill:#d1fae5,stroke:#10b981,stroke-width:3px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/tile_window_5.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Best Practices
|
||||
--------------
|
||||
|
||||
@@ -649,7 +687,7 @@ Key benefits:
|
||||
4. **Composability**: Integrates cleanly with other CK abstractions
|
||||
5. **Intelligence**: LoadStoreTraits analyzes and optimizes every access
|
||||
|
||||
The TileWindow abstraction is essential for building high-performance GPU kernels, providing a clean interface for complex memory access patterns while maintaining peak performance. The sophisticated analysis performed by LoadStoreTraits ensures that every memory operation is as efficient as possible, while the space-filling curve traversal maximizes cache utilization.
|
||||
The TileWindow abstraction is essential for building high-performance GPU kernels, providing a clean interface for complex memory access patterns while maintaining peak performance. The compile-time analysis performed by LoadStoreTraits ensures that every memory operation is as efficient as possible, while the space-filling curve traversal maximizes cache utilization.
|
||||
|
||||
Next Steps
|
||||
----------
|
||||
|
||||
@@ -29,26 +29,34 @@ Zero-Copy Logical Operations
|
||||
- **Data Storage**: The actual tensor data remains stored in memory in linear fashion, exactly as specified by the original tensor shape and strides at creation time (see :ref:`ck_tile_buffer_views` for raw memory access)
|
||||
- **Logical Mapping**: Transforms only change how we interpret and access coordinates - they create different logical views of the same underlying data (building on :ref:`ck_tile_tensor_views`)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Tensor Coordinate Transformation"
|
||||
US["Lower Dimension Space<br/>Source coordinate system"]
|
||||
LS["Upper Dimension Space<br/>Target coordinate system"]
|
||||
|
||||
DATA["Linear Data in Memory<br/>Layout determined by tensor<br/>shape & strides"]
|
||||
end
|
||||
|
||||
US -->|"Forward Transform"| LS
|
||||
LS -->|"Inverse Transform"| US
|
||||
|
||||
DATA -.->|"Same data,<br/>different views"| US
|
||||
DATA -.->|"Same data,<br/>different views"| LS
|
||||
|
||||
style US fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style LS fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "Tensor Coordinate Transformation"
|
||||
US["Lower Dimension Space<br/>Source coordinate system"]
|
||||
LS["Upper Dimension Space<br/>Target coordinate system"]
|
||||
|
||||
DATA["Linear Data in Memory<br/>Layout determined by tensor<br/>shape & strides"]
|
||||
end
|
||||
|
||||
US -->|"Forward Transform"| LS
|
||||
LS -->|"Inverse Transform"| US
|
||||
|
||||
DATA -.->|"Same data,<br/>different views"| US
|
||||
DATA -.->|"Same data,<br/>different views"| LS
|
||||
|
||||
style US fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style LS fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_1.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Index Calculation Operations
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
@@ -63,61 +71,77 @@ These operations enable bidirectional navigation between different coordinate re
|
||||
Transform System Architecture
|
||||
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
|
||||
subgraph "Transform Types"
|
||||
EMB["EmbedTransform<br/>Linear → Multi-D Strided"]
|
||||
UNM["MergeTransform<br/>Multi-D → Linear"]
|
||||
MRG["UnmergeTransform<br/>Linear → Multi-D"]
|
||||
REP["ReplicateTransform<br/>0D → Multi-D Broadcast"]
|
||||
OFF["OffsetTransform<br/>Translation"]
|
||||
PAS["PassThroughTransform<br/>Identity"]
|
||||
PAD["PadTransform<br/>Boundaries"]
|
||||
end
|
||||
|
||||
subgraph "Operations"
|
||||
FWD["Forward<br/>calculate_lower_index()"]
|
||||
BWD["Backward<br/>calculate_upper_index()"]
|
||||
UPD["Update<br/>update_lower_index()"]
|
||||
end
|
||||
|
||||
EMB --> FWD
|
||||
UNM --> FWD
|
||||
MRG --> FWD
|
||||
REP --> FWD
|
||||
OFF --> FWD
|
||||
PAS --> FWD
|
||||
PAD --> FWD
|
||||
|
||||
style FWD fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
|
||||
subgraph "Transform Types"
|
||||
EMB["EmbedTransform<br/>Linear → Multi-D Strided"]
|
||||
UNM["MergeTransform<br/>Multi-D → Linear"]
|
||||
MRG["UnmergeTransform<br/>Linear → Multi-D"]
|
||||
REP["ReplicateTransform<br/>0D → Multi-D Broadcast"]
|
||||
OFF["OffsetTransform<br/>Translation"]
|
||||
PAS["PassThroughTransform<br/>Identity"]
|
||||
PAD["PadTransform<br/>Boundaries"]
|
||||
end
|
||||
|
||||
subgraph "Operations"
|
||||
FWD["Forward<br/>calculate_lower_index()"]
|
||||
BWD["Backward<br/>calculate_upper_index()"]
|
||||
UPD["Update<br/>update_lower_index()"]
|
||||
end
|
||||
|
||||
EMB --> FWD
|
||||
UNM --> FWD
|
||||
MRG --> FWD
|
||||
REP --> FWD
|
||||
OFF --> FWD
|
||||
PAS --> FWD
|
||||
PAD --> FWD
|
||||
|
||||
style FWD fill:#e8f5e9,stroke:#388e3c,stroke-width:2px
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_2.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
MergeTransform
|
||||
--------------
|
||||
|
||||
MergeTransform collapses multiple dimensions from the lower coordinate space into a single dimension in the upper coordinate space, effectively reducing the dimensionality of the tensor representation while preserving all data relationships. This transform is fundamental to the :ref:`tile distribution system <ck_tile_tile_distribution>`.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "MergeTransform: Multi-D → Linear"
|
||||
LS["Lower Coordinate Space<br/>2D: [4, 5]<br/>Coord: (2, 3)"]
|
||||
US["Upper Coordinate Space<br/>1D Linear<br/>Index: 13"]
|
||||
|
||||
DATA["Same Tensor Data<br/>Layout: row-major<br/>Size: 20 elements"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>2×5 + 3 = 13"| US
|
||||
US -->|"Inverse Transform<br/>13÷5=2, 13%5=3"| LS
|
||||
|
||||
DATA -.->|"Multi-dimensional<br/>view"| LS
|
||||
DATA -.->|"Linear<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "MergeTransform: Multi-D → Linear"
|
||||
LS["Lower Coordinate Space<br/>2D: [4, 5]<br/>Coord: (2, 3)"]
|
||||
US["Upper Coordinate Space<br/>1D Linear<br/>Index: 13"]
|
||||
|
||||
DATA["Same Tensor Data<br/>Layout: row-major<br/>Size: 20 elements"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>2×5 + 3 = 13"| US
|
||||
US -->|"Inverse Transform<br/>13÷5=2, 13%5=3"| LS
|
||||
|
||||
DATA -.->|"Multi-dimensional<br/>view"| LS
|
||||
DATA -.->|"Linear<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_3.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -153,26 +177,34 @@ UnmergeTransform
|
||||
|
||||
UnmergeTransform expands coordinates from a single dimension in the lower coordinate space into multiple dimensions in the upper coordinate space, effectively increasing the dimensionality of the tensor representation while preserving all data relationships.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "UnmergeTransform: Linear → Multi-D"
|
||||
LS["Lower Coordinate Space<br/>1D Linear<br/>Index: 14"]
|
||||
US["Upper Coordinate Space<br/>3D: [3, 4, 2]<br/>Coord: (1, 3, 0)"]
|
||||
|
||||
DATA["Same Tensor Data<br/>Layout: row-major<br/>Size: 24 elements"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>14 = 1×8 + 3×2 + 0"| US
|
||||
US -->|"Inverse Transform<br/>linearize back"| LS
|
||||
|
||||
DATA -.->|"Linear<br/>view"| LS
|
||||
DATA -.->|"Multi-dimensional<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "UnmergeTransform: Linear → Multi-D"
|
||||
LS["Lower Coordinate Space<br/>1D Linear<br/>Index: 14"]
|
||||
US["Upper Coordinate Space<br/>3D: [3, 4, 2]<br/>Coord: (1, 3, 0)"]
|
||||
|
||||
DATA["Same Tensor Data<br/>Layout: row-major<br/>Size: 24 elements"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>14 = 1×8 + 3×2 + 0"| US
|
||||
US -->|"Inverse Transform<br/>linearize back"| LS
|
||||
|
||||
DATA -.->|"Linear<br/>view"| LS
|
||||
DATA -.->|"Multi-dimensional<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_4.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -218,26 +250,34 @@ EmbedTransform
|
||||
|
||||
EmbedTransform expands linear indices from the lower coordinate space into multi-dimensional coordinates in the upper coordinate space using configurable strides, enabling flexible strided tensor layouts and sub-tensor views within larger buffers.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "EmbedTransform: Linear → Multi-D Strided"
|
||||
LS["Lower Coordinate Space<br/>1D Linear<br/>Index: 14"]
|
||||
US["Upper Coordinate Space<br/>2D: [2, 3]<br/>Coord: (1, 2)"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform <br/>Strides: [12, 1] <br/>14 ÷ 12 = 1, 14 % 12 = 2"| US
|
||||
US -->|"Inverse Transform<br/>1×12 + 2×1 = 14"| LS
|
||||
|
||||
DATA -.->|"Linear<br/>index view"| LS
|
||||
DATA -.->|"Multi-dimensional<br/>strided view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "EmbedTransform: Linear → Multi-D Strided"
|
||||
LS["Lower Coordinate Space<br/>1D Linear<br/>Index: 14"]
|
||||
US["Upper Coordinate Space<br/>2D: [2, 3]<br/>Coord: (1, 2)"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform <br/>Strides: [12, 1] <br/>14 ÷ 12 = 1, 14 % 12 = 2"| US
|
||||
US -->|"Inverse Transform<br/>1×12 + 2×1 = 14"| LS
|
||||
|
||||
DATA -.->|"Linear<br/>index view"| LS
|
||||
DATA -.->|"Multi-dimensional<br/>strided view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_5.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -272,26 +312,34 @@ ReplicateTransform
|
||||
|
||||
ReplicateTransform creates a higher-dimensional tensor by replicating (broadcasting) a lower-dimensional tensor. It's essentially a broadcasting operation that takes a tensor with fewer dimensions and logically replicates it across new dimensions without data duplication. An example is taking a scalar (0-dimensional) input and broadcasting it across multiple dimensions, enabling efficient broadcasting patterns where a single value appears at every position in a multi-dimensional coordinate space.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "ReplicateTransform: 0D → Multi-D Broadcasting"
|
||||
LS["Lower Coordinate Space<br/>0D: Scalar<br/>Empty coordinate []"]
|
||||
US["Upper Coordinate Space<br/>2D: [3, 4]<br/>All coords: (i, j)"]
|
||||
|
||||
DATA["Single Scalar Value"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>[] → (i,j) for any i,j"| US
|
||||
US -->|"Inverse Transform<br/>(i,j) → [] for any i,j"| LS
|
||||
|
||||
DATA -.->|"One scalar<br/>value"| LS
|
||||
DATA -.->|"Broadcasted view<br/>at all positions"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "ReplicateTransform: 0D → Multi-D Broadcasting"
|
||||
LS["Lower Coordinate Space<br/>0D: Scalar<br/>Empty coordinate []"]
|
||||
US["Upper Coordinate Space<br/>2D: [3, 4]<br/>All coords: (i, j)"]
|
||||
|
||||
DATA["Single Scalar Value"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>[] → (i,j) for any i,j"| US
|
||||
US -->|"Inverse Transform<br/>(i,j) → [] for any i,j"| LS
|
||||
|
||||
DATA -.->|"One scalar<br/>value"| LS
|
||||
DATA -.->|"Broadcasted view<br/>at all positions"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_6.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -337,26 +385,34 @@ OffsetTransform
|
||||
|
||||
OffsetTransform shifts coordinates by a fixed offset, creating a translated view of the coordinate space. It performs simple translation operations where each coordinate in the upper space is mapped to a coordinate in the lower space by adding a constant offset.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "OffsetTransform: 1D → 1D Translation"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 63]<br/>Coord: index + offset"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 47]<br/>Coord: index"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx → idx + 16"| US
|
||||
US -->|"Inverse Transform<br/>idx + 16 → idx"| LS
|
||||
|
||||
DATA -.->|"Lower<br/>view"| LS
|
||||
DATA -.->|"Upper<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "OffsetTransform: 1D → 1D Translation"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 63]<br/>Coord: index + offset"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 47]<br/>Coord: index"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx → idx + 16"| US
|
||||
US -->|"Inverse Transform<br/>idx + 16 → idx"| LS
|
||||
|
||||
DATA -.->|"Lower<br/>view"| LS
|
||||
DATA -.->|"Upper<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_7.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -403,26 +459,34 @@ PassThroughTransform - Identity
|
||||
|
||||
No-op transform that passes coordinates unchanged. The PassThrough transform is the simplest coordinate transformation in CK Tile, implementing a perfect identity mapping where input coordinates are passed through unchanged to the output. This transform is essential as a placeholder in transformation chains and for dimensions that require no modification.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "PassThroughTransform: 1D → 1D Identity"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 59]<br/>Coord: index"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 59]<br/>Coord: index"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -.->|"Perfect Identity<br/>idx → idx"| US
|
||||
US -.->|"Perfect Identity<br/>idx → idx"| LS
|
||||
|
||||
DATA -->|"Same buffer<br/>same view"| LS
|
||||
DATA -->|"Same buffer<br/>same view"| US
|
||||
|
||||
style LS fill:#e8f5e8,stroke:#2e7d32,stroke-width:3px
|
||||
style US fill:#e8f5e8,stroke:#2e7d32,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "PassThroughTransform: 1D → 1D Identity"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 59]<br/>Coord: index"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 59]<br/>Coord: index"]
|
||||
|
||||
DATA["Linear Buffer in Memory"]
|
||||
end
|
||||
|
||||
LS -.->|"Perfect Identity<br/>idx → idx"| US
|
||||
US -.->|"Perfect Identity<br/>idx → idx"| LS
|
||||
|
||||
DATA -->|"Same buffer<br/>same view"| LS
|
||||
DATA -->|"Same buffer<br/>same view"| US
|
||||
|
||||
style LS fill:#e8f5e8,stroke:#2e7d32,stroke-width:3px
|
||||
style US fill:#e8f5e8,stroke:#2e7d32,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_8.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -464,26 +528,34 @@ PadTransform
|
||||
|
||||
PadTransform adds padding to tensor dimensions, mapping coordinates from upper dimension space (with padding) to lower dimension space (original data).
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "PadTransform: 1D → 1D with Padding"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 2] (original data)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 4] (with padding)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx + left_pad"| US
|
||||
US -->|"Inverse Transform<br/>idx - left_pad"| LS
|
||||
|
||||
DATA -.->|"Original view"| LS
|
||||
DATA -.->|"Padded view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "PadTransform: 1D → 1D with Padding"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 2] (original data)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 4] (with padding)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx + left_pad"| US
|
||||
US -->|"Inverse Transform<br/>idx - left_pad"| LS
|
||||
|
||||
DATA -.->|"Original view"| LS
|
||||
DATA -.->|"Padded view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_9.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
**C++ Implementation:**
|
||||
|
||||
.. code-block:: cpp
|
||||
@@ -528,76 +600,100 @@ XorTransform
|
||||
|
||||
XorTransform applies a 2D XOR mapping for specialized memory access patterns. It performs XOR operations on coordinates to create transformed memory layouts for specific algorithmic optimizations, particularly useful for avoiding :ref:`LDS bank conflicts <ck_tile_lds_bank_conflicts>`.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "XorTransform: 2D → 2D XOR Mapping"
|
||||
LS["Lower Coordinate Space<br/>2D: [4, 8]<br/>XOR-transformed coords"]
|
||||
US["Upper Coordinate Space<br/>2D: [4, 8]<br/>Normal coords"]
|
||||
|
||||
DATA["Same Tensor Data"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>apply XOR reverse"| US
|
||||
US -->|"Inverse Transform<br/>apply XOR mapping"| LS
|
||||
|
||||
DATA -.->|"XOR pattern<br/>view"| LS
|
||||
DATA -.->|"Normal<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "XorTransform: 2D → 2D XOR Mapping"
|
||||
LS["Lower Coordinate Space<br/>2D: [4, 8]<br/>XOR-transformed coords"]
|
||||
US["Upper Coordinate Space<br/>2D: [4, 8]<br/>Normal coords"]
|
||||
|
||||
DATA["Same Tensor Data"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>apply XOR reverse"| US
|
||||
US -->|"Inverse Transform<br/>apply XOR mapping"| LS
|
||||
|
||||
DATA -.->|"XOR pattern<br/>view"| LS
|
||||
DATA -.->|"Normal<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_10.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
SliceTransform
|
||||
~~~~~~~~~~~~~~
|
||||
|
||||
SliceTransform extracts a sub-region from a tensor dimension.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "SliceTransform: 1D → 1D Sub-region"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 9] (original range)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 4] (slice range)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx + slice_begin"| US
|
||||
US -->|"Inverse Transform<br/>idx - slice_begin"| LS
|
||||
|
||||
DATA -.->|"Full tensor<br/>view"| LS
|
||||
DATA -.->|"Sub-region<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "SliceTransform: 1D → 1D Sub-region"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 9] (original range)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 4] (slice range)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx + slice_begin"| US
|
||||
US -->|"Inverse Transform<br/>idx - slice_begin"| LS
|
||||
|
||||
DATA -.->|"Full tensor<br/>view"| LS
|
||||
DATA -.->|"Sub-region<br/>view"| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_11.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
ModuloTransform
|
||||
~~~~~~~~~~~~~~~
|
||||
|
||||
ModuloTransform applies cyclic wrapping to coordinates using modulo operations.
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "ModuloTransform: 1D → 1D Cyclic"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 3] (modulus range)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 15] (full range)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx * cycle_count"| US
|
||||
US -->|"Inverse Transform<br/>idx % modulus"| LS
|
||||
|
||||
DATA -.->|" "| LS
|
||||
DATA -.->|" "| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
..
|
||||
Original mermaid diagram (edit here, then run update_diagrams.py)
|
||||
|
||||
.. mermaid::
|
||||
|
||||
graph TB
|
||||
subgraph "ModuloTransform: 1D → 1D Cyclic"
|
||||
LS["Lower Coordinate Space<br/>1D: [0, 3] (modulus range)"]
|
||||
US["Upper Coordinate Space<br/>1D: [0, 15] (full range)"]
|
||||
|
||||
DATA["Tensor Data in Memory"]
|
||||
end
|
||||
|
||||
LS -->|"Forward Transform<br/>idx * cycle_count"| US
|
||||
US -->|"Inverse Transform<br/>idx % modulus"| LS
|
||||
|
||||
DATA -.->|" "| LS
|
||||
DATA -.->|" "| US
|
||||
|
||||
style LS fill:#e3f2fd,stroke:#1976d2,stroke-width:3px
|
||||
style US fill:#fff3e0,stroke:#f57c00,stroke-width:3px
|
||||
style DATA fill:#f0f9ff,stroke:#0284c7,stroke-width:2px,stroke-dasharray: 5 5
|
||||
|
||||
|
||||
|
||||
.. image:: diagrams/transforms_12.svg
|
||||
:alt: Diagram
|
||||
:align: center
|
||||
Summary
|
||||
-------
|
||||
|
||||
|
||||
199
docs/conceptual/ck_tile/update_diagrams.py
Normal file
@@ -0,0 +1,199 @@
|
||||
#!/usr/bin/env python3
|
||||
"""
|
||||
Helper script to update SVG diagrams from commented mermaid sources in RST files.
|
||||
|
||||
This script scans RST files for commented mermaid blocks (created by convert_mermaid_to_svg.py)
|
||||
and regenerates the corresponding SVG files when the source has been modified.
|
||||
|
||||
Usage:
|
||||
python update_diagrams.py # Update all diagrams
|
||||
python update_diagrams.py <file.rst> # Update diagrams in a specific file
|
||||
"""
|
||||
|
||||
import os
|
||||
import re
|
||||
import subprocess
|
||||
import sys
|
||||
import tempfile
|
||||
from pathlib import Path
|
||||
|
||||
# Configuration
|
||||
DOCS_DIR = Path(__file__).parent
|
||||
DIAGRAMS_DIR = DOCS_DIR / 'diagrams'
|
||||
|
||||
# Pattern to find commented mermaid blocks followed by image references
|
||||
COMMENTED_MERMAID_PATTERN = re.compile(
|
||||
r'\.\.\s*\n' # Comment start
|
||||
r'(?: .*\n|\s*\n)*?' # Comment description lines (may have blank lines)
|
||||
r'( \.\. mermaid::\s*\n' # Commented mermaid directive
|
||||
r'(?: \n| .*\n|\s*\n)*?)' # Mermaid content (including blank lines)
|
||||
r'\.\. image:: diagrams/([^\s]+)', # Image reference
|
||||
re.MULTILINE
|
||||
)
|
||||
|
||||
|
||||
def extract_mermaid_from_comment(commented_block):
|
||||
"""Extract mermaid code from a commented block."""
|
||||
# Remove the comment indentation (3 spaces at start of each line)
|
||||
lines = commented_block.split('\n')
|
||||
content_lines = []
|
||||
|
||||
for line in lines:
|
||||
if line.startswith(' '):
|
||||
# Remove the 3-space comment indentation
|
||||
content_lines.append(line[3:])
|
||||
elif line.strip() == '':
|
||||
content_lines.append('')
|
||||
|
||||
# Now we have the mermaid block, extract the actual mermaid code
|
||||
mermaid_content = '\n'.join(content_lines)
|
||||
|
||||
# Remove the ".. mermaid::" directive and extract the indented content
|
||||
mermaid_match = re.search(r'\.\. mermaid::\s*\n((?:(?:\n| .*))*)', mermaid_content)
|
||||
if mermaid_match:
|
||||
mermaid_code = mermaid_match.group(1)
|
||||
# Remove RST indentation from mermaid code
|
||||
code_lines = []
|
||||
for line in mermaid_code.split('\n'):
|
||||
if line.startswith(' '):
|
||||
code_lines.append(line[3:])
|
||||
elif line.strip() == '':
|
||||
code_lines.append('')
|
||||
return '\n'.join(code_lines).strip()
|
||||
|
||||
return None
|
||||
|
||||
|
||||
def convert_mermaid_to_svg(mermaid_code, output_path):
|
||||
"""Convert mermaid code to SVG using mmdc."""
|
||||
# Create a temporary file for the mermaid code
|
||||
with tempfile.NamedTemporaryFile(mode='w', suffix='.mmd', delete=False, encoding='utf-8') as tmp:
|
||||
tmp.write(mermaid_code)
|
||||
tmp_path = tmp.name
|
||||
|
||||
try:
|
||||
# Run mmdc to convert to SVG
|
||||
result = subprocess.run(
|
||||
['mmdc', '-i', tmp_path, '-o', str(output_path), '-t', 'neutral', '-b', 'transparent'],
|
||||
capture_output=True,
|
||||
text=True,
|
||||
check=True,
|
||||
shell=True # Required for Windows .cmd files
|
||||
)
|
||||
return True, None
|
||||
except subprocess.CalledProcessError as e:
|
||||
return False, e.stderr
|
||||
finally:
|
||||
# Clean up temp file
|
||||
os.unlink(tmp_path)
|
||||
|
||||
|
||||
def process_file(file_path, force_update=False):
|
||||
"""Process a single RST file to update diagrams."""
|
||||
print(f"Checking {file_path.name}...")
|
||||
|
||||
with open(file_path, 'r', encoding='utf-8') as f:
|
||||
content = f.read()
|
||||
|
||||
# Find all commented mermaid blocks
|
||||
matches = list(COMMENTED_MERMAID_PATTERN.finditer(content))
|
||||
|
||||
if not matches:
|
||||
print(f" No commented mermaid diagrams found.")
|
||||
return 0, 0
|
||||
|
||||
updated_count = 0
|
||||
error_count = 0
|
||||
|
||||
for match in matches:
|
||||
commented_mermaid = match.group(1)
|
||||
svg_filename = match.group(2)
|
||||
svg_path = DIAGRAMS_DIR / svg_filename
|
||||
|
||||
# Extract mermaid code
|
||||
mermaid_code = extract_mermaid_from_comment(commented_mermaid)
|
||||
if not mermaid_code:
|
||||
print(f" ⚠ Could not extract mermaid code for {svg_filename}")
|
||||
error_count += 1
|
||||
continue
|
||||
|
||||
# Check if SVG needs updating
|
||||
needs_update = force_update or not svg_path.exists()
|
||||
|
||||
if not needs_update:
|
||||
# For a more sophisticated check, we could hash the mermaid code
|
||||
# and compare with a stored hash, but for simplicity we just check existence
|
||||
print(f" ✓ {svg_filename} exists (use --force to regenerate)")
|
||||
continue
|
||||
|
||||
# Generate SVG
|
||||
success, error = convert_mermaid_to_svg(mermaid_code, svg_path)
|
||||
|
||||
if success:
|
||||
print(f" ✓ Updated: {svg_filename}")
|
||||
updated_count += 1
|
||||
else:
|
||||
print(f" ✗ Error updating {svg_filename}: {error}")
|
||||
error_count += 1
|
||||
|
||||
return updated_count, error_count
|
||||
|
||||
|
||||
def find_rst_files():
|
||||
"""Find all RST files in the CK tile docs directory."""
|
||||
return list(DOCS_DIR.glob('*.rst'))
|
||||
|
||||
|
||||
def main():
|
||||
"""Main function."""
|
||||
print("CK Tile Diagram Updater")
|
||||
print("=" * 50)
|
||||
|
||||
# Verify mmdc is available
|
||||
try:
|
||||
subprocess.run(['mmdc', '--version'], capture_output=True, check=True, shell=True)
|
||||
except (subprocess.CalledProcessError, FileNotFoundError):
|
||||
print("Error: mermaid-cli (mmdc) not found. Please install it:")
|
||||
print(" npm install -g @mermaid-js/mermaid-cli")
|
||||
return 1
|
||||
|
||||
# Ensure diagrams directory exists
|
||||
DIAGRAMS_DIR.mkdir(parents=True, exist_ok=True)
|
||||
|
||||
# Parse command line arguments
|
||||
force_update = '--force' in sys.argv or '-f' in sys.argv
|
||||
specific_file = None
|
||||
|
||||
for arg in sys.argv[1:]:
|
||||
if arg not in ['--force', '-f'] and arg.endswith('.rst'):
|
||||
specific_file = DOCS_DIR / arg
|
||||
if not specific_file.exists():
|
||||
print(f"Error: File not found: {arg}")
|
||||
return 1
|
||||
|
||||
# Get files to process
|
||||
if specific_file:
|
||||
files_to_process = [specific_file]
|
||||
else:
|
||||
files_to_process = find_rst_files()
|
||||
|
||||
# Process files
|
||||
total_updated = 0
|
||||
total_errors = 0
|
||||
|
||||
for file_path in files_to_process:
|
||||
updated, errors = process_file(file_path, force_update)
|
||||
total_updated += updated
|
||||
total_errors += errors
|
||||
|
||||
print("\n" + "=" * 50)
|
||||
print(f"✓ Update complete!")
|
||||
print(f" Updated: {total_updated} diagram(s)")
|
||||
if total_errors > 0:
|
||||
print(f" Errors: {total_errors}")
|
||||
|
||||
return 0 if total_errors == 0 else 1
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
exit(main())
|
||||