2.3.1. CUDA Platform

• Added Event Tracing for Windows (ETW) support for CUDA driver activity reporting.

  ETW is a high-performance, low-overhead logging system built into the Windows operating system. This support enables CUDA driver activity to be reported through ETW-based diagnostics for debugging, performance monitoring, and analysis.

• Added mmap() support for DMA-BUF file descriptors exported from CUDA device memory on discrete GPUs.

  This support extends the GPU driver’s existing DMA-BUF mmap() support from Tegra system memory cases to discrete GPU video memory. It provides a low-latency CPU mapping of discrete GPU memory in environments where installing GDRCopy kernel drivers might not be desirable.

  To use this support:

  • Allocate CUDA device memory.

  • Export the desired address range as a DMA-BUF file descriptor through the CUDA DMA-BUF export path, for example by using the following driver API:

    cuMemGetHandleForAddressRange(..., CU_MEM_RANGE_HANDLE_TYPE_DMA_BUF_FD, ...)
    

  • Check whether CU_DEVICE_ATTRIBUTE_DMA_BUF_MMAP_SUPPORTED is supported.

  • Call Linux mmap() on the exported DMA-BUF file descriptor.

  If mmap() fails with ENOTSUPP, the driver and kernel combination does not support this path.

  After the mapping is created, applications can use the CPU pointer for low-latency access to GPU memory from the CPU. Applications must follow the standard DMA-BUF CPU access protocol, including SYNC_START and SYNC_END.

  Manual synchronization or fencing might be required when the same buffer is also used by the GPU, I/O devices, persistent nv-p2p or GDRCopy mappings, or P2PDMA.

• Added the cuStreamBeginRecaptureToGraph() API, which allows applications to initiate stream capture into an existing source graph.

  As the graph is recaptured, updated node parameters are applied to the existing nodes. The nodes must be recaptured in the same order as the original source graph. All recaptures to an existing source graph check for a topological match and fail if the recaptured topology diverges from the original graph.

• Updated Green Contexts so that creation of the default, or NULL, stream is no longer required through the CU_GREEN_CTX_DEFAULT_STREAM flag.

Creating the default stream for a Green Context is now optional.

• Added partial support for error isolation when using MPS.

  This feature addresses a long-standing MPS limitation where some fatal GPU errors could only be contained at a broad affected-device scope, causing collateral termination of clients that did not cause the fault. This support is intended for deployments that rely on MPS for high GPU utilization while also requiring stronger service isolation, such as mixed online and offline workloads, many small database or query workloads, or embedded and robotics pipelines with independent processes.

  Partial error isolation relies on exclusive SM ownership. The MPS control daemon creates SM partitions, the MPS server assigns clients to those partitions, and the client enforces the partition through the launch affinity programmed for its work. Because an SM is owned by only one partition, CUDA can attribute SM error state to the faulting partition or client. When a fault is detected, the CUDA driver terminates work for the faulting client and prevents additional work from being submitted. Faulting or terminated operations may report CUDA_ERROR_LAUNCH_FAILED or a more specific CUDA error.

  Clients in different SM partitions are isolated from each other’s SM-triggered kernel faults. One or more MPS clients may intentionally share the same partition, in which case those clients also share the same fault domain. This feature provides partial isolation and does not guarantee that every possible GPU or system-level failure is isolated per process.

  To use partial error isolation with MPS:

  • Start MPS in static-partition mode:

    nvidia-cuda-mps-control -d -S
    

    Alternatively, use:

    nvidia-cuda-mps-control -d -static-partition
    

  • Create one or more SM partitions:

    echo "sm_partition add <device UUID> <number of chunks>" | nvidia-cuda-mps-control
    

    The command returns the full partition ID.

  • Inspect the configured partitions:

    echo "lspart" | nvidia-cuda-mps-control
    

    To view the client PID, partition, and SM count, use:

    echo "ps -f" | nvidia-cuda-mps-control
    

  • Assign a client process to a partition:

    CUDA_MPS_SM_PARTITION=<device UUID>/<partition ID> ./application
    

  • Remove an idle partition:

    echo "sm_partition rm <device UUID> <partition ID>" | nvidia-cuda-mps-control
    

  When static-partition mode is enabled, a client that starts without CUDA_MPS_SM_PARTITION, uses an invalid partition ID, or tries to assign more than one partition from the same device fails context creation with CUDA_ERROR_INVALID_RESOURCE_CONFIGURATION. Removing a partition fails while clients are still actively using it. Partition IDs are deterministic for the same command sequence, partition sizes, and initial state on matching GPU SKUs.

  Partitions are requested in chunks rather than percentages. For dGPU, a chunk is 4 SMs on Ampere-class GPUs and 8 SMs on Hopper and newer GPUs. For iGPU, the chunk size is 2 SMs to better fit lower-SM-count devices.

  This mode trades some MPS flexibility for isolation. SMs reserved for a partition remain exclusive to the clients assigned to that partition, and other clients cannot automatically borrow idle SMs from it. This mode is best suited for workloads with known resource needs, workloads that can tolerate strict resource limits, or deployments where fault containment is more important than opportunistic sharing.

• Added the nvmlDeviceGetRemappedRows_v2 NVML API. In addition to the information returned by nvmlDeviceGetRemappedRows, nvmlDeviceGetRemappedRows_v2 returns the number of inactive row remappings.

• Added support for using both coherent-memory and non-coherent-memory GPUs in the same process on NVIDIA DGX Station systems.

• The accompanying r610 driver branch improves observability and error classification for Blackwell systems.

  Previously, certain error conditions were reported as XID 119 timeouts, followed by XID 94, and then XID 154 to indicate that the GPU required a reset. This sequence could be misleading because XID 94 can imply error containment, while in this scenario the GPU is not usable without a reset.

  With the updated behavior, the driver logs XID 140 for the unrecoverable ECC poison condition, skips the generic XID 119 timeout report, logs XID 95 for an uncontained robust-channel error from the fatal GSP poison interrupt path, and marks the GPU for reset with XID 154.

  Monitoring rules should treat XID 140 or XID 95 followed by XID 154 as a GSP memory poison and reset-required condition instead of as a general GSP timeout.

• Added Dynamic Boost support for Grace Blackwell systems.

  Note

  This feature is available on Grace Blackwell systems and requires the installation of the NVIDIA r610 driver, corresponding to CUDA Toolkit 13.3, or later.

  Dynamic Boost is a Grace CPU frequency optimization feature that runs as a background service. It uses machine learning to model performance and power impact, and balances them based on workload demands. It continuously monitors the system and adjusts the CPU clock speed in real time to match the needs of the running application.

  When the workload enters phases that are less CPU-bound, Dynamic Boost temporarily lowers the CPU frequency to reduce power consumption. The saved power can then be reallocated to the GPU for improved performance. When the workload becomes more CPU-bound, Dynamic Boost can increase the CPU frequency accordingly.

  Dynamic Boost is available on GB200 and GB300 systems. The service can be managed using standard systemctl commands:

  • Enable the service:

    systemctl enable nvidia-powerd.service
    

  • Start the service:

    systemctl start nvidia-powerd.service
    

  • Stop the service:

    systemctl stop nvidia-powerd.service
    

  • Disable the service:

    systemctl disable nvidia-powerd.service
    

