[Common] Fix pointer arithmatic to generate correct LDS/STS instructions

[kainzhong]

Description

The previous MXFP8 dynamic-SMEM alignment code round-tripped the buffer base through uintptr_t and cast the integer back to a pointer. That breaks the pointer's link to the extern shared object, so the compiler can no longer prove the address is in the shared window and falls back to generic address-space loads/stores (LD.E/ST.E), which resolve the window at runtime.

This PR fixes this behavior so compiler is able to generate the correct LDS / STS instructions.

I have a vibe coded demo repro on https://github.com/kainzhong/ADDR_PTR_REPRO which should explain why.

The pointer works this way:
base + (-base & 127)
= base & ~127 + (base & 127) + (-base & 127)

• base & ~127: set the last 7 digits to 0
• base & 127: take the last 7 digits
• -base & 127: take the last 7 digits of -base
  If the last 7 digits of base is 0000000, then (base & 127) + (-base & 127) is 0
  Otherwise it will become 1 << 8

TODO: try int instead of uint

Fixes # (issue)

Type of change

• Documentation change (change only to the documentation, either a fix or a new content)
• Bug fix (non-breaking change which fixes an issue)
• New feature (non-breaking change which adds functionality)
• Breaking change (fix or feature that would cause existing functionality to not work as expected)
• Infra/Build change
• Code refactoring

Changes

Please list the changes introduced in this PR:

• Compute the aligned pointer directly without going through uintptr_t, in which case nvcc loses track of this pointer's semantics and degrade from ld.shared / st.shared (LDS / STS) to ld.generic / st.generic (LD / ST)

Checklist:

• I have read and followed the contributing guidelines
• The functionality is complete
• I have commented my code, particularly in hard-to-understand areas
• I have made corresponding changes to the documentation
• My changes generate no new warnings
• I have added tests that prove my fix is effective or that my feature works
• New and existing unit tests pass locally with my changes

[greptile-apps]

Greptile Summary

The PR preserves shared-memory pointer provenance while retaining the existing alignment behavior.

• Replaces integer-to-pointer round trips with pointer-relative alignment in MXFP8 and NVFP4 quantization kernels.
• Keeps the aligned offsets and dynamic shared-memory requirements unchanged.
• Enables the compiler to emit shared-memory LDS/STS instructions instead of generic LD/ST instructions.

Confidence Score: 5/5

The PR appears safe to merge.

No blocking failure remains.

Important Files Changed

Filename	Overview
transformer_engine/common/cast/mxfp8/gated_mxfp8.cuh	Changes the dynamic shared-memory alignment expression to preserve pointer provenance without changing the computed offset.
transformer_engine/common/cast/mxfp8/quantize_mxfp8.cuh	Uses pointer-relative alignment for the MXFP8 quantization kernel while retaining the existing 128-byte alignment.
transformer_engine/common/cast/mxfp8/specialized/quantize_mxfp8.cuh	Updates both specialized cast-only kernels to derive aligned addresses directly from their shared-memory objects.
transformer_engine/common/cast/nvfp4/group_quantize_transpose_nvfp4.cuh	Applies the same provenance-preserving alignment calculation to grouped NVFP4 quantization.
transformer_engine/common/cast/nvfp4/quantize_transpose_nvfp4.cuh	Updates both NVFP4 shared-memory alignment sites without changing their layout or allocation requirements.

_{Reviews (3): Last reviewed commit: "[Common] Fix pointer arithmatic to gener..." | Re-trigger Greptile}

[kainzhong]

Vibe coded a benchmark. Looks like the mostly it's a gain or no change on perf numbers except for a few specific shape:

• MXFP8 generic: mostly improvement but regress on gelu fusion. Claude Code's explanation is MIO throttle increased (LDS/STS goes to MIO which is also what special math units for gelu uses)
• MXFP8 specialized bidimensional: pretty much no meaningful change
• MXFP8 gated: overall a win
• NVFP4 2D: slightly better

MXFP8 / NVFP4 shared-memory provenance fix -- results

Aligning the dynamic-SHMEM base through an integer (uintptr_t) and casting the integer back to a pointer loses the link to the extern __shared__ object. ptxas can then no longer prove the address is in the shared window, and falls back to generic address-space accesses (LD.E/ST.E) that resolve the window at runtime. Deriving the aligned address from the original char* by pointer arithmetic restores LDS/STS.

