### Launch Jupyter notebook examples

Source: https://github.com/nvidia/cutlass/blob/main/python/README.md

Start Jupyter Lab to access CUTLASS Python interface examples located in the examples/python directory.

```bash
jupyter-lab ../examples/python
```

--------------------------------

### Install Recommended Development Dependencies

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/pythonDSL/quick_start.rst

Installs additional Python packages recommended for running examples and developing with CUTLASS DSL, including PyTorch, Jupyter, and Mypy.

```bash
pip install torch jupyter mypy==1.19.1
```

--------------------------------

### Install CUTLASS DSL from Source

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/pythonDSL/quick_start.rst

Clones the CUTLASS repository and installs the DSL using the provided setup script, specifying the CUDA Toolkit version.

```bash
git clone https://github.com/NVIDIA/cutlass.git 

# For CUDA Toolkit 12.9:
./cutlass/python/CuTeDSL/setup.sh --cu12

# For CUDA Toolkit 13.1:
./cutlass/python/CuTeDSL/setup.sh --cu13
```

--------------------------------

### Build CUTLASS example with synclog enabled

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/utilities.md

After enabling `synclog` during compilation, navigate to the example directory and use `make` to build the application. This compiles the example with `synclog` instrumentation.

```bash
$ cd examples/54_hopper_fp8_warp_specialized_gemm
$ make
```

--------------------------------

### Install cute-viz for Layout Visualization

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/elementwise_add.ipynb

Installs the `cute-viz` library, which is used to visualize CuTeDSL tensor layouts.

```bash
pip install -U git+https://github.com/NTT123/cute-viz.git
```

--------------------------------

### Install Pandoc for Documentation Building

Source: https://github.com/nvidia/cutlass/blob/main/python/docs/index.html

Install the Pandoc utility, a prerequisite for building the CUTLASS Python interface documentation.

```bash
sudo apt-get install pandoc
```

--------------------------------

### Display Help for Turing TensorOp Conv2dFprop Example

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/implicit_gemm_convolution.md

Run this command to view the available command-line options and usage instructions for configuring the convolution example.

```bash
./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop --help
```

--------------------------------

### Install Test Executable and Result Cache

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This section handles the installation of the test executable and its associated result cache file, if installation of tests is enabled and not explicitly disabled for the executable.

```cmake
  if (NOT __DISABLE_EXECUTABLE_INSTALL_RULE AND CUTLASS_INSTALL_TESTS)

    # file(RELATIVE_PATH CMAKE_CURRENT_BINARY_RELATIVE_DIR ${CMAKE_BINARY_DIR} ${CMAKE_CURRENT_BINARY_DIR})

    install(
      TARGETS ${TARGET}
      RUNTIME DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR}
      )

    if (__RESULT_CACHE_FILE)

     install(
       FILES ${__RESULT_CACHE_FILE}
       DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR}
       )

    endif()

  endif()
```

--------------------------------

### Conditionally Include CMake 'examples' Subdirectory

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

Includes the 'examples' subdirectory if CUTLASS_ENABLE_EXAMPLES is true, and adds example tests as a dependency.

```cmake
if (CUTLASS_ENABLE_EXAMPLES)
  add_subdirectory(examples)
  add_dependencies(test_all test_examples)
endif()
```

--------------------------------

### Installing CUTLASS Headers and Target

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

These commands define the installation rules for CUTLASS, ensuring that its include directories and the CUTLASS target itself are placed in the specified installation paths.

```CMake
install(
  DIRECTORY
  ${CUTLASS_INCLUDE_DIR}/
  ${CMAKE_CURRENT_BINARY_DIR}/include/
  DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
  )

install(
  TARGETS CUTLASS
  EXPORT NvidiaCutlass
  PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
  )
```

--------------------------------

### Profiling Output Example

Source: https://github.com/nvidia/cutlass/blob/main/python/docs/externals/02_pytorch_extension_grouped_gemm.ipynb

Example output showing the measured execution times for grouped and non-grouped GEMM operations, and the calculated speedup.

```text
Grouped:     400.696 us
Non-Grouped: 646.670 us
Speedup: 1.614
```

--------------------------------

### Install CMake Package Configuration Files

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

Installs the generated NvidiaCutlassConfig.cmake and NvidiaCutlassConfigVersion.cmake files to the specified CMake installation directory.

```cmake
install(
  FILES
    ${CMAKE_CURRENT_BINARY_DIR}/NvidiaCutlassConfig.cmake
    ${CMAKE_CURRENT_BINARY_DIR}/NvidiaCutlassConfigVersion.cmake
  DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/NvidiaCutlass
  )
```

--------------------------------

### Build Specific Blackwell GEMM Green Context Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/95_blackwell_gemm_green_context/README.md

Build only the `95_blackwell_gemm_green_context` example after configuring CUTLASS.

```shell
make 95_blackwell_gemm_green_context -j$(nproc)
```

--------------------------------

### Register a CUTLASS Example Executable

Source: https://github.com/nvidia/cutlass/blob/main/examples/48_hopper_warp_specialized_gemm/CMakeLists.txt

This snippet shows how to use the `cutlass_example_add_executable` function to define a new example executable, linking it to its source file.

```CMake
cutlass_example_add_executable(
  48_hopper_warp_specialized_gemm
  48_hopper_warp_specialized_gemm.cu
  )
```

--------------------------------

### Build CUTLASS Turing TensorOp Conv2dFprop Example

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/implicit_gemm_convolution.md

This command compiles the `09_turing_tensorop_conv2dfprop` example executable and then lists its directory to confirm its creation.

```bash
make 09_turing_tensorop_conv2dfprop

ls examples/09_turing_tensorop_conv2dfprop
```

--------------------------------

### Install CUTLASS Python Interface from Source

Source: https://github.com/nvidia/cutlass/blob/main/python/docs_src/source/index.md

