Release Notes#

This section summarizes significant changes, new features, performance improvements, and known issues. Unless otherwise noted, listed issues should not impact functionality. When functionality is affected, a workaround is provided if available.

0.7.1#

Patch release mainly fixing multi-stage TMEM accumulators bug and TRSM NVRTC compatibility.

Resolved Issues#

  • Fixed a bug affecting multi-atom, multi-stage TMEM accumulators (for example M = 256), where the accumulator was not sized for the full multi-M footprint across stages.

  • Fixed a TRSM incompatibility under NVRTC and added a TRSM NVRTC example.

  • Added missing documentation for the reduce_and_store() functionality.

0.7.0#

Release updating the pipelined GEMM API and adding support for flexible-precision floating-point emulation.

New Features#

  • Pipelined GEMM creation now reports recoverable failures through cublasdx::pipeline_error instead of requiring users to rely on assertions. pipeline_error::code identifies cuBLASDx validation failures or cuda for CUDA-only failures, while pipeline_error::get_cuda_error() reports CUDA runtime failures from allocation or preprocessing paths.

  • cuBLASDx now supports flexible-precision floating-point emulation through the RequiredMantissaBits<> operator. Existing FP64 descriptors can enable emulation with no kernel code changes by adding RequiredMantissaBits<> to the descriptor. Even without loss of precision this technique can provide speedups up to 2x on B200 and 16x on RTX Pro SE Blackwell, with further speedups available by decreasing the mantissa requirement.

  • RequiredMantissaBits<> descriptors can use the pipelined GEMM path with either cublasdx::internal_accumulator or cublasdx::reusable_accumulator. The emulation pipeline manages temporary device storage on the host side and passes a regular device pipeline handle to kernels. This storage requires device memory allocations and deallocations on the order of the combined size of the A and B input matrices.

  • Several Blackwell optimizations were added, including 3-way warp specialization and TMEM multi-staging. These optimizations allow efficient persistent kernel construction with staged asynchronous SMEM_LOAD and MMA work before the single epilogue stage for each output tile.

  • Added CUBLASDX_SKIP_IF_NOT_APPLICABLE_SM() macro to skip kernel compilation for unintended device architectures, reducing compilation time and the final binary size. See SM for details.

Breaking Changes#

  • suggest_device_pipeline was renamed to suggest_pipeline.

  • make_device_pipeline was renamed to make_pipeline.

  • The pipeline result-storage option names changed from internal_accumulation and external_accumulation to internal_accumulator and reusable_accumulator. Use reusable_accumulator when user code owns the accumulator.

  • Pipeline creation now returns cublasdx::detail::expected<host_pipeline, pipeline_error>. The host pipeline owns any host-side resources, and kernels should receive pipeline->get_device_handle() instead of pipeline.value().

  • If emulation preprocessing fails after cuBLASDx validation succeeds, pipeline_error::code is pipeline_error_code::cuda and the CUDA status is available through pipeline_error::get_cuda_error().

Updated Usage Patterns#

  • The recommended pipeline creation idiom is now:

    auto pipeline = cublasdx::suggest_pipeline<BLAS>(global_a, global_b);
if(not pipeline) {
    auto const& error = pipeline.error();
    // Handle error.code and error.get_cuda_error().
}

auto shared_memory_size = cublasdx::make_shared_storage_calculator()
                              .add(pipeline->buffer_alignment(), pipeline->buffer_size())
                              .get();

kernel<<<grid, pipeline->get_block_dim(), shared_memory_size>>>(
    pipeline->get_device_handle(), ...);

  • When using the manual reusable_accumulator flow, call tile_pipeline.finish_accumulation() after tile_pipeline.execute(accumulator) and before reading accumulator results. Call it with the same thread scope as tile_pipeline.execute(accumulator); do not guard it with if(accumulator.is_thread_active()) because internal warp-specialization may require broader participation. The tile_pipeline.epilogue(accumulator, functor) convenience path calls finish_accumulation() internally.

  • Descriptors using WithPipeline() must retrieve accumulators from the per-tile pipeline with tile_pipeline.get_accumulator(). Descriptor-level accumulator retrieval is only valid for regular, non-pipelined tile execution.

0.6.0#

Major release adding triangular solve (TRSM), raising the minimum required CUDA Toolkit to 13.0, and removing Volta GPU support.