before	8ffbd1324647 (main, no fixes)
after	c88551a4bbea (all four fixes)
GPU	NVIDIA GB200
timing	pure GPU kernel time, median CUPTI duration via nsys nvtx_kern_sum
dtype	bf16 -> e4m3 (MXFP8), bf16 -> e2m1 (NVFP4)

Summary

fix	file	n	cold median	cold best	warm median	warm best
generic	cast/mxfp8/quantize_mxfp8.cuh	96	1.042x	1.174x	1.040x	1.306x
specialized	cast/mxfp8/specialized/quantize_mxfp8.cuh	12	1.003x	1.012x	1.002x	1.118x
gated	cast/mxfp8/gated_mxfp8.cuh	54	1.017x	1.089x	1.019x	1.095x
nvfp4	cast/nvfp4/quantize_transpose_nvfp4.cuh	18	1.025x	1.059x	1.025x	1.073x
all		180	1.027x	1.174x	1.030x	1.306x

Total kernel time over all 180 workloads: cold 21599 -> 20891 us (1.034x), warm 21357 -> 20610 us (1.036x).

Worst single workload: cold 0.937x, warm 0.951x.

MXFP8 generic quantize kernel

cast/mxfp8/quantize_mxfp8.cuh

fusion	dir	n	cold median	cold range	warm median	warm range
dbias	row	6	1.151x	1.130-1.174	1.155x	1.117-1.175
dbias	col	6	1.089x	1.047-1.142	1.181x	1.110-1.306
dbias_dgelu	both	6	1.080x	1.067-1.089	1.088x	1.081-1.103
dbias	both	6	1.079x	1.069-1.083	1.092x	1.087-1.106
gelu	row	6	1.079x	1.055-1.082	1.076x	1.063-1.083
plain	col	6	1.060x	1.049-1.070	1.176x	1.074-1.265
dgelu	row	6	1.050x	1.045-1.055	1.041x	1.036-1.046
gelu	both	6	1.035x	0.937-1.040	1.038x	0.951-1.041
dbias_dgelu	row	6	1.029x	1.007-1.074	1.031x	1.004-1.084
dbias_dgelu	col	6	1.029x	1.000-1.070	1.032x	1.016-1.068
dgelu	col	6	1.028x	1.025-1.046	1.037x	1.033-1.062
dsilu	row	6	1.027x	0.979-1.044	1.028x	1.026-1.050
dsilu	col	6	1.023x	1.014-1.040	1.017x	0.992-1.023
gelu	col	6	1.004x	0.999-1.047	1.003x	1.001-1.016
dgelu	both	6	1.001x	0.994-1.009	1.004x	0.996-1.011
dsilu	both	6	0.995x	0.988-1.019	0.999x	0.993-1.017

MXFP8 specialized cast-only kernel (bidimensional)

cast/mxfp8/specialized/quantize_mxfp8.cuh

fusion	dir	n	cold median	cold range	warm median	warm range
plain	both	6	1.003x	1.002-1.012	1.006x	0.995-1.118
plain	row	6	1.002x	0.999-1.006	1.000x	0.995-1.002

MXFP8 gated (swiglu / geglu) kernel

cast/mxfp8/gated_mxfp8.cuh

fusion	dir	n	cold median	cold range	warm median	warm range
dswiglu	col	6	1.086x	1.076-1.089	1.090x	1.082-1.091
geglu	row	6	1.078x	1.076-1.082	1.087x	1.062-1.095
geglu	both	6	1.058x	1.052-1.059	1.059x	1.043-1.065
dswiglu	both	6	1.021x	1.013-1.035	1.019x	1.013-1.030
dswiglu	row	6	1.016x	1.013-1.017	1.016x	1.013-1.019
swiglu	col	6	1.015x	1.013-1.019	1.016x	1.013-1.020
swiglu	both	6	1.008x	1.001-1.059	1.005x	1.000-1.052
swiglu	row	6	1.007x	0.984-1.016	1.016x	1.009-1.019
geglu	col	6	1.005x	1.001-1.006	1.005x	1.002-1.019

NVFP4 2D quantize-transpose kernel

cast/nvfp4/quantize_transpose_nvfp4.cuh

fusion	dir	n	cold median	cold range	warm median	warm range
nvfp42d	row	6	1.028x	1.018-1.033	1.026x	1.020-1.035
nvfp42d	col	6	1.025x	1.014-1.059	1.024x	1.013-1.073
nvfp42d	both	6	1.022x	0.999-1.046	1.018x	1.001-1.040