Installs the CUTLASS Python interface directly from its source directory.

```bash
pip install .
```

--------------------------------

### Define installation for cutlass_profiler

Source: https://github.com/nvidia/cutlass/blob/main/tools/profiler/CMakeLists.txt

Specifies the installation target, export set, and runtime destination for the profiler executable.

```cmake
install(
  TARGETS cutlass_profiler
  EXPORT NvidiaCutlass
  RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
  )
```

--------------------------------

### Install TVM FFI

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/pythonDSL/cute_dsl_general/compile_with_tvm_ffi.rst

Install the apache-tvm-ffi package and optionally torch-c-dlpack-ext for improved torch tensor calling performance.

```bash
pip install apache-tvm-ffi
# optional package for improved torch tensor calling performance
pip install torch-c-dlpack-ext
```

--------------------------------

### Running CuTe DSL Programs with JIT and Pre-compilation

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/hello_world.ipynb

Initializes the CUDA context and demonstrates two methods to run the `hello_world` function: immediate Just-In-Time (JIT) compilation and separate pre-compilation for repeated execution, including options to dump PTX/CUBIN files.

```python
# Initialize CUDA context for launching a kernel with error checking
# We make context initialization explicit to allow users to control the context creation
# and avoid potential issues with multiple contexts
cutlass.cuda.initialize_cuda_context()

# Method 1: Just-In-Time (JIT) compilation - compiles and runs the code immediately
print("Running hello_world()...")
hello_world()

# Method 2: Compile first (useful if you want to run the same code multiple times)
print("Compiling...")
hello_world_compiled = cute.compile(hello_world)

# Dump PTX/CUBIN files while compiling
from cutlass.cute import KeepPTX, KeepCUBIN

print("Compiling with PTX/CUBIN dumped...")
hello_world_compiled_ptx_on = cute.compile[KeepPTX, KeepCUBIN](hello_world)

# Run the pre-compiled version
print("Running compiled version...")
hello_world_compiled()
```

--------------------------------

### Build GEMM Softmax Example with CUTLASS

Source: https://github.com/nvidia/cutlass/blob/main/examples/35_gemm_softmax/CMakeLists.txt

This command adds an executable for the GEMM Softmax example to the build system, specifying its name and source file.

```CMake
cutlass_example_add_executable(
  35_gemm_softmax
  gemm_softmax.cu
  )
```

--------------------------------

### Device GEMM Instantiation Example

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/gemm_api.md

Example demonstrating how to instantiate and launch a mixed-precision GEMM targeting Volta Tensor Cores using the cutlass::gemm::device::Gemm operator.

```APIDOC
## Mixed-Precision GEMM Launch Example

### Description
Launch a mixed-precision GEMM targeting Volta Tensor Cores with half-precision inputs and float accumulation.

### Template Parameters
- **ElementA** (cutlass::half_t) - Data type of matrix A
- **LayoutA** (cutlass::layout::ColumnMajor) - Memory layout of matrix A
- **ElementB** (cutlass::half_t) - Data type of matrix B
- **LayoutB** (cutlass::layout::ColumnMajor) - Memory layout of matrix B
- **ElementOutput** (cutlass::half_t) - Data type of output matrix
- **LayoutOutput** (cutlass::layout::ColumnMajor) - Memory layout of output matrix
- **ElementAccumulator** (float) - Data type for accumulation
- **OpClass** (cutlass::arch::OpClassTensorOp) - Indicates Tensor Cores
- **ArchTag** (cutlass::arch::Sm70) - Target GPU compute architecture (Volta)

### Execution Parameters
- **Problem Size**: {m, n, k} - Matrix dimensions
- **Matrix A**: {ptrA, lda} - Pointer and leading dimension
- **Matrix B**: {ptrB, ldb} - Pointer and leading dimension
- **Matrix C**: {ptrC, ldc} - Pointer and leading dimension
- **Matrix D**: {ptrD, ldd} - Pointer and leading dimension (output)
- **Scalars**: {alpha, beta} - Scaling factors

### Return Value
- **cutlass::Status** - Operation status (kSuccess on successful completion)

### Code Example
```cpp
using Gemm = cutlass::gemm::device::Gemm<
  cutlass::half_t,                           // ElementA
  cutlass::layout::ColumnMajor,              // LayoutA
  cutlass::half_t,                           // ElementB
  cutlass::layout::ColumnMajor,              // LayoutB
  cutlass::half_t,                           // ElementOutput
  cutlass::layout::ColumnMajor,              // LayoutOutput
  float,                                     // ElementAccumulator
  cutlass::arch::OpClassTensorOp,            // tag indicating Tensor Cores
  cutlass::arch::Sm70                        // tag indicating target GPU compute architecture
>;

Gemm gemm_op;
cutlass::Status status;

// Launch GEMM on the device
status = gemm_op({
  {m, n, k},
  {ptrA, lda},
  {ptrB, ldb},
  {ptrC, ldc},
  {ptrD, ldd},
  {alpha, beta}
});

if (status != cutlass::Status::kSuccess) {
  return -1;
}
```
```

--------------------------------

### Install NVSHMEM4Py and NVSHMEM Library

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/blackwell/kernel/distributed/README.md

Install the necessary Python bindings and native NVSHMEM library for a specific CUDA version. `nvshmem4py` version >= 0.1.3 is recommended for these examples.

```bash
pip install nvshmem4py-cu12 nvidia-nvshmem-cu12
```

```bash
pip install nvshmem4py-cu13 nvidia-nvshmem-cu13
```

--------------------------------

### GET data()

Source: https://github.com/nvidia/cutlass/blob/main/docs/structcutlass_1_1epilogue_1_1threadblock_1_1EpilogueBase_1_1SharedStorage.html

Returns a pointer to the start of the shared memory buffer used for storage.