• Added support for remote CPU-to-GPU mapping of managed memory.

  CDMM-managed memory that is resident on the CPU can now be mapped directly onto a GPU, allowing the GPU to access the memory without an explicit migration. This capability is reported based on GPU and platform support. It is enabled when cuDeviceGetAttribute() reports CU_DEVICE_ATTRIBUTE_DIRECT_MANAGED_MEM_ACCESS_FROM_HOST = 1 for the device.

• Added System-Allocated Memory (SAM) migration support for CDMM.

  CDMM can now migrate system-allocated memory, such as memory allocated with malloc(), between the CPU and GPU. This support requires a Linux kernel that includes the necessary HMM fixes. Kernel 6.12 or later is recommended. On older kernels, SAM migration remains disabled and existing behavior is preserved.

CUDA Compiler

• For new features from PTX, refer to PTX ISA version 9.3.

• Introduced CUDA Tile C++, which adds support for tile programming in CUDA C++. This feature is supported in both nvcc and NVRTC. For more information about writing tile kernels, see the CUDA Programming Guide and the CUDA Tile C++ API Reference.

  For CUDA Tile C++ API-specific release information, see the CUDA Tile C++ API Reference release notes.

• Added NVRTC support for installing a curated set of CUDA C++ and CCCL headers directly from the NVRTC distribution.

  This support enables runtime compilation of code that uses modern CUDA facilities without requiring a full CUDA Toolkit installation or manual header management.

• Added support for Advanced Control Files (ACFs) in NVIDIA compiler toolchains through the --apply-controls=<file> option.

  This option enables compiler binaries to process encrypted control files as part of the build workflow. Documentation updates, including usage examples for --apply-controls, are available in the CUDA Programming Guide and the NVIDIA CUDA Compiler Driver, NVCC, documentation.

  ACFs are generated using NVIDIA CompileIQ. For more information, see CompileIQ on GitHub.

  Caveats:

  • This feature is not supported in NVRTC.

  • PTXAS support is available only in offline compilation flows and is not supported when using static libraries. For more information, refer to the PTXAS help documentation.

  • Correct compilation behavior is not guaranteed when using ACFs. For more information, see CompileIQ on GitHub.

• Added official C++23 support in nvcc and NVRTC.

  This support enables developers to adopt the latest C++ language features across both ahead-of-time and runtime CUDA compilation workflows. It improves code portability, modernizes GPU application development, and helps align CUDA codebases with contemporary C++ standards.

• Added nvprune functionality to nvcc.

  This support helps developers streamline deployment artifacts and manage builds for targeted GPU architectures without requiring a separate pruning step to remove unnecessary code from the libraries they use.

2.3.2. CUDA Developer Tools

For details on new features, improvements, and bug fixes, see the changelogs for:

• Nsight Systems.

• Nsight Visual Studio Edition.

• CUPTI.

• Nsight Compute.

• Compute Sanitizer.

• NVML.

2.3.3. CUDA C++ Core Libraries (CCCL)

• Added tensor interoperability support in libcu++.

  The cuda::to_device_mdspan(), cuda::to_host_mdspan(), and cuda::to_managed_mdspan() APIs convert a DLPack DLTensor into a typed cuda::std::mdspan view that CUDA kernels can operate on with shape and stride metadata. DLPack is an interchange format used by Python frameworks such as PyTorch, JAX, and CuPy.

  The cuda::to_dlpack_tensor() API converts an mdspan to a DLManagedTensor wrapper.

  The cuda::shared_memory_mdspan API provides a multidimensional view over a CUDA shared memory tile. The accessor guarantees shared memory load and store instructions and adds address space safety checks.

• Added expanded random number distribution support in libcu++.

  The <cuda/std/random> header now includes 17 host- and device-compatible random number distributions, bringing libcu++ closer to parity with the C++ standard library <random> header.

  Supported distributions include:

  • Uniform distributions:

    • cuda::std::uniform_int_distribution

    • cuda::std::uniform_real_distribution

  • Normal-family distributions:

    • cuda::std::normal_distribution

    • cuda::std::lognormal_distribution

  • Discrete distributions:

    • cuda::std::bernoulli_distribution

    • cuda::std::binomial_distribution

    • cuda::std::negative_binomial_distribution

    • cuda::std::geometric_distribution

    • cuda::std::poisson_distribution

  • Continuous distributions:

    • cuda::std::exponential_distribution

    • cuda::std::gamma_distribution

    • cuda::std::weibull_distribution

    • cuda::std::extreme_value_distribution

    • cuda::std::cauchy_distribution

    • cuda::std::chi_squared_distribution

    • cuda::std::fisher_f_distribution

    • cuda::std::student_t_distribution

  CCCL 3.3 also backports the C++26 counter-based engines cuda::std::philox4x32 and cuda::std::philox4x64 to C++17. The <cuda/random> header adds cuda::pcg64 as an NVIDIA extension.

• Added new CUB device-wide algorithms.

  New and updated algorithms include:

  • cub::DeviceFind::FindIf: Improves performance for device-wide searches for the first element that satisfies a predicate.

  • cub::DeviceFind::LowerBound and cub::DeviceFind::UpperBound: Add parallel binary search support to locate multiple values in an ordered sequence in a single device-wide call.

  • cub::DeviceSegmentedScan: Adds segmented prefix-sum and scan support over multiple independent contiguous segments in a single device-wide call.

  • cub::DeviceTransform::Transform: Adds an N-input, M-output form. A single call can read from N input sequences and write to M output sequences, driven by a user-provided operation that returns a cuda::std::tuple<...>.

  • cub::DeviceTopK and cub::DeviceSegmentedTopK: Add device-wide top-k selection support, including a fixed-size segments variant.

• For the complete list of CCCL 3.3.0 changes, see the CCCL 3.3.0 changelog.

2.3.4. CUDA Python

• First stable release of cuda.core.

  As of version 1.0.0, all APIs are considered stable and follow Semantic Versioning (SemVer), with appropriate deprecation periods for breaking changes. See the support policy for details.

• Added green context support for CUDA 12.4 and later.

  New types Context, ContextOptions, SMResource, SMResourceOptions, WorkqueueResource, and WorkqueueResourceOptions enable GPU SM and workqueue resource partitioning. Green contexts can be created with Device.create_context(). Use Context.create_stream() and Context.resources to work within the partitioned resources. #1976

