> ## Documentation Index
> Fetch the complete documentation index at: https://mintlify.com/NVIDIA/cutlass/llms.txt
> Use this file to discover all available pages before exploring further.

# Memory Operations

> CuTe DSL memory management and tensor creation APIs

The CuTe DSL provides comprehensive memory management APIs for creating and manipulating tensors across different memory spaces (global, shared, register).

## Tensor Creation

### make\_tensor

Creates a tensor from a pointer and layout.

```python theme={null}
tensor = cute.make_tensor(ptr, layout)
```

<ParamField path="ptr" type="Pointer" required>
  Pointer to memory (global, shared, or register)
</ParamField>

<ParamField path="layout" type="Layout" required>
  Layout defining shape and stride
</ParamField>

**Example:**

```python theme={null}
# Create tensor from global memory pointer
ptr = cute.make_ptr(data, element_type=cute.Float32)
layout = cute.make_layout(
    shape=(M, N),
    stride=(N, 1)  # Row-major
)
tensor = cute.make_tensor(ptr, layout)
```

### make\_fragment

Allocates a register memory tensor (fragment).

```python theme={null}
fragment = cute.make_fragment(element_type, layout)
```

<ParamField path="element_type" type="Type[Numeric]" required>
  Element data type (e.g., `cute.Float32`, `cute.BFloat16`)
</ParamField>

<ParamField path="layout" type="Layout" required>
  Layout of the fragment
</ParamField>

**Example:**

```python theme={null}
# Allocate 16x8 fragment in registers
layout = cute.make_layout((16, 8))
fragment = cute.make_fragment(cute.Float32, layout)
```

### make\_fragment\_like

Allocates a fragment with the same shape as another tensor.

```python theme={null}
fragment = cute.make_fragment_like(tensor)
```

<ParamField path="tensor" type="Tensor" required>
  Template tensor
</ParamField>

**Example:**

```python theme={null}
# Create fragment matching partitioned tensor
thrA = thr_copy.partition_S(blockA)
fragA = cute.make_fragment_like(thrA)
```

### make\_rmem\_tensor

Allocates register memory tensor.

```python theme={null}
tensor = cute.make_rmem_tensor(element_type, layout)
```

Alias for `make_fragment`.

### make\_rmem\_tensor\_like

Allocates register memory tensor matching another tensor.

```python theme={null}
tensor = cute.make_rmem_tensor_like(template_tensor)
```

Alias for `make_fragment_like`.

## Tensor Initialization

### zeros\_like

Creates a tensor initialized to zero.

```python theme={null}
zero_tensor = cute.zeros_like(tensor)
```

**Example:**

```python theme={null}
# Initialize accumulator to zero
accum = cute.zeros_like(fragC)
```

### ones\_like

Creates a tensor initialized to one.

```python theme={null}
one_tensor = cute.ones_like(tensor)
```

### full

Creates a tensor filled with a specific value.

```python theme={null}
filled_tensor = cute.full(layout, value, element_type)
```

<ParamField path="layout" type="Layout" required>
  Tensor layout
</ParamField>

<ParamField path="value" type="Numeric" required>
  Fill value
</ParamField>

<ParamField path="element_type" type="Type[Numeric]" required>
  Element data type
</ParamField>

### full\_like

Creates a tensor filled with a value, matching another tensor's shape.

```python theme={null}
filled_tensor = cute.full_like(tensor, value)
```

**Example:**

```python theme={null}
# Initialize accumulator to identity (for debugging)
identity = cute.full_like(fragC, 1.0)
```

### empty\_like

Allocates an uninitialized tensor matching another tensor's shape.

```python theme={null}
empty_tensor = cute.empty_like(tensor)
```

## Pointer Operations

### make\_ptr

Creates a typed pointer from raw memory address.

```python theme={null}
ptr = cute.make_ptr(
    address,
    element_type=cute.Float32,
    address_space=cute.AddressSpace.GLOBAL
)
```

<ParamField path="address" type="Union[int, Tensor]" required>
  Memory address or tensor
</ParamField>

<ParamField path="element_type" type="Type[Numeric]" required>
  Element data type
</ParamField>

