Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion examples/59_ampere_gather_scatter_conv/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ This example demonstrates a few super cool features of CUTLASS and CuTe. It show
The most common strategy for implementing high performance convolution kernels on the GPU is to transform
the activation tensor in such a way that we can perform the computation as a GEMM. This is called the
image to column (im2col) transformation. [CUTLASS 2.x implementation of im2col based convolutions is
documented separately](../../media/docs/implicit_gemm_convolution.md), and here we consider a fresh approach for CuTe.
documented separately](../../media/docs/cpp/implicit_gemm_convolution.md), and here we consider a fresh approach for CuTe.

A 3D convolution has the following input tensors:
- Activation tensor (Act): `((N,(D,H,W)), (C,(1,1,1)))`
Expand Down
2 changes: 1 addition & 1 deletion examples/65_distributed_gemm/REQUIREMENTS.md
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@

### Build
Make sure to set up CUTLASS with
support for [Programmatic Dependent Launch (PDL)](../../media/docs/dependent_kernel_launch.md),
support for [Programmatic Dependent Launch (PDL)](../../media/docs/cpp/dependent_kernel_launch.md),
that is with the `CUTLASS_ENABLE_GDC_FOR_SM90` flag.

```bash
Expand Down
2 changes: 1 addition & 1 deletion examples/82_blackwell_distributed_gemm/REQUIREMENTS.md
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@

### Build
Make sure to set up CUTLASS with
support for [Programmatic Dependent Launch (PDL)](../../media/docs/dependent_kernel_launch.md),
support for [Programmatic Dependent Launch (PDL)](../../media/docs/cpp/dependent_kernel_launch.md),
that is with the `CUTLASS_ENABLE_GDC_FOR_SM100` flag.

```bash
Expand Down
6 changes: 3 additions & 3 deletions examples/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@

demonstrates CUTLASS Utilities for allocating and initializing tensors

* [02_dump_reg_smem](02_dump_reg_smem/)
* [02_dump_reg_shmem](02_dump_reg_shmem/)

debugging utilities for printing register and shared memory contents

Expand Down Expand Up @@ -173,7 +173,7 @@

example demonstrating CUTLASS with Python interface

* [41_multi_head_attention](41_multi_head_attention/)
* [41_fused_multi_head_attention](41_fused_multi_head_attention/)

example demonstrating attention example with non-fixed sequence length input

Expand Down Expand Up @@ -206,7 +206,7 @@

Simple tensorop GEMM example using CUTLASS 3.0 APIs targeting NVIDIA Hopper architecture

* [49_hopper_gemm_schedules_with_collective_builder](49_hopper_gemm_schedules_with_collective_builder/)
* [49_hopper_gemm_with_collective_builder](49_hopper_gemm_with_collective_builder/)

Hopper GEMM example leveraging collective operation builders to showcase the builder API and the various kernel scheduled supported in CUTLASS 3.0 such as warp specialized persistent mainloops.

Expand Down
10 changes: 5 additions & 5 deletions examples/python/deprecated/README.md
Original file line number Diff line number Diff line change
@@ -1,23 +1,23 @@
# Examples of using the CUTLASS Python interface

* [00_basic_gemm](/examples/python/00_basic_gemm.ipynb)
* [00_basic_gemm](00_basic_gemm.ipynb)

Shows how declare, configure, compile, and run a CUTLASS GEMM using the Python interface

* [01_epilogue](/examples/python/01_epilogue.ipynb)
* [01_epilogue](01_epilogue.ipynb)

Shows how to fuse elementwise activation functions to GEMMs via the Python interface

* [02_pytorch_extension_grouped_gemm](/examples/python/02_pytorch_extension_grouped_gemm.ipynb)
* [02_pytorch_extension_grouped_gemm](02_pytorch_extension_grouped_gemm.ipynb)

Shows how to declare, compile, and run a grouped GEMM operation via the Python interface,
along with how the emitted kernel can be easily exported to a PyTorch CUDA extension.

* [03_basic_conv2d](/examples/python/03_basic_conv2d.ipynb)
* [03_basic_conv2d](03_basic_conv2d.ipynb)

Shows how to declare, configure, compile, and run a CUTLASS Conv2d using the Python interface

* [04_epilogue_visitor](/examples/python/04_epilogue_visitor.ipynb)
* [04_epilogue_visitor](04_epilogue_visitor.ipynb)

Shows how to fuse elementwise activation functions to GEMMs via the Python Epilogue Visitor interface