• Added the cuda.core.checkpoint module for CUDA process checkpointing.

  The module includes string process state queries, lock, checkpoint, restore, and unlock operations, and GPU UUID remapping support for restore. #1343

• Added an optional cache= keyword argument to Program.compile().

  This argument helps avoid recompilation of identical source, options, and target combinations. Two concrete implementations of the ProgramCacheResource abstract base class are provided:

  • InMemoryProgramCache, a thread-safe, single-process LRU cache.

  • FileStreamProgramCache, a disk-backed, cross-process safe, LRU-evicting cache.

  A standalone make_program_cache_key() function is also exposed for callers who need to include additional content, such as headers or PCH files, in the cache key. #1912

• Added NVIDIA Management Library (NVML) access updates to the cuda.core.system module.

  The following APIs are now available:

  • system.Device.mig for querying and setting MIG mode, enumerating MIG device instances, and navigating parent and child relationships. #1916

  • system.Device.compute_running_processes for querying running compute processes on a device. This API returns ProcessInfo objects with PID, GPU memory usage, and MIG instance IDs. #1917

  • system.Device.get_nvlink() for querying NVLink version and state per link. system.Device.utilization now returns current GPU and memory utilization rates. #1918

  • cuda.core.typing.VirtualMemoryLocationType.

• For the full cuda.core 1.0 changelog, see cuda.core 1.0.0 release notes.

• First stable release of cuda.compute.

  As of version 1.0.0, all APIs are considered stable and follow Semantic Versioning (SemVer), with appropriate deprecation periods for breaking changes. See the support policy for details.

  • Added broad algorithm coverage to cuda.compute.

    Supported algorithms include reduce_into, inclusive_scan, exclusive_scan, unary_transform, binary_transform, lower_bound, upper_bound, segmented_reduce, merge_sort, radix_sort, histogram, select, three_way_partition, and unique_by_key.

  • Added support for Python callables, including plain lambda expressions, as device-wide operators across reductions, scans, and transforms.

    The callable is JIT-compiled to LTO-IR through Numba and linked into the underlying CUB kernel at runtime. Stateful operators that capture mutable globals or device-side state are also supported.

  • Added cuda.compute.lower_bound and cuda.compute.upper_bound for parallel binary search.

    These APIs provide device-wide functionality similar to numpy.searchsorted. Both APIs return insertion indices and accept any iterator type.

  • Added ShuffleIterator for deterministic pseudorandom permutation of an input range.

    ShuffleIterator complements the existing iterators: CountingIterator, ConstantIterator, TransformIterator, ZipIterator, and PermutationIterator.

  • For the full cuda.compute 1.0 changelog, see cuda.compute 1.0 release notes.

2.3.5. CUDA TILE

2.4.1. CUDA Compiler

• Fixed an issue where applications that use NVRTC and are linked with LLVM lld or mold could fail to initialize the supported architecture list. This issue could cause nvrtcGetNumSupportedArchs to return 0 and nvrtcGetSupportedArchs to fail with NVRTC_ERROR_INVALID_INPUT. [5020829]

• Fixed a compiler issue, present since CUDA 12.8, that could cause compiler-inserted thread reconvergence to fail and leave stale or corrupted values in registers, resulting in incorrect program execution. This issue could occur only in kernels that contain two or more nested levels of thread divergence where the compiler elided convergence instructions for one or more divergence levels. Kernels with only a single level of divergence are unaffected.[6156910]

  Minimal illustrative example:

  __global__voidaffected(constint*in,int*out){
  inttid=threadIdx.x;
  
  if(tid<16){// Level 1 divergence
  if(in[tid]>0){// Level 2 divergence, nested
  // If the compiler elides reconvergence for one of these
  // divergence levels, a register written in this region
  // can carry a stale value across the reconvergence point.
  out[tid]=compute(in[tid]);
  }
  
  // Implicit reconvergence: failure can manifest here.
  }
  }
  

2.4.2. CUDA Tools

• Fixed Data Race in WGMMA A/B Register Copy Propagation

  Symptom:

  Kernels that use wgmma.mma_async with wgmma.wait_group.sync.aligned N, where N >= 1, might produce incorrect numerical results when mov instructions are intentionally placed after the wait to protect WGMMA input registers from being overwritten by in-flight WGMMA operations.

  Problem Description:

  In pipelined warp-group MMA loops, mov.b32 instructions are sometimes intentionally placed after wgmma.wait_group.sync.aligned N, where N > 0, to ensure that WGMMA input registers are not overwritten while WGMMA operations from a previous iteration are still in flight. In this specific case, ptxas could incorrectly copy-propagate across wgmma.wait_group.sync.aligned N and eliminate those mov instructions.

  For example, in the following pattern, ldmatrix.x4 feeds two consecutive wgmma calls. The compiler could incorrectly eliminate the second mov.b32 instruction, exposing the race:

  $loop:
   ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r71, %r91, %r81, %r101}, [%r8];
  
   mov.b32 %r70, %r71;
   wgmma.fence.sync.aligned;
   wgmma.mma_async { }, {%r70, ..}, ...;
   wgmma.commit_group.sync.aligned;
   wgmma.wait_group.sync.aligned 1;
  
   mov.b32 %r80, %r81; // compiler eliminates this mov
   wgmma.fence.sync.aligned;
   wgmma.mma_async { }, {%r80, ..}, ...; // uses %r81 directly after copy propagation
   wgmma.commit_group.sync.aligned;
   wgmma.wait_group.sync.aligned 1;
  
   bra $loop;
  
  wgmma.wait_group.sync.aligned 0;
  

  Workaround: For CUDA Toolkits prior to 13.3, use one of the following workarounds:

  • Use wgmma.wait_group.sync.aligned 0 to wait for all in-flight groups before the next ldmatrix instruction.

  • Use individual ldmatrix.x1 loads directly into WGMMA input registers, eliminating the mov instruction that triggers copy propagation.

2.5.1. General CUDA

• Compute Fabric Transport

  • Enabling support for unicast logical endpoints requires Fabric Manager 610 or later.

  • The performance of the fabric.try_pullred instruction is suboptimal. This issue will be addressed in a future CUDA release.

  • If a memory allocation is bound to a logical endpoint using cuLogicalEndpointBindMem or cuLogicalEndpointBindAddr, pointer access to that allocation from multiple GPUs within the same process might not work correctly. This issue will be addressed in a future CUDA release.

2.5.2. CUDA Compiler

