[ROCm] Optimize AMD normalization backward kernel by using tiled and two pass implementations where they are most efficient. by anatoliylitv · Pull Request #3424 · ROCm/pytorch

anatoliylitv · 2026-07-11T22:03:49Z

Hugging Face model dropped 5-10% after switching to tiled kernel (10+ model and tests).

Implemented mixed approach, using combination of Tiled and Two Pass:

Hugging Face (huggingface_bart) performance is back:
Legacy Two pass performance=1652, 1644
Tiled performance =1574, 1568
Mixed performance =1649, 1652

Synthetic reproducer where performance from chess board like became uniform, keeping benefits from both implementation:

Layer Norm Backward Benchmark: Tiled only vs Tiled+Two pass

Device: AMD Instinct MI350X | Warmup: 20 | Iters: 100 | Runs averaged: 3
Tiled only: PID 3151521 (built from PID 3155055)
Tiled+Two pass: PID 3160321

Summary (avg µs over 3 runs)

Benchmark	Op	Shape	dtype	Branch	Tiled only	Tiled+Two pass	Δ (µs)	Speedup	Winner
tile8_small	layer_norm	(32, 512)	float16	Tile-8	7.94	7.97	+0.03	1.00×	~tie
tile64_medium	layer_norm	(96, 768)	float16	Tile-64	20.71	20.72	+0.01	1.00×	~tie
tile128_medium	layer_norm	(192, 1024)	bfloat16	Tile-128	8.20	12.47	+4.27	0.66×	Tiled only
tile256_large	layer_norm	(4096, 1024)	float16	Tile-256	31.07	12.52	-18.55	2.48×	Two pass
tile256_bert	layer_norm	(1024, 768)	float16	Tile-256 BERT	11.15	12.64	+1.49	0.88×	Tiled only
tile256_large_bf16	layer_norm	(4096, 1024)	bfloat16	Tile-256	31.28	12.46	-17.82	2.51×	Two pass
tile256_large_fp32	layer_norm	(4096, 1024)	float32	Tile-256	30.87	15.82	-15.05	1.95×	Two pass
two_pass_huge_M	layer_norm	(131072, 64)	float16	Two-pass M-parallel	26.48	26.69	+0.21	0.99×	~tie
llm_hidden_4096	layer_norm	(32, 4096, 4096)	bfloat16	Tile-256 LLM	509.65	510.26	+0.61	1.00×	~tie
gpt2_style	layer_norm	(8, 1024, 768)	float16	Tile-256 GPT2	58.22	16.62	-41.60	3.50×	Two pass
rms_tile256	rms_norm	(4096, 1024)	float16	Tile-256 rms	22.34	11.60	-10.74	1.93×	Two pass
rms_llm	rms_norm	(32, 4096, 4096)	bfloat16	Tile-256 rms LLM	476.18	473.47	-2.71	1.01×	~tie
rms_two_pass	rms_norm	(131072, 64)	float16	Two-pass rms	16.73	16.61	-0.12	1.01×	~tie

Δ = Tiled only − Tiled+Two pass (negative = Two pass faster). Speedup = Tiled only / Tiled+Two pass.

cc @jeffdaily @sunway513 @jithunnair-amd @pruthvistony @ROCmSupport @jataylo @hongxiayang @naromero77amd @pragupta @jerrymannil @xinyazhang

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

… it accordingly with tiled one where performance is better.

rocm-repo-management-api · 2026-07-11T22:15:16Z

Jenkins build for 20a3ce6fac4bcf46d9846f98270b85ce3346efc5 commit finished as FAILURE
Links: Pipeline Overview / Build artifacts / Test Results

anatoliylitv · 2026-07-30T20:56:32Z

Results of related unit tests run:
root@gbt350-odcdh5-wbc2-b:/myworkspace/pytorch# python test/test_nn.py -v -k "ayerormcuda"
E-001h rocSHMEM Could not open libnuma. Returning NUMAWrapper@src/gda/numa_wrapper.cpp:48
W0730 20:52:27.856000 3209662 /myworkspace/pytorch/torch/_native/cutedsl_utils.py:55] CuTeDSL operators require optional Python packages nvidia-cutlass-dsl and apache-tvm-ffi; missing optional dependency nvidia_cutlass_dsl (importlib.util.find_spec(nvidia_cutlass_dsl) failed)
test_LayerNorm_3d_no_affine_large_feature_cuda (main.TestNN.test_LayerNorm_3d_no_affine_large_feature_cuda) ... /myworkspace/pytorch/torch/backends/cudnn/init.py:175: UserWarning: cuDNN Benchmark limit is not supported in MIOpen and will have no effect. (Triggered internally at /myworkspace/pytorch/torch/csrc/cuda/Module.cpp:1990.)
torch._C._cuda_set_cudnn_benchmark_limit(_benchmark_limit)
ok
test_LayerNorm_3d_no_affine_large_feature_eval_cuda (main.TestNN.test_LayerNorm_3d_no_affine_large_feature_eval_cuda) ... ok
test_LayerNorm_general_cuda (main.TestNNDeviceTypeCUDA.test_LayerNorm_general_cuda) ... ok
test_LayerNorm_numeric_cuda (main.TestNNDeviceTypeCUDA.test_LayerNorm_numeric_cuda) ... ok
test_layernorm_half_precision_cuda (main.TestNNDeviceTypeCUDA.test_layernorm_half_precision_cuda) ... ok
test_layernorm_weight_bias_cuda (main.TestNNDeviceTypeCUDA.test_layernorm_weight_bias_cuda) ... ok