New Features#

  • Triangular solve (function::TRSM): block-level, in-place solver for \(A \times X = B\) or \(X \times A = B\), operating on shared memory.

    • New operators required to describe a TRSM problem:

      • Side — whether the triangular matrix appears on the left or right.

      • FillMode — which triangle (upper or lower) of A holds the data.

      • Diag — whether the diagonal of A is explicitly stored (non_unit) or implicitly ones (unit).

      • BatchesPerBlock — number of independent TRSM instances solved per CUDA block.

    • New TRSM traits: side_of, fill_mode_of, diag_of.

    • TRSM requires linking against the pre-built libcublasdx.fatbin device-code library and separable CUDA compilation. See the Quick Installation Guide for exact CMake setup.

    • New examples 17_trsm/trsm_block and 17_trsm/trsm_thread demonstrate block-level and thread-level TRSM usage.

    • New example 18_tensor_transform/trsm_conj_transpose demonstrates how to use new transform tensor views with TRSM.

    • New examples 19_gemm_batched/batched_gemm_pipeline and 19_gemm_batched/batched_gemm demonstrate how to use GEMM and Pipelined GEMM for batching.

Breaking Changes#

  • Default CMake target now links the libcublasdx.fatbin device-code library. The mathdx::cublasdx target now links the fatbin automatically on CUDA Toolkit ≥ 13.2, enabling both GEMM and TRSM. This means existing projects using mathdx::cublasdx must now set CUDA_SEPARABLE_COMPILATION ON and INTERPROCEDURAL_OPTIMIZATION ON on consuming targets. Projects that use only GEMM and want to avoid this overhead should switch to the new mathdx::cublasdx_no_lto target, which is header-only and produces a static_assert if TRSM is accidentally used. For CUDA Toolkit < 13.2, use mathdx::cublasdx_fatbin which explicitly links libcublasdx.fatbin. See the Quick Installation Guide for details.

  • CUDA Toolkit 12 support removed. The minimum required version is now CUDA Toolkit 13.0.

  • Volta GPU support removed. SM70 (sm_70) and SM72 (sm_72) are no longer supported.

  • Accumulator precision restricted. The C precision in Precision<PA, PB, PC> is now limited to __half, float, double, int32_t, uint32_t, int64_t, uint64_t. Using __nv_bfloat16, cublasdx::tfloat32_t, __nv_fp8_e4m3, __nv_fp8_e5m2, int8_t, uint8_t, int16_t, or uint16_t as the accumulator now produces a static_assert.

Resolved Issues#

  • SM120 now correctly emits TMA instructions where applicable.

  • SM103 (Blackwell Ultra) integral UTCMMA removed. Due to a bug cuBLASDx emitted UTCIMMA instructions for SM103, where they are not available.

  • Fixed default pipeline constructor producing incorrect initial state.

  • Fixed misalignment issues in tensor copy operations.

0.5.1#

Patch release amending heuristics for Ozaki emulation example as well as adding missing operators to performance focused examples (11 and 12)

0.5.0#

Major release adding experimental pipelining API, including support for Hopper and Blackwell features such as TMA, WGMMA and 1SM UTCMMA.

New Features#

  • Experimental Pipelining API supporting fusable asynchronous execution starting from global memory

  • Initial support for WGMMA, 1SM UTCMMA, TMA features

  • Updated Ozaki example, providing significant performance upgrade targeted for Blackwell B200 GPU

  • New accumulator API, opaquely handling fragment operations

  • CUDA 13.1 support

Breaking Changes#

  • Suggested layouts can no longer be mixed with non-suggested fragments, and the other way

  • Partitioner API removed, now merged into Accumulator API

  • StaticBlockDim operator moved from cublasdx::experimental to cublasdx namespace

  • Alignment operator now is also understood to describe alignment of dynamic leading dimensions (if used)

0.4.1#

Patch release adding support for CUDA 13

New Features#

  • Support for CUDA 13.0

  • Deprecation of CUDA 11.X

  • Support for Thor SM renaming from sm_101 to sm_110 starting from CUDA 13.0 * Note: For CUDA 12.9 and older releases, Thor stays labeled as sm_101

0.4.0#

Release introducing wider support for suggested layouts and basic Blackwell support.

New Features#

  • Suggested layouts support wider range of problem sizes, threadblocks and instructions introducing analytical swizzling heuristics

  • Both suggested and non-suggested layouts now generate ld.matrix and st.matrix for either 1 (x1), 2 (x2), or 4 (x4) matrices

  • Support for PTX 8.7 superMMA instructions and fma.f32x2 fused multiply-add instructions.

  • SM100, SM101, SM120 support

  • SM103 and SM121 experimental support

Breaking Changes#

  • The entire shared memory slicing API has been refactored and generalized to support all Dx libraries.

  • SM72 has been deprecated

Known Issues#

CUDA 12.8.0, 12.8.1 and 12.9.0 have been known to miscompile cuBLASDx 0.3.1 and 0.4.X code with high register pressure when

  • Any of the computation types is fp8_e5m2, fp8_e4m3, fp16, bf16, or int8

  • Any of M, N or K (Size Operator) is not a multiple of 16

  • Custom static leading dimension is used (LeadingDimension Operator)