```APIDOC
## GET data()\n\n### Description\nReturns a pointer to the start of the shared memory buffer.\n\n### Method\nCUTLASS_DEVICE\n\n### Endpoint\ncutlass::epilogue::threadblock::EpilogueBase::SharedStorage::data()\n\n### Parameters\n(None)\n\n### Response\n#### Success Response\n- **return** (Element*) - Pointer to the shared memory buffer.
```

--------------------------------

### Shape Division Examples for Layouts

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/cute/02_layout_algebra.md

These examples illustrate how a layout's shape is 'divided out' by an integer, starting from the left. This operation is used to determine the shape of an intermediate layout when computing composition.

```text
(6,2) /  2 => (3,2)
(6,2) /  3 => (2,2)
(6,2) /  6 => (1,2)
(6,2) / 12 => (1,1)
(3,6,2,8) /  3 => (1,6,2,8)
(3,6,2,8) /  6 => (1,3,2,8)
(3,6,2,8) /  9 => (1,2,2,8)
(3,6,2,8) / 72 => (1,1,1,4)
```

--------------------------------

### Shape Modulo Examples for Layouts

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/cute/02_layout_algebra.md

These examples demonstrate how a layout's shape is 'modded out' by an integer, starting from the left. This operation ensures the resulting layout has a shape compatible with the inner layout B.

```text
(6,2) %  2 => (2,1)
(6,2) %  3 => (3,1)
(6,2) %  6 => (6,1)
(6,2) % 12 => (6,2)
(3,6,2,8) %  6 => (3,2,1,1)
(3,6,2,8) %  9 => (3,3,1,1)
(1,2,2,8) %  2 => (1,2,1,1)
(1,2,2,8) % 16 => (1,2,2,4)
```

--------------------------------

### GET cutlass::Array::cbegin() / cend()

Source: https://github.com/nvidia/cutlass/blob/main/docs/functions_c.html

Provides constant iterators to the start and end of a cutlass::Array container for read-only traversal.

```APIDOC
## GET cutlass::Array::cbegin()

### Description
Returns a constant iterator to the first element of the array container.

### Method
GET

### Endpoint
cutlass::Array<T, N, IsPacked>::cbegin()

### Parameters
#### Path Parameters
- **T** (type) - Required - Element type.
- **N** (int) - Required - Number of elements.

### Response
#### Success Response (200)
- **iterator** (const_iterator) - A constant iterator to the beginning of the array.

---

## GET cutlass::Array::cend()

### Description
Returns a constant iterator to the element following the last element of the array.

### Method
GET

### Endpoint
cutlass::Array<T, N, IsPacked>::cend()

### Response
#### Success Response (200)
- **iterator** (const_iterator) - A constant iterator to the end of the array.
```

--------------------------------

### Get Stride for ColumnMajorInterleaved Layout (C++)

Source: https://github.com/nvidia/cutlass/blob/main/docs/layout_2matrix_8h_source.html

Returns the stride of the ColumnMajorInterleaved layout. This value indicates the distance in elements between the start of consecutive rows or columns.

```C++
CUTLASS_HOST_DEVICE Stride stride() const
```

--------------------------------

### Creating a Host Function to Launch a CuTe DSL Kernel

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/hello_world.ipynb

Defines a host function using the `@cute.jit` decorator that prints "hello world" from the CPU and then launches the GPU kernel with specified grid and block dimensions.

```python
@cute.jit
def hello_world():
    # Print hello world from host code
    cute.printf("hello world")

    # Launch kernel
    kernel().launch(
        grid=(1, 1, 1),  # Single thread block
        block=(32, 1, 1),  # One warp (32 threads) per thread block
    )
```

--------------------------------

### Define `cutlass_example_add_executable` CMake Function

Source: https://github.com/nvidia/cutlass/blob/main/examples/CMakeLists.txt

This function encapsulates the logic for adding an executable, linking necessary libraries (CUTLASS, cuBLAS, CUDA), including directories, installing the target, and setting up tests for a CUTLASS example.

```cmake
function(cutlass_example_add_executable NAME)

  set(options)
  set(oneValueArgs DISABLE_TESTS)
  set(multiValueArgs DEPENDS DEPENDEES TEST_COMMAND_OPTIONS)
  cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})

  if (NOT DEFINED __DISABLE_TESTS)
    set(__DISABLE_TESTS OFF)
  endif()

  cutlass_add_executable(${NAME} ${__UNPARSED_ARGUMENTS} BATCH_SOURCES OFF)

  add_dependencies(cutlass_examples ${NAME})

  target_link_libraries(
    ${NAME}
    PRIVATE
    CUTLASS
    cutlass_tools_util_includes
    $<$<BOOL:${CUTLASS_ENABLE_CUBLAS}>:nvidia::cublas>
    cuda
    )

  target_include_directories(
    ${NAME}
    PRIVATE
    ${CUTLASS_EXAMPLES_COMMON_SOURCE_DIR}
    ${CUTLASS_EXAMPLES_UTILS_DIR}
    )

  install(
    TARGETS ${NAME}
    RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
    )

  cutlass_add_executable_tests(
    test_examples_${NAME} ${NAME}
    DEPENDS ${__DEPENDS}
    DEPENDEES test_examples ${__DEPENDEES}
    TEST_COMMAND_OPTIONS ${__TEST_COMMAND_OPTIONS}
    DISABLE_EXECUTABLE_INSTALL_RULE
    DISABLE_TESTS ${__DISABLE_TESTS}
    )

endfunction()
```

--------------------------------

### Parameter Space Sweep with Range Syntax

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/profiler.md

Sweep over a range of GEMM K dimension values using inclusive range syntax (start:end:increment). This example sweeps k from 8 to 4096 in increments of 8.

```bash
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sgemm_128x128_nn --m=4352 --n=4096 --k=8:4096:8
```

--------------------------------

### Header Setup and Namespaces

Source: https://github.com/nvidia/cutlass/blob/main/docs/inner__product_8h_source.html