Full per-workload numbers (us, pure GPU kernel time)

fix	fusion	dir	shape	model tensor	cold before	cold after	cold	warm before	warm after	warm
gated	dswiglu	both	4096x4096	Llama-3-8B hidden	72.4	69.9	1.035x	70.5	68.4	1.030x
gated	dswiglu	both	2048x12288	GPT-3-175B hidden	104.9	101.6	1.033x	103.3	100.8	1.025x
gated	dswiglu	both	8192x8192	Llama-3-70B hidden	262.2	255.8	1.025x	259.6	254.0	1.022x
gated	dswiglu	both	4096x14336	Llama-3-8B FFN	230.0	226.1	1.017x	227.9	224.7	1.014x
gated	dswiglu	both	16384x5120	long-seq hidden	322.7	317.3	1.017x	319.8	314.8	1.016x
gated	dswiglu	both	8192x28672	Llama-3-70B FFN	888.7	877.5	1.013x	881.2	869.8	1.013x
gated	dswiglu	col	8192x28672	Llama-3-70B FFN	460.2	422.4	1.089x	458.7	420.8	1.090x
gated	dswiglu	col	16384x5120	long-seq hidden	169.0	155.3	1.088x	167.5	153.6	1.091x
gated	dswiglu	col	4096x14336	Llama-3-8B FFN	120.9	111.3	1.086x	119.6	109.7	1.091x
gated	dswiglu	col	8192x8192	Llama-3-70B hidden	136.9	126.0	1.086x	135.4	124.4	1.088x
gated	dswiglu	col	2048x12288	GPT-3-175B hidden	56.6	52.4	1.081x	55.2	50.6	1.091x
gated	dswiglu	col	4096x4096	Llama-3-8B hidden	40.3	37.4	1.076x	38.8	35.9	1.082x
gated	dswiglu	row	16384x5120	long-seq hidden	201.8	198.4	1.017x	199.8	196.4	1.017x
gated	dswiglu	row	4096x4096	Llama-3-8B hidden	47.7	46.9	1.017x	45.8	45.0	1.019x
gated	dswiglu	row	8192x8192	Llama-3-70B hidden	163.5	160.8	1.017x	161.5	158.7	1.017x
gated	dswiglu	row	4096x14336	Llama-3-8B FFN	144.3	142.2	1.015x	142.1	140.1	1.015x
gated	dswiglu	row	2048x12288	GPT-3-175B hidden	67.1	66.2	1.015x	64.6	63.6	1.015x
gated	dswiglu	row	8192x28672	Llama-3-70B FFN	548.7	541.5	1.013x	546.3	539.4	1.013x
gated	geglu	both	4096x4096	Llama-3-8B hidden	39.2	37.0	1.059x	37.4	35.8	1.043x
gated	geglu	both	8192x28672	Llama-3-70B FFN	473.5	447.1	1.059x	471.1	445.5	1.058x
gated	geglu	both	16384x5120	long-seq hidden	171.9	162.4	1.058x	170.7	160.4	1.064x
gated	geglu	both	4096x14336	Llama-3-8B FFN	121.8	115.1	1.058x	121.1	114.5	1.058x
gated	geglu	both	8192x8192	Llama-3-70B hidden	138.6	131.1	1.057x	137.9	129.4	1.065x
gated	geglu	both	2048x12288	GPT-3-175B hidden	55.6	52.8	1.052x	54.9	51.8	1.060x
gated	geglu	col	4096x14336	Llama-3-8B FFN	76.8	76.3	1.006x	75.9	75.6	1.003x
gated	geglu	col	2048x12288	GPT-3-175B hidden	35.9	35.7	1.006x	35.2	34.9	1.008x
gated	geglu	col	4096x4096	Llama-3-8B hidden	25.7	25.5	1.005x	23.9	23.5	1.019x
gated	geglu	col	8192x8192	Llama-3-70B hidden	87.0	86.6	1.005x	86.2	85.6	1.007x
gated	geglu	col	8192x28672	Llama-3-70B FFN	289.2	288.5	1.002x	288.2	287.7	1.002x
gated	geglu	col	16384x5120	long-seq hidden	107.0	106.8	1.001x	106.1	105.8	1.003x
gated	geglu	row	8192x28672	Llama-3-70B FFN	285.4	263.8	1.082x	284.2	261.8	1.086x
gated	geglu	row	4096x14336	Llama-3-8B FFN	77.0	71.4	1.079x	75.9	69.7	1.088x
gated	geglu	row	2048x12288	GPT-3-175B hidden	36.8	34.1	1.078x	35.9	32.7	1.095x
gated	geglu	row	8192x8192	Llama-3-70B hidden	87.2	80.8	1.078x	85.6	78.9	1.086x
gated	geglu	row	4096x4096	Llama-3-8B hidden	26.8	24.8	1.078x	24.0	22.6	1.062x
gated	geglu	row	16384x5120	long-seq hidden	106.9	99.3	1.076x	105.5	97.0	1.088x
gated	swiglu	both	4096x4096	Llama-3-8B hidden	43.4	40.9	1.059x	40.4	38.4	1.052x
gated	swiglu	both	2048x12288	GPT-3-175B hidden	61.8	59.0	1.047x	60.1	57.7	1.042x
gated	swiglu	both	8192x8192	Llama-3-70B hidden	149.0	147.3	1.011x	146.7	145.8	1.006x
gated	swiglu	both	4096x14336	Llama-3-8B FFN	130.3	129.7	1.005x	128.4	128.4	1.000x
gated	swiglu	both	16384x5120	long-seq hidden	183.2	182.5	1.004x	181.2	181.1	1.001x
gated	swiglu	both	8192x28672	Llama-3-70B FFN	502.5	502.1	1.001x	500.5	498.6	1.004x
gated	swiglu	col	4096x4096	Llama-3-8B hidden	27.0	26.5	1.019x	25.0	24.5	1.020x
gated	swiglu	col	4096x14336	Llama-3-8B FFN	81.2	79.8	1.018x	80.4	79.2	1.015x
gated	swiglu	col	2048x12288	GPT-3-175B hidden	37.9	37.3	1.015x	37.4	36.8	1.016x
gated	swiglu	col	16384x5120	long-seq hidden	113.3	111.6	1.015x	113.0	111.2	1.017x
gated	swiglu	col	8192x28672	Llama-3-70B FFN	307.2	303.1	1.013x	306.3	302.4	1.013x
gated	swiglu	col	8192x8192	Llama-3-70B hidden	91.9	90.7	1.013x	91.2	89.8	1.015x
gated	swiglu	row	8192x28672	Llama-3-70B FFN	338.5	333.3	1.016x	337.1	330.9	1.019x
gated	swiglu	row	16384x5120	long-seq hidden	126.0	124.7	1.010x	124.2	122.3	1.016x
gated	swiglu	row	8192x8192	Llama-3-70B hidden	102.8	101.7	1.010x	101.0	99.6	1.014x
gated	swiglu	row	4096x14336	Llama-3-8B FFN	90.7	90.3	1.004x	89.5	88.0	1.017x
gated	swiglu	row	2048x12288	GPT-3-175B hidden	43.7	44.1	0.992x	42.3	41.9	1.009x
gated	swiglu	row	4096x4096	Llama-3-8B hidden	30.9	31.4	0.984x	26.8	26.3	1.019x
generic	dbias	both	16384x5120	long-seq hidden	68.7	63.5	1.083x	68.8	63.0	1.092x
generic	dbias	both	4096x14336	Llama-3-8B FFN	50.2	46.3	1.083x	49.8	45.8	1.087x
generic	dbias	both	4096x4096	Llama-3-8B hidden	18.9	17.6	1.079x	16.8	15.2	1.104x
generic	dbias	both	8192x28672	Llama-3-70B FFN	177.1	164.3	1.078x	178.0	163.2	1.091x
generic	dbias	both	8192x8192	Llama-3-70B hidden	56.4	52.4	1.077x	56.5	52.0	1.087x
generic	dbias	both	2048x12288	GPT-3-175B hidden	24.5	22.9	1.069x	23.4	21.1	1.106x
generic	dbias	col	4096x4096	Llama-3-8B hidden	13.1	11.5	1.142x	10.9	9.0	1.213x
generic	dbias	col	2048x12288	GPT-3-175B hidden	17.9	16.2	1.106x	15.6	11.9	1.306x
generic	dbias	col	4096x14336	Llama-3-8B FFN	33.6	30.8	1.094x	32.7	27.4	1.195x
generic	dbias	col	8192x28672	Llama-3-70B FFN	115.9	106.9	1.084x	115.3	103.8	1.110x
generic	dbias	col	16384x5120	long-seq hidden	45.3	41.8	1.082x	44.9	38.8	1.158x
generic	dbias	col	8192x8192	Llama-3-70B hidden	37.3	35.7	1.047x	36.9	31.6	1.168x
generic	dbias	row	8192x28672	Llama-3-70B FFN	160.1	136.4	1.174x	157.5	134.0	1.175x
generic	dbias	row	4096x14336	Llama-3-8B FFN	45.2	39.2	1.153x	42.7	36.9	1.157x
generic	dbias	row	4096x4096	Llama-3-8B hidden	16.7	14.5	1.152x	12.8	11.5	1.117x
generic	dbias	row	16384x5120	long-seq hidden	61.3	53.2	1.151x	59.1	50.8	1.163x
generic	dbias	row	8192x8192	Llama-3-70B hidden	50.8	44.3	1.147x	48.3	41.8	1.154x
generic	dbias	row	2048x12288	GPT-3-175B hidden	22.4	19.9	1.130x	17.1	14.9	1.149x
generic	dbias_dgelu	both	4096x14336	Llama-3-8B FFN	147.2	135.2	1.089x	147.5	135.6	1.087x
generic	dbias_dgelu	both	16384x5120	long-seq hidden	207.0	191.4	1.082x	207.6	191.4	1.085x
generic	dbias_dgelu	both	4096x4096	Llama-3-8B hidden	46.7	43.2	1.081x	46.8	42.4	1.103x
generic	dbias_dgelu	both	8192x8192	Llama-3-70B hidden	165.9	153.7	1.079x	168.4	154.1	1.093x
generic	dbias_dgelu	both	8192x28672	Llama-3-70B FFN	571.0	531.2	1.075x	572.6	529.6	1.081x
generic	dbias_dgelu	both	2048x12288	GPT-3-175B hidden	65.7	61.5	1.067x	66.4	61.0	1.089x
generic	dbias_dgelu	col	8192x28672	Llama-3-70B FFN	381.7	356.6	1.070x	379.7	355.7	1.068x
generic	dbias_dgelu	col	16384x5120	long-seq hidden	134.9	130.6	1.033x	134.9	129.7	1.040x
generic	dbias_dgelu	col	8192x8192	Llama-3-70B hidden	109.3	105.8	1.033x	108.3	104.7	1.035x
generic	dbias_dgelu	col	4096x14336	Llama-3-8B FFN	95.5	93.2	1.024x	95.4	92.7	1.029x
generic	dbias_dgelu	col	2048x12288	GPT-3-175B hidden	43.9	43.4	1.011x	43.2	42.5	1.016x
generic	dbias_dgelu	col	4096x4096	Llama-3-8B hidden	31.4	31.4	1.000x	29.2	28.4	1.028x
generic	dbias_dgelu	row	2048x12288	GPT-3-175B hidden	47.7	44.4	1.074x	47.1	43.6	1.080x
generic	dbias_dgelu	row	4096x4096	Llama-3-8B hidden	33.9	32.1	1.058x	31.9	29.4	1.084x
generic	dbias_dgelu	row	4096x14336	Llama-3-8B FFN	101.2	98.3	1.030x	100.6	97.6	1.031x
generic	dbias_dgelu	row	16384x5120	long-seq hidden	141.5	137.8	1.027x	140.3	139.2	1.008x
generic	dbias_dgelu	row	8192x8192	Llama-3-70B hidden	114.8	111.9	1.026x	114.6	111.0	1.032x
generic	dbias_dgelu	row	8192x28672	Llama-3-70B FFN	382.8	380.3	1.007x	382.4	380.7	1.004x
generic	dgelu	both	8192x8192	Llama-3-70B hidden	162.5	161.0	1.009x	162.3	161.7	1.003x
generic	dgelu	both	8192x28672	Llama-3-70B FFN	552.4	547.4	1.009x	553.6	547.5	1.011x
generic	dgelu	both	2048x12288	GPT-3-175B hidden	64.8	64.6	1.003x	64.1	63.9	1.003x
generic	dgelu	both	4096x14336	Llama-3-8B FFN	142.8	143.0	0.999x	143.8	142.6	1.009x
generic	dgelu	both	16384x5120	long-seq hidden	201.1	201.4	0.998x	201.5	200.6	1.005x
generic	dgelu	both	4096x4096	Llama-3-8B hidden	45.5	45.7	0.994x	44.5	44.7	0.996x
generic	dgelu	col	8192x28672	Llama-3-70B FFN	354.4	338.8	1.046x	354.3	333.5	1.062x
generic	dgelu	col	4096x14336	Llama-3-8B FFN	91.3	88.6	1.031x	90.5	87.2	1.038x
generic	dgelu	col	8192x8192	Llama-3-70B hidden	103.6	100.7	1.029x	102.6	99.4	1.033x
generic	dgelu	col	2048x12288	GPT-3-175B hidden	42.7	41.6	1.026x	42.1	40.6	1.037x
generic	dgelu	col	4096x4096	Llama-3-8B hidden	30.8	30.0	1.026x	29.1	27.7	1.050x
generic	dgelu	col	16384x5120	long-seq hidden	127.9	124.8	1.025x	127.0	122.6	1.036x
generic	dgelu	row	2048x12288	GPT-3-175B hidden	44.3	42.0	1.055x	43.4	41.7	1.041x
generic	dgelu	row	4096x4096	Llama-3-8B hidden	31.7	30.0	1.055x	29.3	28.0	1.046x
generic	dgelu	row	16384x5120	long-seq hidden	134.1	127.6	1.051x	127.9	122.6	1.043x
generic	dgelu	row	8192x28672	Llama-3-70B FFN	346.4	330.4	1.048x	344.4	330.4	1.042x
generic	dgelu	row	8192x8192	Llama-3-70B hidden	105.3	100.8	1.045x	103.8	99.8	1.040x
generic	dgelu	row	4096x14336	Llama-3-8B FFN	93.0	89.0	1.045x	91.5	88.3	1.036x
generic	dsilu	both	2048x12288	GPT-3-175B hidden	63.3	62.1	1.019x	62.2	61.2	1.017x
generic	dsilu	both	4096x4096	Llama-3-8B hidden	44.3	43.8	1.011x	41.4	41.5	0.997x
generic	dsilu	both	4096x14336	Llama-3-8B FFN	135.9	136.4	0.997x	135.4	135.2	1.001x
generic	dsilu	both	8192x8192	Llama-3-70B hidden	153.9	154.9	0.994x	153.2	152.5	1.005x
generic	dsilu	both	16384x5120	long-seq hidden	190.5	191.9	0.993x	188.8	189.6	0.996x
generic	dsilu	both	8192x28672	Llama-3-70B FFN	522.2	528.7	0.988x	521.6	525.1	0.993x
generic	dsilu	col	8192x28672	Llama-3-70B FFN	389.9	375.0	1.040x	386.2	380.5	1.015x
generic	dsilu	col	16384x5120	long-seq hidden	140.3	136.4	1.028x	139.3	140.5	0.992x
generic	dsilu	col	2048x12288	GPT-3-175B hidden	46.8	45.7	1.024x	45.3	44.3	1.023x
generic	dsilu	col	4096x14336	Llama-3-8B FFN	99.5	97.3	1.022x	98.2	96.5	1.018x
generic	dsilu	col	8192x8192	Llama-3-70B hidden	112.9	110.6	1.021x	111.2	109.7	1.014x
generic	dsilu	col	4096x4096	Llama-3-8B hidden	32.4	32.0	1.014x	28.0	27.5	1.019x
generic	dsilu	row	8192x28672	Llama-3-70B FFN	364.0	348.8	1.044x	364.0	346.6	1.050x
generic	dsilu	row	16384x5120	long-seq hidden	132.9	129.4	1.027x	131.5	127.7	1.029x
generic	dsilu	row	2048x12288	GPT-3-175B hidden	45.2	44.0	1.027x	44.0	42.8	1.028x
generic	dsilu	row	8192x8192	Llama-3-70B hidden	107.6	104.8	1.027x	106.5	103.8	1.026x
generic	dsilu	row	4096x14336	Llama-3-8B FFN	95.3	93.1	1.024x	94.2	91.5	1.030x
generic	dsilu	row	4096x4096	Llama-3-8B hidden	31.6	32.3	0.979x	27.5	26.8	1.026x
generic	gelu	both	16384x5120	long-seq hidden	143.5	138.0	1.040x	143.5	138.0	1.040x
generic	gelu	both	4096x14336	Llama-3-8B FFN	101.8	98.3	1.036x	101.5	97.6	1.040x
generic	gelu	both	8192x28672	Llama-3-70B FFN	392.1	378.9	1.035x	392.4	378.7	1.036x
generic	gelu	both	8192x8192	Llama-3-70B hidden	115.6	111.6	1.035x	115.8	111.2	1.041x
generic	gelu	both	2048x12288	GPT-3-175B hidden	46.6	45.9	1.014x	46.4	45.0	1.031x
generic	gelu	both	4096x4096	Llama-3-8B hidden	31.1	33.2	0.937x	29.9	31.5	0.951x
generic	gelu	col	16384x5120	long-seq hidden	95.2	90.9	1.047x	90.8	90.6	1.002x
generic	gelu	col	4096x4096	Llama-3-8B hidden	22.7	22.5	1.011x	21.2	20.8	1.016x
generic	gelu	col	8192x8192	Llama-3-70B hidden	74.2	73.8	1.006x	73.7	73.5	1.002x
generic	gelu	col	2048x12288	GPT-3-175B hidden	31.4	31.3	1.003x	29.9	29.7	1.005x
generic	gelu	col	8192x28672	Llama-3-70B FFN	244.9	244.8	1.000x	244.4	244.2	1.001x
generic	gelu	col	4096x14336	Llama-3-8B FFN	65.5	65.6	0.999x	65.2	64.9	1.004x
generic	gelu	row	8192x28672	Llama-3-70B FFN	261.0	241.3	1.082x	260.5	240.5	1.083x
generic	gelu	row	16384x5120	long-seq hidden	96.8	89.6	1.081x	96.1	89.2	1.078x
generic	gelu	row	8192x8192	Llama-3-70B hidden	78.5	72.7	1.079x	77.8	72.4	1.074x
generic	gelu	row	4096x14336	Llama-3-8B FFN	69.5	64.5	1.078x	68.9	63.8	1.080x
generic	gelu	row	4096x4096	Llama-3-8B hidden	23.4	22.0	1.060x	21.7	20.4	1.063x
generic	gelu	row	2048x12288	GPT-3-175B hidden	32.4	30.7	1.055x	31.1	29.1	1.069x
generic	plain	col	8192x8192	Llama-3-70B hidden	35.2	32.9	1.070x	34.4	29.7	1.158x
generic	plain	col	8192x28672	Llama-3-70B FFN	111.2	104.2	1.067x	110.0	102.4	1.074x
generic	plain	col	4096x4096	Llama-3-8B hidden	11.8	11.2	1.062x	9.9	7.9	1.247x
generic	plain	col	2048x12288	GPT-3-175B hidden	15.9	15.0	1.059x	13.7	10.8	1.265x
generic	plain	col	16384x5120	long-seq hidden	42.8	40.4	1.058x	42.0	37.3	1.126x
generic	plain	col	4096x14336	Llama-3-8B FFN	31.3	29.8	1.049x	30.7	25.7	1.195x
nvfp4	nvfp42d	both	4096x4096	Llama-3-8B hidden	15.3	14.7	1.046x	15.1	14.6	1.040x
nvfp4	nvfp42d	both	2048x12288	GPT-3-175B hidden	21.9	21.2	1.037x	21.2	20.4	1.037x
nvfp4	nvfp42d	both	4096x14336	Llama-3-8B FFN	45.4	44.4	1.022x	44.9	44.2	1.016x
nvfp4	nvfp42d	both	8192x8192	Llama-3-70B hidden	51.0	49.9	1.021x	51.0	49.9	1.021x
nvfp4	nvfp42d	both	16384x5120	long-seq hidden	65.0	63.7	1.020x	64.7	63.8	1.014x
nvfp4	nvfp42d	both	8192x28672	Llama-3-70B FFN	188.5	188.7	0.999x	188.2	188.0	1.001x
nvfp4	nvfp42d	col	4096x4096	Llama-3-8B hidden	11.7	11.1	1.059x	11.5	10.8	1.065x
nvfp4	nvfp42d	col	2048x12288	GPT-3-175B hidden	16.5	15.8	1.042x	16.0	14.9	1.073x
nvfp4	nvfp42d	col	16384x5120	long-seq hidden	48.5	47.3	1.025x	48.7	47.4	1.027x
nvfp4	nvfp42d	col	4096x14336	Llama-3-8B FFN	34.0	33.1	1.025x	33.4	32.7	1.020x
nvfp4	nvfp42d	col	8192x8192	Llama-3-70B hidden	40.5	39.6	1.023x	40.1	39.6	1.013x
nvfp4	nvfp42d	col	8192x28672	Llama-3-70B FFN	149.6	147.6	1.014x	149.4	147.2	1.015x
nvfp4	nvfp42d	row	16384x5120	long-seq hidden	37.5	36.3	1.033x	37.5	36.3	1.035x
nvfp4	nvfp42d	row	8192x28672	Llama-3-70B FFN	95.9	92.8	1.033x	96.0	93.2	1.031x
nvfp4	nvfp42d	row	2048x12288	GPT-3-175B hidden	14.8	14.4	1.031x	13.9	13.6	1.025x
nvfp4	nvfp42d	row	8192x8192	Llama-3-70B hidden	31.5	30.7	1.026x	31.4	30.7	1.025x
nvfp4	nvfp42d	row	4096x4096	Llama-3-8B hidden	10.2	10.0	1.019x	10.0	9.8	1.020x
nvfp4	nvfp42d	row	4096x14336	Llama-3-8B FFN	28.5	28.0	1.018x	27.9	27.1	1.028x
specialized	plain	both	8192x8192	Llama-3-70B hidden	47.2	46.6	1.012x	44.2	44.0	1.005x
specialized	plain	both	2048x12288	GPT-3-175B hidden	20.1	20.0	1.006x	16.2	15.0	1.083x
specialized	plain	both	16384x5120	long-seq hidden	57.3	57.1	1.004x	54.3	53.9	1.007x
specialized	plain	both	8192x28672	Llama-3-70B FFN	150.0	149.5	1.003x	145.4	146.1	0.995x
specialized	plain	both	4096x14336	Llama-3-8B FFN	42.3	42.2	1.003x	39.6	39.4	1.004x
specialized	plain	both	4096x4096	Llama-3-8B hidden	15.9	15.9	1.002x	11.7	10.4	1.118x
specialized	plain	row	4096x4096	Llama-3-8B hidden	10.0	10.0	1.006x	7.4	7.3	1.002x
specialized	plain	row	16384x5120	long-seq hidden	38.8	38.6	1.004x	36.7	36.7	0.999x
specialized	plain	row	4096x14336	Llama-3-8B FFN	28.2	28.1	1.003x	25.1	25.0	1.001x
specialized	plain	row	8192x28672	Llama-3-70B FFN	101.1	101.1	1.001x	98.7	98.7	1.000x
specialized	plain	row	8192x8192	Llama-3-70B hidden	31.3	31.4	0.999x	28.6	28.5	1.002x
specialized	plain	row	2048x12288	GPT-3-175B hidden	14.0	14.0	0.999x	10.1	10.1	0.995x