Ran 6 tests in 6.364s

OK
root@gbt350-odcdh5-wbc2-b:/myworkspace/pytorch# python test/test_nn.py -v -k layer_norm_backwards_eps
E-001h rocSHMEM Could not open libnuma. Returning NUMAWrapper@src/gda/numa_wrapper.cpp:48
W0730 20:53:14.651000 3210323 /myworkspace/pytorch/torch/_native/cutedsl_utils.py:55] CuTeDSL operators require optional Python packages nvidia-cutlass-dsl and apache-tvm-ffi; missing optional dependency nvidia_cutlass_dsl (importlib.util.find_spec(nvidia_cutlass_dsl) failed)
test_layer_norm_backwards_eps (main.TestNN.test_layer_norm_backwards_eps) ... ok

Ran 1 test in 3.070s

OK
root@gbt350-odcdh5-wbc2-b:/myworkspace/pytorch# python test/test_ops.py -v -k layer_norm -k rms_norm
E-001h rocSHMEM Could not open libnuma. Returning NUMAWrapper@src/gda/numa_wrapper.cpp:48
W0730 20:53:32.777000 3210963 /myworkspace/pytorch/torch/_native/cutedsl_utils.py:55] CuTeDSL operators require optional Python packages nvidia-cutlass-dsl and apache-tvm-ffi; missing optional dependency nvidia_cutlass_dsl (importlib.util.find_spec(nvidia_cutlass_dsl) failed)
test_dtypes__refs_native_layer_norm_cpu (main.TestCommonCPU.test_dtypes__refs_native_layer_norm_cpu) ... /myworkspace/pytorch/torch/_refs/init.py:3340: UserWarning: var_mean(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
biased_var, mean = torch.var_mean(
ok
test_dtypes__refs_nn_functional_layer_norm_cpu (main.TestCommonCPU.test_dtypes__refs_nn_functional_layer_norm_cpu) ... ok
test_dtypes_native_layer_norm_cpu (main.TestCommonCPU.test_dtypes_native_layer_norm_cpu) ... ok
test_dtypes_nn_functional_layer_norm_cpu (main.TestCommonCPU.test_dtypes_nn_functional_layer_norm_cpu) ... ok
test_dtypes_nn_functional_rms_norm_cpu (main.TestCommonCPU.test_dtypes_nn_functional_rms_norm_cpu) ... ok
test_errors_native_layer_norm_cpu (main.TestCommonCPU.test_errors_native_layer_norm_cpu) ... ok
test_errors_nn_functional_rms_norm_cpu (main.TestCommonCPU.test_errors_nn_functional_rms_norm_cpu) ... ok
test_multiple_devices_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_multiple_devices_native_layer_norm_cpu_float32) ... skipped "Only runs on ['cuda', 'xpu']"
test_multiple_devices_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_multiple_devices_nn_functional_layer_norm_cpu_float32) ... skipped "Only runs on ['cuda', 'xpu']"
test_multiple_devices_nn_functional_rms_norm_cpu_float32 (main.TestCommonCPU.test_multiple_devices_nn_functional_rms_norm_cpu_float32) ... skipped "Only runs on ['cuda', 'xpu']"
test_noncontiguous_samples_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_noncontiguous_samples_native_layer_norm_cpu_float32) ... /myworkspace/pytorch/torch/backends/cudnn/init.py:175: UserWarning: cuDNN Benchmark limit is not supported in MIOpen and will have no effect. (Triggered internally at /myworkspace/pytorch/torch/csrc/cuda/Module.cpp:1990.)
torch._C._cuda_set_cudnn_benchmark_limit(_benchmark_limit)
ok
test_noncontiguous_samples_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_noncontiguous_samples_nn_functional_layer_norm_cpu_float32) ... ok
test_noncontiguous_samples_nn_functional_rms_norm_cpu_complex64 (main.TestCommonCPU.test_noncontiguous_samples_nn_functional_rms_norm_cpu_complex64) ... ok
test_noncontiguous_samples_nn_functional_rms_norm_cpu_float32 (main.TestCommonCPU.test_noncontiguous_samples_nn_functional_rms_norm_cpu_float32) ... ok
test_numpy_ref_native_layer_norm_cpu_float64 (main.TestCommonCPU.test_numpy_ref_native_layer_norm_cpu_float64) ... skipped 'XXX: raises tensor-likes are not close.'
test_numpy_ref_nn_functional_layer_norm_cpu_float64 (main.TestCommonCPU.test_numpy_ref_nn_functional_layer_norm_cpu_float64) ... skipped 'XXX: raises tensor-likes are not close.'
test_numpy_ref_nn_functional_rms_norm_cpu_complex128 (main.TestCommonCPU.test_numpy_ref_nn_functional_rms_norm_cpu_complex128) ... skipped 'XXX: raises tensor-likes are not close.'
test_numpy_ref_nn_functional_rms_norm_cpu_float64 (main.TestCommonCPU.test_numpy_ref_nn_functional_rms_norm_cpu_float64) ... skipped 'XXX: raises tensor-likes are not close.'
test_out__refs_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_out__refs_native_layer_norm_cpu_float32) ... ok
test_out__refs_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_out__refs_nn_functional_layer_norm_cpu_float32) ... ok
test_out_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_out_native_layer_norm_cpu_float32) ... ok
test_out_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_out_nn_functional_layer_norm_cpu_float32) ... ok
test_out_nn_functional_rms_norm_cpu_float32 (main.TestCommonCPU.test_out_nn_functional_rms_norm_cpu_float32) ... ok
test_out_warning__refs_native_layer_norm_cpu (main.TestCommonCPU.test_out_warning__refs_native_layer_norm_cpu) ... ok
test_out_warning__refs_nn_functional_layer_norm_cpu (main.TestCommonCPU.test_out_warning__refs_nn_functional_layer_norm_cpu) ... ok
test_out_warning_native_layer_norm_cpu (main.TestCommonCPU.test_out_warning_native_layer_norm_cpu) ... ok
test_out_warning_nn_functional_layer_norm_cpu (main.TestCommonCPU.test_out_warning_nn_functional_layer_norm_cpu) ... ok
test_out_warning_nn_functional_rms_norm_cpu (main.TestCommonCPU.test_out_warning_nn_functional_rms_norm_cpu) ... ok
test_python_ref__refs_native_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref__refs_native_layer_norm_cpu_bfloat16) ... /myworkspace/pytorch/torch/_prims/init.py:2403: UserWarning: var(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
return torch.var(input, dim=dim, correction=correction, **kwargs)
ok
test_python_ref__refs_native_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref__refs_native_layer_norm_cpu_float16) ... ok
test_python_ref__refs_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref__refs_native_layer_norm_cpu_float32) ... skipped 'Skipped!'
test_python_ref__refs_native_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref__refs_native_layer_norm_cpu_float64) ... ok
test_python_ref__refs_nn_functional_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref__refs_nn_functional_layer_norm_cpu_bfloat16) ... ok
test_python_ref__refs_nn_functional_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref__refs_nn_functional_layer_norm_cpu_float16) ... ok
test_python_ref__refs_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref__refs_nn_functional_layer_norm_cpu_float32) ... skipped 'Skipped!'
test_python_ref__refs_nn_functional_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref__refs_nn_functional_layer_norm_cpu_float64) ... ok
test_python_ref_errors__refs_native_layer_norm_cpu (main.TestCommonCPU.test_python_ref_errors__refs_native_layer_norm_cpu) ... ok
test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_bfloat16) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float16 (main.TestCommonCPU.test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float16) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float32 (main.TestCommonCPU.test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float32) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float64 (main.TestCommonCPU.test_python_ref_executor__refs_native_layer_norm_executor_aten_cpu_float64) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_bfloat16) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float16 (main.TestCommonCPU.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float16) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float32 (main.TestCommonCPU.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float32) ... skipped 'Only runs on cuda'
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float64 (main.TestCommonCPU.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cpu_float64) ... skipped 'Only runs on cuda'
test_python_ref_meta__refs_native_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_meta__refs_native_layer_norm_cpu_bfloat16) ... /myworkspace/pytorch/torch/_refs/init.py:3340: UserWarning: var_mean(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
biased_var, mean = torch.var_mean(
ok
test_python_ref_meta__refs_native_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref_meta__refs_native_layer_norm_cpu_float16) ... ok
test_python_ref_meta__refs_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref_meta__refs_native_layer_norm_cpu_float32) ... ok
test_python_ref_meta__refs_native_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref_meta__refs_native_layer_norm_cpu_float64) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_meta__refs_nn_functional_layer_norm_cpu_bfloat16) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float16) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float32) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref_meta__refs_nn_functional_layer_norm_cpu_float64) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_native_layer_norm_cpu_bfloat16) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float16) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float32) ... skipped 'Skipped!'
test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_native_layer_norm_cpu_float64) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_bfloat16 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_bfloat16) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float16 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float16) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float32) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float64 (main.TestCommonCPU.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cpu_float64) ... ok
test_variant_consistency_eager_native_layer_norm_cpu_float32 (main.TestCommonCPU.test_variant_consistency_eager_native_layer_norm_cpu_float32) ... ok
test_variant_consistency_eager_nn_functional_layer_norm_cpu_float32 (main.TestCommonCPU.test_variant_consistency_eager_nn_functional_layer_norm_cpu_float32) ... ok
test_variant_consistency_eager_nn_functional_rms_norm_cpu_complex64 (main.TestCommonCPU.test_variant_consistency_eager_nn_functional_rms_norm_cpu_complex64) ... ok
test_variant_consistency_eager_nn_functional_rms_norm_cpu_float32 (main.TestCommonCPU.test_variant_consistency_eager_nn_functional_rms_norm_cpu_float32) ... ok
test_dtypes__refs_native_layer_norm_cuda (main.TestCommonCUDA.test_dtypes__refs_native_layer_norm_cuda) ... ok
test_dtypes__refs_nn_functional_layer_norm_cuda (main.TestCommonCUDA.test_dtypes__refs_nn_functional_layer_norm_cuda) ... Some dtypes for _refs.nn.functional.layer_norm on device type cuda are only partially supported!
The following dtypes only worked on some samples during forward: {torch.int16, torch.int32, torch.int64, torch.uint8, torch.bool, torch.int8}.