Required includes and namespace hierarchy for the CUTLASS reference inner product implementation.

```cpp
#pragma once

#include "cutlass/cutlass.h"
#include "cutlass/array.h"

namespace cutlass {
namespace reference {
namespace detail {
```

--------------------------------

### Run Blackwell Distributed GEMM Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/82_blackwell_distributed_gemm/README.md

Provides a sample command to execute the Blackwell Distributed GEMM example with specified M, N, K dimensions and iteration counts.

```bash
./82_blackwell_distributed_gemm --m=16384 --n=106496 --k=16384 --warmup-iterations=10 --iterations=100
```

--------------------------------

### Complete Multi-GPU Sharded Vector Add with CUTLASS and JAX

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/dsl_tutorials/jax/cute_dsl_jax.ipynb

This comprehensive example shows the full setup for sharding a CUTLASS vector addition kernel across multiple GPUs using `jax.shard_map`, including mesh creation, sharding specification, and execution with assertion.

```python
import warnings
from functools import partial
from jax.sharding import PartitionSpec as P
with warnings.catch_warnings():
    warnings.simplefilter("ignore", DeprecationWarning)
    from jax.experimental.shard_map import shard_map

num_devices = len(jax.devices())
print(f"Number of devices: {num_devices}")

if num_devices > 1:
    mesh = jax.make_mesh((num_devices,), "x")
    # Kernel expects 3-D tensors: (elems_per_thread, threads, blocks)
    # Shard along the blocks axis (last dim)
    sharding = P(None, None, "x")

    @jax.jit
    def sharded_vector_add(a, b):
        @partial(
            shard_map,
            mesh=mesh,
            in_specs=(sharding, sharding),
            out_specs=sharding,
        )
        def _add(a_shard, b_shard):
            call = cjax.cutlass_call(
                launch_vector_add,
                output_shape_dtype=jax.ShapeDtypeStruct(
                    a_shard.shape, a_shard.dtype
                ),
                use_static_tensors=True,
            )
            return call(a_shard, b_shard)
        return _add(a, b)

    # Create 3-D tensors: (1, 256, total_blocks) with total_blocks divisible by device count
    blocks_per_device = 16
    total_blocks = blocks_per_device * num_devices
    shape = (1, BLOCK, total_blocks)
    a_m = jax.random.normal(jax.random.PRNGKey(10), shape, dtype=jnp.float32)
    b_m = jax.random.normal(jax.random.PRPRNGKey(11), shape, dtype=jnp.float32)

    c_m = sharded_vector_add(a_m, b_m)
    np.testing.assert_allclose(np.array(c_m), np.array(a_m + b_m), rtol=1e-5)
    N_total = int(np.prod(shape))
    print(f"Sharded Vector Add PASSED across {num_devices} devices (N={N_total})")
else:
    print("Only 1 device detected. Skipping multi-GPU example.")
    print("On a multi-GPU system, shard_map distributes CUTLASS kernels across devices.")
```

--------------------------------

### Defining and Installing Test Directories

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This snippet sets various CMake cache variables to control the installation of test executables and defines the directory structure for installing tests relative to the main installation prefix.

```CMake
set(CUTLASS_INSTALL_TESTS ON CACHE BOOL "Install test executables")
set(CUTLASS_TEST_EXECUTION_ENVIRONMENT "" CACHE BOOL "Environment in which to invoke unit test executables")

set(CMAKE_TEST_INSTALL_PREFIX test CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_PREFIX ${CMAKE_TEST_INSTALL_PREFIX}/cutlass CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_BINDIR ${CUTLASS_TEST_INSTALL_PREFIX}/${CMAKE_INSTALL_BINDIR} CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_LIBDIR ${CUTLASS_TEST_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR} CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")

install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_LIBDIR})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/ctest)
```

--------------------------------

### Run CUTLASS example and capture synclog output

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/utilities.md

Execute the compiled CUTLASS example with `synclog` enabled, redirecting its output to a file. Setting `--iterations=0` ensures `synclog` information is printed for the reference run.

```bash
$ ./54_hopper_fp8_warp_specialized_gemm --iterations=0 &> synclog.txt
```

--------------------------------

### POST /cutlass/gemm/device/GemmBatched/initialize

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html

Initializes the GEMM state using provided arguments, workspace, and CUDA stream.

```APIDOC
## POST /cutlass/gemm/device/GemmBatched/initialize

### Description
Initializes GEMM state from arguments.

### Method
POST

### Endpoint
cutlass::gemm::device::GemmBatched::initialize(Arguments const &args, void *workspace, cudaStream_t stream)

### Parameters
#### Request Body
- **args** (Arguments) - Required - GEMM configuration arguments including matrix dimensions and pointers.
- **workspace** (void*) - Optional - Pointer to device workspace memory (defaults to nullptr).
- **stream** (cudaStream_t) - Optional - CUDA stream for execution (defaults to nullptr).

### Response
#### Success Response (200)
- **status** (Status) - Returns a cutlass::Status code indicating success or the specific error encountered during initialization.
```

--------------------------------

### Execute Turing TensorOp Conv2dFprop Example with Various Parameters

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/implicit_gemm_convolution.md

These commands demonstrate running the convolution example with different tensor dimensions and an optional host-side reference check for correctness verification.

```bash
./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=32 --h=224 --w=224 --c=128 --k=256 --r=1 --s=1
```

```bash
./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=1 --h=224 --w=224 --c=32 --k=32 --r=3 --s=3 --ref-check
```

--------------------------------

### Configure CUTLASS for Blackwell GEMM Green Context Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/95_blackwell_gemm_green_context/README.md

Configure the CUTLASS build system to enable examples and target the SM100a architecture for Blackwell GPUs.