• CUDA Tile C++

  The Tile C++ compiler may incorrectly issue a diagnostic when printf() is invoked from a tile function if GLIBC fortification level 2 or higher is enabled by the host compiler. This level of GLIBC fortification is enabled by default on Ubuntu distributions of GCC when -O1 or higher is present. It can also be enabled explicitly when -O1 or higher is present and -D_FORTIFY_SOURCE=2 or higher is specified.

  The following example demonstrates the issue:

  #include<cstdio>
  
  __tile_global__voidkernel(){
  printf("hello\n");
  }
  

  Compile the example with the following command:

  nvcc --enable-tile -arch sm_80 -std=c++20 -O1 -D_FORTIFY_SOURCE=2 main.cu
  

  The compiler may report the following diagnostic:

  error: calling a __host__ __device__ function("printf") from a __tile_global__ function("kernel") is not allowed
  

  As a workaround, disable GLIBC fortification by adding -U_FORTIFY_SOURCE to the nvcc command line. A fix for this issue will be available in a future CUDA release.

• For semantic clarifications, refer to Changes in PTX ISA version 9.3.

2.6.1. CUDA Developer Tools

• Legacy Nsight Eclipse Edition plugins are no longer delivered in CUDA Toolkit packages beginning with CUDA 13.3.

2.6.2. Architectures

• None

2.6.3. Operating Systems

• None

2.6.4. CUDA Toolchains

• None

3.1.1. cuBLAS: Release 13.3

• New Features

• Enabled memory-parsimonious tiling for FP64 emulated matrix multiplications. This improvement ensures that the workspace memory budget no longer exceeds 8 GB.

• Added support for CUDA Green contexts.

• Improved FP4 matrix multiplication performance on Blackwell Ultra GPUs by a geometric mean of 5% across a wide range of problems, with up to 7% speedup for some small problems.

• Improved TF32 matrix multiplication performance on Blackwell and Blackwell Ultra GPUs by a geometric mean of 27% across a wide range of problems and layouts, with up to 3.5x speedup for some small problems.

• Improved TF32 TN matrix multiplication performance on Hopper GPUs by a geometric mean of 11% across a wide range of problems, with up to 40% speedup for some small problems.

• Improved SYMV performance with TMA-based acceleration for Hopper, Blackwell, and Blackwell Ultra kernels with up to 27% geomean speedup.

3.1.2. cuBLAS: Release 13.2 Update 1

• New Features

  • Extended the experimental Grouped GEMM API in cuBLASLt to support NVFP4 inputs and bias epilogues on Blackwell GPUs with Compute Capability 10.x and 11.0.

  • Extended the experimental Grouped GEMM API in cuBLASLt to support BF16, FP16, and FP8 input data types with BF16, FP16, and FP32 output data types on Hopper GPUs. For FP8 inputs, tensorwide scaling and block scaling (VEC128 and BLK128x128) are supported.

  • Improved Grouped GEMM performance on Blackwell GPUs, providing up to 20% higher performance for large problem sizes.

• Resolved Issues

  • Fixed an issue in cublasLtMatmulAlgoGetHeuristic() that could result in no algorithm candidates being returned for Grouped GEMM on Blackwell GPUs. [CUB-9657]

3.1.3. cuBLAS: Release 13.2

• New Features

  • Extended the experimental Grouped GEMM API in cuBLASLt to support MXFP8 inputs on GPUs with Compute Capability 10.x and 11.0.

  • Added control over special-case handling in FP32 emulation via the environment variable CUBLAS_EMULATION_SPECIAL_VALUES_SUPPORT_MASK. Setting CUBLAS_EMULATION_SPECIAL_VALUES_SUPPORT_MASK=0 can improve performance for applications that do not require preservation of infinity and NaN values, without requiring code changes. For more information, see the cudaEmulationSpecialValuesSupport_t documentation.

  • Added FP64 fixed-point emulation support to the cublas[D|Z]syrk, cublas[D|Z]syr2k, cublasZherk, and cublasZher2k routines. When the math mode is set to CUBLAS_FP64_EMULATED_FIXEDPOINT_MATH, cuBLAS will automatically use FP64 emulation for sufficiently large SYRK and HERK problems.

  • Improved performance on RTX PRO 6000 GPUs, delivering up to 20% speedup for FP8, FP16/BF16, TF32, and INT8 precisions.

  • Improved GEMM performance on DGX Spark systems for MXFP8 and NVFP4 data types in large M and N problem sizes, with up to 3× performance improvement for selected matrix shapes.

• Known Issues

  • On Blackwell GPUs, FP64 fixed-point emulation kernels may produce incorrect results or experience data corruption when executed concurrently with third-party kernels that allocate tensor memory.[CUB-9633]

• Resolved Issues

  • Fixed an issue in cublasLtMatmul that could lead to incorrect results when it ran concurrently with another kernel that uses Tensor Memory. This issue only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with Compute Capability 10.x and 11.x, and existed since cuBLAS 12.8. [5807900]

  • Fixed an issue in cublasLtMatmul that could lead to incorrect results or invalid memory access errors for large leading dimensions, specifically when the product of the data type size and the leading dimension of a matrix exceeded the bounds of a signed 32 bit integer. This issue affected GPUs with Compute Capability 9.0, 10.x, or 11.0, existed since cuBLAS 12.6 Update 2, and only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66. [CUB-9572]

  • Fixed an issue in the cuBLASLt Matmul API that could cause FP8 kernels to hang on GPUs with Compute Capability 9.0 when beta != 0 and scale_C = 0. This issue only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66. [CUB-9627]

  • Fixed an issue in the cuBLASLt Grouped GEMM API that ignored groups with k = 0, leading to incorrect results. This issue existed since CUDA 13.1. [CUB-9529]

  • Fixed an issue in the cuBLASLt Matmul API that could cause incorrect results when C broadcasting was used (LDC = 0). [5845724]

  • Added missing checks for matrix pointer alignment in the cublasLtMatmul API. [CUB-9577, CUB-9599, CUB-9585]

  • Fixed an issue in cublasLtMatmul that could lead to incorrect results for NVFP4 precision on B300 and GB300 GPUs when the m dimension was not a multiple of 64. [CUB-9577]

  • Fixed an issue in cublasLtMatmul that could lead to incorrect results for NVFP4 precision on future GPUs, impacting future hardware compatibility. [CUB-9570]

  • Fixed an issue in GEMM and Matmul APIs with BF16 and FP16 inputs on DGX Spark and FP8 inputs on GeForce that could potentially cause illegal memory accesses. [5846563]

  • Fixed an issue in cuBLASLt to enable CUBLASLT_EPILOGUE_BGRADA and CUBLASLT_EPILOGUE_BGRADB epilogues when the C matrix CUBLASLT_MATRIX_LAYOUT_ORDER was set to CUBLASLT_ORDER_ROW. [4617436]

  • Fixed an integer overflow bug in complex, emulated FP64 matrix multiplication. The affected routines include cublasZgemm, cublasZtrsm, cublasGemmEx, and cublasLtMatmul. The overflow occurred when 2*m*n + m exceeded UINT_MAX, where m is the number of rows of op(A) and C, and n is the number of columns of op(B) and C. [5720478]

  • Improved GB200 and B200 performance for MXFP8 and NVFP4 precisions when M and N were less than or equal to 32. [CUB-9646]

