diff --git a/CHANGELOG.md b/CHANGELOG.md index fe5c4d7a0..c19e56d83 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -2,6 +2,13 @@ # CUTLASS 4.x +## [4.3.6](https://github.com/NVIDIA/cutlass/releases/tag/v4.3.6) (2026-06-17) + +### CUTLASS C++ +* Make [version.h](https://github.com/NVIDIA/cutlass/blob/release/4.4/include/cutlass/version.h) NVRTC JIT compilation compatible. +* Allow linking large cutlass library on 64bit platform. +* Fix alignment-related miscalculation for pipeline stages of Blackwell blockscaled GEMM. + ## [4.3.5](https://github.com/NVIDIA/cutlass/releases/tag/v4.3.5) (2026-01-09) ### CuTe DSL diff --git a/README.md b/README.md index af692c802..959f55c30 100644 --- a/README.md +++ b/README.md @@ -1,9 +1,9 @@ ![ALT](./media/images/gemm-hierarchy-with-epilogue-no-labels.png "Complete CUDA GEMM decomposition") # Overview -# CUTLASS 4.3.5 +# CUTLASS 4.3.6 -_CUTLASS 4.3.5 - Jan 2026_ +_CUTLASS 4.3.6 - June 2026_ CUTLASS is a collection of abstractions for implementing high-performance matrix-matrix multiplication (GEMM) and related computations at all levels and scales within CUDA. It incorporates strategies for diff --git a/include/cutlass/gemm/collective/builders/sm100_blockscaled_umma_builder.inl b/include/cutlass/gemm/collective/builders/sm100_blockscaled_umma_builder.inl index 78b66196d..2c28c2ca4 100644 --- a/include/cutlass/gemm/collective/builders/sm100_blockscaled_umma_builder.inl +++ b/include/cutlass/gemm/collective/builders/sm100_blockscaled_umma_builder.inl @@ -73,7 +73,7 @@ sm100_compute_stage_count_or_override_blockscaled(StageCountAutoCarveout::SharedStorage); + constexpr auto mainloop_pipeline_bytes = cutlass::round_up(sizeof(typename cutlass::PipelineTmaUmmaAsync<1>::SharedStorage), 128);; constexpr auto a_bits = cute::sizeof_bits_v; constexpr auto b_bits = cute::sizeof_bits_v; constexpr auto stage_sfa_bytes = size(filter_zeros(TileShapeSFA{})); diff --git a/include/cutlass/version.h b/include/cutlass/version.h index 658933e68..9f87549fd 100644 --- a/include/cutlass/version.h +++ b/include/cutlass/version.h @@ -31,12 +31,16 @@ #pragma once +#if defined(__CUDACC_RTC__) +#include CUDA_STD_HEADER(cstdint) +#else #include #include +#endif #define CUTLASS_MAJOR 4 #define CUTLASS_MINOR 3 -#define CUTLASS_PATCH 5 +#define CUTLASS_PATCH 6 #ifdef CUTLASS_VERSIONS_GENERATED #include "cutlass/version_extended.h" @@ -65,6 +69,7 @@ namespace cutlass { return CUTLASS_BUILD + 0; } +#if !defined(__CUDACC_RTC__) inline std::string getVersionString() { std::string version = "@CUTLASS_VERSION@"; if (getVersionBuild()) { @@ -72,9 +77,10 @@ namespace cutlass { } return version; } - + inline std::string getGitRevision() { return "@CUTLASS_REVISION@"; } +#endif // !defined(__CUDACC_RTC__) } // namespace cutlass diff --git a/python/cutlass_cppgen/__init__.py b/python/cutlass_cppgen/__init__.py index 26085c3e1..fa2398f6f 100644 --- a/python/cutlass_cppgen/__init__.py +++ b/python/cutlass_cppgen/__init__.py @@ -133,7 +133,7 @@ def get_option_registry(): this._option_registry = OptionRegistry(device_cc()) return this._option_registry -this.__version__ = '4.3.5' +this.__version__ = '4.3.6' from cutlass_cppgen.backend import create_memory_pool from cutlass_cppgen.emit.pytorch import pytorch diff --git a/python/setup_cutlass.py b/python/setup_cutlass.py index d057de160..7b6a65ad6 100644 --- a/python/setup_cutlass.py +++ b/python/setup_cutlass.py @@ -51,7 +51,7 @@ setup_pycute.perform_setup() setup( name='cutlass_cppgen', - version='4.3.5', + version='4.3.6', description='CUTLASS Pythonic Interface', package_dir={'': '.'}, packages=[ diff --git a/python/setup_library.py b/python/setup_library.py index c8f0bc95d..f4442e5d6 100644 --- a/python/setup_library.py +++ b/python/setup_library.py @@ -36,7 +36,7 @@ from setuptools import setup def perform_setup(): setup( name='cutlass_library', - version='4.3.5', + version='4.3.6', description='CUTLASS library generation scripts', packages=['cutlass_library'] ) diff --git a/python/setup_pycute.py b/python/setup_pycute.py index c13638782..64a703b5f 100644 --- a/python/setup_pycute.py +++ b/python/setup_pycute.py @@ -36,7 +36,7 @@ from setuptools import setup def perform_setup(): setup( name='pycute', - version='4.3.5', + version='4.3.6', description='Python implementation of CuTe', packages=['pycute'], ) diff --git a/tools/library/CMakeLists.txt b/tools/library/CMakeLists.txt index b512aa9c4..c8dd78ef8 100644 --- a/tools/library/CMakeLists.txt +++ b/tools/library/CMakeLists.txt @@ -123,6 +123,19 @@ function(cutlass_add_cutlass_library) PRIVATE cutlass_library_internal_interface ) + # Use medium code model to handle large libraries (>2GB) on Linux x86_64 + # Conservative approach: only apply on Linux with GCC/Clang/Intel compilers + # This addresses the "relocation truncated to fit" linker error when libraries exceed 2GB + if(CMAKE_SYSTEM_NAME STREQUAL "Linux" AND + CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|AMD64" AND + CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang|IntelLLVM") + message(STATUS "CUTLASS: Applying -mcmodel=medium for large library support (Linux x86_64, ${CMAKE_CXX_COMPILER_ID})") + target_compile_options(${__NAME}_objs PRIVATE + $<$:-mcmodel=medium> + $<$:-Xcompiler=-mcmodel=medium> + ) + endif() + if (CUTLASS_BUILD_MONO_LIBRARY AND __SUFFIX) # If we're only building a single monolithic library then we