```shell
cmake <cutlass_root> \
  -DCUTLASS_NVCC_ARCHS=100a \
  -DCUTLASS_ENABLE_EXAMPLES=ON \
  -DCUTLASS_ENABLE_TESTS=OFF \
  -DCUTLASS_ENABLE_LIBRARY=OFF \
  -DCUTLASS_ENABLE_PROFILER=OFF \
  -DCMAKE_BUILD_TYPE=Release
```

--------------------------------

### Add CUTLASS Example Executable with CMake

Source: https://github.com/nvidia/cutlass/blob/main/examples/02_dump_reg_shmem/CMakeLists.txt

Use this CMake function to register a new executable example within the CUTLASS project. It specifies the example name, source file, and disables testing for this particular example.

```CMake
cutlass_example_add_executable(
  02_dump_reg_shmem 
  dump_reg_shmem.cu
  DISABLE_TESTS ON
)
```

--------------------------------

### Set Default Installation Prefix in CMake

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This snippet sets the default installation prefix to 'install' if it hasn't been initialized, making it a cached path variable.

```CMake
if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
  set(CMAKE_INSTALL_PREFIX install CACHE PATH "Default installation location." FORCE)
endif()
```

--------------------------------

### initialize(Arguments const &args, void *workspace)

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1gemm_1_1device_1_1GemmSplitKParallel.html

Initializes the GEMM state using provided arguments and workspace.

```APIDOC
## METHOD initialize

### Description
Initializes GEMM state from arguments.

### Method
MEMBER FUNCTION

### Endpoint
initialize(Arguments const &args, void *workspace)

### Parameters
#### Path Parameters
- No path parameters.

#### Query Parameters
- No query parameters.

#### Request Body
- **args** (Arguments const &) - Required - Structure containing GEMM configuration arguments.
- **workspace** (void *) - Required - Pointer to a workspace memory region.

### Request Example
{
  "args": {
    "//": "Details of Arguments struct"
  },
  "workspace": "0x..."
}

### Response
#### Success Response (Status)
- Returns a `Status` indicating the success or failure of the initialization.

#### Response Example
{
  "status": "Success"
}
```

--------------------------------

### Generate CTest File for Installation

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This snippet configures a CTest file specifically for installation, using a non-extended format compatible with CTest's add_test when installed.

```CMake
set(TEST_EXE_PATH $<TARGET_FILE_NAME:${TARGET}>)
set(TEST_USE_EXTENDED_FORMAT OFF) # ctest does not support extended add_test format.
configure_file("${CUTLASS_CTEST_TEMPLATE_FILE}" "${TEST_GEN_DIR}/CTestTestfile.${TEST_NAME}.install.cmake.in" @ONLY)
```

--------------------------------

### Run Ampere Gather/Scatter Convolution Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/59_ampere_gather_scatter_conv/README.md

This command executes the Ampere gather/scatter convolution example, demonstrating performance for both dense and gather/scatter forward propagation kernels on Ampere and Ada GPUs. The output shows TFLOP/s for each kernel.

```Shell
$> ./examples/59_ampere_gather_scatter_conv/59_ampere_gather_scatter_conv --n=131072 --i=128 --no-check
Ampere convolution forward propagation kernel supporting both affine and gather/scatter tensors.

Allocating tensors ... done.
Initializing data ... done.
Initializing gather/scatter index buffers ... done.

Running dense fprop kernel
Conv TFLOP count = 0.927713
Conv dense perf: 31.027376ms | TFLOP/s = 29.899819

Running gather/scatter fprop kernel
Conv TFLOP count = 0.927713
Conv gather/scatter perf: 28.973721ms | TFLOP/s = 32.019117
```

--------------------------------

### cutlass::library::Manifest::initialize

Source: https://github.com/nvidia/cutlass/blob/main/docs/manifest_8h_source.html

Performs top-level initialization for the manifest.

```APIDOC
## Function: cutlass::library::Manifest::initialize

### Description
Top-level initialization.

### Function Signature
Status initialize()

### Parameters
- No parameters.

### Return Value
- **Status** - A status code indicating the success or failure of the initialization.
```

--------------------------------

### GET get()

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__eb7d20f8b9d69e0ae5e7ef51dc480867.html

Retrieves a pointer to the current tile access location.

```APIDOC
## GET get()\n\n### Description\nReturns a pointer to the current tile access location for the iterator.\n\n### Method\nGET (C++ Member Function)\n\n### Endpoint\ncutlass::transform::threadblock::RegularTileAccessIterator::get\n\n### Response\n#### Success Response (200)\n- **AccessType*** (pointer) - A pointer to the memory location of the current tile access.
```

--------------------------------

### Define JIT Function with Dynamic and Static Arguments (Python)

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/pythonDSL/cute_dsl_general/dsl_jit_arg_generation.rst

Illustrates how to define a JIT function using `@cute.jit` with both dynamic (e.g., `cutlass.Int32`) and static (`cutlass.Constexpr`) arguments, showing how `printf` handles them.

```python
import cutlass
import cutlass.cute as cute

@cute.jit
def foo(x: cutlass.Int32, y: cutlass.Constexpr):
    print("x = ", x)
    print("y = ", y)
    cute.printf("x: {}", x)
    cute.printf("y: {}", y)

foo(2, 2)
```

--------------------------------

### Use Zipped, Tiled, and Flat Products in CuTe

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/cute_layout_algebra.ipynb

Illustrates `cute.zipped_product`, `cute.tiled_product`, and `cute.flat_product` as variations of `logical_product` that rearrange modes into specific forms for different tiling strategies.

