I encountered a strange warning when compiling a GEMM kernel for Hopper cards. ptxas info : (C7514) Potential Performance Loss: wgmma.mma_async instructions are serialized due to non wgmma instructions reading accumulator registers of a wgmma between start and end of the pipeline stage in the f…

Hi， thanks for reporting NVBUG 5431330 . Below the latest engineering team update . Thank you for reporting this case to us. As you’ve already discovered in your other report, changing the first wgmma.wait_group.sync.aligned 1 in the loop to wgmma.wait_group.sync.aligned 0 can serve as a workaround…

See Log in | NVIDIA Developer for details.

Now I believe this is a bug of NVCC. Here is a simple program that can reproduce this issue: #include <cstdint> __launch_bounds__(128, 1) __global__ void test(unsigned *result, unsigned iters) { __shared__ unsigned buffer[8192]; unsigned *buf1 = &buffer[0], *buf2 = &buffer[4096]; unsigned *b…

Generally speaking, the NVCC compiler fails to analyze dependencies of wgmma instructions correctly when encountered simple loop branchings, which is annoying. This issue does not exist for other asynchronous instructions like “cp.async(.wait_group N)”.

This issue still exist for CUDA-13.0

[image] jiachen17713: Now I believe this is a bug of NVCC. You might want to file a bug and refer to this thread.

A bug has already been filed as indicated in comment #2 in this thread. The development team is looking at it internally as recently as August 6th. It doesn’t appear to be resolved yet. However there is no indication (that I can see) in the bug referring back to this thread. I’ve added a note to…

PTXAS: mysterious warning for wgmma.mma_async instruction serialization

Accelerated Computing CUDA CUDA NVCC Compiler

rs277 August 9, 2025, 5:05am 6

You might want to file a bug and refer to this thread.

Topic		Replies	Views
Incorrect PTX complier warning? Potential performance loss for WGMMA pipeline crossing func boundary CUDA Programming and Performance cuda , compile	0	126	October 17, 2025
PTXAS says registers not enough for m64n256k32 integer wgmma after setting regcount from 128 to 200 by setmaxnreg.inc CUDA Programming and Performance	7	158	April 21, 2026
Sm_90a: only one WGMMA of a 2-mma group overlaps the accumulator XOR-reduction — why? CUDA Programming and Performance	5	98	June 24, 2026
Example with wgmma.mma_async CUDA Programming and Performance	4	2660	November 23, 2024
Question about wgmma instruction in Hopper CUDA Programming and Performance	3	1583	October 25, 2024
How many tensor cores to execute the wmma.mma.sync.aligned.{alayout}.{blayout}.m16n16k16 instruction？ CUDA Programming and Performance cuda	23	478	December 12, 2025
Throughput and latency of mma.sync instruction CUDA Programming and Performance	4	670	August 19, 2024
Wmma vs Wgmma On H100 GPU CUDA Programming and Performance cublas	4	421	December 15, 2025
Fastest Tiled WMMA for Matrices of Any Size? CUDA Programming and Performance	3	529	October 26, 2024
PTX instruction `mma` not lowered to tensor core related SASS instruction TensorRT	2	1452	March 22, 2022

PTXAS: mysterious warning for wgmma.mma_async instruction serialization

Related topics