> ## 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.

# cutlass::Array

> Statically sized array for all CUTLASS-supported numeric types, safe for use in unions

## Overview

`cutlass::Array` is a statically sized array template that accommodates all CUTLASS-supported numeric types and is safe to use in a union. It provides STL-compatible interfaces including iterators, element access, and various operations.

**Header:** `cutlass/array.h`

## Template Signature

```cpp theme={null}
template <
  typename T,
  int N,
  bool RegisterSized = sizeof_bits<T>::value >= 32
>
struct Array;
```

## Template Parameters

<ParamField path="T" type="typename">
  Element type. Can be any CUTLASS numeric type including `float`, `half_t`, `int8_t`, etc.
</ParamField>

<ParamField path="N" type="int">
  Number of elements in the array
</ParamField>

<ParamField path="RegisterSized" type="bool" default="sizeof_bits<T>::value >= 32">
  Internal optimization flag. Set to true for types with 32 or more bits.
</ParamField>

## Member Types

```cpp theme={null}
using Storage = T;
using Element = T;

static constexpr size_t kStorageElements = N;
static constexpr size_t kElements = N;

// C++ standard container types
typedef T value_type;
typedef size_t size_type;
typedef ptrdiff_t difference_type;
typedef value_type &reference;
typedef value_type const & const_reference;
typedef value_type *pointer;
typedef value_type const * const_pointer;
```

## Data Members

```cpp theme={null}
Storage storage[kElements];
```

Internal storage for array elements.

## Member Functions

### Element Access

#### at

```cpp theme={null}
CUTLASS_HOST_DEVICE
reference at(size_type pos);

CUTLASS_HOST_DEVICE
const_reference at(size_type pos) const;
```

Accesses element at specified position with bounds checking.

#### operator\[]

```cpp theme={null}
CUTLASS_HOST_DEVICE
reference operator[](size_type pos);

CUTLASS_HOST_DEVICE
const_reference operator[](size_type pos) const;
```

Accesses element at specified position without bounds checking.

#### front

```cpp theme={null}
CUTLASS_HOST_DEVICE
reference front();

CUTLASS_HOST_DEVICE
const_reference front() const;
```

Returns reference to the first element.

#### back

```cpp theme={null}
CUTLASS_HOST_DEVICE
reference back();

CUTLASS_HOST_DEVICE
const_reference back() const;
```

Returns reference to the last element.

#### data

```cpp theme={null}
CUTLASS_HOST_DEVICE
pointer data();

CUTLASS_HOST_DEVICE
const_pointer data() const;
```

Returns pointer to the underlying array.

#### raw\_data

```cpp theme={null}
CUTLASS_HOST_DEVICE
pointer raw_data();

CUTLASS_HOST_DEVICE
const_pointer raw_data() const;
```

Returns pointer to the underlying storage.

### Capacity

#### empty

```cpp theme={null}
CUTLASS_HOST_DEVICE
constexpr bool empty() const;
```

Returns true if the array is empty (N == 0).

#### size

```cpp theme={null}
CUTLASS_HOST_DEVICE
constexpr size_type size() const;
```

Returns the number of elements (N).

#### max\_size

```cpp theme={null}
CUTLASS_HOST_DEVICE
constexpr size_type max_size() const;
```

Returns the maximum number of elements (N).

### Operations

#### clear

```cpp theme={null}
CUTLASS_HOST_DEVICE
void clear();
```

Sets all elements to zero using `fill(T(0))`.

#### fill

```cpp theme={null}
CUTLASS_HOST_DEVICE
void fill(T const &value);
```

Fills all elements with the specified value.

**Implementation:**

```cpp theme={null}
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < int(kElements); ++i) {
  storage[i] = static_cast<Storage>(value);
}
```

### Iterators

#### begin / end

```cpp theme={null}
CUTLASS_HOST_DEVICE
iterator begin();

CUTLASS_HOST_DEVICE
const_iterator begin() const;

CUTLASS_HOST_DEVICE
const_iterator cbegin() const;

CUTLASS_HOST_DEVICE
iterator end();

CUTLASS_HOST_DEVICE
const_iterator end() const;

CUTLASS_HOST_DEVICE
const_iterator cend() const;
```

Returns iterators to the beginning and end of the array.

#### rbegin / rend

```cpp theme={null}
CUTLASS_HOST_DEVICE
reverse_iterator rbegin();

CUTLASS_HOST_DEVICE
const_reverse_iterator rbegin() const;

CUTLASS_HOST_DEVICE
const_reverse_iterator crbegin() const;

CUTLASS_HOST_DEVICE
reverse_iterator rend();

CUTLASS_HOST_DEVICE
const_reverse_iterator rend() const;

CUTLASS_HOST_DEVICE
const_reverse_iterator crend() const;
```

Returns reverse iterators.

## Factory Functions

### make\_Array

```cpp theme={null}
template <typename Element>
CUTLASS_HOST_DEVICE
Array<Element, 1> make_Array(Element x);

template <typename Element>
CUTLASS_HOST_DEVICE
Array<Element, 2> make_Array(Element x, Element y);

template <typename Element>
CUTLASS_HOST_DEVICE
Array<Element, 3> make_Array(Element x, Element y, Element z);

template <typename Element>
CUTLASS_HOST_DEVICE
Array<Element, 4> make_Array(Element x, Element y, Element z, Element w);
```