```python
@cute.jit
def zipped_tiled_flat_product_example():
    """
    Demonstrates zipped, tiled, and flat products
    Layout Shape : (M, N, L, ...)
    Tiler Shape  : <TileM, TileN>

    zipped_product  : ((M,N), (TileM,TileN,L,...))
    tiled_product   : ((M,N), TileM, TileN, L, ...)
    flat_product    : (M, N, TileM, TileN, L, ...)
    """
    # Define the original layout
    layout = cute.make_layout((2, 5), stride=(5, 1))

    # Define the tiler
    tiler = cute.make_layout((3, 4), stride=(1, 3))

    # Apply zipped product
    zipped_result = cute.zipped_product(layout, tiler=tiler)

    # Apply tiled product
    tiled_result = cute.tiled_product(layout, tiler=tiler)

    # Apply flat product
    flat_result = cute.flat_product(layout, tiler=tiler)

    # Print results
    print(">>> Layout:", layout)
    print(">>> Tiler :", tiler)
    print(">>> Zipped Product Result:", zipped_result)
    print(">>> Tiled Product Result:", tiled_result)
    print(">>> Flat Product Result:", flat_result)
    cute.printf(">?? Zipped Product Result: {}", zipped_result)
    cute.printf(">?? Tiled Product Result: {}", tiled_result)
    cute.printf(">?? Flat Product Result: {}", flat_result)


z_t_f_product_example()
```

--------------------------------

### GET get

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__a3c11cf1f00ef7a1efb8389ac6e4c6e0.html

Returns a pointer to the current access type location.

```APIDOC
## GET get

### Description
Returns a pointer to the current memory location based on the iterator state.

### Method
GET

### Endpoint
get()

### Response
#### Success Response (200)
- **pointer** (AccessType*) - A pointer to the current element access type.
```

--------------------------------

### Perform 1D Logical Product with CuTe

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/cute_layout_algebra.ipynb

Use `cute.logical_product` to reproduce a 1D layout based on a tiler, demonstrating how modes are combined and restrided.

```python
@cute.jit
def logical_product_1d_example():
    """
    Demonstrates 1D logical product
    """
    # Define the original layout
    layout = cute.make_layout((2, 2), stride=(4, 1))  # (2,2):(4,1)

    # Define the tiler
    tiler = cute.make_layout(6, stride=1)  # Apply to layout 6:1

    # Apply logical product
    result = cute.logical_product(layout, tiler=tiler)

    # Print results
    print(">>> Layout:", layout)
    print(">>> Tiler :", tiler)
    print(">>> Logical Product Result:", result)
    cute.printf(">?? Logical Product Result: {}", result)


logical_product_1d_example()
```

--------------------------------

### Install Generated Kernel List File in CMake

Source: https://github.com/nvidia/cutlass/blob/main/tools/library/CMakeLists.txt

Installs the generated kernel list file to the CMake info directory under cutlass. This makes the kernel manifest available after installation.

```cmake
install(
  FILES ${CUTLASS_LIBRARY_GENERATED_KERNEL_LIST_FILE}
  DESTINATION ${CMAKE_INSTALL_INFODIR}/cutlass
  )
```

--------------------------------

### Install Profiler Regressions List File

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This command installs the specified profiler regression list file to the CMake install info directory, renaming it to 'profiler_regressions.csv' for consistent access.

```CMake
  install(
    FILES ${CUTLASS_PROFILER_REGRESSION_LIST_FILE}
    DESTINATION ${CMAKE_INSTALL_INFODIR}/cutlass
    RENAME profiler_regressions.csv
    )
```

--------------------------------

### Conditionally Install CTest Testfile

Source: https://github.com/nvidia/cutlass/blob/main/CMakeLists.txt

This block conditionally generates and installs a CTest test file if the CUTLASS_INSTALL_TESTS option is enabled. It renames the file during installation to a standard CTestTestfile name.

```CMake
    if (CUTLASS_INSTALL_TESTS)

    file(GENERATE
      OUTPUT "${TEST_GEN_DIR}/CTestTestfile.${TEST_NAME}.install.cmake"
      INPUT "${TEST_GEN_DIR}/CTestTestfile.${TEST_NAME}.install.cmake.in"
      )

    install(
      FILES "${TEST_GEN_DIR}/CTestTestfile.${TEST_NAME}.install.cmake"
      DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/ctest/${TEST_NAME}
      RENAME CTestTestfile.${TEST_NAME}.cmake
      )

    endif()
```

--------------------------------

### Manifest::initialize()

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1library_1_1Manifest-members.html

Initializes the manifest and prepares it for use. This method must be called before adding operations to the manifest.

```APIDOC
## initialize()

### Description
Initializes the manifest and prepares it for use.

### Method
Member Function

### Signature
```cpp
void initialize()
```
```

--------------------------------

### Pipeline Example - Multi-threaded Synchronization

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/pipeline.md

Complete example demonstrating a 4-stage asynchronous pipeline with 2 producer threads and 1 consumer thread. Shows the typical pattern of producer acquisition, work, commit, consumer wait, work, and release operations.

```APIDOC
## Pipeline Example - 4-Stage with 2 Producers and 1 Consumer

### Overview
Demonstrates basic asynchronous pipeline usage with multiple producer threads and a single consumer thread.

### Configuration
- **NumStages**: 4
- **Producer Threads**: 2 (thread_idx 0 and 1)
- **Consumer Threads**: 1 (thread_idx 2)

### Producer Thread Pattern
```cpp
if (thread_idx == 0 or thread_idx == 1) {
  PipelineState smem_pipe_write = cutlass::make_producer_start_state<MainloopPipeline>();
  for (; iter > 0; --iter) {
    pipeline.producer_acquire(smem_pipe_write);
    // Producer operations (e.g., shared memory writes)
    pipeline.producer_commit(smem_pipe_write);
    ++smem_pipe_write;
  }
}
```

### Consumer Thread Pattern
```cpp
else if (thread_idx == 2) {
  PipelineState smem_pipe_read;
  for (; iter > 0; --iter) {
    pipeline.consumer_wait(smem_pipe_read);
    // Consumer operations
    pipeline.consumer_release(smem_pipe_read);
    ++smem_pipe_read;
  }
}
```

### Key Points
- Producer threads must guarantee memory visibility for consumer threads
- Multiple producer threads can work on different stages simultaneously
- Consumer thread blocks until producer commits to a stage
- Pipeline stages are reused in circular fashion
- Refer to unit tests and pipeline classes for advanced variations
```