Expand Down
10 changes: 5 additions & 5 deletions media/docs/cpp/blackwell_functionality.md
Original file line number Diff line number Diff line change
Expand Up @@ -654,16 +654,16 @@ auto val_a_mk = tensor_sfa(make_coord(m,k,0));
# Blackwell SM120 GEMMs
The NVIDIA RTX 5000 Series GPUs introduce support for new narrow precision (4bit and 6bit) block-scaled and non-block-scaled tensor cores. The PTX ISA has extended the `mma` instructions to support these data formats which are 1x to 4x faster than Ada architecture's fp8 tensor cores. For more detailed information see [`mma` PTX documentation](https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-for-mma).

CUTLASS 4.0 has added support for these newly introduced narrow precision GEMMs. Similar to the Blackwell SM100 GEMMs, the SM120 GEMMs can be built using the collective builder interface. See examples in [examples/79_blackwell_geforce_gemm/](../../examples/79_blackwell_geforce_gemm/) and unit tests listed below.
CUTLASS 4.0 has added support for these newly introduced narrow precision GEMMs. Similar to the Blackwell SM100 GEMMs, the SM120 GEMMs can be built using the collective builder interface. See examples in [examples/79_blackwell_geforce_gemm/](../../../examples/79_blackwell_geforce_gemm/) and unit tests listed below.

The data types supported and tensor alignment requirements are the same as the Blackwell SM100 GEMMs. The scale factor layout is also the same as SM100 mentioned above. `OpClassTensorOp` is used for non-blockscaled narrow precision GEMMs and `OpClassBlockScaledTensorOp` is used for blockscaled narrow precision GEMMs.

| Ptx Instruction | Throughput | Notes | Unit Test |
|---------------------------------------------------------------------|----------------------------|-------|-----------|
|mma.sync.aligned.kind::f8f6f4 | 1x Ada Fp8 Tensor Core(2x for FP32 accumulator) | Mixed precision MMA with A={f4,f6,f8} x B={f4,f6,f8} TN layouts | [unit test](../../test/unit/gemm/device/sm120_tensorop_gemm/) |
|mma.sync.aligned.kind::mxf8f6f4.block_scale | 1x Ada Fp8 Tensor Core(2x for FP32 accumulator) | Block scaled mixed precision MMA with A={mxf4,mxf6,mxf8} x B={mxf4,mxf6,mxf8} with TN layouts | [unit test](../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_mxf6_mxf8_f32_f32.cu) |
|mma.sync.aligned.kind::mxf4.block_scale | 2x Ada Fp8 Tensor Core(4x for FP32 accumulator) | Block scaled MMA with A={mxf4} x B={mxf4} with TN layouts | [unit test](../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_mxf4_mxf4_f32_f32.cu) |
|mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::[2X\|4X] | 2x Ada Fp8 Tensor Core(4x for FP32 accumulator) | Block scaled MMA with A={mxf4} x B={mxf4} or A={nvf4} x B={nvf4} with TN layouts | [unit test](../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_nvf4_nvf4_f32_f32.cu) |
|mma.sync.aligned.kind::f8f6f4 | 1x Ada Fp8 Tensor Core(2x for FP32 accumulator) | Mixed precision MMA with A={f4,f6,f8} x B={f4,f6,f8} TN layouts | [unit test](../../../test/unit/gemm/device/sm120_tensorop_gemm/) |
|mma.sync.aligned.kind::mxf8f6f4.block_scale | 1x Ada Fp8 Tensor Core(2x for FP32 accumulator) | Block scaled mixed precision MMA with A={mxf4,mxf6,mxf8} x B={mxf4,mxf6,mxf8} with TN layouts | [unit test](../../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_mxf6_mxf8_f32_f32.cu) |
|mma.sync.aligned.kind::mxf4.block_scale | 2x Ada Fp8 Tensor Core(4x for FP32 accumulator) | Block scaled MMA with A={mxf4} x B={mxf4} with TN layouts | [unit test](../../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_mxf4_mxf4_f32_f32.cu) |
|mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::[2X\|4X] | 2x Ada Fp8 Tensor Core(4x for FP32 accumulator) | Block scaled MMA with A={mxf4} x B={mxf4} or A={nvf4} x B={nvf4} with TN layouts | [unit test](../../../test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/sm120_bs_gemm_nvf4_nvf4_f32_f32.cu) |

Besides the similarities, there are some key differences from the Blackwell SM100 GEMMs:

Expand Down