3.1.4. cuBLAS: Release 13.1 Update 1

• Known Issues

  • The cuBLASLt Grouped GEMM API ignores groups with k = 0, which can lead to incorrect results. As a workaround, initialize output matrices D with beta*C for all groups, and then compute Grouped GEMM as D += A*B so the result for groups with k = 0 is computed properly. This issue applies to the experimental cuBLASLt Grouped GEMM API introduced in CUDA 13.1. [CUB-9529]

  • Complex FP64 GEMM routines using fixed-point emulation can produce incorrect results when matrix dimensions are large enough that m*n > 2^31 due to integer overflow in an address calculation. [5720478]

• Resolved issues

  • Fixed an issue where fixed point emulation with 7 mantissa bits or less could trigger unspecified launch failures. [5692684]

  • Fixed an issue where cublasLtMatmul with FP8 arguments and CUBLASLT_MATMUL_MATRIX_SCALE_SCALAR_32F scaling mode (default) incorrectly required scaling factor addresses to be 16-byte aligned. This issue existed since cuBLAS 12.9. [5728938]

3.1.5. cuBLAS: Release 13.1

• New Features

  • Introduced experimental support for grouped GEMM in cuBLASLt. Users can create a matrix with grouped layout using cublasLtGroupedMatrixLayoutCreate or cublasLtGroupedMatrixLayoutInit, where matrix shapes are passed as device arrays. cublasLtMatmul now accepts matrices with grouped layout, in which case matrices are passed as a device array of pointers, where each pointer is a separate matrix that represents a group with its own shapes. Initial support covers A/B types FP8 (E4M3/E5M2), FP16, and BF16, with C/D types FP16, BF16, and FP32; column-major only, default epilogue, 16-byte alignment; requires GPUs with compute capability 10.x or 11.0.

    In addition, the following experimental features were added as part of grouped GEMM:

    • Per-batch tensor-wide scaling for FP8 inputs, enabled by the new cublasLtMatmulDescAttributes_t entry CUBLASLT_MATMUL_MATRIX_SCALE_PER_BATCH_SCALAR_32F.

    • Per-batch device-side alpha and beta, enabled by the new cublasLtMatmulDescAttributes_t entries CUBLASLT_MATMUL_DESC_ALPHA_BATCH_STRIDE and CUBLASLT_MATMUL_DESC_BETA_BATCH_STRIDE.

  • Improved performance on NVIDIA DGX Spark for CFP32 GEMMs. [5514146]

  • Added SM121 DriveOS support.

  • Improved performance on Blackwell (sm_100 and sm_103) via heuristics tuning for FP32 GEMMs whose shapes satisfy M, N >> K. [CUB-8572]

  • Improved performance of FP16, FP32, and CFP32 GEMMs on Blackwell Thor.

• Resolved Issues

  • Fixed missing memory initialization in cublasCreate() that could result in emulation environment variables being ignored. [CUB-9302]

  • Removed unnecessary overhead related to loading kernels on GPUs with compute capability 10.3. [5547886]

  • Fixed FP8 matmuls potentially failing to launch on multi-device Blackwell GeForce systems. [CUB-9487]

  • Added stricter checks for in-place matmul to prevent invalid use cases (C == D is allowed if and only if Cdesc == Ddesc). As a side effect, users are no longer able to use D as a dummy pointer for C when using CUBLASLT_POINTER_MODE_DEVICE with beta = 0. However, a distinct dummy pointer may still be passed. The stricter checking was added in CUDA Toolkit 13.0 Update 2. [5471880]

  • Fixed cublasLtMatmul with INT8 inputs, INT32 accumulation, and INT32 outputs potentially returning CUBLAS_STATUS_NOT_SUPPORTED when dimension N is larger than 65,536 or when batch count is larger than 1. [5541380]

  • Added validation for batched matmul to reject invalid configurations where the batch counts differ (Adesc batch count != Bdesc batch count). [5645772]

• Known Issues

  • The Grouped GEMM cuBLASLt API ignores groups with k = 0, which can lead to incorrect results. As a workaround, initialize each output matrix D with beta * C for all groups before the call, then compute Grouped GEMM as D += A * B so that the result for groups with k = 0 is preserved. This issue applies to the experimental Grouped GEMM cuBLASLt API released in CUDA 13.1. [CUB-9529]

3.1.6. cuBLAS: Release 13.0 Update 2

• New Features

  • Enabled opt-in fixed-point emulation for FP64 matmuls (D/ZGEMM) which improves performance and power-efficiency. The implementation follows the Ozaki-1 Scheme and leverages an automatic dynamic precision framework to ensure FP64-level accuracy. See here for more details on fixed-point emulation along with the table of supported compute-capabilities and the CUDA library samples for example usages.

  • Improved performance on NVIDIA DGX Spark for FP16/BF16 and FP8 GEMMs.

  • Added support for BF16x9 FP32 emulation to cublas[SC]syr[2]k and cublasCher[2]k routines. With the math mode set to CUBLAS_FP32_EMULATED_BF16X9_MATH, for large enough problems, cuBLAS will automatically dispatch SYRK and HERK to BF16x9-accelerated algorithms.

• Resolved Issues

  • Fixed undefined behavior caused by dereferencing a nullptr when passing an uninitialized matrix layout descriptor for Cdesc in cublasLtMatmul. [CUB-8911]

  • Improved performance of cublas[SCDZ]syr[2]k and cublas[CZ]her[2]k on Hopper GPUs when dimension N is large. [CUB-8293, 5384826]

• Known Issues

  • cublasLtMatmul with INT8 inputs, INT32 accumulation, and INT32 outputs might return CUBLAS_STATUS_NOT_SUPPORTED when dimension N is larger than 65,536 or when the batch count is larger than 1. The issue has existed since CUDA Toolkit 13.0 Update 1 and will be fixed in a later release. [5541380]

3.1.7. cuBLAS: Release 13.0 Update 1

• New Features

  • Improved performance:

    • Block-scaled FP4 GEMMs on NVIDIA Blackwell and Blackwell Ultra GPUs

    • SYMV on NVIDIA Blackwell GPUs [5171345]

    • cublasLtMatmul for small cases when run concurrently with other CUDA kernels [5238629]

    • TF32 GEMMs on Thor GPUs [5313616]

    • Programmatic Dependent Launch (PDL) is now supported in some cuBLAS kernels for architectures sm_90 and above, decreasing kernel launch latencies when executed alongside other PDL kernels.

