diff --git a/CHANGELOG.md b/CHANGELOG.md index f027e0d37..3df2d095c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -2,6 +2,21 @@ # CUTLASS 4.x +## [4.5.2](https://github.com/NVIDIA/cutlass/releases/tag/v4.5.2) (2026-05-22) + +### CuTe DSL +* New features + - Python 3.14t is now supported with GIL enabled + +* Bug fixing and improvements + - Fixed following issues: + https://github.com/NVIDIA/cutlass/issues/3240 + https://github.com/NVIDIA/cutlass/issues/3241 + +### CUTLASS C++ +* Fix missing convert fucntion in EVT for fp4 kernels. +* Avoid instantiate 2sm tma kernels where ctaN is none power of 64 when ctaN > 128 in profiler. + ## [4.5.1](https://github.com/NVIDIA/cutlass/releases/tag/v4.5.1) (2026-05-15) ### CuTe DSL diff --git a/README.md b/README.md index 28e2fe498..fdbb6f3db 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.5.1 +# CUTLASS 4.5.2 -_CUTLASS 4.5.1 - May 2026_ +_CUTLASS 4.5.2 - May 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 @@ -56,6 +56,7 @@ To get started quickly - please refer : - Initial linter support: Improved type hints on CuTe DSL APIs to support static type checkers like MyPy - dataclasses.dataclass is now supported for JIT compilaton and cute.compile for both plain and tvm-ffi path - cute.copy now supports user specified loop unrolling + - Python 3.14t is now supported with GIL enabled * Bug fixing and improvements - Improved source code correlation for profiling/debugging @@ -69,6 +70,8 @@ To get started quickly - please refer : https://github.com/NVIDIA/cutlass/issues/3208 https://github.com/NVIDIA/cutlass/issues/3201 https://github.com/NVIDIA/cutlass/issues/3227 + https://github.com/NVIDIA/cutlass/issues/3240 + https://github.com/NVIDIA/cutlass/issues/3241 - Fixed Jax int64 stride divisibility issue - Fixed issues for SM120 blockscaled MMAs - added missing MXFP8MMAOP and MXF8F6F4MMAOP for sm120. @@ -110,8 +113,10 @@ To get started quickly - please refer : - Fix CUTLASS clang build issues - Remove `PipelineStorage` shadowing in SM100 complex epilogue - Fix build issue in SM90 epilogue fusion visitor TMA warpspecialized + - Fix missing convert fucntion in EVT for fp4 kernels * Fix some profiler issues: - Add missing reference kernels for blockwise GEMM profiler. + - Avoid instantiate 2sm tma kernels where ctaN is none power of 64 when ctaN > 128 in profiler. Note: CUTLASS 4.x builds are known to be down on Windows platforms for all CUDA toolkits. CUTLASS team is working on a fix. diff --git a/examples/python/CuTeDSL/cute_ext/blackwell/dense_gemm.py b/examples/python/CuTeDSL/cute_ext/blackwell/dense_gemm.py index b717cb76d..cbb8a3314 100644 --- a/examples/python/CuTeDSL/cute_ext/blackwell/dense_gemm.py +++ b/examples/python/CuTeDSL/cute_ext/blackwell/dense_gemm.py @@ -50,8 +50,7 @@ import cutlass.cute.testing as testing # - D has shape (M, N, L) and is the output in global memory # - L is the batch dimension # -# The kernel uses the LIR (Low-level Intermediate Representation) DSL which is a Python DSL -# for writing high-performance, Blackwell (SM100)-targeted kernels on top of CuTe abstractions. +# The kernel uses Python DSL for writing high-performance, Blackwell (SM100)-targeted kernels on top of CuTe abstractions. # # KEY CONCEPTS: # - TMA (Tensor Memory Accelerator): Hardware feature for high-bandwidth GMEM <-> SMEM transfers @@ -89,7 +88,7 @@ class DenseGemmKernel: This class encapsulates all the configuration and logic for a high-performance batched matrix multiplication: D = A @ B (with optional epilogue operation). - The design follows LIR conventions: + The design follows conventions: 1. __init__: Store configuration parameters 2. __call__: JIT-decorated host launcher that computes grid and calls kernel 3. kernel: Device kernel that performs the actual computation @@ -223,7 +222,7 @@ class DenseGemmKernel: - Contains all SMEM/TMEM/RMEM allocations, pipeline setup, and compute logic - Is compiled to PTX and executed by each thread in the grid - This kernel follows the standard LIR GEMM structure: + This kernel follows the standard GEMM structure: 1. Create tiled_mma configuration 2. Compute tiler and divide tensors 3. Allocate SMEM, TMEM, and RMEM buffers @@ -307,7 +306,7 @@ class DenseGemmKernel: # ======================================================================================== # STEP 3: DIVIDE GLOBAL TENSORS INTO TILES (zipped_divide) # ======================================================================================== - # cute.zipped_divide is the PRIMARY tiling operation in LIR kernels. + # cute.zipped_divide is the PRIMARY tiling operation in kernels. # # CUTE ALGEBRA EXPLANATION - zipped_divide: # ------------------------------------------ diff --git a/include/cutlass/epilogue/collective/builders/sm100_builder.inl b/include/cutlass/epilogue/collective/builders/sm100_builder.inl index e5fafd520..b3c10aa1c 100644 --- a/include/cutlass/epilogue/collective/builders/sm100_builder.inl +++ b/include/cutlass/epilogue/collective/builders/sm100_builder.inl @@ -1209,6 +1209,27 @@ private: using CtaTileShape_MNK = decltype(cta_tile_shape()); using TmemWarpShape_MN = decltype(detail::sm100_tmem_warps()); + // SM100 2SM TMA epilogue with 16-bit elements and EpilogueTileAuto requires CtaN divisible + // by 64 when CtaN > 128. When MaxBits==16 and CtaN>128, N_perf=64 is selected, but if + // CtaN%64!=0 the epilogue tile falls back to CtaN, producing non-64-aligned strides in + // the SMEM swizzle layout that break upcast<64>. + static constexpr int EpiSmemMaxBits_ = + (sizeof_bits_v >= sizeof_bits_v) + ? sizeof_bits_v + : sizeof_bits_v; + static_assert( + !Is2SmMma || + !cute::is_same_v || + FusionOp::IsPerColScaleSupported || + EpiSmemMaxBits_ != 16 || + size<1>(CtaTileShape_MNK{}) <= 128 || + size<1>(CtaTileShape_MNK{}) % 64 == 0, + "SM100 2SM TMA epilogue: 16-bit element types (f16/bf16) with CtaN > 128 require " + "CtaN to be divisible by 64. CtaN=160 and CtaN=224 are unsupported. " + "Use a CtaN that is a multiple of 64 (e.g. 128, 192, 256) " + "or use a 32-bit output type (f32)." + ); + // Attempts to compute a reasonably performant epilogue tile or allows the user to provide one. static constexpr auto epilogue_tile() { diff --git a/include/cutlass/epilogue/collective/sm100_epilogue_array_nosmem.hpp b/include/cutlass/epilogue/collective/sm100_epilogue_array_nosmem.hpp index 176e7ad9c..d108bcb5f 100644 --- a/include/cutlass/epilogue/collective/sm100_epilogue_array_nosmem.hpp +++ b/include/cutlass/epilogue/collective/sm100_epilogue_array_nosmem.hpp @@ -42,6 +42,7 @@ #include "cute/tensor.hpp" #include "cute/numeric/numeric_types.hpp" #include "cutlass/cuda_host_adapter.hpp" +#include "cutlass/numeric_conversion.h" ///////////////////////////////////////////////////////////////////////////////////////////////// @@ -838,7 +839,12 @@ public: CUTLASS_PRAGMA_UNROLL for (int epi_v = 0; epi_v < size(tTR_rAcc_frg); ++epi_v) { - tTR_rD_frg(epi_v) = cst_callbacks.visit(tTR_rAcc_frg(epi_v), epi_v, epi_m, epi_n); + auto frg_visited = cst_callbacks.visit(tTR_rAcc_frg(epi_v), epi_v, epi_m, epi_n); + // For 4-bit output types (e.g. float_e2m1_t), the visitor returns ElementCompute + // (scaled float values) rather than ElementD. Explicitly convert here so the + // assignment is always well-typed regardless of visitor return type. + using VisitedElement = typename decltype(frg_visited)::Element; + tTR_rD_frg(epi_v) = NumericArrayConverter{}(frg_visited); } Tensor reduction_buffer = make_tensor( diff --git a/include/cutlass/epilogue/collective/sm100_epilogue_nosmem.hpp b/include/cutlass/epilogue/collective/sm100_epilogue_nosmem.hpp index c8944f157..924f19be9 100644 --- a/include/cutlass/epilogue/collective/sm100_epilogue_nosmem.hpp +++ b/include/cutlass/epilogue/collective/sm100_epilogue_nosmem.hpp @@ -759,7 +759,12 @@ public: CUTLASS_PRAGMA_UNROLL for (int epi_v = 0; epi_v < size(tTR_rAcc_frg); ++epi_v) { - tTR_rD_frg(epi_v) = cst_callbacks.visit(tTR_rAcc_frg(epi_v), epi_v, epi_m, epi_n); + auto frg_visited = cst_callbacks.visit(tTR_rAcc_frg(epi_v), epi_v, epi_m, epi_n); + // For 4-bit output types (e.g. float_e2m1_t), the visitor returns ElementCompute + // (scaled float values) rather than ElementD. Explicitly convert here so the + // assignment is always well-typed regardless of visitor return type. + using VisitedElement = typename decltype(frg_visited)::Element; + tTR_rD_frg(epi_v) = NumericArrayConverter{}(frg_visited); } Tensor reduction_buffer = make_tensor( diff --git a/include/cutlass/float8.h b/include/cutlass/float8.h index 6019de643..507d6858d 100644 --- a/include/cutlass/float8.h +++ b/include/cutlass/float8.h @@ -60,7 +60,7 @@ #if defined(__CUDA_ARCH__) # if (__CUDA_ARCH__ >= 1000) # if (__CUDACC_VER_MAJOR__ > 13) || ((__CUDACC_VER_MAJOR__ >= 13) && (__CUDACC_VER_MINOR__ >= 2)) -# define CUDA_PTX_FP8_BF16_CVT_ENABLED 1 +# define CUDA_PTX_FP8_BF16_CVT_SUPPORTED 1 # endif // (__CUDACC_VER_MAJOR__ > 13) || ((__CUDACC_VER_MAJOR__ >= 13) && (__CUDACC_VER_MINOR__ >= 2)) # endif // (__CUDA_ARCH__ >= 1000) #endif // defined(__CUDA_ARCH__) @@ -70,12 +70,18 @@ defined(CUTLASS_ARCH_MMA_SM103A_ENABLED) || defined(CUTLASS_ARCH_MMA_SM110A_ENABLED) ||\ defined(CUTLASS_ARCH_MMA_SM120A_ENABLED) || defined(CUTLASS_ARCH_MMA_SM121A_ENABLED)) # define CUDA_PTX_UE8M0_CVT_ENABLED 1 +# if CUDA_PTX_FP8_BF16_CVT_SUPPORTED +# define CUDA_PTX_FP8_BF16_CVT_ENABLED 1 +# endif #endif #if (defined(CUTLASS_ARCH_MMA_SM100F_ENABLED) || defined(CUTLASS_ARCH_MMA_SM101F_ENABLED) ||\ defined(CUTLASS_ARCH_MMA_SM103F_ENABLED) || defined(CUTLASS_ARCH_MMA_SM110F_ENABLED) ||\ defined(CUTLASS_ARCH_MMA_SM120F_ENABLED) || defined(CUTLASS_ARCH_MMA_SM121F_ENABLED)) # define CUDA_PTX_UE8M0_CVT_ENABLED 1 +# if CUDA_PTX_FP8_BF16_CVT_SUPPORTED +# define CUDA_PTX_FP8_BF16_CVT_ENABLED 1 +# endif #endif #ifdef __GNUC__ diff --git a/include/cutlass/version.h b/include/cutlass/version.h index 08596546e..f22a08a24 100644 --- a/include/cutlass/version.h +++ b/include/cutlass/version.h @@ -36,7 +36,7 @@ #define CUTLASS_MAJOR 4 #define CUTLASS_MINOR 5 -#define CUTLASS_PATCH 1 +#define CUTLASS_PATCH 2 #ifdef CUTLASS_VERSIONS_GENERATED #include "cutlass/version_extended.h" diff --git a/python/CuTeDSL/cutlass/cute/__init__.py b/python/CuTeDSL/cutlass/cute/__init__.py index aaa7787e6..eadb458bd 100644 --- a/python/CuTeDSL/cutlass/cute/__init__.py +++ b/python/CuTeDSL/cutlass/cute/__init__.py @@ -145,6 +145,7 @@ from .tuple import ( unwrap, wrap, ) +from cutlass._mlir.dialects.cute_nvgpu import ReductionKind as ReductionKind from .tensor import ( TensorSSA, ReductionOp, @@ -256,6 +257,7 @@ __all__ = [ "ThrCopy", "TensorSSA", "ReductionOp", + "ReductionKind", "SymInt", # Basic utility functions "assume", diff --git a/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py b/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py index 8adfccc5b..21b6b3e60 100644 --- a/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py +++ b/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py @@ -10,6 +10,7 @@ # is strictly prohibited. import enum +import warnings from dataclasses import dataclass from typing import Any, Optional, Type from typing_extensions import deprecated @@ -18,9 +19,13 @@ from cutlass.base_dsl.arch import Arch from cutlass.cutlass_dsl import BaseDSL import cutlass._mlir.dialects.cute_nvgpu as _cute_nvgpu_ir -from cutlass._mlir.dialects.nvvm import ReductionOp as ReductionOp +from cutlass._mlir.dialects.cute_nvgpu import ReductionKind as ReductionKind +from cutlass._mlir.dialects.cute import ReductionOp as _CuteReductionOp +from cutlass._mlir.dialects.nvvm import ReductionOp as _NvvmReductionOp from cutlass._mlir import ir +ReductionOp = ReductionKind + from ...atom import CopyOp, Trait, make_atom from ...typing import Int16, Int32, Int64, Pointer, Integer, Numeric from ..common import OpError, LoadCacheMode as LoadCacheMode_ @@ -952,9 +957,41 @@ class CopyReduceBulkTensorTileS2GOp(TmaCopyOp): This Operation uses TMA in the ``.tile`` mode. """ - reduction_kind: ReductionOp = ReductionOp.ADD + reduction_kind: ReductionKind = ReductionKind.ADD def __post_init__(self) -> None: + # Canonicalize reduction_kind to cute.ReductionKind by NAME (not value) + # to (1) eliminate the silent IntEnum cross-class miscompile + # (cute.ReductionOp.ADD == nvvm.ReductionOp.AND because both are + # integer 0) and (2) decouple from upstream NVVM enum stability. + kind = self.reduction_kind + if isinstance(kind, ReductionKind): + pass # native path — already the contract type + elif isinstance(kind, _NvvmReductionOp): + # Hardware-correct upstream enum; convert silently by name. + self.reduction_kind = ReductionKind[kind.name] + elif isinstance(kind, _CuteReductionOp): + # Validate the name first so an unknown member (e.g. MUL) raises + # a clean TypeError. Doing it before warnings.warn guarantees the + # invalid-name case isn't masked by the DeprecationWarning being + # escalated to an exception under -W error / warnings-as-errors. + if kind.name not in ReductionKind.__members__: + raise TypeError( + f"cute.ReductionOp.{kind.name} is not a valid TMA " + f"reduction kind. Valid kinds: " + f"{[m.name for m in ReductionKind]}." + ) + self.reduction_kind = ReductionKind[kind.name] + warnings.warn( + "Passing cute.ReductionOp to CopyReduceBulkTensorTileS2GOp " + "is deprecated: cute.ReductionOp models compute reductions " + "(for cute.reduce). TMA hardware reduction uses a separate " + "enum -- use cute.ReductionKind instead.", + DeprecationWarning, + stacklevel=3, + ) + # else: leave as-is; _to_ir raises TypeError on unknown types. + # Arch verification arch = BaseDSL._get_dsl().get_arch_enum() if not arch >= Arch.sm_90: @@ -979,25 +1016,18 @@ class CopyReduceBulkTensorTileS2GOp(TmaCopyOp): "Use cpasync.make_tiled_tma_atom to obtain a copy Atom for TMA" ) - def _to_ir(self) -> _cute_nvgpu_ir.ReductionKind: - if self.reduction_kind == ReductionOp.ADD: - return _cute_nvgpu_ir.ReductionKind.ADD - elif self.reduction_kind == ReductionOp.MIN: - return _cute_nvgpu_ir.ReductionKind.MIN - elif self.reduction_kind == ReductionOp.MAX: - return _cute_nvgpu_ir.ReductionKind.MAX - elif self.reduction_kind == ReductionOp.INC: - return _cute_nvgpu_ir.ReductionKind.INC - elif self.reduction_kind == ReductionOp.DEC: - return _cute_nvgpu_ir.ReductionKind.DEC - elif self.reduction_kind == ReductionOp.AND: - return _cute_nvgpu_ir.ReductionKind.AND - elif self.reduction_kind == ReductionOp.OR: - return _cute_nvgpu_ir.ReductionKind.OR - elif self.reduction_kind == ReductionOp.XOR: - return _cute_nvgpu_ir.ReductionKind.XOR - else: - assert False, "unrecognized self.reduction_kind" + def _to_ir(self) -> ReductionKind: + # __post_init__ canonicalizes reduction_kind to ReductionKind, so this + # is normally a pass-through. Raise (rather than assert) so the error + # survives ``python -O`` and we never silently return None. + kind = self.reduction_kind + if isinstance(kind, ReductionKind): + return kind + raise TypeError( + f"reduction_kind must be cute.ReductionKind; got " + f"{type(kind).__module__}.{type(kind).__name__} ({kind!r}). " + f"Valid kinds: {[m.name for m in ReductionKind]}." + ) class CopyReduceBulkTensorTileS2GNonExecTrait(Trait): diff --git a/python/CuTeDSL/requirements-cu13.txt b/python/CuTeDSL/requirements-cu13.txt index 6434aeb2d..8486beced 100644 --- a/python/CuTeDSL/requirements-cu13.txt +++ b/python/CuTeDSL/requirements-cu13.txt @@ -1,3 +1,3 @@ # Use `pip install -r requirements-cu13.txt` with the present file to install a # wheel consistent with the present state of the github repository -nvidia-cutlass-dsl[cu13]==4.5.1 +nvidia-cutlass-dsl[cu13]==4.5.2 diff --git a/python/CuTeDSL/requirements.txt b/python/CuTeDSL/requirements.txt index 7e73af326..83175d1de 100644 --- a/python/CuTeDSL/requirements.txt +++ b/python/CuTeDSL/requirements.txt @@ -1,3 +1,3 @@ # Use `pip install -r requirements.txt` with the present file to install a # wheel consistent with the present state of the github repository -nvidia-cutlass-dsl==4.5.1 +nvidia-cutlass-dsl==4.5.2 diff --git a/python/cutlass_cppgen/__init__.py b/python/cutlass_cppgen/__init__.py index 93c2d0fd2..2928b3b71 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.5.1' +this.__version__ = '4.5.2' from cutlass_cppgen.backend import create_memory_pool from cutlass_cppgen.emit.pytorch import pytorch diff --git a/python/cutlass_library/generator.py b/python/cutlass_library/generator.py index f5c9df283..37665e906 100644 --- a/python/cutlass_library/generator.py +++ b/python/cutlass_library/generator.py @@ -7057,8 +7057,19 @@ def GenerateSM100_TensorOp_16b_UMMA_gemm(manifest, cuda_version, gemm_kind=GemmK for layout in layouts: layout[2][1] = 128 // DataTypeSize[data_types_mixed[0]["d_type"]] - CreateGemmUniversal3xOperator(manifest, layouts, tile_descriptions, data_types_mixed, - [[kernel_schedule, epi_schedule]], tile_schedulers=tile_schedulers, gemm_kind=gemm_kind) + # SM100 2SM TMA epilogue with 16-bit output requires CtaN divisible by 64 when CtaN > 128. + # N=160 and N=224 fail the upcast<64> stride-divisibility check in the SMEM swizzle + # layout. Skip tile shapes where CtaN is not 64-aligned. + if DataTypeSize[data_types_mixed[0]["d_type"]] == 16: + tile_descs_for_mixed = [ + td for td in tile_descriptions if td.threadblock_shape[1] <= 128 or td.threadblock_shape[1] % 64 == 0 + ] + else: + tile_descs_for_mixed = tile_descriptions + + if tile_descs_for_mixed: + CreateGemmUniversal3xOperator(manifest, layouts, tile_descs_for_mixed, data_types_mixed, + [[kernel_schedule, epi_schedule]], tile_schedulers=tile_schedulers, gemm_kind=gemm_kind) def GenerateSM100_TensorOp_16b_UMMA_alignx_gemm(manifest, cuda_version, gemm_kind=GemmKind.Universal3x): if not CudaToolkitVersionSatisfies(cuda_version, 12, 8): @@ -7448,6 +7459,18 @@ def GenerateSM100_TensorOp_fp8_UMMA_gemm(manifest, cuda_version, gemm_kind=GemmK ( data_type["d_type"] == DataType.e5m2 ): continue + # SM100 2SM TMA epilogue with 16-bit output requires CtaN divisible by 64 when CtaN > 128. + # N=160 and N=224 fail the upcast<64> stride-divisibility check in the SMEM swizzle + # layout. Skip tile shapes where CtaN is not 64-aligned. + if DataTypeSize[data_type["d_type"]] == 16: + tile_descs_for_dtype = [ + td for td in tile_descriptions if td.threadblock_shape[1] <= 128 or td.threadblock_shape[1] % 64 == 0 + ] + else: + tile_descs_for_dtype = tile_descriptions + if not tile_descs_for_dtype: + continue + if grouped: epi_schedule = EpilogueScheduleType.PtrArrayTmaWarpSpecialized2Sm elif math_inst.instruction_shape[0] == 128: @@ -7456,7 +7479,7 @@ def GenerateSM100_TensorOp_fp8_UMMA_gemm(manifest, cuda_version, gemm_kind=GemmK epi_schedule = EpilogueScheduleType.ScheduleAuto kernel_schedule = to_grouped_schedule(KernelScheduleType.TmaWarpSpecialized2SmSm100, grouped) - CreateGemmUniversal3xOperator(manifest, layouts, tile_descriptions, data_type, + CreateGemmUniversal3xOperator(manifest, layouts, tile_descs_for_dtype, data_type, [[kernel_schedule, epi_schedule]], tile_schedulers=tile_schedulers, gemm_kind=gemm_kind) def GenerateSM100_TensorOp_fp8_UMMA_alignx_gemm(manifest, cuda_version, gemm_kind=GemmKind.Universal3x): diff --git a/python/setup_cutlass.py b/python/setup_cutlass.py index e78aaa1ca..31fa5b0f7 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.5.1', + version='4.5.2', description='CUTLASS Pythonic Interface', package_dir={'': '.'}, packages=[ diff --git a/python/setup_library.py b/python/setup_library.py index 229e7f145..534723ca2 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.5.1', + version='4.5.2', description='CUTLASS library generation scripts', packages=['cutlass_library'] ) diff --git a/python/setup_pycute.py b/python/setup_pycute.py index 21ee8e059..de04dd92c 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.5.1', + version='4.5.2', description='Python implementation of CuTe', packages=['pycute'], )