<ParamField path="address_space" type="AddressSpace" default="GLOBAL">
  Memory address space
</ParamField>

**Example:**

```python theme={null}
# Create pointer from DLPack tensor
from cutlass.cute.runtime import from_dlpack

tensor_desc = from_dlpack(torch_tensor)
ptr = cute.make_ptr(
    tensor_desc,
    element_type=cute.Float16,
    address_space=cute.AddressSpace.GLOBAL
)
```

### recast\_ptr

Recasts a pointer to a different element type.

```python theme={null}
new_ptr = cute.recast_ptr(ptr, new_element_type)
```

<ParamField path="ptr" type="Pointer" required>
  Original pointer
</ParamField>

<ParamField path="new_element_type" type="Type[Numeric]" required>
  New element data type
</ParamField>

**Example:**

```python theme={null}
# Recast FP16 pointer to INT8 for quantization
f16_ptr = cute.make_ptr(data, element_type=cute.Float16)
i8_ptr = cute.recast_ptr(f16_ptr, cute.Int8)
```

## Layout Operations

### make\_layout

Creates a layout from shape and stride.

```python theme={null}
layout = cute.make_layout(
    shape=(M, N),
    stride=(N, 1)  # Row-major
)
```

<ParamField path="shape" type="Shape" required>
  Shape tuple (can be nested)
</ParamField>

<ParamField path="stride" type="Stride" default="compact">
  Stride tuple (defaults to compact row-major)
</ParamField>

**Examples:**

```python theme={null}
# Row-major 2D
layout = cute.make_layout((M, N), stride=(N, 1))

# Column-major 2D
layout = cute.make_layout((M, N), stride=(1, M))

# Hierarchical layout
layout = cute.make_layout(
    shape=((4, 8), (2, 16)),
    stride=((128, 1), (64, 8))
)
```

### make\_identity\_layout

Creates an identity layout (coordinate layout).

```python theme={null}
layout = cute.make_identity_layout(shape)
```

**Example:**

```python theme={null}
# Create coordinate layout for predication
coord_layout = cute.make_identity_layout((M, N))
```

### make\_ordered\_layout

Creates a layout with specified dimension ordering.

```python theme={null}
layout = cute.make_ordered_layout(
    shape=(M, N),
    order=(1, 0)  # Column-major
)
```

### make\_layout\_like

Creates a layout matching another layout's shape.

```python theme={null}
new_layout = cute.make_layout_like(template_layout)
```

### recast\_layout

Recasts a layout to a different element type.

```python theme={null}
new_layout = cute.recast_layout(layout, old_type, new_type)
```

**Example:**

```python theme={null}
# Recast layout from FP32 to FP16 (doubles logical size)
f16_layout = cute.recast_layout(f32_layout, cute.Float32, cute.Float16)
```

## Identity and Coordinate Tensors

### make\_identity\_tensor

Creates a coordinate tensor for bounds checking.

```python theme={null}
coord_tensor = cute.make_identity_tensor(shape)
```

<ParamField path="shape" type="Shape" required>
  Tensor shape
</ParamField>

**Example:**

```python theme={null}
# Create coordinate tensor for predication
shape = (M, N)
coord = cute.make_identity_tensor(shape)

# Partition to threads
thr_coord = thr_copy.partition_S(coord)

# Create predicate for out-of-bounds
pred = thr_coord < (M, N)

# Predicated copy
cute.basic_copy_if(pred, src, dst)
```

## Tensor Recasting

### recast\_tensor

Recasts a tensor to a different element type.

```python theme={null}
new_tensor = cute.recast_tensor(tensor, new_element_type)
```

<ParamField path="tensor" type="Tensor" required>
  Original tensor
</ParamField>

<ParamField path="new_element_type" type="Type[Numeric]" required>
  New element data type
</ParamField>

**Example:**

```python theme={null}
# Recast for type conversion
f16_tensor = cute.recast_tensor(f32_tensor, cute.Float16)
```

## Memory Space Management

### Address Spaces