Patched but not measurable

Three further sites carry the identical anti-pattern but cannot be reached, so they are fixed for consistency rather than measured. Leaving them would let the pattern be reintroduced if those paths are ever enabled.

site	why unreachable
cast/mxfp8/specialized/quantize_mxfp8.cuh (warp-specialized variant)	guarded by enable_if_t<CastTraits::_use_warp_specialization>, and that flag is false; the library contains only the 8 non-warp-specialized instantiations
cast/nvfp4/quantize_transpose_nvfp4.cuh (1D kernel)	quantize_transpose<false> is only called for bf16, and its first statement diverts every bf16 1D call to quantize_transpose_tuned_1D
cast/nvfp4/group_quantize_transpose_nvfp4.cuh	the PyTorch grouped path requires RHT ("graph safe grouped quant kernel for non-RHT path is not ready yet"), which routes to the hadamard fusion kernels instead

hadamard_transform/*.cu carries the same pattern at 3 more sites but was left alone: across all 108 hadamard kernels there are only 72 generic LD/ST in total (worst kernel: 6), so there is nothing measurable to win.

Method

• Both arms are full builds of the same tree; each profile records the md5 of the libtransformer_engine.so actually mapped into the process (read from /proc/self/maps after load), so a stale build cannot masquerade as a fresh one.
• cold L2: 512 MB scratch buffer zeroed before every iteration. warm L2: no flush, back-to-back steady state. Both are pure kernel time either way; the cache state only changes what the kernel is bottlenecked on.
• Numerics are bitwise identical across the two builds (72 MXFP8 and 18 NVFP4 combo/shape/direction configurations hashed). This is a pure codegen change.
• Gains track how scalar the shared-memory access pattern is, because the generic-address penalty is charged per instruction: the MXFP8 colwise path moves 1.5 bytes/instruction and gains most, while already-vectorized paths (10 bytes/instruction) barely move.
• Reproduce with ./run_final.sh && python gen_final_report.py.