These code corruptions may manifest as either incorrect results or illegal memory access errors.

To highlight this issue cuBLASDx 0.4.X will overprotectively hard fail if these conditions are met.

If you are using cuBLASDx, and this happens to your code, you can:

  • Update to the latest CUDA Toolkit and NVCC compiler (12.9.1 is known to work)

  • define the CUBLASDX_IGNORE_NVBUG_5218000_ASSERT macro to ignore these assertions and verify correctness of the results.

  • if the case is indeed affected by the bug, adding the -Xptxas -O1 flag to the compilation command will limit PTX optimization phase and produce correct binary, although potentially slower.

0.3.1#

Minor release with quality of life changes and minor heuristic improvements.

New Features#

  • Improved usability, readability and reference in device_gemm_performance example.

  • Improved GEMM heuristics for non-suggested executions.

  • Improved safety of mixing suggested and non-suggested layouts and accumulators in execution.

0.3.0#

The 3rd early access (EA) release of cuBLASDx library brings support for selected integer types, matrix multiplication with results stored in registers, and decoupling of computation types from input/output types.

New Features#

  • Register fragment tensor support:

    • Added partitioning, predication and transformation tools.

    • Added copying and partitioning utilities for moving data from / into register file buffers.

  • New GEMM Register APIs allowing for better performance.

    • Using registers to store C matrix (result matrix) allows to save on shared memory size and transfers by performing accumulation and input / output in registers.

  • Integral types support, including MMA instruction support for integral types.

  • Compute precision decoupled from input precision.

    • A GEMM function can now accept any input and convert it in registers, saving on memory usage.

  • More robust shared memory management tools have been added, enabling construction of efficient pipelined execution.

  • The library no longer statically asserts on whether the GEMM problem will fit in shared memory, due to:

    • A and B can alias or overlap themselves,

    • C can be entirely in register file, and

    • input precision can be arbitrary, and is not defined by compute precision.

Breaking Changes#

  • Shared memory slicing and shared memory size utilities are no longer available as BLAS methods.

  • Shared memory size trait has been removed from BLAS type.

  • is_supported trait has been removed, and is_supported_smem_restrict, is_supported_rmem_restrict added in its place.

  • The library no longer asserts on whether a size will fit in shared memory, it’s the user’s responsibility now.

Resolved Issues#

  • cuBLASDx internal versioning definitions have been fixed: * CUBLASDX_VERSION_MAJOR is now defined as a multiple of 10000 in CUBLASDX_VERSION instead of 1000, i.e. CUBLASDX_VERSION_MAJOR = CUBLASDX_VERSION / 10000. * CUBLASDX_VERSION_MINOR now can have a new maximum of two digits. * Since there was no major release there is no extra gap between this and previous version. * All definitions can be checked in the file cublasdx_version.hpp.

  • Missing traits have been added to follow C++ metaprogramming convention

Known Issues#

  • It’s recommended to use the latest CUDA Toolkit and NVCC compiler.

    • CUDA 12.4 has known edge cases producing incorrect FP8 MMA emulation code on SM90 with register APIs.

    • CUDA 12.1 has known edge cases of crashing when passing previously named values into 3-value operators.

      • e.g. Alignment<varn_name_1, var_name_2, var_name_3> may cause compilation hang, while Alignment<8, 8, 8> will always work.

0.2.0#

The second early access (EA) release of cuBLASDx library brings tensor API, mixed precision, and performance improvements.

New Features#

Known Issues#

  • It’s recommended to use the latest CUDA Toolkit and NVCC compiler.

0.1.0#

The first early access (EA) release of cuBLASDx library.

New Features#

  • Support for general matrix multiply.

    • Tensor cores support for fp16, fp64, complex fp64 calculations.

  • Support for SM70 - SM90 CUDA architectures.

  • Multiple examples included.

Known Issues#

  • Since CUDA Toolkit 12.2, the NVCC compiler in certain situations reports an incorrect compilation error when the value_type type of a GEMM description type is used. The problematic code with possible workarounds is presented below:

    // Any GEMM description type
using GEMM = decltype(Size<32 /* M */, 32 /* N */, 32 /* K */>()
            + Precision<double>()
            + Type<type::real>()
            + TransposeMode<t_mode /* A */, t_mode /* B */>()
            + Function<function::MM>()
            + SM<890>()
            + Block());
using type = typename GEMM::value_type; // compilation error

// Workaround #1
using type = typename decltype(GEMM())::value_type;

// Workaround #2 (used in cuBLASDx examples)
template <typename T>
using value_type_t = typename T::value_type;

using type = value_type_t<GEMM>;