Skip to content

Commit 35ae515

Browse files
authored
[MLIR][NVVM][Docs] Explain memory spaces (#168059)
1 parent 83fc85c commit 35ae515

File tree

1 file changed

+34
-0
lines changed

1 file changed

+34
-0
lines changed

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -79,6 +79,40 @@ def NVVM_Dialect : Dialect {
7979
sequence must be expressed directly, NVVM provides an `nvvm.inline_ptx` op to
8080
embed PTX inline as a last-resort escape hatch, with explicit operands and
8181
results.
82+
83+
84+
**Memory Spaces:** The NVVM dialect introduces the following memory spaces,
85+
each with distinct scopes and lifetimes:
86+
```
87+
| Memory Space | Address Space | Scope | Lifetime |
88+
|-------------------|---------------|----------------------|-------------------|
89+
| `generic` | 0 | All threads | Context-dependent |
90+
| `global` | 1 | All threads (device) | Application |
91+
| `shared` | 3 | Thread block (CTA) | Kernel execution |
92+
| `constant` | 4 | All threads (RO) | Application |
93+
| `local` | 5 | Single thread | Kernel execution |
94+
| `tensor` | 6 | Thread block (CTA) | Kernel execution |
95+
| `shared_cluster` | 7 | Thread block cluster | Kernel execution |
96+
```
97+
**Memory Space Details:**
98+
- **generic**: Can point to any memory space; requires runtime resolution of
99+
actual address space. Use when pointer origin is unknown at compile time.
100+
Performance varies based on the underlying memory space.
101+
- **global**: Accessible by all threads across all blocks; persists across
102+
kernel launches. Highest latency but largest capacity (device memory). Best
103+
for large data and inter-kernel communication.
104+
- **shared**: Shared within a thread block (CTA); very fast on-chip memory for
105+
cooperation between threads in the same block. Limited capacity. Ideal for
106+
block-level collaboration, caching, and reducing global memory traffic.
107+
- **constant**: Read-only memory cached per SM. Size typically limited to
108+
64KB. Best for read-only data and uniform values accessed by all threads.
109+
- **local**: Private to each thread. Use for per-thread private data and
110+
automatic variables that don't fit in registers.
111+
- **tensor**: Special memory space for tensor core operations. Used by
112+
`tcgen05` instructions on SM 100+ for tensor input/output operations.
113+
- **shared_cluster**: Distributed shared memory across thread blocks within
114+
a cluster (SM 90+). Enables collaboration beyond single-block scope with
115+
fast access across cluster threads.
82116
}];
83117

84118
let name = "nvvm";

0 commit comments

Comments
 (0)