```python theme={null}
# Global memory
ptr_global = cute.make_ptr(
    address,
    element_type=cute.Float32,
    address_space=cute.AddressSpace.GLOBAL
)

# Shared memory
ptr_shared = cute.make_ptr(
    address,
    element_type=cute.Float32,
    address_space=cute.AddressSpace.SHARED
)

# Register memory (fragments)
fragment = cute.make_fragment(cute.Float32, layout)
```

## Tensor Partitioning

### Partition by Tiled Copy

```python theme={null}
# Partition source tensor
thrA_src = thr_copy.partition_S(block_tensor)

# Partition destination tensor
thrA_dst = thr_copy.partition_D(fragment)
```

### Partition by Tiled MMA

```python theme={null}
# Partition A operand
thrA = thr_mma.partition_A(block_A)

# Partition B operand
thrB = thr_mma.partition_B(block_B)

# Partition C (accumulator)
thrC = thr_mma.partition_C(block_C)
```

## Complete Example: Memory Management

```python theme={null}
import cutlass.cute as cute
from cutlass.cute.runtime import from_dlpack

@cute.kernel
def memory_example_kernel(
    gA: cute.Tensor,  # Global memory input
    gB: cute.Tensor,  # Global memory output
    sA_layout: cute.Layout,  # Shared memory layout
):
    tidx, _, _ = cute.arch.thread_idx()
    bidx, _, _ = cute.arch.block_idx()
    
    # Allocate shared memory
    sA = cute.make_tensor(
        cute.make_ptr(
            cute.arch.smem_ptr(0),
            element_type=gA.element_type,
            address_space=cute.AddressSpace.SHARED
        ),
        sA_layout
    )
    
    # Partition global memory block
    blkA = gA[((None, None), bidx)]
    blkB = gB[((None, None), bidx)]
    
    # Create tiled copy
    copy_atom = cute.make_copy_atom(
        cute.nvgpu.CopyUniversalOp(),
        gA.element_type
    )
    thr_layout = cute.make_layout((4, 32), stride=(32, 1))
    val_layout = cute.make_layout((1, 4), stride=(4, 1))
    tiled_copy = cute.make_tiled_copy_tv(
        copy_atom,
        thr_layout,
        val_layout
    )
    thr_copy = tiled_copy.get_slice(tidx)
    
    # Partition tensors
    thrA_gmem = thr_copy.partition_S(blkA)
    thrA_smem = thr_copy.partition_D(sA)
    
    # Allocate register fragment
    fragA = cute.make_fragment_like(thrA_gmem)
    
    # Load: Global -> Register
    cute.copy(copy_atom, thrA_gmem, fragA)
    
    # Store: Register -> Shared
    cute.copy(copy_atom, fragA, thrA_smem)
    
    # Synchronize threads
    cute.arch.syncthreads()
    
    # Partition output
    thrB = thr_copy.partition_D(blkB)
    
    # Process data (example: simple copy)
    fragB = cute.make_fragment_like(thrB)
    cute.basic_copy(thrA_smem, fragB)
    
    # Store: Register -> Global
    cute.copy(copy_atom, fragB, thrB)


@cute.jit
def memory_example(
    torch_A,
    torch_B,
    stream,
):
    # Convert from DLPack
    A_desc = from_dlpack(torch_A)
    B_desc = from_dlpack(torch_B)
    
    # Create tensors
    M, N = torch_A.shape
    gA = cute.make_tensor(
        cute.make_ptr(A_desc, element_type=cute.Float32),
        cute.make_layout((M, N), stride=(N, 1))
    )
    gB = cute.make_tensor(
        cute.make_ptr(B_desc, element_type=cute.Float32),
        cute.make_layout((M, N), stride=(N, 1))
    )
    
    # Shared memory layout
    tile_M, tile_N = 128, 128
    sA_layout = cute.make_layout((tile_M, tile_N))
    
    # Launch kernel
    num_blocks = (M + tile_M - 1) // tile_M
    memory_example_kernel(gA, gB, sA_layout).launch(
        grid=(num_blocks, 1, 1),
        block=(128, 1, 1),
        stream=stream
    )
```

## See Also

* [Kernel Definition](/api/dsl/kernel-definition) - Defining GPU kernels
* [Operations Reference](/api/dsl/operations) - DSL operations
* [Atom Operations](/api/dsl/atoms) - Copy and MMA atoms