• Resolved Issues

  • Fixed an issue where some cublasSsyrkx kernels produced incorrect results when beta = 0 on NVIDIA Blackwell GPUs. [CUB-8846]

  • Resolved issues in cublasLtMatmul with INT8 inputs, INT32 accumulation, and INT32 outputs where:

    • cublasLtMatmul could have produced incorrect results when A and B matrices used regular ordering (CUBLASLT_ORDER_COL or CUBLASLT_ORDER_ROW). [CUB-8874]

    • cublasLtMatmul could have been run with unsupported configurations of alpha/ beta, which must be 0 or 1. [CUB-8873]

3.1.8. cuBLAS: Release 13.0

• New Features

  • The cublasGemmEx, cublasGemmBatchedEx, and cublasGemmStridedBatchedEx functions now accept CUBLAS_GEMM_AUTOTUNE as a valid value for the algo parameter. When this option is used, the library benchmarks a selection of available algorithms internally and chooses the optimal one based on the given problem configuration. The selected algorithm is cached within the current cublasHandle_t, so subsequent calls with the same problem descriptor will reuse the cached configuration for improved performance.

    This is an experimental feature. Users are encouraged to transition to the cuBLASLt API, which provides fine-grained control over algorithm selection through the heuristics API and includes support for additional data types such as FP8 and block-scaled formats, as well as kernel fusion. (see autotuning example in cuBLASLt).

  • Improved performance of BLAS Level 3 non-GEMM kernels (SYRK, HERK, TRMM, SYMM, HEMM) for FP32 and CF32 precisions on NVIDIA Blackwell GPUs.

  • This release adds support of SM110 GPUs for arm64-sbsa on Linux.

• Known Issues

  • cublasLtMatmul previously ignored user-specified auxiliary (Aux) data types for ReLU epilogues and defaulted to using a bitmask. The correct behavior is now enforced: an error is returned if an invalid Aux data type is specified for ReLU epilogues. [CUB-7984]

• Deprecations

  • The experimental feature for atomic synchronization along the rows (CUBLASLT_MATMUL_DESC_ATOMIC_SYNC_NUM_CHUNKS_D_ROWS) and columns (CUBLASLT_MATMUL_DESC_ATOMIC_SYNC_NUM_CHUNKS_D_COLS) of the output matrix which was deprecated in 12.8 has now been removed.

  • Starting with this release, cuBLAS will return CUBLAS_STATUS_NOT_SUPPORTED if any of the following descriptor attributes are set but the corresponding scale is not supported:

    • CUBLASLT_MATMUL_DESC_A_SCALE_POINTER

    • CUBLASLT_MATMUL_DESC_B_SCALE_POINTER

    • CUBLASLT_MATMUL_DESC_D_SCALE_POINTER

    • CUBLASLT_MATMUL_DESC_D_OUT_SCALE_POINTER

    • CUBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER

  • Previously, this restriction applied only to non-narrow precision matmuls. It now also applies to narrow precision matmuls when a scale is set for a non-narrow precision tensor.

3.2.1. cuFFT: Release 13.3

• New Features

  • Expanded LTO support to include transform sizes divisible by primes larger than 127, along with increased callback support.

• Resolved Issues

  • Fixed an issue where cufftXtQueryPlan could result in floating-point exceptions when querying multi-GPU plans that are not single-batch one-dimensional FFTs.[5923044]

3.2.2. cuFFT: Release 13.2

• New Features

  • Using cuFFT link-time optimized (LTO) kernels now requires NVRTC.

• Deprecations

  • cufftDebug is deprecated and will be removed in a future release.

3.2.3. cuFFT: Release 13.1

• New Features

  • Improved performance for transforms whose sizes are powers of 2, 3, 5, and 7 on Blackwell GPUs, in both single and double precision.

  • Improved performance for selected power-of-two sizes in 2D and 3D transforms, in both single and double precision.

  • Introduced an experimental cuFFT device API that provides host functions to query or generate device function code and exposes database metadata through a C++ header for use with the cuFFTDx library.

• Resolved Issues

  • Fixed a correctness issue, identified in CUDA 13.0, that affected a very specific subset of kernels: half- and bfloat16-precision strided R2C and C2R FFTs of size 1.

3.2.4. cuFFT: Release 13.0 Update 1

• Known Issues

  • In CUDA 13.0, a correctness issue affects a specific subset of kernels, namely half and bfloat precision size 1 strided R2C and C2R kernels. A fix will be included in a future CUDA release.

3.2.5. cuFFT: Release 13.0

• New Features

  • Added new error codes:

    • CUFFT_MISSING_DEPENDENCY

    • CUFFT_NVRTC_FAILURE

    • CUFFT_NVJITLINK_FAILURE

    • CUFFT_NVSHMEM_FAILURE

  • Introduced CUFFT_PLAN_NULL, a value that can be assigned to a cufftHandle to indicate a null handle. It is safe to call cufftDestroy on a null handle.

  • Improved performance for single-precision C2C multi-dimensional FFTs and large power-of-2 FFTs.

• Known Issues

  • An issue identified in CUDA 13.0 affects the correctness of a specific subset of cuFFT kernels, specifically half-precision and bfloat16 size-1 strided R2C and C2R transforms. A fix will be included in a future CUDA release.

• Deprecations

  • Removed support for Maxwell, Pascal, and Volta GPUs, corresponding to compute capabilities earlier than Turing.

  • Removed legacy cuFFT error codes:

    • CUFFT_INCOMPLETE_PARAMETER_LIST

    • CUFFT_PARSE_ERROR

    • CUFFT_LICENSE_ERROR

  • Removed the libcufft../_static_nocallback.a static library. Users should link against libcufft../_static.a instead, as both are functionally equivalent.

3.3.1. cuSOLVER: Release 13.3

• New Features

  • Improved cusolverDnXgeev performance when computing eigenvectors by moving eigenvector post-processing from the host to the device.

• Known Issues

  • The cusolverDn{C,Z}sytrf and cusolverDnXsytrs APIs assume that the complex input matrix A is Hermitian instead of symmetric when devIpiv is set to NULL. This issue exists starting with CUDA Toolkit 13.1. [5797471]

3.3.2. cuSOLVER: Release 13.2 Update 1

• New Features

  • Improved performance of cusolverDnXgeqrf() and cusolverDn<S,D,C,Z>geqrf() on sm_90, sm_100, sm_103, and sm_120 for matrices with m <= 65536.

  • Added the new public 64-bit interface cusolverDnXpolar(), which exposes the QDWH algorithm implementation for polar decomposition in cuSOLVERDn.

  • Added the new public 64-bit interface cusolverDnXstedc(), which computes the eigenvalues and, optionally, eigenvectors of a symmetric tridiagonal matrix using the divide-and-conquer method.

3.3.3. cuSOLVER: Release 13.2