Factory functions to construct arrays with 1-4 elements.

## Numeric Specializations

Array provides specialized implementations of numeric operations for SIMD efficiency on supported types like `half_t`.

### Arithmetic Operations

```cpp theme={null}
template <typename T, int N>
struct plus<Array<T, N>>;

template <typename T, int N>
struct minus<Array<T, N>>;

template <typename T, int N>
struct multiplies<Array<T, N>>;

template <typename T, int N>
struct divides<Array<T, N>>;

template <typename T, int N>
struct negate<Array<T, N>>;
```

### Comparison Operations

```cpp theme={null}
template <typename T, int N, bool PropagateNaN>
struct maximum<Array<T, N>, PropagateNaN>;

template <typename T, int N, bool PropagateNaN>
struct minimum<Array<T, N>, PropagateNaN>;
```

### Fused Operations

```cpp theme={null}
template <typename T, int N>
struct multiply_add<Array<T, N>, Array<T, N>, Array<T, N>>;

template <typename T, int N>
struct multiply_add_relu0<Array<T, N>, Array<T, N>, Array<T, N>>;
```

### Mathematical Functions

```cpp theme={null}
template <typename T, int N>
struct absolute_value_op<Array<T, N>>;

template <typename T, int N>
struct inverse_square_root<Array<T, N>>;

template <typename T, int N>
struct reciprocal_approximate<Array<T, N>>;
```

## Usage Examples

### Basic Array Operations

```cpp theme={null}
// Create and initialize an array
cutlass::Array<float, 4> arr;
arr.fill(1.0f);

// Access elements
arr[0] = 2.0f;
float val = arr.at(1);

// Iterate over elements
for (auto it = arr.begin(); it != arr.end(); ++it) {
  *it *= 2.0f;
}

// Range-based for loop
for (float& elem : arr) {
  elem += 1.0f;
}
```

### Factory Functions

```cpp theme={null}
// Create arrays with specific values
auto vec3 = cutlass::make_Array(1.0f, 2.0f, 3.0f);
auto vec4 = cutlass::make_Array(1.0f, 2.0f, 3.0f, 4.0f);
```

### Numeric Operations

```cpp theme={null}
cutlass::Array<float, 4> a, b, c;
a.fill(2.0f);
b.fill(3.0f);

// Element-wise addition
cutlass::plus<cutlass::Array<float, 4>> add_op;
c = add_op(a, b);  // c = {5, 5, 5, 5}

// Multiply-add: c = a * b + c
cutlass::multiply_add<cutlass::Array<float, 4>> ma_op;
c = ma_op(a, b, c);
```

### With Half-Precision Types

```cpp theme={null}
// Array of half-precision values
cutlass::Array<cutlass::half_t, 8> half_arr;
half_arr.fill(cutlass::half_t(1.5f));

// SIMD-optimized operations on half_t arrays
cutlass::multiplies<cutlass::Array<cutlass::half_t, 8>> mul_op;
auto result = mul_op(half_arr, cutlass::half_t(2.0f));
```

### In Device Code

```cpp theme={null}
__global__ void kernel() {
  // Arrays work in both host and device code
  cutlass::Array<float, 4> fragment;
  
  // Initialize
  fragment.clear();
  
  // Compute
  CUTLASS_PRAGMA_UNROLL
  for (int i = 0; i < fragment.size(); ++i) {
    fragment[i] = threadIdx.x * i;
  }
}
```

## SIMD Optimization

For `half_t` arrays, operations are automatically SIMD-optimized on SM\_53 and later:

```cpp theme={null}
// From include/cutlass/array.h:1180-1212
template <int N>
struct plus<Array<half_t, N>> {
  CUTLASS_HOST_DEVICE
  Array<half_t, N> operator()(Array<half_t, N> const & lhs, 
                              Array<half_t, N> const &rhs) const {
    Array<half_t, N> result;
    #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 530)

    __half2 *result_ptr = reinterpret_cast<__half2 *>(&result);
    __half2 const *lhs_ptr = reinterpret_cast<__half2 const *>(&lhs);
    __half2 const *rhs_ptr = reinterpret_cast<__half2 const *>(&rhs);

    CUTLASS_PRAGMA_UNROLL
    for (int i = 0; i < N / 2; ++i) {
      result_ptr[i] = __hadd2(lhs_ptr[i], rhs_ptr[i]);
    }

    if constexpr (N % 2) {
      __half const *a_residual_ptr = reinterpret_cast<__half const *>(&lhs);
      __half const *b_residual_ptr = reinterpret_cast<__half const *>(&rhs);
      __half d_residual = __hadd(a_residual_ptr[N - 1], b_residual_ptr[N - 1]);
      result[N - 1] = reinterpret_cast<half_t const &>(d_residual);
    }

    #else
    // Fallback for host code or older architectures
    CUTLASS_PRAGMA_UNROLL
    for (int i = 0; i < N; ++i) {
      result[i] = lhs[i] + rhs[i];
    }
    #endif

    return result;
  }
};
```

## See Also

* [TensorRef](/api/cpp/tensor-ref) - Uses Array for fragment operations
* [Gemm](/api/cpp/gemm) - Uses Array internally for accumulation
* [Epilogue Operators](/api/cpp/epilogue-operators) - Operate on Array fragments