--------------------------------

### Install cutlass_library package

Source: https://github.com/nvidia/cutlass/blob/main/python/README.md

Install the cutlass_library package for enumerating and emitting CUTLASS C++ kernels. This is used by the CMake system and can also be installed automatically with the main CUTLASS Python interface.

```bash
python setup_library.py develop --user
```

--------------------------------

### Build CUTLASS Python Interface Documentation

Source: https://github.com/nvidia/cutlass/blob/main/python/docs/index.html

Execute these commands sequentially to generate the HTML documentation for the CUTLASS Python interface after installing prerequisites.

```bash
sphinx-apidoc -o docs_src/source/ cutlass/ cutlass/backend\*
```

```bash
cd docs_src
```

```bash
make html
```

```bash
mv _build/\* ../docs
```

--------------------------------

### Test CUTLASS Python interface installation

Source: https://github.com/nvidia/cutlass/blob/main/python/README.md

Verify successful installation by importing cutlass and running a simple GEMM operation with float16 matrices. Requires numpy and a working CUDA installation.

```python
import cutlass
import numpy as np

plan = cutlass.op.Gemm(element=np.float16, layout=cutlass.LayoutType.RowMajor)
A, B, C, D = [np.ones((128, 128), dtype=np.float16) for i in range(4)]
plan.run(A, B, C, D)
```

--------------------------------

### Execute Ampere Gather/Scatter Convolution Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/59_ampere_gather_scatter_conv/README.md

These commands demonstrate how to run the Ampere gather/scatter convolution example with various input sizes and options. Different parameters like 'n' (number of images) and 'i' (iterations) can be adjusted.

```sh
./59_ampere_gather_scatter_conv
```

```sh
./59_ampere_gather_scatter_conv --n=108
```

```sh
./59_ampere_gather_scatter_conv --n=4096 --i=1
```

```sh
./59_ampere_gather_scatter_conv --n=1080 --i=1000
```

```sh
./59_ampere_gather_scatter_conv --n=131072 --i=1000 --no-check
```

--------------------------------

### Applying Logical and Zipped Divides in C++

Source: https://github.com/nvidia/cutlass/blob/main/media/docs/cpp/cute/02_layout_algebra.md

Demonstrates how to use logical_divide and zipped_divide with a specific layout and tiler.

```cpp
// A: shape is (9,32)
auto layout_a = make_layout(make_shape (Int< 9>{}, make_shape (Int< 4>{}, Int<8>{})),
                            make_stride(Int<59>{}, make_stride(Int<13>{}, Int<1>{})));
// B: shape is (3,8)
auto tiler = make_tile(Layout<_3,_3>{},           // Apply     3:3     to mode-0
                       Layout<Shape <_2,_4>,      // Apply (2,4):(1,8) to mode-1
                              Stride<_1,_8>>{});

// ((TileM,RestM), (TileN,RestN)) with shape ((3,3), (8,4))
auto ld = logical_divide(layout_a, tiler);
// ((TileM,TileN), (RestM,RestN)) with shape ((3,8), (3,4))
auto zd = zipped_divide(layout_a, tiler);
```

--------------------------------

### CMake Conditional Build for Distributed GEMM Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/65_distributed_gemm/CMakeLists.txt

Registers the distributed GEMM example executable only when CUTLASS_NVCC_ARCHS includes the 90a architecture. This ensures the example is built only on compatible GPU targets.

```cmake
if (CUTLASS_NVCC_ARCHS MATCHES 90a)
cutlass_example_add_executable(
  65_distributed_gemm
  65_distributed_gemm.cu
  )
endif()
```

--------------------------------

### GET get()

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen809793e785fb4211888c6b4e5dcfcb39.html

Retrieves a pointer to the current tile access location in memory.

```APIDOC
## GET get()

### Description
Returns a pointer to the current tile access location in memory for the specialized RowMajorInterleaved layout.

### Method
GET

### Endpoint
cutlass::transform::threadblock::PredicatedTileAccessIterator::get()

### Response
#### Success Response (200)
- **AccessType*** (pointer) - Pointer to the current memory access location.

### Response Example
{
  "pointer": "0x7f8c12345678"
}
```

--------------------------------

### Install CUTLASS Python interface in developer mode

Source: https://github.com/nvidia/cutlass/blob/main/python/docs/_sources/install.md.txt

Installs the package in developer mode so that source changes are reflected immediately. Ensure the installed cuda-python version matches the CUDA version in CUDA_INSTALL_PATH.

```bash
python setup.py develop --user
```

--------------------------------

### CUTLASS Tiled Copy Tutorial Executables

Source: https://github.com/nvidia/cutlass/blob/main/examples/cute/tutorial/CMakeLists.txt

Registers two tiled copy tutorial example executables: a basic tiled copy implementation and a conditional variant.

```cmake
cutlass_example_add_executable(
  cute_tutorial_tiled_copy
  tiled_copy.cu
)
```

```cmake
cutlass_example_add_executable(
  cute_tutorial_tiled_copy_if
  tiled_copy_if.cu
)
```

--------------------------------

### GET get()

Source: https://github.com/nvidia/cutlass/blob/main/docs/classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemenab63a1e105bf37f6371516cb9e2c5a7a.html

Retrieves a pointer to the current tile's access type for the iterator.

```APIDOC
## GET get()

### Description
Returns a pointer to the current tile's access type. This is a constant inline method used to access the underlying data.

### Method
GET

### Endpoint
cutlass::transform::threadblock::PredicatedTileAccessIterator::get

### Response
#### Success Response (200)
- **AccessType*** (pointer) - A pointer to the memory location of the current tile.
```