• New Features

  • Added FP64 fixed-point emulation support to cuSOLVERDn. The following new APIs are available:

  • cusolverDnSetFixedPointEmulationMantissaControl()

  • cusolverDnGetFixedPointEmulationMantissaControl()

  • cusolverDnSetFixedPointEmulationMaxMantissaBitCount()

  • cusolverDnGetFixedPointEmulationMaxMantissaBitCount()

  • cusolverDnSetFixedPointEmulationMantissaBitOffset()

  • cusolverDnGetFixedPointEmulationMantissaBitOffset()

  • cusolverDnSetEmulationSpecialValuesSupport()

  • cusolverDnGetEmulationSpecialValuesSupport()

  • Added the cusolverDnXsygvd API to support larger problem sizes.

• Known Issues

  • Starting with CUDA Toolkit 13.1, cusolverDn{C,Z}sytrf and cusolverDnXsytrs assume the complex input matrix A is Hermitian (instead of symmetric) when devIpiv == NULL.[5797471]

3.3.4. cuSOLVER: Release 13.1

• Resolved Issues

  • Fixed a bug that prevented users from changing the algorithm for cusolverDnXsyevBatched by using cusolverDnSetAdvOptions.[5539844]

3.3.5. cuSOLVER: Release 13.0 Update 1

• Resolved Issues

  • Fixed a race condition in cusolverDnXgeev that could occur when using multiple host threads with either separate handles per thread or a shared handle, which caused execution to abort and returned CUSOLVER_STATUS_INTERNAL_ERROR.

3.3.6. cuSOLVER: Release 13.0

• New Features

  • cuSOLVER offers a new math mode to leverage improved performance of emulated FP32 arithmetic on Nvidia Blackwell GPUs.

    To enable and control this feature, the following new APIs have been added:

    • cusolverDnSetMathMode()

    • cusolverDnGetMathMode()

    • cusolverDnSetEmulationStrategy()

    • cusolverDnGetEmulationStrategy()

  • Performance improvements for cusolverDnXsyevBatched() have been made by introducing an internal algorithm switch on Blackwell GPUs for matrices of size n <= 32.

    To revert to the previous algorithm for all problem sizes, use cusolverDnSetAdvOptions().

    For more details, refer to the cusolverDnXsyevBatched() documentation.

• Deprecations

  • cuSOLVERMg is deprecated and may be removed in an upcoming major release. Users are encouraged to use cuSOLVERMp for multi-GPU functionality across both single and multi-node environments. To disable the deprecation warning, add the compiler flag -DDISABLE_CUSOLVERMG_DEPRECATED.

  • cuSOLVERSp and cuSOLVERRf are fully deprecated and may be removed in an upcoming major release. Users are encouraged to use the cuDSS library for better performance and ongoing support.

    For help with the transition, refer to the cuDSS samples or CUDA samples for migrating from cuSOLVERSp to cuDSS.

    To disable the deprecation warning, add the compiler flag: -DDISABLE_CUSOLVER_DEPRECATED.

• Resolved Issues

  • The supported input matrix size for cusolverDnXsyevd, cusolverDnXsyevdx, cusolverDnXsyevBatched, cusolverDn<t>syevd, and cusolverDn<t>syevdx is no longer limited to n <= 32768.

    This update also applies to routines that share the same internal implementation: cusolverDnXgesvdr, cusolverDnXgesvdp, cusolverDn<t>sygvd, cusolverDn<t>sygvdx, and cusolverDn<t>gesvdaStridedBatched.

3.4.1. cuSPARSE: Release 13.3

• New Features

  • Added support for the CSC format in SpSV and SpSM.

  • Improved CSR SpMV ALG2 performance by an average of 11%.

  • Added the Generic API SpGEAM for sparse matrix-matrix addition.

  • Added SpMVOp ALG1 with reduced preprocessing overhead.

  • Added support for mixed index types in SpMVOp computation for CSR matrices with 64-bit offsets and 32-bit indices.

  • Added support for the FP32 data type in SpMVOp.

  • Avoided recompilation for the same epilogue in SpMVOp.

  • Added mixed-precision support in SpMV for 32-bit input matrices and 64-bit input vectors.

  • Added support for updating matrix values after preprocessing in SpMVOp ALG1.

• Resolved Issues

  • Fixed a memory leak in SpMVOp when destroy_lrb() was called. [5974043]

3.4.2. cuSPARSE: Release 13.2 Update 1

• New Features

  • Improved cusparseSpMVOp_createDescr() performance by up to 2.5x.

  • Reduced cusparseSpMVOp_createPlan() planning latency for default epilogues through ahead-of-time compilation, avoiding JIT compilation in this case.

• Resolved Issues

  • Fixed an issue that caused performance regressions in BSR SpMM for certain block sizes. [5860241]

• Deprecation

  • Deprecated the SpMMOp and SpGEMMreuse APIs.

3.4.3. cuSPARSE: Release 13.2

• New Features

  • Improved the runtime of the SpMVOp::buffer_size_estimate API.

3.4.4. cuSPARSE: Release 13.1 Update 1

• New Features

  • Added a new cusparseSpMVOp_bufferSize API that returns the size of the workspace buffer required for SpMVOp computations. Users provide this buffer when creating cusparseSpMVOpDescr_t, removing internal memory allocations.

  • Improved SpMVOp performance on B200. [CUSPARSE-2931] [CUSPARSE-2932] [CUSPARSE-2933]

• Resolved Issues

  • Fixed an accuracy issue in mixed-precision CSR/COO SpMM computations. [CUSPARSE-2349]

  • Fixed an issue in CSR SpMM computations when the input dense matrix has a high number of columns. [CUSPARSE-2301]

3.4.5. cuSPARSE: Release 13.1

• New Features

  • Introduced an experimental Sparse Matrix-Vector Multiplication (SpMVOp) API that provides improved performance compared with the existing generic CsrMV API. This API supports CSR format with 32-bit indices, double precision, and user-defined epilogues.

  • The nvJitLink shared library is now loaded dynamically at runtime.

  • Improved cusparseXcsrsort with reduced memory usage and higher performance. [CUSPARSE-2630]

• Known Issues

  • When using 32-bit indexing, cusparseSpSV and cusparseSpSM may crash if the number of nonzero elements (nnz) approaches 2^31 - 1.[CUSPARSE-2211]

• Resolved Issues

  • Fixed potential issues when input and output pointers are not 16-byte aligned in cusparseCsr2cscEx2, cusparseSparseToDense, and CSR/COO cusparseSpMM. [CUSPARSE-2380]

  • Fixed a determinism issue in CSR cusparseSpMM ALG3. [CUSPARSE-2612]

  • All routines now support matrices with up to 2^31 - 1 nonzero elements (nnz) when using 32-bit indexing, with the exception of cusparseSpSV and cusparseSpSM. [CUSPARSE-2153]

  • Fixed a potential race condition that could occur when dynamically loading driver APIs. [CUSPARSE-2764]

3.4.6. cuSPARSE: Release 13.0 Update 1