ok
test_dtypes_native_layer_norm_cuda (main.TestCommonCUDA.test_dtypes_native_layer_norm_cuda) ... Some dtypes for native_layer_norm on device type cuda are only partially supported!
The following dtypes only worked on some samples during forward: {torch.int16, torch.int32, torch.int64, torch.uint8, torch.bool, torch.int8}.

ok
test_dtypes_nn_functional_layer_norm_cuda (main.TestCommonCUDA.test_dtypes_nn_functional_layer_norm_cuda) ... Some dtypes for nn.functional.layer_norm on device type cuda are only partially supported!
The following dtypes only worked on some samples during forward: {torch.int16, torch.int32, torch.int64, torch.uint8, torch.bool, torch.int8}.

ok
test_dtypes_nn_functional_rms_norm_cuda (main.TestCommonCUDA.test_dtypes_nn_functional_rms_norm_cuda) ... Some dtypes for nn.functional.rms_norm on device type cuda are only partially supported!
The following dtypes only worked on some samples during forward: {torch.int16, torch.int32, torch.int64, torch.uint8, torch.bool, torch.int8}.

ok
test_errors_native_layer_norm_cuda (main.TestCommonCUDA.test_errors_native_layer_norm_cuda) ... ok
test_errors_nn_functional_rms_norm_cuda (main.TestCommonCUDA.test_errors_nn_functional_rms_norm_cuda) ... ok
test_multiple_devices_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_multiple_devices_native_layer_norm_cuda_float32) ... ok
test_multiple_devices_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_multiple_devices_nn_functional_layer_norm_cuda_float32) ... ok
test_multiple_devices_nn_functional_rms_norm_cuda_float32 (main.TestCommonCUDA.test_multiple_devices_nn_functional_rms_norm_cuda_float32) ... ok
test_noncontiguous_samples_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_noncontiguous_samples_native_layer_norm_cuda_float32) ... ok
test_noncontiguous_samples_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_noncontiguous_samples_nn_functional_layer_norm_cuda_float32) ... ok
test_noncontiguous_samples_nn_functional_rms_norm_cuda_complex64 (main.TestCommonCUDA.test_noncontiguous_samples_nn_functional_rms_norm_cuda_complex64) ... ok
test_noncontiguous_samples_nn_functional_rms_norm_cuda_float32 (main.TestCommonCUDA.test_noncontiguous_samples_nn_functional_rms_norm_cuda_float32) ... ok
test_numpy_ref_native_layer_norm_cuda_float64 (main.TestCommonCUDA.test_numpy_ref_native_layer_norm_cuda_float64) ... ok
test_numpy_ref_nn_functional_layer_norm_cuda_float64 (main.TestCommonCUDA.test_numpy_ref_nn_functional_layer_norm_cuda_float64) ... ok
test_numpy_ref_nn_functional_rms_norm_cuda_complex128 (main.TestCommonCUDA.test_numpy_ref_nn_functional_rms_norm_cuda_complex128) ... ok
test_numpy_ref_nn_functional_rms_norm_cuda_float64 (main.TestCommonCUDA.test_numpy_ref_nn_functional_rms_norm_cuda_float64) ... ok
test_out__refs_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_out__refs_native_layer_norm_cuda_float32) ... ok
test_out__refs_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_out__refs_nn_functional_layer_norm_cuda_float32) ... ok
test_out_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_out_native_layer_norm_cuda_float32) ... ok
test_out_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_out_nn_functional_layer_norm_cuda_float32) ... ok
test_out_nn_functional_rms_norm_cuda_float32 (main.TestCommonCUDA.test_out_nn_functional_rms_norm_cuda_float32) ... ok
test_out_warning__refs_native_layer_norm_cuda (main.TestCommonCUDA.test_out_warning__refs_native_layer_norm_cuda) ... ok
test_out_warning__refs_nn_functional_layer_norm_cuda (main.TestCommonCUDA.test_out_warning__refs_nn_functional_layer_norm_cuda) ... ok
test_out_warning_native_layer_norm_cuda (main.TestCommonCUDA.test_out_warning_native_layer_norm_cuda) ... ok
test_out_warning_nn_functional_layer_norm_cuda (main.TestCommonCUDA.test_out_warning_nn_functional_layer_norm_cuda) ... ok
test_out_warning_nn_functional_rms_norm_cuda (main.TestCommonCUDA.test_out_warning_nn_functional_rms_norm_cuda) ... ok
test_python_ref__refs_native_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref__refs_native_layer_norm_cuda_bfloat16) ... /myworkspace/pytorch/torch/_prims/init.py:2403: UserWarning: var(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
return torch.var(input, dim=dim, correction=correction, **kwargs)
ok
test_python_ref__refs_native_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref__refs_native_layer_norm_cuda_float16) ... ok
test_python_ref__refs_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref__refs_native_layer_norm_cuda_float32) ... ok
test_python_ref__refs_native_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref__refs_native_layer_norm_cuda_float64) ... ok
test_python_ref__refs_nn_functional_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref__refs_nn_functional_layer_norm_cuda_bfloat16) ... ok
test_python_ref__refs_nn_functional_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref__refs_nn_functional_layer_norm_cuda_float16) ... ok
test_python_ref__refs_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref__refs_nn_functional_layer_norm_cuda_float32) ... ok
test_python_ref__refs_nn_functional_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref__refs_nn_functional_layer_norm_cuda_float64) ... ok
test_python_ref_errors__refs_native_layer_norm_cuda (main.TestCommonCUDA.test_python_ref_errors__refs_native_layer_norm_cuda) ... ok
test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_bfloat16) ... ok
test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float16 (main.TestCommonCUDA.test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float16) ... ok
test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float32 (main.TestCommonCUDA.test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float32) ... ok
test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float64 (main.TestCommonCUDA.test_python_ref_executor__refs_native_layer_norm_executor_aten_cuda_float64) ... ok
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_bfloat16) ... ok
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float16 (main.TestCommonCUDA.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float16) ... ok
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float32 (main.TestCommonCUDA.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float32) ... ok
test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float64 (main.TestCommonCUDA.test_python_ref_executor__refs_nn_functional_layer_norm_executor_aten_cuda_float64) ... ok
test_python_ref_meta__refs_native_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_meta__refs_native_layer_norm_cuda_bfloat16) ... /myworkspace/pytorch/torch/_refs/init.py:3340: UserWarning: var_mean(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
biased_var, mean = torch.var_mean(
ok
test_python_ref_meta__refs_native_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref_meta__refs_native_layer_norm_cuda_float16) ... ok
test_python_ref_meta__refs_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref_meta__refs_native_layer_norm_cuda_float32) ... ok
test_python_ref_meta__refs_native_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref_meta__refs_native_layer_norm_cuda_float64) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_meta__refs_nn_functional_layer_norm_cuda_bfloat16) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float16) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float32) ... ok
test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref_meta__refs_nn_functional_layer_norm_cuda_float64) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_native_layer_norm_cuda_bfloat16) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float16) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float32) ... ok
test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_native_layer_norm_cuda_float64) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_bfloat16 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_bfloat16) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float16 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float16) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float32) ... ok
test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float64 (main.TestCommonCUDA.test_python_ref_torch_fallback__refs_nn_functional_layer_norm_cuda_float64) ... ok
test_variant_consistency_eager_native_layer_norm_cuda_float32 (main.TestCommonCUDA.test_variant_consistency_eager_native_layer_norm_cuda_float32) ... ok
test_variant_consistency_eager_nn_functional_layer_norm_cuda_float32 (main.TestCommonCUDA.test_variant_consistency_eager_nn_functional_layer_norm_cuda_float32) ... ok
test_variant_consistency_eager_nn_functional_rms_norm_cuda_complex64 (main.TestCommonCUDA.test_variant_consistency_eager_nn_functional_rms_norm_cuda_complex64) ... ok
test_variant_consistency_eager_nn_functional_rms_norm_cuda_float32 (main.TestCommonCUDA.test_variant_consistency_eager_nn_functional_rms_norm_cuda_float32) ... ok
test_backward_native_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_backward_native_layer_norm_cpu_float32) ... ok
test_backward_nn_functional_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_backward_nn_functional_layer_norm_cpu_float32) ... ok
test_backward_nn_functional_rms_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_backward_nn_functional_rms_norm_cpu_float32) ... ok
test_cow_input_native_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_cow_input_native_layer_norm_cpu_float32) ... ok
test_cow_input_nn_functional_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_cow_input_nn_functional_layer_norm_cpu_float32) ... ok
test_cow_input_nn_functional_rms_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_cow_input_nn_functional_rms_norm_cpu_float32) ... ok
test_forward_ad_native_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_forward_ad_native_layer_norm_cpu_float32) ... skipped 'Does not support forward_ad'
test_forward_ad_nn_functional_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_forward_ad_nn_functional_layer_norm_cpu_float32) ... /myworkspace/pytorch/torch/jit/_script.py:1488: DeprecationWarning: torch.jit.script is deprecated. Please switch to torch.compile or torch.export.
warnings.warn(
ok
test_forward_ad_nn_functional_rms_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_forward_ad_nn_functional_rms_norm_cpu_float32) ... ok
test_operator_native_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_operator_native_layer_norm_cpu_float32) ... ok
test_operator_nn_functional_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_operator_nn_functional_layer_norm_cpu_float32) ... ok
test_operator_nn_functional_rms_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_operator_nn_functional_rms_norm_cpu_float32) ... ok
test_view_replay_native_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_view_replay_native_layer_norm_cpu_float32) ... ok
test_view_replay_nn_functional_layer_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_view_replay_nn_functional_layer_norm_cpu_float32) ... ok
test_view_replay_nn_functional_rms_norm_cpu_float32 (main.TestCompositeComplianceCPU.test_view_replay_nn_functional_rms_norm_cpu_float32) ... ok
test_backward_native_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_backward_native_layer_norm_cuda_float32) ... ok
test_backward_nn_functional_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_backward_nn_functional_layer_norm_cuda_float32) ... ok
test_backward_nn_functional_rms_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_backward_nn_functional_rms_norm_cuda_float32) ... ok
test_cow_input_native_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_cow_input_native_layer_norm_cuda_float32) ... ok
test_cow_input_nn_functional_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_cow_input_nn_functional_layer_norm_cuda_float32) ... ok
test_cow_input_nn_functional_rms_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_cow_input_nn_functional_rms_norm_cuda_float32) ... ok
test_forward_ad_native_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_forward_ad_native_layer_norm_cuda_float32) ... skipped 'Does not support forward_ad'
test_forward_ad_nn_functional_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_forward_ad_nn_functional_layer_norm_cuda_float32) ... ok
test_forward_ad_nn_functional_rms_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_forward_ad_nn_functional_rms_norm_cuda_float32) ... ok
test_operator_native_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_operator_native_layer_norm_cuda_float32) ... ok
test_operator_nn_functional_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_operator_nn_functional_layer_norm_cuda_float32) ... ok
test_operator_nn_functional_rms_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_operator_nn_functional_rms_norm_cuda_float32) ... ok
test_view_replay_native_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_view_replay_native_layer_norm_cuda_float32) ... ok
test_view_replay_nn_functional_layer_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_view_replay_nn_functional_layer_norm_cuda_float32) ... ok
test_view_replay_nn_functional_rms_norm_cuda_float32 (main.TestCompositeComplianceCUDA.test_view_replay_nn_functional_rms_norm_cuda_float32) ... ok
test_fake_autocast_native_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_autocast_native_layer_norm_cpu_float32) ... ok
test_fake_autocast_nn_functional_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_autocast_nn_functional_layer_norm_cpu_float32) ... ok
test_fake_autocast_nn_functional_rms_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_autocast_nn_functional_rms_norm_cpu_float32) ... ok
test_fake_crossref_backward_amp_native_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_amp_native_layer_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_crossref_backward_amp_nn_functional_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_amp_nn_functional_layer_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_crossref_backward_amp_nn_functional_rms_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_amp_nn_functional_rms_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_crossref_backward_no_amp_native_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_no_amp_native_layer_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_crossref_backward_no_amp_nn_functional_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_no_amp_nn_functional_layer_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_crossref_backward_no_amp_nn_functional_rms_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_crossref_backward_no_amp_nn_functional_rms_norm_cpu_float32) ... skipped 'Only runs on cuda'
test_fake_native_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_native_layer_norm_cpu_float32) ... ok
test_fake_nn_functional_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_nn_functional_layer_norm_cpu_float32) ... ok
test_fake_nn_functional_rms_norm_cpu_float32 (main.TestFakeTensorCPU.test_fake_nn_functional_rms_norm_cpu_float32) ... ok
test_pointwise_ops_native_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_pointwise_ops_native_layer_norm_cpu_float32) ... ok
test_pointwise_ops_nn_functional_layer_norm_cpu_float32 (main.TestFakeTensorCPU.test_pointwise_ops_nn_functional_layer_norm_cpu_float32) ... ok
test_pointwise_ops_nn_functional_rms_norm_cpu_float32 (main.TestFakeTensorCPU.test_pointwise_ops_nn_functional_rms_norm_cpu_float32) ... ok
test_fake_autocast_native_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_autocast_native_layer_norm_cuda_float32) ... ok
test_fake_autocast_nn_functional_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_autocast_nn_functional_layer_norm_cuda_float32) ... ok
test_fake_autocast_nn_functional_rms_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_autocast_nn_functional_rms_norm_cuda_float32) ... ok
test_fake_crossref_backward_amp_native_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_amp_native_layer_norm_cuda_float32) ... /myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
ok
test_fake_crossref_backward_amp_nn_functional_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_amp_nn_functional_layer_norm_cuda_float32) ... /myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
ok
test_fake_crossref_backward_amp_nn_functional_rms_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_amp_nn_functional_rms_norm_cuda_float32) ... /myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
/myworkspace/pytorch/test/test_ops.py:2967: FutureWarning: torch.cuda.amp.autocast(args...) is deprecated. Please use torch.amp.autocast('cuda', args...) instead.
context(),
ok
test_fake_crossref_backward_no_amp_native_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_no_amp_native_layer_norm_cuda_float32) ... ok
test_fake_crossref_backward_no_amp_nn_functional_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_no_amp_nn_functional_layer_norm_cuda_float32) ... ok
test_fake_crossref_backward_no_amp_nn_functional_rms_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_crossref_backward_no_amp_nn_functional_rms_norm_cuda_float32) ... ok
test_fake_native_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_native_layer_norm_cuda_float32) ... ok
test_fake_nn_functional_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_nn_functional_layer_norm_cuda_float32) ... ok
test_fake_nn_functional_rms_norm_cuda_float32 (main.TestFakeTensorCUDA.test_fake_nn_functional_rms_norm_cuda_float32) ... ok
test_pointwise_ops_native_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_pointwise_ops_native_layer_norm_cuda_float32) ... ok
test_pointwise_ops_nn_functional_layer_norm_cuda_float32 (main.TestFakeTensorCUDA.test_pointwise_ops_nn_functional_layer_norm_cuda_float32) ... ok
test_pointwise_ops_nn_functional_rms_norm_cuda_float32 (main.TestFakeTensorCUDA.test_pointwise_ops_nn_functional_rms_norm_cuda_float32) ... ok
test_conj_view_nn_functional_rms_norm_cpu_complex64 (main.TestMathBitsCPU.test_conj_view_nn_functional_rms_norm_cpu_complex64) ... ok
test_neg_conj_view_nn_functional_rms_norm_cpu_complex128 (main.TestMathBitsCPU.test_neg_conj_view_nn_functional_rms_norm_cpu_complex128) ... ok
test_neg_view__refs_native_layer_norm_cpu_float64 (main.TestMathBitsCPU.test_neg_view__refs_native_layer_norm_cpu_float64) ... /myworkspace/pytorch/torch/_refs/init.py:3340: UserWarning: var_mean(): degrees of freedom is <= 0. Correction should be strictly less than the reduction factor (input numel divided by output numel). (Triggered internally at /myworkspace/pytorch/aten/src/ATen/native/ReduceOps.cpp:1879.)
biased_var, mean = torch.var_mean(
ok
test_neg_view__refs_nn_functional_layer_norm_cpu_float64 (main.TestMathBitsCPU.test_neg_view__refs_nn_functional_layer_norm_cpu_float64) ... ok
test_neg_view_native_layer_norm_cpu_float64 (main.TestMathBitsCPU.test_neg_view_native_layer_norm_cpu_float64) ... ok
test_neg_view_nn_functional_layer_norm_cpu_float64 (main.TestMathBitsCPU.test_neg_view_nn_functional_layer_norm_cpu_float64) ... ok
test_neg_view_nn_functional_rms_norm_cpu_float64 (main.TestMathBitsCPU.test_neg_view_nn_functional_rms_norm_cpu_float64) ... ok
test_conj_view_nn_functional_rms_norm_cuda_complex64 (main.TestMathBitsCUDA.test_conj_view_nn_functional_rms_norm_cuda_complex64) ... ok
test_neg_conj_view_nn_functional_rms_norm_cuda_complex128 (main.TestMathBitsCUDA.test_neg_conj_view_nn_functional_rms_norm_cuda_complex128) ... ok
test_neg_view__refs_native_layer_norm_cuda_float64 (main.TestMathBitsCUDA.test_neg_view__refs_native_layer_norm_cuda_float64) ... ok
test_neg_view__refs_nn_functional_layer_norm_cuda_float64 (main.TestMathBitsCUDA.test_neg_view__refs_nn_functional_layer_norm_cuda_float64) ... ok
test_neg_view_native_layer_norm_cuda_float64 (main.TestMathBitsCUDA.test_neg_view_native_layer_norm_cuda_float64) ... ok
test_neg_view_nn_functional_layer_norm_cuda_float64 (main.TestMathBitsCUDA.test_neg_view_nn_functional_layer_norm_cuda_float64) ... ok
test_neg_view_nn_functional_rms_norm_cuda_float64 (main.TestMathBitsCUDA.test_neg_view_nn_functional_rms_norm_cuda_float64) ... ok
test_refs_are_in_decomp_table_op__refs_native_layer_norm_cpu (main.TestRefsOpsInfoCPU.test_refs_are_in_decomp_table_op__refs_native_layer_norm_cpu) ... ok
test_refs_are_in_decomp_table_op__refs_nn_functional_layer_norm_cpu (main.TestRefsOpsInfoCPU.test_refs_are_in_decomp_table_op__refs_nn_functional_layer_norm_cpu) ... ok
test_refs_are_in_python_ref_db_op__refs_native_layer_norm_cpu (main.TestRefsOpsInfoCPU.test_refs_are_in_python_ref_db_op__refs_native_layer_norm_cpu) ... ok
test_refs_are_in_python_ref_db_op__refs_nn_functional_layer_norm_cpu (main.TestRefsOpsInfoCPU.test_refs_are_in_python_ref_db_op__refs_nn_functional_layer_norm_cpu) ... ok
test_tags__refs_native_layer_norm_cpu_float32 (main.TestTagsCPU.test_tags__refs_native_layer_norm_cpu_float32) ... ok
test_tags__refs_nn_functional_layer_norm_cpu_float32 (main.TestTagsCPU.test_tags__refs_nn_functional_layer_norm_cpu_float32) ... ok
test_tags_native_layer_norm_cpu_float32 (main.TestTagsCPU.test_tags_native_layer_norm_cpu_float32) ... ok
test_tags_nn_functional_layer_norm_cpu_float32 (main.TestTagsCPU.test_tags_nn_functional_layer_norm_cpu_float32) ... ok
test_tags_nn_functional_rms_norm_cpu_float32 (main.TestTagsCPU.test_tags_nn_functional_rms_norm_cpu_float32) ... ok
test_tags__refs_native_layer_norm_cuda_float32 (main.TestTagsCUDA.test_tags__refs_native_layer_norm_cuda_float32) ... skipped 'Only runs on cpu'
test_tags__refs_nn_functional_layer_norm_cuda_float32 (main.TestTagsCUDA.test_tags__refs_nn_functional_layer_norm_cuda_float32) ... skipped 'Only runs on cpu'
test_tags_native_layer_norm_cuda_float32 (main.TestTagsCUDA.test_tags_native_layer_norm_cuda_float32) ... skipped 'Only runs on cpu'
test_tags_nn_functional_layer_norm_cuda_float32 (main.TestTagsCUDA.test_tags_nn_functional_layer_norm_cuda_float32) ... skipped 'Only runs on cpu'
test_tags_nn_functional_rms_norm_cuda_float32 (main.TestTagsCUDA.test_tags_nn_functional_rms_norm_cuda_float32) ... skipped 'Only runs on cpu'

Ran 218 tests in 15.669s

OK (skipped=31)

anatoliylitv · 2026-07-30T21:36:39Z

Test	Result	Duration
(unidentified — command scrolled out of buffer, likely `test_distributed_spawn.py` run with `WORLD_SIZE=8`)	PASS (`OK`, 1 test)	17.542s
`test_autograd.py::TestAutogradDeviceTypeCUDA::test_reentrant_parent_error_on_cpu_cuda`	PASS (`OK`, 1 test)	0.829s
`test_autograd.py::TestAutogradDeviceTypeCUDA::test_reentrant_parent_error_on_cpu_cuda` (rerun)	PASS (`OK`, 1 test)	0.720s
`test_autograd.py::TestAutogradDeviceTypeCUDA::test_reentrant_parent_error_on_cpu_cuda` (with `PYTORCH_TEST_WITH_ROCM=1`)	PASS (`OK`, 1 test)	0.734s

Total: 4/4 tests passed, 0 failed.

anatoliylitv · 2026-08-14T18:17:07Z

Replaced with:
#3564
#3565

anatoliylitv added 2 commits July 10, 2026 21:15

Return back two-pass gamma/beta backward kernel for LayerNorm and use…

b31fe4a

… it accordingly with tiled one where performance is better.

Update function names for clarity

20a3ce6

anatoliylitv changed the title ~~Anatoliylitv/tiled and 2 pass layer norm~~ [ROCm] Optimize AMD normalization backward kernel by using tiled and two pass implementations where they are most efficient. Jul 21, 2026

Remove unreachable ROCm tuning blocks

aedb098

anatoliylitv closed this Aug 14, 2026

Navigation Menu

Uh oh!

[ROCm] Optimize AMD normalization backward kernel by using tiled and two pass implementations where they are most efficient. - #3424

[ROCm] Optimize AMD normalization backward kernel by using tiled and two pass implementations where they are most efficient.#3424
anatoliylitv wants to merge 3 commits into
release/2.12from
anatoliylitv/tiled-and-2-pass-layer-norm

anatoliylitv commented Jul 11, 2026 •

edited

Loading

Uh oh!

rocm-repo-management-api Bot commented Jul 11, 2026 •

edited

Loading

Uh oh!

anatoliylitv commented Jul 30, 2026

Uh oh!

anatoliylitv commented Jul 30, 2026 •

edited

Loading

Uh oh!

anatoliylitv commented Aug 14, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

anatoliylitv commented Jul 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Layer Norm Backward Benchmark: Tiled only vs Tiled+Two pass

Summary (avg µs over 3 runs)

Submission Checklist

Uh oh!

rocm-repo-management-api Bot commented Jul 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

anatoliylitv commented Jul 30, 2026

Uh oh!

anatoliylitv commented Jul 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

anatoliylitv commented Aug 14, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

anatoliylitv commented Jul 11, 2026 •

edited

Loading

rocm-repo-management-api Bot commented Jul 11, 2026 •

edited

Loading

anatoliylitv commented Jul 30, 2026 •

edited

Loading