--------------------------------

### initialize(Arguments const &args, void *workspace, cudaStream_t stream)

Source: https://github.com/nvidia/cutlass/blob/main/docs/include_2cutlass_2gemm_2device_2gemm_8h_source.html

Prepares the GEMM object for execution by setting up internal parameters and associating the workspace and CUDA stream.

```APIDOC
## INITIALIZE cutlass::gemm::device::Gemm::initialize

### Description
Initializes the GEMM state with the provided arguments. This method must be called before launching the kernel to ensure the grid shape and internal parameters are correctly configured.

### Method
INSTANCE

### Endpoint
cutlass::gemm::device::Gemm::initialize

### Parameters
#### Request Body
- **args** (Arguments) - Required - GEMM configuration parameters including problem size and tensor references.
- **workspace** (void*) - Optional - Pointer to allocated device memory workspace (default: nullptr).
- **stream** (cudaStream_t) - Optional - The CUDA stream in which the GEMM operation will be executed (default: nullptr).

### Response
#### Success Response (200)
- **status** (Status) - Returns Status::kSuccess upon successful initialization of the GEMM operator.

### Response Example
{
  "status": "Status::kSuccess"
}
```

--------------------------------

### Demonstrate Static and Dynamic Value Printing in CuTe

Source: https://github.com/nvidia/cutlass/blob/main/examples/python/CuTeDSL/cute/notebooks/print.ipynb

Defines a JIT-compiled function `print_example` to illustrate how Python's `print` handles static values at compile time and `cute.printf` displays both static and dynamic values at runtime, including layout representations.

```python
@cute.jit
def print_example(a: cutlass.Int32, b: cutlass.Constexpr[int]):
    """
    Demonstrates different printing methods in CuTe and how they handle static vs dynamic values.

    This example shows:
    1. How Python's `print` function works with static values at compile time but can't show dynamic values
    2. How `cute.printf` can display both static and dynamic values at runtime
    3. The difference between types in static vs dynamic contexts
    4. How layouts are represented in both printing methods

    Args:
        a: A dynamic Int32 value that will be determined at runtime
        b: A static (compile-time constant) integer value
    """
    # Use Python `print` to print static information
    print(">>>", b)  # => 2
    # `a` is dynamic value
    print(">>>", a)  # => ?

    # Use `cute.printf` to print dynamic information
    cute.printf(">?? {}", a)  # => 8
    cute.printf(">?? {}", b)  # => 2

    print(">>>", type(a))  # => <class 'cutlass.Int32'>
    print(">>>", type(b))  # => <class 'int'>

    layout = cute.make_layout((a, b))
    print(">>>", layout)  # => (?,2):(1,?)
cute.printf(">?? {}", layout)  # => (8,2):(1,8)
```

--------------------------------

### Conditional Blackwell Sparse GEMM Example Build

Source: https://github.com/nvidia/cutlass/blob/main/examples/83_blackwell_sparse_gemm/CMakeLists.txt

CMake configuration that conditionally compiles the Blackwell sparse GEMM example only when CUTLASS_NVCC_ARCHS includes the 100a architecture. Uses the cutlass_example_add_executable macro to register the example.

```CMake
if (CUTLASS_NVCC_ARCHS MATCHES 100a)

cutlass_example_add_executable(
  83_blackwell_sparse_gemm
  83_blackwell_sparse_gemm.cu
)

endif()
```

--------------------------------

### CMake Conditional Build for Blackwell GEMM Example

Source: https://github.com/nvidia/cutlass/blob/main/examples/78_blackwell_emulated_bf16x9_gemm/CMakeLists.txt

Registers a CUTLASS example executable only when targeting Blackwell GPU architecture (CUTLASS_NVCC_ARCHS matches 100a). Use this pattern to conditionally compile architecture-specific examples.

```CMake
if (CUTLASS_NVCC_ARCHS MATCHES 100a)
cutlass_example_add_executable(
  78_blackwell_emulated_bf16x9_gemm
  78_blackwell_emulated_bf16x9_gemm.cu
  )
endif()
```

--------------------------------

### Execute Hopper GEMM with L2 Weight Prefetch

Source: https://github.com/nvidia/cutlass/blob/main/examples/63_hopper_gemm_with_weight_prefetch/README.md

Run the example executable with various configurations for matrix dimensions, overlap ratio, and prefetch ratio to observe performance differences.

```bash
./63_hopper_gemm_with_weight_prefetch --m=8192 --n=1 --k=8192
```

```bash
echo "Without overlap and prefetch"
./63_hopper_gemm_with_weight_prefetch --o=-1.0 --p=-1.0
```

```bash
echo "Overlap ratio of 0.5, best effort prefetch"
./63_hopper_gemm_with_weight_prefetch --o=0.5 --p=-1.0
```

```bash
echo "Overlap ratio of 0.8, prefetch ratio of 0.7"
./63_hopper_gemm_with_weight_prefetch --o=0.8 --p=0.7
```

--------------------------------

### CMake executable target for CUTLASS SYRK example

Source: https://github.com/nvidia/cutlass/blob/main/examples/31_basic_syrk/CMakeLists.txt

Registers a CMake executable target for the basic SYRK example using the cutlass_example_add_executable helper function. The target compiles basic_syrk.cu and is identified as example 31.

```CMake
cutlass_example_add_executable(
  31_basic_syrk
  basic_syrk.cu
)
```

--------------------------------

### Define Main Custom Target for Two-Tensor Fusion Examples

Source: https://github.com/nvidia/cutlass/blob/main/examples/13_two_tensor_op_fusion/CMakeLists.txt

This command creates a top-level custom target that depends on both the fused GEMM and fused CONV example targets, ensuring all fusion examples are built together.

```cmake
add_custom_target(13_two_tensor_op_fusion
DEPENDS 13_fused_two_gemms
        13_fused_two_convs
)
```