• New Features

  • Added support for the BSR format in the generic SpMV API (CUSPARSE-2518).

• Deprecation

  • Deprecated the legacy BSR SpMV API (replaced by the generic SpMV API).

• Resolved Issues

  • Enabled all generic APIs to support zero-dimension matrices/vectors (m, n, k = 0) (CUSPARSE-2378).

  • Enabled all generic APIs to support small-dimension matrices/vectors (small m, n, or k) (CUSPARSE-2379).

  • Fixed incorrect results in mixed-precision CSR/COO SpMV computations (CUSPARSE-2349).

3.4.7. cuSPARSE: Release 13.0

• New Features

  • Added support for 64-bit index matrices in SpGEMM computation. (CUSPARSE-2365)

• Known Issues

  • cuSPARSE logging APIs can crash on Windows.

  • CUSPARSE_SPMM_CSR_ALG3 does not return deterministic results as stated in the documentation.

• Deprecation

  • Dropped support for pre-Turing architectures (Maxwell, Volta, and Pascal).

• Resolved Issues

  • Fixed a bug in cusparseSparseToDense_bufferSize that caused it to request up to 16× more memory than required. [CUSPARSE-2352]

  • Fixed unwanted 16-byte alignment requirements on the external buffer. Most routines will now work with any alignment. In the generic API, only cusparseSpGEMM routines are still affected. [CUSPARSE-2352]

  • Fixed incorrect results from cusparseCsr2cscEx2 when any of the input matrix dimensions are zero, such as when m = 0 or n = 0. [CUSPARSE-2319]

  • Fixed incorrect results from CSR SpMV when any of the input matrix dimensions are zero, such as when m = 0 or n = 0. [CUSPARSE-1800]

3.5.1. CUDA Math: Release 13.3

• Resolved Issues

  • Fixed an issue where silent data corruption could occur when the CUDA Math API __mul24() intrinsic was called with compile-time constant inputs due to undefined behavior from compiler optimizations applied to overflowing signed integer multiplication. This issue was introduced in CUDA Toolkit 11.1 and resolved in CUDA Toolkit 13.3. [5807344]

3.5.2. CUDA Math: Release 13.2 Update 1

• Known Issues

  • Silent data corruption can occur when the CUDA Math API __mul24() intrinsic is called with compile-time constant inputs. Compiler optimizations applied to overflowing signed integer multiplication can expose the program to undefined behavior. This issue was introduced in CUDA Toolkit 11.1 and will be fixed in a future release. [5807344]

3.5.3. CUDA Math: Release 13.2

• New Features

  • Accuracy and performance improvements were made to the following libdevice single-precision math functions:

    • expm1f(): up to 20% faster, with minor accuracy improvements.

    • erff(): 5% to 10% faster, with minor accuracy improvements.

  These gains come from algorithmic simplifications, reduced branching, and tighter approximations.[5480287]

3.5.4. CUDA Math: Release 13.0

• New Features

  • Single and double precision math functions received targeted performance and accuracy improvements through algorithmic simplifications, reduced branching, and tighter approximations.

    • atan2f, atan2: Up to 10% faster with minor improvements in accuracy.

    • sinhf, coshf, acoshf, asinhf, asinh: Up to 50% speedups with minor improvements in accuracy.

    • cbrtf, rcbrtf: 15% faster with minor improvements in accuracy.

    • erfinvf, erfcinvf, normcdfinvf: Minor accuracy improvements, performance neutral.

    • ldexpf, ldexp: Up to 3x faster in single precision and 30% faster in double precision, with no accuracy loss.

    • modff, modf: Up to 50% faster in single precision and 10% faster in double precision, with no accuracy loss.

3.6.1. nvJPEG: Release 13.3

• New Features

  • Added support for region-of-interest decoding with nvjpegDecodeBatchedEx when using the NVJPEG_BACKEND_LOSSLESS_JPEG backend.

Resolved Issues

• Fixed an issue with boundary handling when decoding a region of interest with NVJPEG_FLAGS_UPSAMPLING_WITH_INTERPOLATION enabled.

3.6.2. nvJPEG: Release 13.2 Update 1

• New Features

  • Added the NVJPEG_OUTPUT_UNCHANGEDI enum value to nvjpegOutputFormat_t for unchanged interleaved output. For chroma subsampling formats other than 4:4:4, chroma values are duplicated so that the chroma and luma dimensions match.

3.6.3. nvJPEG: Release 13.1

Resolved Issues

• nvJPEG’s lossless JPEG 92 (lj92) implementation can now correctly handle lj92 files that contain a comment marker in the header. [5484797]

3.6.4. nvJPEG: Release 13.0 Update 1

Resolved Issues

• Fixed a race condition in certain cases during progressive encoding (5307748).

• Fixed an uninitialized read when encoding images as 4:1:0 JPEG bitstreams (5308008).

3.6.5. nvJPEG: Release 13.0

• Deprecations

  • Removed the nvjpegEncoderParamsCopyHuffmanTables API.

Resolved Issues

• nvJPEG is now more robust and no longer crashes or exhibits undefined behavior when decoding malformed or truncated bitstreams. [5168024, 5133845, 5143450]

• nvjpegEncodeYUV now avoids reading outside of allocated device memory in certain cases. [5133826]

• Optimized memory usage when encoding RGB inputs using the hardware encoder.

• Fixed issues related to rounding in various transform, sampling, and conversion steps, improving image quality for both encoder and decoder. [5064901, 3976092]

• Various bug fixes for improved security.

3.7.1. NPP: Release 13.1 Update 1

• Resolved Issues

  • Reduced nvJPEG Encoder initialization time on Thor. [5533951]

3.7.2. NPP: Release 13.1

• Resolved Issues

  • Fixed an issue in nppiCFAToRGB_8u_C1C3R() affecting SSIM validation for NPPI_BAYER_GBRG patterns. [5192648]

3.7.3. NPP: Release 13.0

• Deprecations

  • Removal of Legacy Non-Context APIs

    All legacy NPP APIs without the _Ctx suffix have been deprecated and are now removed starting with this release. Developers should transition to the context-aware (_Ctx) versions to ensure continued support and compatibility with the latest CUDA releases.

  • Deprecation of ``nppGetStreamContext()``

    The nppGetStreamContext() API has been deprecated and removed. Developers are strongly encouraged to adopt application-managed stream contexts by explicitly managing the NppStreamContext structure. For guidance, refer to the NPP Documentation – General Conventions and the usage demonstrated in the StreamContexts example.

• Resolved Issues

  • Fixed an issue in nppiFloodFillRange_8u_C1IR_Ctx where the flood fill operation did not correctly fill the full target area. [5141474]

  • Resolved a bug in the nppiDebayer() API that affected proper reconstruction of color data during Bayer pattern conversion. [5138782]