From 04072e382b305ff5af8294004ae59c783f7cb126 Mon Sep 17 00:00:00 2001 From: Jeremy Berchtold Date: Mon, 15 Dec 2025 16:04:56 -0800 Subject: [PATCH 1/5] HLO FFI tests Signed-off-by: Jeremy Berchtold --- tests/jax/ffi_hlo/transformer_stablehlo.txt | 2591 +++++++++++++++++++ tests/jax/test_custom_call_compute.py | 179 ++ 2 files changed, 2770 insertions(+) create mode 100644 tests/jax/ffi_hlo/transformer_stablehlo.txt diff --git a/tests/jax/ffi_hlo/transformer_stablehlo.txt b/tests/jax/ffi_hlo/transformer_stablehlo.txt new file mode 100644 index 00000000000..b4a5376f0da --- /dev/null +++ b/tests/jax/ffi_hlo/transformer_stablehlo.txt @@ -0,0 +1,2591 @@ +#loc = loc(unknown) +#loc1 = loc("var_collect['collection']['mask1']") +#loc2 = loc("var_collect['params']['TransformerLayer_0']['attention']['out']['kernel'].value") +#loc3 = loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['kernel'].value") +#loc4 = loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['ln_bias'].value") +#loc5 = loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['scale'].value") +#loc6 = loc("var_collect['params']['TransformerLayer_0']['mlp']['ln_bias'].value") +#loc7 = loc("var_collect['params']['TransformerLayer_0']['mlp']['scale'].value") +#loc8 = loc("var_collect['params']['TransformerLayer_0']['mlp']['wi_kernel'].value") +#loc9 = loc("var_collect['params']['TransformerLayer_0']['mlp']['wo_kernel'].value") +#loc10 = loc("var_collect['params']['TransformerLayer_0']['relpos_bias']['rel_embedding'].value") +#loc11 = loc("x") +#loc12 = loc("grouped_kernel") +#loc15 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":2007:12 to :54) +#loc16 = loc("/usr/local/lib/python3.12/dist-packages/_pytest/python.py":166:13 to :37) +#loc17 = loc("/usr/local/lib/python3.12/dist-packages/pluggy/_callers.py":121:26 to :51) +#loc18 = loc("/usr/local/lib/python3.12/dist-packages/pluggy/_manager.py":120:15 to :76) +#loc19 = loc("/usr/local/lib/python3.12/dist-packages/pluggy/_hooks.py":512:15 to :85) +#loc24 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1960:24 to 1965:24) +#loc25 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1985:20 to :87) +#loc31 = loc("/opt/transformerengine/transformer_engine/jax/quantize/helper.py":673:12 to :91) +#loc32 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":402:20 to 404:9) +#loc33 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":738:24 to 740:9) +#loc34 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1333:35 to 1356:27) +#loc35 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":2088:20 to 2120:96) +#loc36 = loc("/opt/transformerengine/transformer_engine/jax/quantize/helper.py":637:21 to :76) +#loc37 = loc("/opt/transformerengine/transformer_engine/jax/quantize/helper.py":641:27 to 643:9) +#loc77 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":1295:15 to 1312:9) +#loc78 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":352:16 to 371:13) +#loc79 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":818:16 to 841:13) +#loc80 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1570:12 to 1585:61) +#loc92 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":1092:37 to 1111:5) +#loc93 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":923:13 to 941:5) +#loc97 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":3441:37 to 3448:5) +#loc140 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":447:19 to 452:9) +#loc141 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":386:16 to 398:5) +#loc142 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":356:13 to 368:5) +#loc143 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1995:20 to 2001:17) +#loc147 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1227:22 to 1229:9) +#loc148 = loc("scatter") +#loc149 = loc("scatter-add") +#loc151 = loc("scatter-max") +#loc171 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1982:26 to :62) +#loc174 = loc("reduce_window_sum") +#loc177 = loc("TestFFICompatibility.test_generate_hlo"(#loc15)) +#loc178 = loc("pytest_pyfunc_call"(#loc16)) +#loc179 = loc("_multicall"(#loc17)) +#loc180 = loc("PluginManager._hookexec"(#loc18)) +#loc181 = loc("HookCaller.__call__"(#loc19)) +#loc186 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc24)) +#loc187 = loc("TestFFICompatibility.test_generate_hlo..f"(#loc25)) +#loc193 = loc("NVFP4ScalingQuantizeConfig.get_quantize_flax_meta"(#loc31)) +#loc194 = loc("TransformerEngineBase.generate_quantizer_set"(#loc32)) +#loc195 = loc("LayerNormDenseGeneral.__call__"(#loc33)) +#loc196 = loc("MultiHeadAttention.__call__"(#loc34)) +#loc197 = loc("TransformerLayer.__call__"(#loc35)) +#loc198 = loc("NVFP4ScalingQuantizeConfig._make_stochastic_rounding_rng_state"(#loc36)) +#loc199 = loc("NVFP4ScalingQuantizeConfig._make_stochastic_rounding_rng_state"(#loc37)) +#loc239 = loc("fused_attn"(#loc77)) +#loc240 = loc("_FusedDotProductAttention.__call__"(#loc78)) +#loc241 = loc("DotProductAttention.__call__"(#loc79)) +#loc242 = loc("MultiHeadAttention.__call__"(#loc80)) +#loc254 = loc("_fused_attn_fwd_rule"(#loc92)) +#loc255 = loc("_legacy_fused_attn"(#loc93)) +#loc259 = loc("fused_attn_fwd"(#loc97)) +#loc302 = loc("_grouped_dense_fwd_rule"(#loc140)) +#loc303 = loc("_grouped_dense"(#loc141)) +#loc304 = loc("grouped_dense"(#loc142)) +#loc305 = loc("TestFFICompatibility.test_generate_hlo..train_step"(#loc143)) +#loc309 = loc("grouped_quantize"(#loc147)) +#loc330 = loc("TestFFICompatibility.test_generate_hlo"(#loc171)) +#loc333 = loc(callsite(#loc180 at #loc181)) +#loc335 = loc(callsite(#loc186 at #loc187)) +#loc344 = loc(callsite(#loc178 at #loc179)) +#loc347 = loc(callsite(#loc179 at #loc333)) +#loc349 = loc(callsite(#loc197 at #loc335)) +#loc362 = loc(callsite(#loc330 at #loc344)) +#loc365 = loc(callsite(#loc178 at #loc347)) +#loc371 = loc(callsite(#loc242 at #loc349)) +#loc383 = loc(callsite(#loc186 at #loc362)) +#loc386 = loc(callsite(#loc177 at #loc365)) +#loc392 = loc(callsite(#loc241 at #loc371)) +#loc407 = loc(callsite(#loc197 at #loc383)) +#loc417 = loc(callsite(#loc240 at #loc392)) +#loc433 = loc(callsite(#loc305 at #loc386)) +#loc435 = loc(callsite(#loc196 at #loc407)) +#loc447 = loc(callsite(#loc239 at #loc417)) +#loc474 = loc(callsite(#loc304 at #loc433)) +#loc476 = loc(callsite(#loc195 at #loc435)) +#loc491 = loc(callsite(#loc255 at #loc447)) +#loc524 = loc(callsite(#loc303 at #loc474)) +#loc527 = loc(callsite(#loc194 at #loc476)) +#loc545 = loc(callsite(#loc254 at #loc491)) +#loc587 = loc(callsite(#loc302 at #loc524)) +#loc592 = loc(callsite(#loc193 at #loc527)) +#loc642 = loc(callsite(#loc259 at #loc545)) +#loc750 = loc(callsite(#loc309 at #loc587)) +#loc785 = loc(callsite(#loc198 at #loc592)) +#loc786 = loc(callsite(#loc199 at #loc592)) +module @jit_train_step attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} { + func.func public @main(%arg0: tensor<1x1x128x512xbf16> loc("var_collect['collection']['mask1']"), %arg1: tensor<512x512xbf16> loc("var_collect['params']['TransformerLayer_0']['attention']['out']['kernel'].value"), %arg2: tensor<512x3x512xbf16> loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['kernel'].value"), %arg3: tensor<512xbf16> loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['ln_bias'].value"), %arg4: tensor<512xbf16> loc("var_collect['params']['TransformerLayer_0']['attention']['qkv']['scale'].value"), %arg5: tensor<512xbf16> loc("var_collect['params']['TransformerLayer_0']['mlp']['ln_bias'].value"), %arg6: tensor<512xbf16> loc("var_collect['params']['TransformerLayer_0']['mlp']['scale'].value"), %arg7: tensor<512x1x2048xbf16> loc("var_collect['params']['TransformerLayer_0']['mlp']['wi_kernel'].value"), %arg8: tensor<2048x512xbf16> loc("var_collect['params']['TransformerLayer_0']['mlp']['wo_kernel'].value"), %arg9: tensor<8x32xbf16> loc("var_collect['params']['TransformerLayer_0']['relpos_bias']['rel_embedding'].value"), %arg10: tensor<1x128x512xbf16> loc("x"), %arg11: tensor<1x512x512xbf16> loc("grouped_kernel")) -> (tensor {jax.result_info = "result[0]"}, tensor<1x1x128x512xbf16> {jax.result_info = "result[1]['collection']['mask1']"}, tensor<1x1x128x512xbf16> {jax.result_info = "result[1]['collection']['mask2']"}, tensor<512x512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['attention']['out']['kernel'].value"}, tensor<512x3x512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['attention']['qkv']['kernel'].value"}, tensor<512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['attention']['qkv']['ln_bias'].value"}, tensor<512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['attention']['qkv']['scale'].value"}, tensor<512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['mlp']['ln_bias'].value"}, tensor<512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['mlp']['scale'].value"}, tensor<512x1x2048xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['mlp']['wi_kernel'].value"}, tensor<2048x512xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['mlp']['wo_kernel'].value"}, tensor<8x32xbf16> {jax.result_info = "result[1]['params']['TransformerLayer_0']['relpos_bias']['rel_embedding'].value"}) { + %cst = stablehlo.constant dense<4.480000e+02> : tensor loc(#loc) + %c = stablehlo.constant dense<1> : tensor loc(#loc) + %cst_0 = stablehlo.constant dense<6.553600e+04> : tensor loc(#loc) + %c_1 = stablehlo.constant dense<19484454> : tensor loc(#loc) + %c_2 = stablehlo.constant dense<2364379198> : tensor loc(#loc) + %c_3 = stablehlo.constant dense<3284580519> : tensor loc(#loc) + %c_4 = stablehlo.constant dense<128> : tensor loc(#loc) + %cst_5 = stablehlo.constant dense<2.688000e+03> : tensor loc(#loc) + %cst_6 = stablehlo.constant dense<0xFF80> : tensor loc(#loc) + %cst_7 = stablehlo.constant dense<0.000000e+00> : tensor loc(#loc) + %c_8 = stablehlo.constant dense<0> : tensor loc(#loc) + %cst_9 = stablehlo.constant dense<1.600000e+01> : tensor loc(#loc) + %cst_10 = stablehlo.constant dense<0.000000e+00> : tensor loc(#loc) + %cst_11 = stablehlo.constant dense<1.000000e+00> : tensor loc(#loc) + %c_12 = stablehlo.constant dense<2147483647> : tensor loc(#loc) + %c_13 = stablehlo.constant dense<726121952> : tensor loc(#loc) + %c_14 = stablehlo.constant dense<38174041> : tensor loc(#loc) + %c_15 = stablehlo.constant dense<-1> : tensor loc(#loc788) + %c_16 = stablehlo.constant dense<32> : tensor loc(#loc788) + %c_17 = stablehlo.constant dense<0> : tensor loc(#loc) + %c_18 = stablehlo.constant dense<128> : tensor<1xi32> loc(#loc) + %c_19 = stablehlo.constant dense<[1, 1, 1, -1, 1, -1, -1, -1, -1, -1, -1, 1, -1, 1, -1, -1]> : tensor<16xi32> loc(#loc) + %c_20 = stablehlo.constant dense<"0x0100000001000000010000000100000001000000010000000100000001000000010000000100000001000000010000000100000001000000010000000100000001000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF01000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF01000000010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF01000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF01000000010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF010000000100000001000000FFFFFFFFFFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFF0100000001000000010000000100000001000000010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF010000000100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF010000000100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFF01000000010000000100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF0100000001000000010000000100000001000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF01000000FFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF0100000001000000FFFFFFFFFFFFFFFF01000000010000000100000001000000FFFFFFFFFFFFFFFF01000000FFFFFFFFFFFFFFFF01000000FFFFFFFF0100000001000000FFFFFFFFFFFFFFFF0100000001000000FFFFFFFF01000000FFFFFFFFFFFFFFFF01000000"> : tensor<16x16xi32> loc(#loc) + %c_21 = stablehlo.constant dense<"0x000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F00000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F00000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F00000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F00000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F00000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000001F0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000001F0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000001F0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000001F0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000001F0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000001F0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001F0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000001E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001D0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000001C0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000001C0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000001C0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000001C0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000001C0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000001C0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001C0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000001B0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000001A0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000001900000019000000190000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A0000000900000009000000090000000900000008000000080000000800000008000000070000000600000005000000040000000300000002000000010000000000000011000000120000001300000014000000150000001600000017000000180000001800000018000000180000001900000019000000190000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000110000001200000013000000140000001500000016000000170000001800000018000000180000001800000019000000190000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000001800000018000000190000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A0000000900000009000000090000000900000008000000080000000800000008000000070000000600000005000000040000000300000002000000010000000000000011000000120000001300000014000000150000001600000017000000180000001800000018000000180000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000110000001200000013000000140000001500000016000000170000001800000018000000180000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000001700000018000000180000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A0000000900000009000000090000000900000008000000080000000800000008000000070000000600000005000000040000000300000002000000010000000000000011000000120000001300000014000000150000001600000017000000180000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000110000001200000013000000140000001500000016000000170000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000001400000015000000160000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A0000000900000009000000090000000900000008000000080000000800000008000000070000000600000005000000040000000300000002000000010000000000000011000000120000001300000014000000150000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000110000001200000013000000140000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A000000090000000900000009000000090000000800000008000000080000000800000007000000060000000500000004000000030000000200000001000000000000001100000012000000130000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A0000000900000009000000090000000900000008000000080000000800000008000000070000000600000005000000040000000300000002000000010000000000000011000000120000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000110000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000F0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000E0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000D0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000C0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000B0000000A0000000A0000000A0000000A0000000A0000000A0000000A00000009000000090000000900000009000000080000000800000008000000080000000700000006000000050000000400000003000000020000000100000000000000"> : tensor<128x128xi32> loc(#loc) + %0 = stablehlo.broadcast_in_dim %c_21, dims = [1, 2] : (tensor<128x128xi32>) -> tensor<1x128x128xi32> loc(#loc) + %1 = stablehlo.shift_right_logical %c_17, %c_16 : tensor loc(#loc789) + %2 = stablehlo.convert %1 : (tensor) -> tensor loc(#loc790) + %3 = stablehlo.broadcast_in_dim %2, dims = [] : (tensor) -> tensor<1xui32> loc(#loc791) + %4 = stablehlo.and %c_17, %c_15 : tensor loc(#loc792) + %5 = stablehlo.convert %4 : (tensor) -> tensor loc(#loc790) + %6 = stablehlo.broadcast_in_dim %5, dims = [] : (tensor) -> tensor<1xui32> loc(#loc791) + %7 = stablehlo.concatenate %3, %6, dim = 0 : (tensor<1xui32>, tensor<1xui32>) -> tensor<2xui32> loc(#loc793) + %8 = stablehlo.iota dim = 0 : tensor<32x1x1xi32> loc(#loc794) + %9 = stablehlo.broadcast_in_dim %8, dims = [0, 1, 2] : (tensor<32x1x1xi32>) -> tensor<32x128x128xi32> loc(#loc795) + %10 = stablehlo.broadcast_in_dim %0, dims = [0, 1, 2] : (tensor<1x128x128xi32>) -> tensor<32x128x128xi32> loc(#loc795) + %11 = stablehlo.compare EQ, %9, %10, SIGNED : (tensor<32x128x128xi32>, tensor<32x128x128xi32>) -> tensor<32x128x128xi1> loc(#loc795) + %12 = stablehlo.convert %11 : (tensor<32x128x128xi1>) -> tensor<32x128x128xbf16> loc(#loc796) + %13 = stablehlo.dot_general %arg9, %12, contracting_dims = [1] x [0] : (tensor<8x32xbf16>, tensor<32x128x128xbf16>) -> tensor<8x128x128xbf16> loc(#loc797) + %14 = stablehlo.broadcast_in_dim %13, dims = [1, 2, 3] : (tensor<8x128x128xbf16>) -> tensor<1x8x128x128xbf16> loc(#loc798) + %15 = call @_threefry_fold_in(%7, %c_14) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc799) + %16 = call @fold_in(%15, %c_13) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc800) + %17 = call @_randint(%16, %c_17, %c_12) : (tensor<2xui32>, tensor, tensor) -> tensor<1x4xi32> loc(#loc801) + %18 = stablehlo.bitcast_convert %17 : (tensor<1x4xi32>) -> tensor<1x4xui32> loc(#loc802) + %19 = stablehlo.broadcast_in_dim %cst_11, dims = [] : (tensor) -> tensor<1xf32> loc(#loc803) + %20 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc804) + %21:8 = stablehlo.custom_call @te_norm_forward_ffi(%arg10, %19, %20, %arg4, %arg3) {mhlo.backend_config = {epsilon = 9.9999999999999995E-7 : f64, norm_type = 0 : i64, output_amax_when_no_scaling = true, quantize_layout = 0 : i64, scaling_mode = 0 : i64, sm_margin = 0 : i64, zero_centered_gamma = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<512xbf16>, tensor<512xbf16>) -> (tensor<1x128x512xbf16>, tensor<1xbf16>, tensor<0xf32>, tensor<1xf32>, tensor<1xf32>, tensor<1x128xf32>, tensor<1x128xf32>, tensor<1xui8>) loc(#loc805) + %22 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc806) + %23 = call @_diag(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc807) + %24 = stablehlo.dot_general %23, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc808) + %25 = stablehlo.sqrt %cst_9 : tensor loc(#loc809) + %26 = stablehlo.convert %24 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc810) + %27 = stablehlo.broadcast_in_dim %25, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc811) + %28 = stablehlo.divide %26, %27 : tensor<16x16xf32> loc(#loc811) + %29 = stablehlo.convert %28 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc812) + %30:2 = stablehlo.custom_call @te_rht_amax_ffi(%21#0) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = false, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc813) + %31 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc814) + %32:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%21#0, %22, %21#4, %31, %30#1, %29) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = false, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x512xf4E2M1FN>, tensor<512x1x128xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc815) + %33 = stablehlo.slice %32#2 [0:1, 0:128, 0:32] : (tensor<1x128x32xf8E4M3FN>) -> tensor<1x128x32xf8E4M3FN> loc(#loc816) + %34 = stablehlo.slice %32#3 [0:512, 0:1, 0:8] : (tensor<512x1x8xf8E4M3FN>) -> tensor<512x1x8xf8E4M3FN> loc(#loc816) + %35 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc817) + %36 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc818) + %37 = stablehlo.abs %arg2 : tensor<512x3x512xbf16> loc(#loc819) + %38 = stablehlo.reduce(%37 init: %cst_6) applies stablehlo.maximum across dimensions = [0, 1, 2] : (tensor<512x3x512xbf16>, tensor) -> tensor loc(#loc820) + %39 = stablehlo.broadcast_in_dim %38, dims = [] : (tensor) -> tensor<1x1x1xbf16> loc(#loc821) + %40 = stablehlo.convert %39 : (tensor<1x1x1xbf16>) -> tensor<1x1x1xf32> loc(#loc822) + %41 = stablehlo.reshape %40 : (tensor<1x1x1xf32>) -> tensor<1xf32> loc(#loc823) + %42 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc824) + %43 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc825) + %44:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%arg2, %35, %41, %42, %43, %36) {mhlo.backend_config = {flatten_axis = -2 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 5 : i64, stochastic_rounding = false, use_rht = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x3x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<1x1xbf16>) -> (tensor<512x3x512xf4E2M1FN>, tensor<3x512x512xf4E2M1FN>, tensor<512x3x32xf8E4M3FN>, tensor<3x512x32xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc826) + %45 = stablehlo.slice %44#2 [0:512, 0:3, 0:32] : (tensor<512x3x32xf8E4M3FN>) -> tensor<512x3x32xf8E4M3FN> loc(#loc827) + %46 = stablehlo.slice %44#3 [0:3, 0:512, 0:32] : (tensor<3x512x32xf8E4M3FN>) -> tensor<3x512x32xf8E4M3FN> loc(#loc827) + %47 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc828) + %48 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc829) + %49 = stablehlo.divide %32#4, %48 : tensor<1xf32> loc(#loc829) + %50 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc830) + %51 = stablehlo.divide %44#4, %50 : tensor<1xf32> loc(#loc830) + %52 = stablehlo.multiply %49, %51 : tensor<1xf32> loc(#loc831) + %53 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc832) + %54 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc833) + %55:4 = stablehlo.custom_call @te_gemm_ffi(%32#0, %33, %44#1, %46, %53, %54, %52, %47) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 2 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<3x512x512xf4E2M1FN>, tensor<3x512x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x3x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33607936xui8>) loc(#loc834) + %56 = stablehlo.reshape %55#0 : (tensor<1x128x3x512xbf16>) -> tensor<1x128x3x8x64xbf16> loc(#loc835) + %57 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xi32> loc(#loc836) + %58 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xi32> loc(#loc837) + %59 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc838) + %60 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc839) + %61 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc840) + %62 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc841) + %63 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc842) + %64 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc843) + %65 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc844) + %66 = stablehlo.broadcast_in_dim %65, dims = [0] : (tensor<2xui32>) -> tensor<2x1xui32> loc(#loc845) + %67 = stablehlo.reshape %66 : (tensor<2x1xui32>) -> tensor<2xui32> loc(#loc846) + %68 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc847) + %69 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc848) + %70 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc849) + %71 = stablehlo.compare LT, %57, %70, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc849) + %72 = call @_where_79(%71, %c_17, %57) : (tensor<1xi1>, tensor, tensor<1xi32>) -> tensor<1xi32> loc(#loc850) + %73 = call @_cumsum_with_promotion(%72) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc851) + %74 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc852) + %75 = stablehlo.concatenate %74, %73, dim = 0 : (tensor<1xi32>, tensor<1xi32>) -> tensor<2xi32> loc(#loc853) + %76 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc849) + %77 = stablehlo.compare LT, %58, %76, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc849) + %78 = call @_where_79(%77, %c_17, %58) : (tensor<1xi1>, tensor, tensor<1xi32>) -> tensor<1xi32> loc(#loc850) + %79 = call @_cumsum_with_promotion(%78) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc851) + %80 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc852) + %81 = stablehlo.concatenate %80, %79, dim = 0 : (tensor<1xi32>, tensor<1xi32>) -> tensor<2xi32> loc(#loc853) + %82:4 = stablehlo.custom_call @te_fused_attn_forward_ffi(%56, %68, %68, %14, %69, %67, %75, %81, %59, %60, %61, %62, %63, %64) {mhlo.backend_config = {attn_heads = 8 : i64, bias_batch = 1 : i64, bias_heads = 8 : i64, bias_type = 2 : i64, deterministic = false, dropout_probability = 0.000000e+00 : f64, input_batch = 1 : i64, is_training = true, kv_max_seqlen = 128 : i64, mask_type = 2 : i64, max_segments_per_seq = 1 : i64, num_gqa_groups = 8 : i64, q_max_seqlen = 128 : i64, qk_head_dim = 64 : i64, qkv_layout = 5 : i64, scaling_factor = 1.000000e+00 : f64, softmax_type = 0 : i64, v_head_dim = 64 : i64, window_size_left = -1 : i64, window_size_right = -1 : i64}, operand_layouts = [dense<[4, 3, 2, 1, 0]> : tensor<5xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x3x8x64xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<1x8x128x128xbf16>, tensor<0xf32>, tensor<2xui32>, tensor<2xi32>, tensor<2xi32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>) -> (tensor<1x128x8x64xbf16>, tensor<1x8x128x1xf32>, tensor<2x4xui32>, tensor<1xui8>) loc(#loc854) + %83 = stablehlo.reshape %82#0 : (tensor<1x128x8x64xbf16>) -> tensor<1x128x512xbf16> loc(#loc855) + %84 = call @_threefry_fold_in(%7, %c_3) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc856) + %85 = call @fold_in(%84, %c_13) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc857) + %86 = call @_randint(%85, %c_17, %c_12) : (tensor<2xui32>, tensor, tensor) -> tensor<1x4xi32> loc(#loc858) + %87 = stablehlo.bitcast_convert %86 : (tensor<1x4xi32>) -> tensor<1x4xui32> loc(#loc859) + %88 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc860) + %89 = call @_diag(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc861) + %90 = stablehlo.dot_general %89, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc862) + %91 = stablehlo.sqrt %cst_9 : tensor loc(#loc863) + %92 = stablehlo.convert %90 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc864) + %93 = stablehlo.broadcast_in_dim %91, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc865) + %94 = stablehlo.divide %92, %93 : tensor<16x16xf32> loc(#loc865) + %95 = stablehlo.convert %94 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc866) + %96:2 = stablehlo.custom_call @te_rht_amax_ffi(%83) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = true, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc867) + %97 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc868) + %98:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%83, %88, %96#0, %97, %96#1, %95) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = false, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x512xf4E2M1FN>, tensor<512x1x128xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc869) + %99 = stablehlo.slice %98#2 [0:1, 0:128, 0:32] : (tensor<1x128x32xf8E4M3FN>) -> tensor<1x128x32xf8E4M3FN> loc(#loc870) + %100 = stablehlo.slice %98#3 [0:512, 0:1, 0:8] : (tensor<512x1x8xf8E4M3FN>) -> tensor<512x1x8xf8E4M3FN> loc(#loc870) + %101 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc871) + %102 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc872) + %103 = stablehlo.abs %arg1 : tensor<512x512xbf16> loc(#loc873) + %104 = stablehlo.reduce(%103 init: %cst_6) applies stablehlo.maximum across dimensions = [0, 1] : (tensor<512x512xbf16>, tensor) -> tensor loc(#loc874) + %105 = stablehlo.broadcast_in_dim %104, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc875) + %106 = stablehlo.convert %105 : (tensor<1x1xbf16>) -> tensor<1x1xf32> loc(#loc876) + %107 = stablehlo.reshape %106 : (tensor<1x1xf32>) -> tensor<1xf32> loc(#loc877) + %108 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc878) + %109 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc879) + %110:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%arg1, %101, %107, %108, %109, %102) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 5 : i64, stochastic_rounding = false, use_rht = false}, operand_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<1x1xbf16>) -> (tensor<512x512xf4E2M1FN>, tensor<512x512xf4E2M1FN>, tensor<512x32xf8E4M3FN>, tensor<512x32xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc880) + %111 = stablehlo.slice %110#2 [0:512, 0:32] : (tensor<512x32xf8E4M3FN>) -> tensor<512x32xf8E4M3FN> loc(#loc881) + %112 = stablehlo.slice %110#3 [0:512, 0:32] : (tensor<512x32xf8E4M3FN>) -> tensor<512x32xf8E4M3FN> loc(#loc881) + %113 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc882) + %114 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc883) + %115 = stablehlo.divide %98#4, %114 : tensor<1xf32> loc(#loc883) + %116 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc884) + %117 = stablehlo.divide %110#4, %116 : tensor<1xf32> loc(#loc884) + %118 = stablehlo.multiply %115, %117 : tensor<1xf32> loc(#loc885) + %119 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc886) + %120 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc887) + %121:4 = stablehlo.custom_call @te_gemm_ffi(%98#0, %99, %110#1, %112, %119, %120, %118, %113) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x512xf4E2M1FN>, tensor<512x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33575168xui8>) loc(#loc888) + %122 = stablehlo.add %121#0, %arg10 : tensor<1x128x512xbf16> loc(#loc889) + %123 = call @_threefry_fold_in(%7, %c_2) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc890) + %124 = call @fold_in(%123, %c_13) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc891) + %125 = call @_randint(%124, %c_17, %c_12) : (tensor<2xui32>, tensor, tensor) -> tensor<1x4xi32> loc(#loc892) + %126 = stablehlo.bitcast_convert %125 : (tensor<1x4xi32>) -> tensor<1x4xui32> loc(#loc893) + %127 = call @_threefry_fold_in(%7, %c_1) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc894) + %128 = call @fold_in(%127, %c_13) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc895) + %129 = call @_randint(%128, %c_17, %c_12) : (tensor<2xui32>, tensor, tensor) -> tensor<1x4xi32> loc(#loc896) + %130 = stablehlo.bitcast_convert %129 : (tensor<1x4xi32>) -> tensor<1x4xui32> loc(#loc897) + %131 = stablehlo.broadcast_in_dim %cst_11, dims = [] : (tensor) -> tensor<1xf32> loc(#loc898) + %132 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc899) + %133:8 = stablehlo.custom_call @te_norm_forward_ffi(%122, %131, %132, %arg6, %arg5) {mhlo.backend_config = {epsilon = 9.9999999999999995E-7 : f64, norm_type = 0 : i64, output_amax_when_no_scaling = true, quantize_layout = 0 : i64, scaling_mode = 0 : i64, sm_margin = 0 : i64, zero_centered_gamma = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<512xbf16>, tensor<512xbf16>) -> (tensor<1x128x512xbf16>, tensor<1xbf16>, tensor<0xf32>, tensor<1xf32>, tensor<1xf32>, tensor<1x128xf32>, tensor<1x128xf32>, tensor<1xui8>) loc(#loc900) + %134 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc901) + %135 = call @_diag(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc902) + %136 = stablehlo.dot_general %135, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc903) + %137 = stablehlo.sqrt %cst_9 : tensor loc(#loc904) + %138 = stablehlo.convert %136 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc905) + %139 = stablehlo.broadcast_in_dim %137, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc906) + %140 = stablehlo.divide %138, %139 : tensor<16x16xf32> loc(#loc906) + %141 = stablehlo.convert %140 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc907) + %142:2 = stablehlo.custom_call @te_rht_amax_ffi(%133#0) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = false, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc908) + %143 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc909) + %144:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%133#0, %134, %133#4, %143, %142#1, %141) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = false, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x512xf4E2M1FN>, tensor<512x1x128xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc910) + %145 = stablehlo.slice %144#2 [0:1, 0:128, 0:32] : (tensor<1x128x32xf8E4M3FN>) -> tensor<1x128x32xf8E4M3FN> loc(#loc911) + %146 = stablehlo.slice %144#3 [0:512, 0:1, 0:8] : (tensor<512x1x8xf8E4M3FN>) -> tensor<512x1x8xf8E4M3FN> loc(#loc911) + %147 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc912) + %148 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc913) + %149 = stablehlo.abs %arg7 : tensor<512x1x2048xbf16> loc(#loc914) + %150 = stablehlo.reduce(%149 init: %cst_6) applies stablehlo.maximum across dimensions = [0, 1, 2] : (tensor<512x1x2048xbf16>, tensor) -> tensor loc(#loc915) + %151 = stablehlo.broadcast_in_dim %150, dims = [] : (tensor) -> tensor<1x1x1xbf16> loc(#loc916) + %152 = stablehlo.convert %151 : (tensor<1x1x1xbf16>) -> tensor<1x1x1xf32> loc(#loc917) + %153 = stablehlo.reshape %152 : (tensor<1x1x1xf32>) -> tensor<1xf32> loc(#loc918) + %154 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc919) + %155 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc920) + %156:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%arg7, %147, %153, %154, %155, %148) {mhlo.backend_config = {flatten_axis = -2 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 5 : i64, stochastic_rounding = false, use_rht = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x1x2048xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<1x1xbf16>) -> (tensor<512x1x2048xf4E2M1FN>, tensor<1x2048x512xf4E2M1FN>, tensor<512x1x128xf8E4M3FN>, tensor<1x2048x32xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc921) + %157 = stablehlo.slice %156#2 [0:512, 0:1, 0:128] : (tensor<512x1x128xf8E4M3FN>) -> tensor<512x1x128xf8E4M3FN> loc(#loc922) + %158 = stablehlo.slice %156#3 [0:1, 0:2048, 0:32] : (tensor<1x2048x32xf8E4M3FN>) -> tensor<1x2048x32xf8E4M3FN> loc(#loc922) + %159 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc923) + %160 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc924) + %161 = stablehlo.divide %144#4, %160 : tensor<1xf32> loc(#loc924) + %162 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc925) + %163 = stablehlo.divide %156#4, %162 : tensor<1xf32> loc(#loc925) + %164 = stablehlo.multiply %161, %163 : tensor<1xf32> loc(#loc926) + %165 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc927) + %166 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc928) + %167:4 = stablehlo.custom_call @te_gemm_ffi(%144#0, %145, %156#1, %158, %165, %166, %164, %159) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 2 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<1x2048x512xf4E2M1FN>, tensor<1x2048x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x1x2048xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33624320xui8>) loc(#loc929) + %168 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc930) + %169 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc931) + %170:5 = stablehlo.custom_call @te_act_lu_ffi(%167#0, %168, %169) {mhlo.backend_config = {act_enum = 0 : i64, act_params = {clamped_swiglu = {alpha = 1.702000e+00 : f32, limit = 7.000000e+00 : f32}}, output_amax_when_no_scaling = true, quantize_layout = 0 : i64, scaling_mode = 0 : i64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x1x2048xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x2048xbf16>, tensor<1xbf16>, tensor<0xf32>, tensor<1xf32>, tensor<1xf32>) loc(#loc932) + %171 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc933) + %172 = call @_diag(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc934) + %173 = stablehlo.dot_general %172, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc935) + %174 = stablehlo.sqrt %cst_9 : tensor loc(#loc936) + %175 = stablehlo.convert %173 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc937) + %176 = stablehlo.broadcast_in_dim %174, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc938) + %177 = stablehlo.divide %175, %176 : tensor<16x16xf32> loc(#loc938) + %178 = stablehlo.convert %177 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc939) + %179:2 = stablehlo.custom_call @te_rht_amax_ffi(%170#0) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = false, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x2048xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc940) + %180 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc941) + %181:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%170#0, %171, %170#4, %180, %179#1, %178) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = false, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x2048xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x2048xf4E2M1FN>, tensor<2048x1x128xf4E2M1FN>, tensor<1x128x128xf8E4M3FN>, tensor<2048x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc942) + %182 = stablehlo.slice %181#2 [0:1, 0:128, 0:128] : (tensor<1x128x128xf8E4M3FN>) -> tensor<1x128x128xf8E4M3FN> loc(#loc943) + %183 = stablehlo.slice %181#3 [0:2048, 0:1, 0:8] : (tensor<2048x1x8xf8E4M3FN>) -> tensor<2048x1x8xf8E4M3FN> loc(#loc943) + %184 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc944) + %185 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc945) + %186 = stablehlo.abs %arg8 : tensor<2048x512xbf16> loc(#loc946) + %187 = stablehlo.reduce(%186 init: %cst_6) applies stablehlo.maximum across dimensions = [0, 1] : (tensor<2048x512xbf16>, tensor) -> tensor loc(#loc947) + %188 = stablehlo.broadcast_in_dim %187, dims = [] : (tensor) -> tensor<1x1xbf16> loc(#loc948) + %189 = stablehlo.convert %188 : (tensor<1x1xbf16>) -> tensor<1x1xf32> loc(#loc949) + %190 = stablehlo.reshape %189 : (tensor<1x1xf32>) -> tensor<1xf32> loc(#loc950) + %191 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<0xui32> loc(#loc951) + %192 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc952) + %193:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%arg8, %184, %190, %191, %192, %185) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 5 : i64, stochastic_rounding = false, use_rht = false}, operand_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<2048x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<0xui32>, tensor<1xf32>, tensor<1x1xbf16>) -> (tensor<2048x512xf4E2M1FN>, tensor<512x2048xf4E2M1FN>, tensor<2048x32xf8E4M3FN>, tensor<512x128xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc953) + %194 = stablehlo.slice %193#2 [0:2048, 0:32] : (tensor<2048x32xf8E4M3FN>) -> tensor<2048x32xf8E4M3FN> loc(#loc954) + %195 = stablehlo.slice %193#3 [0:512, 0:128] : (tensor<512x128xf8E4M3FN>) -> tensor<512x128xf8E4M3FN> loc(#loc954) + %196 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc955) + %197 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc956) + %198 = stablehlo.divide %181#4, %197 : tensor<1xf32> loc(#loc956) + %199 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc957) + %200 = stablehlo.divide %193#4, %199 : tensor<1xf32> loc(#loc957) + %201 = stablehlo.multiply %198, %200 : tensor<1xf32> loc(#loc958) + %202 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc959) + %203 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc960) + %204:4 = stablehlo.custom_call @te_gemm_ffi(%181#0, %182, %193#1, %195, %202, %203, %201, %196) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x2048xf4E2M1FN>, tensor<1x128x128xf8E4M3FN>, tensor<512x2048xf4E2M1FN>, tensor<512x128xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33636608xui8>) loc(#loc961) + %205 = stablehlo.add %204#0, %122 : tensor<1x128x512xbf16> loc(#loc962) + %206 = stablehlo.reshape %205 : (tensor<1x128x512xbf16>) -> tensor<1x1x128x512xbf16> loc(#loc963) + %207 = stablehlo.custom_call @te_scaled_softmax_forward_ffi(%206) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<1x1x128x512xbf16>) -> tensor<1x1x128x512xbf16> loc(#loc964) + %208 = stablehlo.convert %arg0 : (tensor<1x1x128x512xbf16>) -> tensor<1x1x128x512xui8> loc(#loc965) + %209 = stablehlo.custom_call @te_scaled_masked_softmax_forward_ffi(%207, %208) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<1x1x128x512xbf16>, tensor<1x1x128x512xui8>) -> tensor<1x1x128x512xbf16> loc(#loc966) + %210 = stablehlo.reshape %209 : (tensor<1x1x128x512xbf16>) -> tensor<64x1x32x32xbf16> loc(#loc967) + %211 = stablehlo.custom_call @te_scaled_upper_triang_masked_softmax_forward_ffi(%210) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<64x1x32x32xbf16>) -> tensor<64x1x32x32xbf16> loc(#loc968) + %212 = stablehlo.convert %211 : (tensor<64x1x32x32xbf16>) -> tensor<64x1x32x32xf32> loc(#loc969) + %213 = stablehlo.reduce(%212 init: %cst_10) applies stablehlo.add across dimensions = [0, 1, 2, 3] : (tensor<64x1x32x32xf32>, tensor) -> tensor loc(#loc970) + %214 = stablehlo.divide %213, %cst_0 : tensor loc(#loc971) + %215 = stablehlo.convert %214 : (tensor) -> tensor loc(#loc969) + %216 = stablehlo.divide %cst_11, %cst_0 : tensor loc(#loc972) + %217 = stablehlo.broadcast_in_dim %216, dims = [] : (tensor) -> tensor<64x1x32x32xf32> loc(#loc973) + %218 = stablehlo.convert %217 : (tensor<64x1x32x32xf32>) -> tensor<64x1x32x32xbf16> loc(#loc974) + %219 = stablehlo.custom_call @te_scaled_upper_triang_masked_softmax_backward_ffi(%218, %211) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<64x1x32x32xbf16>, tensor<64x1x32x32xbf16>) -> tensor<64x1x32x32xbf16> loc(#loc975) + %220 = stablehlo.reshape %219 : (tensor<64x1x32x32xbf16>) -> tensor<1x1x128x512xbf16> loc(#loc976) + %221 = stablehlo.custom_call @te_scaled_masked_softmax_backward_ffi(%220, %209) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<1x1x128x512xbf16>, tensor<1x1x128x512xbf16>) -> tensor<1x1x128x512xbf16> loc(#loc977) + %222 = stablehlo.custom_call @te_scaled_softmax_backward_ffi(%221, %207) {mhlo.backend_config = {scale_factor = 1.000000e+00 : f64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>]} : (tensor<1x1x128x512xbf16>, tensor<1x1x128x512xbf16>) -> tensor<1x1x128x512xbf16> loc(#loc978) + %223 = stablehlo.reshape %222 : (tensor<1x1x128x512xbf16>) -> tensor<1x128x512xbf16> loc(#loc979) + %224 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %225 = call @_diag_156(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc981) + %226 = stablehlo.dot_general %225, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc982) + %227 = stablehlo.sqrt %cst_9 : tensor loc(#loc983) + %228 = stablehlo.convert %226 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc984) + %229 = stablehlo.broadcast_in_dim %227, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc985) + %230 = stablehlo.divide %228, %229 : tensor<16x16xf32> loc(#loc985) + %231 = stablehlo.convert %230 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc984) + %232:2 = stablehlo.custom_call @te_rht_amax_ffi(%223) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = true, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc986) + %233:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%223, %224, %232#0, %126, %232#1, %231) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = true, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<1x4xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x512xf4E2M1FN>, tensor<512x1x128xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc987) + %234 = stablehlo.slice %233#2 [0:1, 0:128, 0:32] : (tensor<1x128x32xf8E4M3FN>) -> tensor<1x128x32xf8E4M3FN> loc(#loc988) + %235 = stablehlo.slice %233#3 [0:512, 0:1, 0:8] : (tensor<512x1x8xf8E4M3FN>) -> tensor<512x1x8xf8E4M3FN> loc(#loc988) + %236 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %237 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %238 = stablehlo.divide %233#4, %237 : tensor<1xf32> loc(#loc985) + %239 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %240 = stablehlo.divide %193#4, %239 : tensor<1xf32> loc(#loc985) + %241 = stablehlo.multiply %238, %240 : tensor<1xf32> loc(#loc989) + %242 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %243 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %244:4 = stablehlo.custom_call @te_gemm_ffi(%233#0, %234, %193#0, %194, %242, %243, %241, %236) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<2048x512xf4E2M1FN>, tensor<2048x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x2048xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33624320xui8>) loc(#loc990) + %245 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %246 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %247 = stablehlo.divide %179#1, %246 : tensor<1xf32> loc(#loc985) + %248 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %249 = stablehlo.divide %232#1, %248 : tensor<1xf32> loc(#loc985) + %250 = stablehlo.multiply %247, %249 : tensor<1xf32> loc(#loc989) + %251 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %252 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %253:4 = stablehlo.custom_call @te_gemm_ffi(%181#1, %183, %233#1, %235, %251, %252, %250, %245) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 1 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<2048x1x128xf4E2M1FN>, tensor<2048x1x8xf8E4M3FN>, tensor<512x1x128xf4E2M1FN>, tensor<512x1x8xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<2048x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33575168xui8>) loc(#loc990) + %254 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %255 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %256:7 = stablehlo.custom_call @te_dact_dbias_quantize_ffi(%244#0, %167#0, %254, %255) {mhlo.backend_config = {act_enum = 0 : i64, act_params = {clamped_swiglu = {alpha = 1.702000e+00 : f32, limit = 7.000000e+00 : f32}}, is_dbias = false, output_amax_when_no_scaling = true, quantize_layout = 0 : i64, scaling_mode = 0 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x2048xbf16>, tensor<1x128x1x2048xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x1x2048xbf16>, tensor<1xbf16>, tensor<0xf32>, tensor<1xf32>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc991) + %257 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %258 = call @_diag_156(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc981) + %259 = stablehlo.dot_general %258, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc982) + %260 = stablehlo.sqrt %cst_9 : tensor loc(#loc983) + %261 = stablehlo.convert %259 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc984) + %262 = stablehlo.broadcast_in_dim %260, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc985) + %263 = stablehlo.divide %261, %262 : tensor<16x16xf32> loc(#loc985) + %264 = stablehlo.convert %263 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc984) + %265:2 = stablehlo.custom_call @te_rht_amax_ffi(%256#0) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = false, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x1x2048xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc986) + %266:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%256#0, %257, %256#4, %130, %265#1, %264) {mhlo.backend_config = {flatten_axis = -2 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = true, use_rht = true}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x1x2048xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<1x4xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x1x2048xf4E2M1FN>, tensor<1x2048x1x128xf4E2M1FN>, tensor<1x128x1x128xf8E4M3FN>, tensor<1x2048x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc987) + %267 = stablehlo.slice %266#2 [0:1, 0:128, 0:1, 0:128] : (tensor<1x128x1x128xf8E4M3FN>) -> tensor<1x128x1x128xf8E4M3FN> loc(#loc988) + %268 = stablehlo.slice %266#3 [0:1, 0:2048, 0:1, 0:8] : (tensor<1x2048x1x8xf8E4M3FN>) -> tensor<1x2048x1x8xf8E4M3FN> loc(#loc988) + %269 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %270 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %271 = stablehlo.divide %266#4, %270 : tensor<1xf32> loc(#loc985) + %272 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %273 = stablehlo.divide %156#4, %272 : tensor<1xf32> loc(#loc985) + %274 = stablehlo.multiply %271, %273 : tensor<1xf32> loc(#loc989) + %275 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %276 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %277:4 = stablehlo.custom_call @te_gemm_ffi(%266#0, %267, %156#0, %157, %275, %276, %274, %269) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x1x2048xf4E2M1FN>, tensor<1x128x1x128xf8E4M3FN>, tensor<512x1x2048xf4E2M1FN>, tensor<512x1x128xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33636608xui8>) loc(#loc990) + %278 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc980) + %279 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %280 = stablehlo.divide %142#1, %279 : tensor<1xf32> loc(#loc985) + %281 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc985) + %282 = stablehlo.divide %265#1, %281 : tensor<1xf32> loc(#loc985) + %283 = stablehlo.multiply %280, %282 : tensor<1xf32> loc(#loc989) + %284 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %285 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc980) + %286:4 = stablehlo.custom_call @te_gemm_ffi(%144#1, %146, %266#1, %268, %284, %285, %283, %278) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 1 : i64, lhs_transposed = false, rhs_axis_boundary = 2 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x1x128xf4E2M1FN>, tensor<512x1x8xf8E4M3FN>, tensor<1x2048x1x128xf4E2M1FN>, tensor<1x2048x1x8xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<512x1x2048xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33575168xui8>) loc(#loc990) + %287:4 = stablehlo.custom_call @te_norm_backward_ffi(%277#0, %122, %133#5, %133#6, %arg6) {mhlo.backend_config = {norm_type = 0 : i64, sm_margin = 0 : i64, zero_centered_gamma = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1x128x512xbf16>, tensor<1x128xf32>, tensor<1x128xf32>, tensor<512xbf16>) -> (tensor<1x128x512xbf16>, tensor<512xbf16>, tensor<512xbf16>, tensor<131072xui8>) loc(#loc992) + %288 = stablehlo.add %223, %287#0 : tensor<1x128x512xbf16> loc(#loc993) + %289 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc994) + %290 = call @_diag_156(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc995) + %291 = stablehlo.dot_general %290, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc996) + %292 = stablehlo.sqrt %cst_9 : tensor loc(#loc997) + %293 = stablehlo.convert %291 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc998) + %294 = stablehlo.broadcast_in_dim %292, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc999) + %295 = stablehlo.divide %293, %294 : tensor<16x16xf32> loc(#loc999) + %296 = stablehlo.convert %295 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc998) + %297:2 = stablehlo.custom_call @te_rht_amax_ffi(%288) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = true, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc1000) + %298:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%288, %289, %297#0, %87, %297#1, %296) {mhlo.backend_config = {flatten_axis = -1 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = true, use_rht = true}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<1x4xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x512xf4E2M1FN>, tensor<512x1x128xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc1001) + %299 = stablehlo.slice %298#2 [0:1, 0:128, 0:32] : (tensor<1x128x32xf8E4M3FN>) -> tensor<1x128x32xf8E4M3FN> loc(#loc1002) + %300 = stablehlo.slice %298#3 [0:512, 0:1, 0:8] : (tensor<512x1x8xf8E4M3FN>) -> tensor<512x1x8xf8E4M3FN> loc(#loc1002) + %301 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc994) + %302 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc999) + %303 = stablehlo.divide %298#4, %302 : tensor<1xf32> loc(#loc999) + %304 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc999) + %305 = stablehlo.divide %110#4, %304 : tensor<1xf32> loc(#loc999) + %306 = stablehlo.multiply %303, %305 : tensor<1xf32> loc(#loc1003) + %307 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc994) + %308 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc994) + %309:4 = stablehlo.custom_call @te_gemm_ffi(%298#0, %299, %110#0, %111, %307, %308, %306, %301) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xf4E2M1FN>, tensor<1x128x32xf8E4M3FN>, tensor<512x512xf4E2M1FN>, tensor<512x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33575168xui8>) loc(#loc1004) + %310 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc994) + %311 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc999) + %312 = stablehlo.divide %96#1, %311 : tensor<1xf32> loc(#loc999) + %313 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc999) + %314 = stablehlo.divide %297#1, %313 : tensor<1xf32> loc(#loc999) + %315 = stablehlo.multiply %312, %314 : tensor<1xf32> loc(#loc1003) + %316 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc994) + %317 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc994) + %318:4 = stablehlo.custom_call @te_gemm_ffi(%98#1, %100, %298#1, %300, %316, %317, %315, %310) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 1 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x1x128xf4E2M1FN>, tensor<512x1x8xf8E4M3FN>, tensor<512x1x128xf4E2M1FN>, tensor<512x1x8xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<512x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33562880xui8>) loc(#loc1004) + %319 = stablehlo.reshape %309#0 : (tensor<1x128x512xbf16>) -> tensor<1x128x8x64xbf16> loc(#loc1005) + %320 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc1006) + %321 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<0xf32> loc(#loc1006) + %322 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1007) + %323 = stablehlo.compare LT, %57, %322, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc1007) + %324 = call @_where_79(%323, %c_17, %57) : (tensor<1xi1>, tensor, tensor<1xi32>) -> tensor<1xi32> loc(#loc1008) + %325 = call @_cumsum_with_promotion(%324) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1009) + %326 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1006) + %327 = stablehlo.concatenate %326, %325, dim = 0 : (tensor<1xi32>, tensor<1xi32>) -> tensor<2xi32> loc(#loc1010) + %328 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1007) + %329 = stablehlo.compare LT, %58, %328, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc1007) + %330 = call @_where_79(%329, %c_17, %58) : (tensor<1xi1>, tensor, tensor<1xi32>) -> tensor<1xi32> loc(#loc1008) + %331 = call @_cumsum_with_promotion(%330) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1009) + %332 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1006) + %333 = stablehlo.concatenate %332, %331, dim = 0 : (tensor<1xi32>, tensor<1xi32>) -> tensor<2xi32> loc(#loc1010) + %334:6 = stablehlo.custom_call @te_fused_attn_backward_ffi(%56, %320, %320, %14, %321, %82#1, %82#2, %82#0, %319, %327, %333, %59, %60, %61, %62, %63, %64) {mhlo.backend_config = {attn_heads = 8 : i64, bias_batch = 1 : i64, bias_heads = 8 : i64, bias_type = 2 : i64, deterministic = false, dropout_probability = 0.000000e+00 : f64, input_batch = 1 : i64, is_training = true, kv_max_seqlen = 128 : i64, mask_type = 2 : i64, max_segments_per_seq = 1 : i64, num_gqa_groups = 8 : i64, q_max_seqlen = 128 : i64, qk_head_dim = 64 : i64, qkv_layout = 5 : i64, scaling_factor = 1.000000e+00 : f64, softmax_type = 0 : i64, v_head_dim = 64 : i64, window_size_left = -1 : i64, window_size_right = -1 : i64}, operand_layouts = [dense<[4, 3, 2, 1, 0]> : tensor<5xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[4, 3, 2, 1, 0]> : tensor<5xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x3x8x64xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<1x8x128x128xbf16>, tensor<0xf32>, tensor<1x8x128x1xf32>, tensor<2x4xui32>, tensor<1x128x8x64xbf16>, tensor<1x128x8x64xbf16>, tensor<2xi32>, tensor<2xi32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>, tensor<0xf32>) -> (tensor<1x128x3x8x64xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<1x8x128x128xbf16>, tensor<0xf32>, tensor<266368xui8>) loc(#loc1011) + %335 = stablehlo.reshape %334#0 : (tensor<1x128x3x8x64xbf16>) -> tensor<1x128x3x512xbf16> loc(#loc1012) + %336 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1013) + %337 = call @_diag_156(%c_19) : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1014) + %338 = stablehlo.dot_general %337, %c_20, contracting_dims = [1] x [0] : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi32> loc(#loc1015) + %339 = stablehlo.sqrt %cst_9 : tensor loc(#loc1016) + %340 = stablehlo.convert %338 : (tensor<16x16xi32>) -> tensor<16x16xf32> loc(#loc1017) + %341 = stablehlo.broadcast_in_dim %339, dims = [] : (tensor) -> tensor<16x16xf32> loc(#loc1018) + %342 = stablehlo.divide %340, %341 : tensor<16x16xf32> loc(#loc1018) + %343 = stablehlo.convert %342 : (tensor<16x16xf32>) -> tensor<16x16xbf16> loc(#loc1017) + %344:2 = stablehlo.custom_call @te_rht_amax_ffi(%335) {mhlo.backend_config = {flatten_axis = 2 : i64, produce_regular_amax = true, rht_matrix_random_sign_mask_t = 55272 : i64}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x3x512xbf16>) -> (tensor<1xf32>, tensor<1xf32>) loc(#loc1019) + %345:7 = stablehlo.custom_call @te_dbias_quantize_ffi(%335, %336, %344#0, %18, %344#1, %343) {mhlo.backend_config = {flatten_axis = -2 : i64, is_dbias = false, q_layout = 2 : i64, scaling_mode = 4 : i64, stochastic_rounding = true, use_rht = true}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<[1, 0]> : tensor<2xindex>], output_operand_aliases = [#stablehlo.output_operand_alias], result_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x3x512xbf16>, tensor<1xf32>, tensor<1xf32>, tensor<1x4xui32>, tensor<1xf32>, tensor<16x16xbf16>) -> (tensor<1x128x3x512xf4E2M1FN>, tensor<3x512x1x128xf4E2M1FN>, tensor<1x128x3x32xf8E4M3FN>, tensor<3x512x1x8xf8E4M3FN>, tensor<1xf32>, tensor<1xbf16>, tensor<1xf32>) loc(#loc1020) + %346 = stablehlo.slice %345#2 [0:1, 0:128, 0:3, 0:32] : (tensor<1x128x3x32xf8E4M3FN>) -> tensor<1x128x3x32xf8E4M3FN> loc(#loc1021) + %347 = stablehlo.slice %345#3 [0:3, 0:512, 0:1, 0:8] : (tensor<3x512x1x8xf8E4M3FN>) -> tensor<3x512x1x8xf8E4M3FN> loc(#loc1021) + %348 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1013) + %349 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1018) + %350 = stablehlo.divide %345#4, %349 : tensor<1xf32> loc(#loc1018) + %351 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1018) + %352 = stablehlo.divide %44#4, %351 : tensor<1xf32> loc(#loc1018) + %353 = stablehlo.multiply %350, %352 : tensor<1xf32> loc(#loc1022) + %354 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc1013) + %355 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc1013) + %356:4 = stablehlo.custom_call @te_gemm_ffi(%345#0, %346, %44#0, %45, %354, %355, %353, %348) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 2 : i64, lhs_transposed = false, rhs_axis_boundary = 1 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x3x512xf4E2M1FN>, tensor<1x128x3x32xf8E4M3FN>, tensor<512x3x512xf4E2M1FN>, tensor<512x3x32xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<1x128x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33616128xui8>) loc(#loc1023) + %357 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1013) + %358 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1018) + %359 = stablehlo.divide %30#1, %358 : tensor<1xf32> loc(#loc1018) + %360 = stablehlo.broadcast_in_dim %cst_5, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1018) + %361 = stablehlo.divide %344#1, %360 : tensor<1xf32> loc(#loc1018) + %362 = stablehlo.multiply %359, %361 : tensor<1xf32> loc(#loc1022) + %363 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc1013) + %364 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<0xbf16> loc(#loc1013) + %365:4 = stablehlo.custom_call @te_gemm_ffi(%32#1, %34, %345#1, %347, %363, %364, %362, %357) {mhlo.backend_config = {collective_op = 0 : i64, fuse_bias = false, fuse_gelu = false, grad = false, lhs_axis_boundary = 1 : i64, lhs_transposed = false, rhs_axis_boundary = 2 : i64, rhs_transposed = true, scaling_mode = 4 : i64, use_split_accumulator = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<[3, 2, 1, 0]> : tensor<4xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<512x1x128xf4E2M1FN>, tensor<512x1x8xf8E4M3FN>, tensor<3x512x1x128xf4E2M1FN>, tensor<3x512x1x8xf8E4M3FN>, tensor<0xbf16>, tensor<0xbf16>, tensor<1xf32>, tensor<1xf32>) -> (tensor<512x3x512xbf16>, tensor<0xbf16>, tensor<0xbf16>, tensor<33571072xui8>) loc(#loc1023) + %366:4 = stablehlo.custom_call @te_norm_backward_ffi(%356#0, %arg10, %21#5, %21#6, %arg4) {mhlo.backend_config = {norm_type = 0 : i64, sm_margin = 0 : i64, zero_centered_gamma = false}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<[2, 1, 0]> : tensor<3xindex>, dense<[1, 0]> : tensor<2xindex>, dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x128x512xbf16>, tensor<1x128x512xbf16>, tensor<1x128xf32>, tensor<1x128xf32>, tensor<512xbf16>) -> (tensor<1x128x512xbf16>, tensor<512xbf16>, tensor<512xbf16>, tensor<131072xui8>) loc(#loc1024) + %367 = stablehlo.reduce(%334#3 init: %cst_7) applies stablehlo.add across dimensions = [0] : (tensor<1x8x128x128xbf16>, tensor) -> tensor<8x128x128xbf16> loc(#loc1025) + %368 = stablehlo.dot_general %367, %12, contracting_dims = [1, 2] x [1, 2] : (tensor<8x128x128xbf16>, tensor<32x128x128xbf16>) -> tensor<8x32xbf16> loc(#loc1026) + %369 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1x128x512xbf16> loc(#loc1027) + %370 = stablehlo.broadcast_in_dim %cst_7, dims = [] : (tensor) -> tensor<1x1x128x512xbf16> loc(#loc1027) + %371 = stablehlo.reshape %arg10 : (tensor<1x128x512xbf16>) -> tensor<128x512xbf16> loc(#loc1028) + %372 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1029) + %373 = stablehlo.abs %371 : tensor<128x512xbf16> loc(#loc1030) + %374 = stablehlo.reduce(%373 init: %cst_6) applies stablehlo.maximum across dimensions = [1] : (tensor<128x512xbf16>, tensor) -> tensor<128xbf16> loc(#loc1031) + %375 = stablehlo.iota dim = 0 : tensor<1xi32> loc(#loc1032) + %376 = call @_roll_static(%c_18) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1033) + %377 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1034) + %378 = "stablehlo.scatter"(%376, %377, %c_17) <{indices_are_sorted = true, scatter_dimension_numbers = #stablehlo.scatter, unique_indices = true}> ({ + ^bb0(%arg12: tensor loc("scatter"), %arg13: tensor loc("scatter")): + stablehlo.return %arg13 : tensor loc(#loc1035) + }) : (tensor<1xi32>, tensor<1xi32>, tensor) -> tensor<1xi32> loc(#loc1035) + %379 = call @cumsum_213(%378) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1036) + %380 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<128xi32> loc(#loc1034) + %381 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1037) + %382 = stablehlo.compare LT, %379, %381, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc1037) + %383 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1038) + %384 = stablehlo.add %379, %383 : tensor<1xi32> loc(#loc1038) + %385 = stablehlo.select %382, %384, %379 : tensor<1xi1>, tensor<1xi32> loc(#loc1039) + %386 = stablehlo.broadcast_in_dim %385, dims = [0] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1034) + %387 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1034) + %388 = "stablehlo.scatter"(%380, %386, %387) <{scatter_dimension_numbers = #stablehlo.scatter}> ({ + ^bb0(%arg12: tensor loc("scatter-add"), %arg13: tensor loc("scatter-add")): + %458 = stablehlo.add %arg12, %arg13 : tensor loc(#loc1038) + stablehlo.return %458 : tensor loc(#loc1040) + }) : (tensor<128xi32>, tensor<1x1xi32>, tensor<1xi32>) -> tensor<128xi32> loc(#loc1040) + %389 = call @cumsum_218(%388) : (tensor<128xi32>) -> tensor<128xi32> loc(#loc1036) + %390 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<128xi32> loc(#loc1041) + %391 = stablehlo.subtract %389, %390 : tensor<128xi32> loc(#loc1041) + %392 = call @_take(%375, %391) : (tensor<1xi32>, tensor<128xi32>) -> tensor<128xi32> loc(#loc1042) + %393 = stablehlo.broadcast_in_dim %cst_6, dims = [] : (tensor) -> tensor<1xbf16> loc(#loc1043) + %394 = stablehlo.broadcast_in_dim %392, dims = [0] : (tensor<128xi32>) -> tensor<128x1xi32> loc(#loc1043) + %395 = "stablehlo.scatter"(%393, %394, %374) <{scatter_dimension_numbers = #stablehlo.scatter}> ({ + ^bb0(%arg12: tensor loc("scatter-max"), %arg13: tensor loc("scatter-max")): + %458 = stablehlo.maximum %arg12, %arg13 : tensor loc(#loc1045) + stablehlo.return %458 : tensor loc(#loc1044) + }) : (tensor<1xbf16>, tensor<128x1xi32>, tensor<128xbf16>) -> tensor<1xbf16> loc(#loc1044) + %396 = stablehlo.slice %395 [0:1] : (tensor<1xbf16>) -> tensor<1xbf16> loc(#loc1046) + %397 = stablehlo.reshape %396 : (tensor<1xbf16>) -> tensor loc(#loc1047) + %398 = stablehlo.broadcast_in_dim %cst_11, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1048) + %399 = stablehlo.convert %397 : (tensor) -> tensor loc(#loc1049) + %400 = stablehlo.divide %cst, %399 : tensor loc(#loc1050) + %401 = stablehlo.divide %400, %cst_11 : tensor loc(#loc1051) + %402 = stablehlo.compare GT, %397, %cst_7, FLOAT : (tensor, tensor) -> tensor loc(#loc1052) + %403 = call @_where_237(%402, %401, %398) : (tensor, tensor, tensor<1xf32>) -> tensor<1xf32> loc(#loc1053) + %404 = stablehlo.is_finite %397 : (tensor) -> tensor loc(#loc1054) + %405 = call @_where_239(%404, %403, %398) : (tensor, tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> loc(#loc1055) + %406 = stablehlo.slice %405 [0:1] : (tensor<1xf32>) -> tensor<1xf32> loc(#loc1056) + %407 = stablehlo.reshape %406 : (tensor<1xf32>) -> tensor loc(#loc1057) + %408 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1058) + %409 = "stablehlo.scatter"(%372, %408, %407) <{indices_are_sorted = true, scatter_dimension_numbers = #stablehlo.scatter, unique_indices = true}> ({ + ^bb0(%arg12: tensor loc("scatter"), %arg13: tensor loc("scatter")): + stablehlo.return %arg13 : tensor loc(#loc1059) + }) : (tensor<1xf32>, tensor<1xi32>, tensor) -> tensor<1xf32> loc(#loc1059) + %410:5 = stablehlo.custom_call @te_grouped_quantize_ffi(%371, %409, %c_18) {mhlo.backend_config = {flatten_axis = -1 : i64, q_layout = 0 : i64, scaling_mode = 3 : i64}, operand_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<128x512xbf16>, tensor<1xf32>, tensor<1xi32>) -> (tensor<65536xf8E4M3FN>, tensor<1xf8E4M3FN>, tensor<1xf32>, tensor<1xf32>, tensor<1xf32>) loc(#loc1060) + %411 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1061) + %412 = stablehlo.broadcast_in_dim %cst_10, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1062) + %413 = stablehlo.abs %arg11 : tensor<1x512x512xbf16> loc(#loc1063) + %414 = stablehlo.reduce(%413 init: %cst_6) applies stablehlo.maximum across dimensions = [1, 2] : (tensor<1x512x512xbf16>, tensor) -> tensor<1xbf16> loc(#loc1064) + %415 = stablehlo.iota dim = 0 : tensor<1xi32> loc(#loc1065) + %416 = call @_roll_static(%411) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1066) + %417 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1067) + %418 = "stablehlo.scatter"(%416, %417, %c_17) <{indices_are_sorted = true, scatter_dimension_numbers = #stablehlo.scatter, unique_indices = true}> ({ + ^bb0(%arg12: tensor loc("scatter"), %arg13: tensor loc("scatter")): + stablehlo.return %arg13 : tensor loc(#loc1068) + }) : (tensor<1xi32>, tensor<1xi32>, tensor) -> tensor<1xi32> loc(#loc1068) + %419 = call @cumsum_213(%418) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1069) + %420 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1067) + %421 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1070) + %422 = stablehlo.compare LT, %419, %421, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc1070) + %423 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1071) + %424 = stablehlo.add %419, %423 : tensor<1xi32> loc(#loc1071) + %425 = stablehlo.select %422, %424, %419 : tensor<1xi1>, tensor<1xi32> loc(#loc1072) + %426 = stablehlo.broadcast_in_dim %425, dims = [0] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1067) + %427 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1067) + %428 = "stablehlo.scatter"(%420, %426, %427) <{scatter_dimension_numbers = #stablehlo.scatter}> ({ + ^bb0(%arg12: tensor loc("scatter-add"), %arg13: tensor loc("scatter-add")): + %458 = stablehlo.add %arg12, %arg13 : tensor loc(#loc1071) + stablehlo.return %458 : tensor loc(#loc1073) + }) : (tensor<1xi32>, tensor<1x1xi32>, tensor<1xi32>) -> tensor<1xi32> loc(#loc1073) + %429 = call @cumsum_213(%428) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1069) + %430 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1074) + %431 = stablehlo.subtract %429, %430 : tensor<1xi32> loc(#loc1074) + %432 = call @_take_248(%415, %431) : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi32> loc(#loc1075) + %433 = stablehlo.broadcast_in_dim %cst_6, dims = [] : (tensor) -> tensor<1xbf16> loc(#loc1076) + %434 = stablehlo.broadcast_in_dim %432, dims = [0] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1076) + %435 = "stablehlo.scatter"(%433, %434, %414) <{scatter_dimension_numbers = #stablehlo.scatter}> ({ + ^bb0(%arg12: tensor loc("scatter-max"), %arg13: tensor loc("scatter-max")): + %458 = stablehlo.maximum %arg12, %arg13 : tensor loc(#loc1078) + stablehlo.return %458 : tensor loc(#loc1077) + }) : (tensor<1xbf16>, tensor<1x1xi32>, tensor<1xbf16>) -> tensor<1xbf16> loc(#loc1077) + %436 = stablehlo.slice %435 [0:1] : (tensor<1xbf16>) -> tensor<1xbf16> loc(#loc1079) + %437 = stablehlo.reshape %436 : (tensor<1xbf16>) -> tensor loc(#loc1080) + %438 = stablehlo.broadcast_in_dim %cst_11, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1081) + %439 = stablehlo.convert %437 : (tensor) -> tensor loc(#loc1082) + %440 = stablehlo.divide %cst, %439 : tensor loc(#loc1083) + %441 = stablehlo.divide %440, %cst_11 : tensor loc(#loc1084) + %442 = stablehlo.compare GT, %437, %cst_7, FLOAT : (tensor, tensor) -> tensor loc(#loc1085) + %443 = call @_where_237(%442, %441, %438) : (tensor, tensor, tensor<1xf32>) -> tensor<1xf32> loc(#loc1086) + %444 = stablehlo.is_finite %437 : (tensor) -> tensor loc(#loc1087) + %445 = call @_where_239(%444, %443, %438) : (tensor, tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> loc(#loc1088) + %446 = stablehlo.slice %445 [0:1] : (tensor<1xf32>) -> tensor<1xf32> loc(#loc1089) + %447 = stablehlo.reshape %446 : (tensor<1xf32>) -> tensor loc(#loc1090) + %448 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1091) + %449 = "stablehlo.scatter"(%412, %448, %447) <{indices_are_sorted = true, scatter_dimension_numbers = #stablehlo.scatter, unique_indices = true}> ({ + ^bb0(%arg12: tensor loc("scatter"), %arg13: tensor loc("scatter")): + stablehlo.return %arg13 : tensor loc(#loc1092) + }) : (tensor<1xf32>, tensor<1xi32>, tensor) -> tensor<1xf32> loc(#loc1092) + %450:5 = stablehlo.custom_call @te_grouped_quantize_ffi(%arg11, %449, %411) {mhlo.backend_config = {flatten_axis = -1 : i64, q_layout = 0 : i64, scaling_mode = 3 : i64}, operand_layouts = [dense<[2, 1, 0]> : tensor<3xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>]} : (tensor<1x512x512xbf16>, tensor<1xf32>, tensor<1xi32>) -> (tensor<262144xf8E4M3FN>, tensor<1xf8E4M3FN>, tensor<1xf32>, tensor<1xf32>, tensor<1xf32>) loc(#loc1093) + %451 = stablehlo.broadcast_in_dim %c_17, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1094) + %452:2 = stablehlo.custom_call @te_grouped_gemm_ffi(%410#0, %410#2, %450#0, %450#2, %cst_10, %c_18, %451) {mhlo.backend_config = {K = 512 : i64, M = 128 : i64, N = 512 : i64, has_bias = false, is_grouped_dense_wgrad = false, lhs_is_trans = false, rhs_is_trans = false, scaling_mode = 3 : i64, use_async_d2h_group_sizes = false}, operand_layouts = [dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>, dense<> : tensor<0xindex>, dense<0> : tensor<1xindex>, dense<0> : tensor<1xindex>], result_layouts = [dense<[1, 0]> : tensor<2xindex>, dense<0> : tensor<1xindex>]} : (tensor<65536xf8E4M3FN>, tensor<1xf32>, tensor<262144xf8E4M3FN>, tensor<1xf32>, tensor, tensor<1xi32>, tensor<1xi32>) -> (tensor<128x512xbf16>, tensor<134218016xui8>) loc(#loc1095) + %453 = stablehlo.convert %452#0 : (tensor<128x512xbf16>) -> tensor<128x512xf32> loc(#loc1096) + %454 = stablehlo.reduce(%453 init: %cst_10) applies stablehlo.add across dimensions = [0, 1] : (tensor<128x512xf32>, tensor) -> tensor loc(#loc1097) + %455 = stablehlo.divide %454, %cst_0 : tensor loc(#loc1098) + %456 = stablehlo.convert %455 : (tensor) -> tensor loc(#loc1096) + %457 = stablehlo.add %215, %456 : tensor loc(#loc1099) + return %457, %369, %370, %318#0, %365#0, %366#2, %366#1, %287#2, %287#1, %286#0, %253#0, %368 : tensor, tensor<1x1x128x512xbf16>, tensor<1x1x128x512xbf16>, tensor<512x512xbf16>, tensor<512x3x512xbf16>, tensor<512xbf16>, tensor<512xbf16>, tensor<512xbf16>, tensor<512xbf16>, tensor<512x1x2048xbf16>, tensor<2048x512xbf16>, tensor<8x32xbf16> loc(#loc) + } loc(#loc) + func.func private @_threefry_fold_in(%arg0: tensor<2xui32> loc(unknown), %arg1: tensor loc(unknown)) -> tensor<2xui32> { + %c = stablehlo.constant dense<4294967295> : tensor loc(#loc1195) + %c_0 = stablehlo.constant dense<32> : tensor loc(#loc1195) + %0 = stablehlo.shift_right_logical %arg1, %c_0 : tensor loc(#loc1101) + %1 = stablehlo.broadcast_in_dim %0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1102) + %2 = stablehlo.and %arg1, %c : tensor loc(#loc1103) + %3 = stablehlo.broadcast_in_dim %2, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1102) + %4 = stablehlo.concatenate %1, %3, dim = 0 : (tensor<1xui32>, tensor<1xui32>) -> tensor<2xui32> loc(#loc1104) + %5 = stablehlo.slice %arg0 [0:1] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1105) + %6 = stablehlo.reshape %5 : (tensor<1xui32>) -> tensor loc(#loc1106) + %7 = stablehlo.slice %arg0 [1:2] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1105) + %8 = stablehlo.reshape %7 : (tensor<1xui32>) -> tensor loc(#loc1106) + %9 = stablehlo.slice %4 [0:1] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1107) + %10 = stablehlo.slice %4 [1:2] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1107) + %11:2 = call @threefry2x32(%6, %8, %9, %10) : (tensor, tensor, tensor<1xui32>, tensor<1xui32>) -> (tensor<1xui32>, tensor<1xui32>) loc(#loc785) + %12 = stablehlo.concatenate %11#0, %11#1, dim = 0 : (tensor<1xui32>, tensor<1xui32>) -> tensor<2xui32> loc(#loc1104) + return %12 : tensor<2xui32> loc(#loc1195) + } loc(#loc1195) + func.func private @threefry2x32(%arg0: tensor loc(callsite(#loc198 at #loc592)), %arg1: tensor loc(callsite(#loc198 at #loc592)), %arg2: tensor<1xui32> loc(callsite(#loc198 at #loc592)), %arg3: tensor<1xui32> loc(callsite(#loc198 at #loc592))) -> (tensor<1xui32>, tensor<1xui32>) { + %c = stablehlo.constant dense<5> : tensor loc(#loc172) + %c_0 = stablehlo.constant dense<4> : tensor loc(#loc172) + %c_1 = stablehlo.constant dense<2> : tensor loc(#loc172) + %c_2 = stablehlo.constant dense<8> : tensor loc(#loc) + %c_3 = stablehlo.constant dense<24> : tensor loc(#loc) + %c_4 = stablehlo.constant dense<16> : tensor loc(#loc) + %c_5 = stablehlo.constant dense<3> : tensor loc(#loc) + %c_6 = stablehlo.constant dense<29> : tensor loc(#loc) + %c_7 = stablehlo.constant dense<1> : tensor loc(#loc172) + %c_8 = stablehlo.constant dense<6> : tensor loc(#loc) + %c_9 = stablehlo.constant dense<26> : tensor loc(#loc) + %c_10 = stablehlo.constant dense<17> : tensor loc(#loc) + %c_11 = stablehlo.constant dense<15> : tensor loc(#loc) + %c_12 = stablehlo.constant dense<19> : tensor loc(#loc) + %c_13 = stablehlo.constant dense<13> : tensor loc(#loc) + %c_14 = stablehlo.constant dense<466688986> : tensor loc(#loc172) + %0 = stablehlo.xor %arg0, %arg1 : tensor loc(#loc1108) + %1 = stablehlo.xor %0, %c_14 : tensor loc(#loc1108) + %2 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %3 = stablehlo.add %arg2, %2 : tensor<1xui32> loc(#loc1109) + %4 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %5 = stablehlo.add %arg3, %4 : tensor<1xui32> loc(#loc1109) + %6 = stablehlo.add %3, %5 : tensor<1xui32> loc(#loc1109) + %7 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %8 = stablehlo.shift_left %5, %7 : tensor<1xui32> loc(#loc1110) + %9 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %10 = stablehlo.shift_right_logical %5, %9 : tensor<1xui32> loc(#loc1101) + %11 = stablehlo.or %8, %10 : tensor<1xui32> loc(#loc1111) + %12 = stablehlo.xor %6, %11 : tensor<1xui32> loc(#loc1108) + %13 = stablehlo.add %6, %12 : tensor<1xui32> loc(#loc1109) + %14 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %15 = stablehlo.shift_left %12, %14 : tensor<1xui32> loc(#loc1110) + %16 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %17 = stablehlo.shift_right_logical %12, %16 : tensor<1xui32> loc(#loc1101) + %18 = stablehlo.or %15, %17 : tensor<1xui32> loc(#loc1111) + %19 = stablehlo.xor %13, %18 : tensor<1xui32> loc(#loc1108) + %20 = stablehlo.add %13, %19 : tensor<1xui32> loc(#loc1109) + %21 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %22 = stablehlo.shift_left %19, %21 : tensor<1xui32> loc(#loc1110) + %23 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %24 = stablehlo.shift_right_logical %19, %23 : tensor<1xui32> loc(#loc1101) + %25 = stablehlo.or %22, %24 : tensor<1xui32> loc(#loc1111) + %26 = stablehlo.xor %20, %25 : tensor<1xui32> loc(#loc1108) + %27 = stablehlo.add %20, %26 : tensor<1xui32> loc(#loc1109) + %28 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %29 = stablehlo.shift_left %26, %28 : tensor<1xui32> loc(#loc1110) + %30 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %31 = stablehlo.shift_right_logical %26, %30 : tensor<1xui32> loc(#loc1101) + %32 = stablehlo.or %29, %31 : tensor<1xui32> loc(#loc1111) + %33 = stablehlo.xor %27, %32 : tensor<1xui32> loc(#loc1108) + %34 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %35 = stablehlo.add %27, %34 : tensor<1xui32> loc(#loc1109) + %36 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %37 = stablehlo.add %33, %36 : tensor<1xui32> loc(#loc1109) + %38 = stablehlo.broadcast_in_dim %c_7, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %39 = stablehlo.add %37, %38 : tensor<1xui32> loc(#loc1109) + %40 = stablehlo.add %35, %39 : tensor<1xui32> loc(#loc1109) + %41 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %42 = stablehlo.shift_left %39, %41 : tensor<1xui32> loc(#loc1110) + %43 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %44 = stablehlo.shift_right_logical %39, %43 : tensor<1xui32> loc(#loc1101) + %45 = stablehlo.or %42, %44 : tensor<1xui32> loc(#loc1111) + %46 = stablehlo.xor %40, %45 : tensor<1xui32> loc(#loc1108) + %47 = stablehlo.add %40, %46 : tensor<1xui32> loc(#loc1109) + %48 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %49 = stablehlo.shift_left %46, %48 : tensor<1xui32> loc(#loc1110) + %50 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %51 = stablehlo.shift_right_logical %46, %50 : tensor<1xui32> loc(#loc1101) + %52 = stablehlo.or %49, %51 : tensor<1xui32> loc(#loc1111) + %53 = stablehlo.xor %47, %52 : tensor<1xui32> loc(#loc1108) + %54 = stablehlo.add %47, %53 : tensor<1xui32> loc(#loc1109) + %55 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %56 = stablehlo.shift_left %53, %55 : tensor<1xui32> loc(#loc1110) + %57 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %58 = stablehlo.shift_right_logical %53, %57 : tensor<1xui32> loc(#loc1101) + %59 = stablehlo.or %56, %58 : tensor<1xui32> loc(#loc1111) + %60 = stablehlo.xor %54, %59 : tensor<1xui32> loc(#loc1108) + %61 = stablehlo.add %54, %60 : tensor<1xui32> loc(#loc1109) + %62 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %63 = stablehlo.shift_left %60, %62 : tensor<1xui32> loc(#loc1110) + %64 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %65 = stablehlo.shift_right_logical %60, %64 : tensor<1xui32> loc(#loc1101) + %66 = stablehlo.or %63, %65 : tensor<1xui32> loc(#loc1111) + %67 = stablehlo.xor %61, %66 : tensor<1xui32> loc(#loc1108) + %68 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %69 = stablehlo.add %61, %68 : tensor<1xui32> loc(#loc1109) + %70 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %71 = stablehlo.add %67, %70 : tensor<1xui32> loc(#loc1109) + %72 = stablehlo.broadcast_in_dim %c_1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %73 = stablehlo.add %71, %72 : tensor<1xui32> loc(#loc1109) + %74 = stablehlo.add %69, %73 : tensor<1xui32> loc(#loc1109) + %75 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %76 = stablehlo.shift_left %73, %75 : tensor<1xui32> loc(#loc1110) + %77 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %78 = stablehlo.shift_right_logical %73, %77 : tensor<1xui32> loc(#loc1101) + %79 = stablehlo.or %76, %78 : tensor<1xui32> loc(#loc1111) + %80 = stablehlo.xor %74, %79 : tensor<1xui32> loc(#loc1108) + %81 = stablehlo.add %74, %80 : tensor<1xui32> loc(#loc1109) + %82 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %83 = stablehlo.shift_left %80, %82 : tensor<1xui32> loc(#loc1110) + %84 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %85 = stablehlo.shift_right_logical %80, %84 : tensor<1xui32> loc(#loc1101) + %86 = stablehlo.or %83, %85 : tensor<1xui32> loc(#loc1111) + %87 = stablehlo.xor %81, %86 : tensor<1xui32> loc(#loc1108) + %88 = stablehlo.add %81, %87 : tensor<1xui32> loc(#loc1109) + %89 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %90 = stablehlo.shift_left %87, %89 : tensor<1xui32> loc(#loc1110) + %91 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %92 = stablehlo.shift_right_logical %87, %91 : tensor<1xui32> loc(#loc1101) + %93 = stablehlo.or %90, %92 : tensor<1xui32> loc(#loc1111) + %94 = stablehlo.xor %88, %93 : tensor<1xui32> loc(#loc1108) + %95 = stablehlo.add %88, %94 : tensor<1xui32> loc(#loc1109) + %96 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %97 = stablehlo.shift_left %94, %96 : tensor<1xui32> loc(#loc1110) + %98 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %99 = stablehlo.shift_right_logical %94, %98 : tensor<1xui32> loc(#loc1101) + %100 = stablehlo.or %97, %99 : tensor<1xui32> loc(#loc1111) + %101 = stablehlo.xor %95, %100 : tensor<1xui32> loc(#loc1108) + %102 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %103 = stablehlo.add %95, %102 : tensor<1xui32> loc(#loc1109) + %104 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %105 = stablehlo.add %101, %104 : tensor<1xui32> loc(#loc1109) + %106 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %107 = stablehlo.add %105, %106 : tensor<1xui32> loc(#loc1109) + %108 = stablehlo.add %103, %107 : tensor<1xui32> loc(#loc1109) + %109 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %110 = stablehlo.shift_left %107, %109 : tensor<1xui32> loc(#loc1110) + %111 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %112 = stablehlo.shift_right_logical %107, %111 : tensor<1xui32> loc(#loc1101) + %113 = stablehlo.or %110, %112 : tensor<1xui32> loc(#loc1111) + %114 = stablehlo.xor %108, %113 : tensor<1xui32> loc(#loc1108) + %115 = stablehlo.add %108, %114 : tensor<1xui32> loc(#loc1109) + %116 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %117 = stablehlo.shift_left %114, %116 : tensor<1xui32> loc(#loc1110) + %118 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %119 = stablehlo.shift_right_logical %114, %118 : tensor<1xui32> loc(#loc1101) + %120 = stablehlo.or %117, %119 : tensor<1xui32> loc(#loc1111) + %121 = stablehlo.xor %115, %120 : tensor<1xui32> loc(#loc1108) + %122 = stablehlo.add %115, %121 : tensor<1xui32> loc(#loc1109) + %123 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %124 = stablehlo.shift_left %121, %123 : tensor<1xui32> loc(#loc1110) + %125 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %126 = stablehlo.shift_right_logical %121, %125 : tensor<1xui32> loc(#loc1101) + %127 = stablehlo.or %124, %126 : tensor<1xui32> loc(#loc1111) + %128 = stablehlo.xor %122, %127 : tensor<1xui32> loc(#loc1108) + %129 = stablehlo.add %122, %128 : tensor<1xui32> loc(#loc1109) + %130 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %131 = stablehlo.shift_left %128, %130 : tensor<1xui32> loc(#loc1110) + %132 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %133 = stablehlo.shift_right_logical %128, %132 : tensor<1xui32> loc(#loc1101) + %134 = stablehlo.or %131, %133 : tensor<1xui32> loc(#loc1111) + %135 = stablehlo.xor %129, %134 : tensor<1xui32> loc(#loc1108) + %136 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %137 = stablehlo.add %129, %136 : tensor<1xui32> loc(#loc1109) + %138 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %139 = stablehlo.add %135, %138 : tensor<1xui32> loc(#loc1109) + %140 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %141 = stablehlo.add %139, %140 : tensor<1xui32> loc(#loc1109) + %142 = stablehlo.add %137, %141 : tensor<1xui32> loc(#loc1109) + %143 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %144 = stablehlo.shift_left %141, %143 : tensor<1xui32> loc(#loc1110) + %145 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %146 = stablehlo.shift_right_logical %141, %145 : tensor<1xui32> loc(#loc1101) + %147 = stablehlo.or %144, %146 : tensor<1xui32> loc(#loc1111) + %148 = stablehlo.xor %142, %147 : tensor<1xui32> loc(#loc1108) + %149 = stablehlo.add %142, %148 : tensor<1xui32> loc(#loc1109) + %150 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %151 = stablehlo.shift_left %148, %150 : tensor<1xui32> loc(#loc1110) + %152 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %153 = stablehlo.shift_right_logical %148, %152 : tensor<1xui32> loc(#loc1101) + %154 = stablehlo.or %151, %153 : tensor<1xui32> loc(#loc1111) + %155 = stablehlo.xor %149, %154 : tensor<1xui32> loc(#loc1108) + %156 = stablehlo.add %149, %155 : tensor<1xui32> loc(#loc1109) + %157 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %158 = stablehlo.shift_left %155, %157 : tensor<1xui32> loc(#loc1110) + %159 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %160 = stablehlo.shift_right_logical %155, %159 : tensor<1xui32> loc(#loc1101) + %161 = stablehlo.or %158, %160 : tensor<1xui32> loc(#loc1111) + %162 = stablehlo.xor %156, %161 : tensor<1xui32> loc(#loc1108) + %163 = stablehlo.add %156, %162 : tensor<1xui32> loc(#loc1109) + %164 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1110) + %165 = stablehlo.shift_left %162, %164 : tensor<1xui32> loc(#loc1110) + %166 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1101) + %167 = stablehlo.shift_right_logical %162, %166 : tensor<1xui32> loc(#loc1101) + %168 = stablehlo.or %165, %167 : tensor<1xui32> loc(#loc1111) + %169 = stablehlo.xor %163, %168 : tensor<1xui32> loc(#loc1108) + %170 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %171 = stablehlo.add %163, %170 : tensor<1xui32> loc(#loc1109) + %172 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %173 = stablehlo.add %169, %172 : tensor<1xui32> loc(#loc1109) + %174 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xui32> loc(#loc1109) + %175 = stablehlo.add %173, %174 : tensor<1xui32> loc(#loc1109) + return %171, %175 : tensor<1xui32>, tensor<1xui32> loc(#loc785) + } loc(#loc785) + func.func private @fold_in(%arg0: tensor<2xui32> loc(unknown), %arg1: tensor loc(unknown)) -> tensor<2xui32> { + %0 = stablehlo.convert %arg1 : (tensor) -> tensor loc(#loc1113) + %1 = call @_threefry_fold_in(%arg0, %0) : (tensor<2xui32>, tensor) -> tensor<2xui32> loc(#loc1114) + return %1 : tensor<2xui32> loc(#loc1196) + } loc(#loc1196) + func.func private @_randint(%arg0: tensor<2xui32> loc(unknown), %arg1: tensor loc(unknown), %arg2: tensor loc(unknown)) -> tensor<1x4xi32> { + %c = stablehlo.constant dense<65536> : tensor loc(#loc1197) + %c_0 = stablehlo.constant dense<1> : tensor loc(#loc) + %c_1 = stablehlo.constant dense<32> : tensor loc(#loc) + %c_2 = stablehlo.constant dense<1> : tensor loc(#loc) + %c_3 = stablehlo.constant dense<4> : tensor loc(#loc) + %c_4 = stablehlo.constant dense<2147483647> : tensor loc(#loc) + %c_5 = stablehlo.constant dense<-2147483648> : tensor loc(#loc) + %0 = call @clip(%c_4, %c_5, %c_4) : (tensor, tensor, tensor) -> tensor loc(#loc1116) + %1 = stablehlo.compare GT, %arg2, %0, SIGNED : (tensor, tensor) -> tensor loc(#loc1117) + %2 = call @clip_9(%arg1, %c_5, %c_4) : (tensor, tensor, tensor) -> tensor loc(#loc1116) + %3 = stablehlo.convert %2 : tensor loc(#loc1118) + %4 = call @clip_9(%arg2, %c_5, %c_4) : (tensor, tensor, tensor) -> tensor loc(#loc1116) + %5 = stablehlo.convert %4 : tensor loc(#loc1118) + %6 = stablehlo.broadcast_in_dim %3, dims = [] : (tensor) -> tensor<1x1xi32> loc(#loc1119) + %7 = stablehlo.broadcast_in_dim %5, dims = [] : (tensor) -> tensor<1x1xi32> loc(#loc1119) + %8 = call @_threefry_split(%arg0) : (tensor<2xui32>) -> tensor<2x2xui32> loc(#loc1120) + %9 = stablehlo.slice %8 [0:1, 0:2] : (tensor<2x2xui32>) -> tensor<1x2xui32> loc(#loc1121) + %10 = stablehlo.reshape %9 : (tensor<1x2xui32>) -> tensor<2xui32> loc(#loc1122) + %11 = stablehlo.slice %8 [1:2, 0:2] : (tensor<2x2xui32>) -> tensor<1x2xui32> loc(#loc1121) + %12 = stablehlo.reshape %11 : (tensor<1x2xui32>) -> tensor<2xui32> loc(#loc1122) + %13 = stablehlo.slice %10 [0:1] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %14 = stablehlo.reshape %13 : (tensor<1xui32>) -> tensor loc(#loc1122) + %15 = stablehlo.slice %10 [1:2] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %16 = stablehlo.reshape %15 : (tensor<1xui32>) -> tensor loc(#loc1122) + %17 = stablehlo.iota dim = 0 : tensor<1x4xui64> loc(#loc1123) + %18 = stablehlo.iota dim = 1 : tensor<1x4xui64> loc(#loc1123) + %19 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %20 = stablehlo.multiply %19, %17 : tensor<1x4xui64> loc(#loc1123) + %21 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %22 = stablehlo.multiply %21, %18 : tensor<1x4xui64> loc(#loc1123) + %23 = stablehlo.add %20, %22 : tensor<1x4xui64> loc(#loc1123) + %24 = stablehlo.broadcast_in_dim %c_1, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %25 = stablehlo.shift_right_logical %23, %24 : tensor<1x4xui64> loc(#loc1123) + %26 = stablehlo.convert %23 : (tensor<1x4xui64>) -> tensor<1x4xui32> loc(#loc1123) + %27 = stablehlo.convert %25 : (tensor<1x4xui64>) -> tensor<1x4xui32> loc(#loc1123) + %28:2 = call @threefry2x32_27(%14, %16, %27, %26) : (tensor, tensor, tensor<1x4xui32>, tensor<1x4xui32>) -> (tensor<1x4xui32>, tensor<1x4xui32>) loc(#loc1124) + %29 = stablehlo.xor %28#0, %28#1 : tensor<1x4xui32> loc(#loc1125) + %30 = stablehlo.slice %12 [0:1] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %31 = stablehlo.reshape %30 : (tensor<1xui32>) -> tensor loc(#loc1122) + %32 = stablehlo.slice %12 [1:2] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %33 = stablehlo.reshape %32 : (tensor<1xui32>) -> tensor loc(#loc1122) + %34 = stablehlo.iota dim = 0 : tensor<1x4xui64> loc(#loc1123) + %35 = stablehlo.iota dim = 1 : tensor<1x4xui64> loc(#loc1123) + %36 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %37 = stablehlo.multiply %36, %34 : tensor<1x4xui64> loc(#loc1123) + %38 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %39 = stablehlo.multiply %38, %35 : tensor<1x4xui64> loc(#loc1123) + %40 = stablehlo.add %37, %39 : tensor<1x4xui64> loc(#loc1123) + %41 = stablehlo.broadcast_in_dim %c_1, dims = [] : (tensor) -> tensor<1x4xui64> loc(#loc1123) + %42 = stablehlo.shift_right_logical %40, %41 : tensor<1x4xui64> loc(#loc1123) + %43 = stablehlo.convert %40 : (tensor<1x4xui64>) -> tensor<1x4xui32> loc(#loc1123) + %44 = stablehlo.convert %42 : (tensor<1x4xui64>) -> tensor<1x4xui32> loc(#loc1123) + %45:2 = call @threefry2x32_27(%31, %33, %44, %43) : (tensor, tensor, tensor<1x4xui32>, tensor<1x4xui32>) -> (tensor<1x4xui32>, tensor<1x4xui32>) loc(#loc1124) + %46 = stablehlo.xor %45#0, %45#1 : tensor<1x4xui32> loc(#loc1125) + %47 = stablehlo.subtract %7, %6 : tensor<1x1xi32> loc(#loc1126) + %48 = stablehlo.convert %47 : (tensor<1x1xi32>) -> tensor<1x1xui32> loc(#loc1118) + %49 = stablehlo.compare LE, %7, %6, SIGNED : (tensor<1x1xi32>, tensor<1x1xi32>) -> tensor<1x1xi1> loc(#loc1127) + %50 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<1x1xui32> loc(#loc1119) + %51 = stablehlo.select %49, %50, %48 : tensor<1x1xi1>, tensor<1x1xui32> loc(#loc1128) + %52 = stablehlo.compare GT, %7, %6, SIGNED : (tensor<1x1xi32>, tensor<1x1xi32>) -> tensor<1x1xi1> loc(#loc1117) + %53 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1x1xi1> loc(#loc1129) + %54 = stablehlo.and %53, %52 : tensor<1x1xi1> loc(#loc1129) + %55 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<1x1xui32> loc(#loc1130) + %56 = stablehlo.add %51, %55 : tensor<1x1xui32> loc(#loc1130) + %57 = stablehlo.select %54, %56, %51 : tensor<1x1xi1>, tensor<1x1xui32> loc(#loc1128) + %58 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1x1xui32> loc(#loc1131) + %59 = stablehlo.remainder %58, %57 : tensor<1x1xui32> loc(#loc1131) + %60 = stablehlo.multiply %59, %59 : tensor<1x1xui32> loc(#loc1132) + %61 = stablehlo.remainder %60, %57 : tensor<1x1xui32> loc(#loc1131) + %62 = stablehlo.broadcast_in_dim %57, dims = [0, 1] : (tensor<1x1xui32>) -> tensor<1x4xui32> loc(#loc1131) + %63 = stablehlo.remainder %29, %62 : tensor<1x4xui32> loc(#loc1131) + %64 = stablehlo.broadcast_in_dim %61, dims = [0, 1] : (tensor<1x1xui32>) -> tensor<1x4xui32> loc(#loc1132) + %65 = stablehlo.multiply %63, %64 : tensor<1x4xui32> loc(#loc1132) + %66 = stablehlo.broadcast_in_dim %57, dims = [0, 1] : (tensor<1x1xui32>) -> tensor<1x4xui32> loc(#loc1131) + %67 = stablehlo.remainder %46, %66 : tensor<1x4xui32> loc(#loc1131) + %68 = stablehlo.add %65, %67 : tensor<1x4xui32> loc(#loc1130) + %69 = stablehlo.broadcast_in_dim %57, dims = [0, 1] : (tensor<1x1xui32>) -> tensor<1x4xui32> loc(#loc1131) + %70 = stablehlo.remainder %68, %69 : tensor<1x4xui32> loc(#loc1131) + %71 = stablehlo.convert %70 : (tensor<1x4xui32>) -> tensor<1x4xi32> loc(#loc1118) + %72 = stablehlo.broadcast_in_dim %6, dims = [0, 1] : (tensor<1x1xi32>) -> tensor<1x4xi32> loc(#loc1130) + %73 = stablehlo.add %72, %71 : tensor<1x4xi32> loc(#loc1130) + return %73 : tensor<1x4xi32> loc(#loc1197) + } loc(#loc1197) + func.func private @clip(%arg0: tensor loc(unknown), %arg1: tensor loc(unknown), %arg2: tensor loc(unknown)) -> tensor { + %0 = stablehlo.maximum %arg1, %arg0 : tensor loc(#loc1134) + %1 = stablehlo.minimum %arg2, %0 : tensor loc(#loc1135) + return %1 : tensor loc(#loc1198) + } loc(#loc1198) + func.func private @clip_9(%arg0: tensor loc(unknown), %arg1: tensor loc(unknown), %arg2: tensor loc(unknown)) -> tensor { + %0 = stablehlo.maximum %arg1, %arg0 : tensor loc(#loc1134) + %1 = stablehlo.minimum %arg2, %0 : tensor loc(#loc1135) + return %1 : tensor loc(#loc1198) + } loc(#loc1198) + func.func private @_threefry_split(%arg0: tensor<2xui32> loc(unknown)) -> tensor<2x2xui32> { + %c = stablehlo.constant dense<32> : tensor loc(#loc1123) + %c_0 = stablehlo.constant dense<1> : tensor loc(#loc1123) + %0 = stablehlo.slice %arg0 [0:1] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %1 = stablehlo.reshape %0 : (tensor<1xui32>) -> tensor loc(#loc1122) + %2 = stablehlo.slice %arg0 [1:2] : (tensor<2xui32>) -> tensor<1xui32> loc(#loc1121) + %3 = stablehlo.reshape %2 : (tensor<1xui32>) -> tensor loc(#loc1122) + %4 = stablehlo.iota dim = 0 : tensor<2xui64> loc(#loc1123) + %5 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<2xui64> loc(#loc1123) + %6 = stablehlo.multiply %5, %4 : tensor<2xui64> loc(#loc1123) + %7 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<2xui64> loc(#loc1123) + %8 = stablehlo.shift_right_logical %6, %7 : tensor<2xui64> loc(#loc1123) + %9 = stablehlo.convert %6 : (tensor<2xui64>) -> tensor<2xui32> loc(#loc1123) + %10 = stablehlo.convert %8 : (tensor<2xui64>) -> tensor<2xui32> loc(#loc1123) + %11:2 = call @threefry2x32_14(%1, %3, %10, %9) : (tensor, tensor, tensor<2xui32>, tensor<2xui32>) -> (tensor<2xui32>, tensor<2xui32>) loc(#loc786) + %12 = stablehlo.broadcast_in_dim %11#0, dims = [0] : (tensor<2xui32>) -> tensor<2x1xui32> loc(#loc1119) + %13 = stablehlo.broadcast_in_dim %11#1, dims = [0] : (tensor<2xui32>) -> tensor<2x1xui32> loc(#loc1119) + %14 = stablehlo.concatenate %12, %13, dim = 1 : (tensor<2x1xui32>, tensor<2x1xui32>) -> tensor<2x2xui32> loc(#loc1136) + return %14 : tensor<2x2xui32> loc(#loc1198) + } loc(#loc1198) + func.func private @threefry2x32_14(%arg0: tensor loc(callsite(#loc199 at #loc592)), %arg1: tensor loc(callsite(#loc199 at #loc592)), %arg2: tensor<2xui32> loc(callsite(#loc199 at #loc592)), %arg3: tensor<2xui32> loc(callsite(#loc199 at #loc592))) -> (tensor<2xui32>, tensor<2xui32>) { + %c = stablehlo.constant dense<5> : tensor loc(#loc172) + %c_0 = stablehlo.constant dense<4> : tensor loc(#loc172) + %c_1 = stablehlo.constant dense<2> : tensor loc(#loc172) + %c_2 = stablehlo.constant dense<8> : tensor loc(#loc) + %c_3 = stablehlo.constant dense<24> : tensor loc(#loc) + %c_4 = stablehlo.constant dense<16> : tensor loc(#loc) + %c_5 = stablehlo.constant dense<3> : tensor loc(#loc) + %c_6 = stablehlo.constant dense<29> : tensor loc(#loc) + %c_7 = stablehlo.constant dense<1> : tensor loc(#loc172) + %c_8 = stablehlo.constant dense<6> : tensor loc(#loc) + %c_9 = stablehlo.constant dense<26> : tensor loc(#loc) + %c_10 = stablehlo.constant dense<17> : tensor loc(#loc) + %c_11 = stablehlo.constant dense<15> : tensor loc(#loc) + %c_12 = stablehlo.constant dense<19> : tensor loc(#loc) + %c_13 = stablehlo.constant dense<13> : tensor loc(#loc) + %c_14 = stablehlo.constant dense<466688986> : tensor loc(#loc172) + %0 = stablehlo.xor %arg0, %arg1 : tensor loc(#loc1125) + %1 = stablehlo.xor %0, %c_14 : tensor loc(#loc1125) + %2 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %3 = stablehlo.add %arg2, %2 : tensor<2xui32> loc(#loc1130) + %4 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %5 = stablehlo.add %arg3, %4 : tensor<2xui32> loc(#loc1130) + %6 = stablehlo.add %3, %5 : tensor<2xui32> loc(#loc1130) + %7 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %8 = stablehlo.shift_left %5, %7 : tensor<2xui32> loc(#loc1137) + %9 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %10 = stablehlo.shift_right_logical %5, %9 : tensor<2xui32> loc(#loc1138) + %11 = stablehlo.or %8, %10 : tensor<2xui32> loc(#loc1139) + %12 = stablehlo.xor %6, %11 : tensor<2xui32> loc(#loc1125) + %13 = stablehlo.add %6, %12 : tensor<2xui32> loc(#loc1130) + %14 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %15 = stablehlo.shift_left %12, %14 : tensor<2xui32> loc(#loc1137) + %16 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %17 = stablehlo.shift_right_logical %12, %16 : tensor<2xui32> loc(#loc1138) + %18 = stablehlo.or %15, %17 : tensor<2xui32> loc(#loc1139) + %19 = stablehlo.xor %13, %18 : tensor<2xui32> loc(#loc1125) + %20 = stablehlo.add %13, %19 : tensor<2xui32> loc(#loc1130) + %21 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %22 = stablehlo.shift_left %19, %21 : tensor<2xui32> loc(#loc1137) + %23 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %24 = stablehlo.shift_right_logical %19, %23 : tensor<2xui32> loc(#loc1138) + %25 = stablehlo.or %22, %24 : tensor<2xui32> loc(#loc1139) + %26 = stablehlo.xor %20, %25 : tensor<2xui32> loc(#loc1125) + %27 = stablehlo.add %20, %26 : tensor<2xui32> loc(#loc1130) + %28 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %29 = stablehlo.shift_left %26, %28 : tensor<2xui32> loc(#loc1137) + %30 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %31 = stablehlo.shift_right_logical %26, %30 : tensor<2xui32> loc(#loc1138) + %32 = stablehlo.or %29, %31 : tensor<2xui32> loc(#loc1139) + %33 = stablehlo.xor %27, %32 : tensor<2xui32> loc(#loc1125) + %34 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %35 = stablehlo.add %27, %34 : tensor<2xui32> loc(#loc1130) + %36 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %37 = stablehlo.add %33, %36 : tensor<2xui32> loc(#loc1130) + %38 = stablehlo.broadcast_in_dim %c_7, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %39 = stablehlo.add %37, %38 : tensor<2xui32> loc(#loc1130) + %40 = stablehlo.add %35, %39 : tensor<2xui32> loc(#loc1130) + %41 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %42 = stablehlo.shift_left %39, %41 : tensor<2xui32> loc(#loc1137) + %43 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %44 = stablehlo.shift_right_logical %39, %43 : tensor<2xui32> loc(#loc1138) + %45 = stablehlo.or %42, %44 : tensor<2xui32> loc(#loc1139) + %46 = stablehlo.xor %40, %45 : tensor<2xui32> loc(#loc1125) + %47 = stablehlo.add %40, %46 : tensor<2xui32> loc(#loc1130) + %48 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %49 = stablehlo.shift_left %46, %48 : tensor<2xui32> loc(#loc1137) + %50 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %51 = stablehlo.shift_right_logical %46, %50 : tensor<2xui32> loc(#loc1138) + %52 = stablehlo.or %49, %51 : tensor<2xui32> loc(#loc1139) + %53 = stablehlo.xor %47, %52 : tensor<2xui32> loc(#loc1125) + %54 = stablehlo.add %47, %53 : tensor<2xui32> loc(#loc1130) + %55 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %56 = stablehlo.shift_left %53, %55 : tensor<2xui32> loc(#loc1137) + %57 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %58 = stablehlo.shift_right_logical %53, %57 : tensor<2xui32> loc(#loc1138) + %59 = stablehlo.or %56, %58 : tensor<2xui32> loc(#loc1139) + %60 = stablehlo.xor %54, %59 : tensor<2xui32> loc(#loc1125) + %61 = stablehlo.add %54, %60 : tensor<2xui32> loc(#loc1130) + %62 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %63 = stablehlo.shift_left %60, %62 : tensor<2xui32> loc(#loc1137) + %64 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %65 = stablehlo.shift_right_logical %60, %64 : tensor<2xui32> loc(#loc1138) + %66 = stablehlo.or %63, %65 : tensor<2xui32> loc(#loc1139) + %67 = stablehlo.xor %61, %66 : tensor<2xui32> loc(#loc1125) + %68 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %69 = stablehlo.add %61, %68 : tensor<2xui32> loc(#loc1130) + %70 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %71 = stablehlo.add %67, %70 : tensor<2xui32> loc(#loc1130) + %72 = stablehlo.broadcast_in_dim %c_1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %73 = stablehlo.add %71, %72 : tensor<2xui32> loc(#loc1130) + %74 = stablehlo.add %69, %73 : tensor<2xui32> loc(#loc1130) + %75 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %76 = stablehlo.shift_left %73, %75 : tensor<2xui32> loc(#loc1137) + %77 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %78 = stablehlo.shift_right_logical %73, %77 : tensor<2xui32> loc(#loc1138) + %79 = stablehlo.or %76, %78 : tensor<2xui32> loc(#loc1139) + %80 = stablehlo.xor %74, %79 : tensor<2xui32> loc(#loc1125) + %81 = stablehlo.add %74, %80 : tensor<2xui32> loc(#loc1130) + %82 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %83 = stablehlo.shift_left %80, %82 : tensor<2xui32> loc(#loc1137) + %84 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %85 = stablehlo.shift_right_logical %80, %84 : tensor<2xui32> loc(#loc1138) + %86 = stablehlo.or %83, %85 : tensor<2xui32> loc(#loc1139) + %87 = stablehlo.xor %81, %86 : tensor<2xui32> loc(#loc1125) + %88 = stablehlo.add %81, %87 : tensor<2xui32> loc(#loc1130) + %89 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %90 = stablehlo.shift_left %87, %89 : tensor<2xui32> loc(#loc1137) + %91 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %92 = stablehlo.shift_right_logical %87, %91 : tensor<2xui32> loc(#loc1138) + %93 = stablehlo.or %90, %92 : tensor<2xui32> loc(#loc1139) + %94 = stablehlo.xor %88, %93 : tensor<2xui32> loc(#loc1125) + %95 = stablehlo.add %88, %94 : tensor<2xui32> loc(#loc1130) + %96 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %97 = stablehlo.shift_left %94, %96 : tensor<2xui32> loc(#loc1137) + %98 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %99 = stablehlo.shift_right_logical %94, %98 : tensor<2xui32> loc(#loc1138) + %100 = stablehlo.or %97, %99 : tensor<2xui32> loc(#loc1139) + %101 = stablehlo.xor %95, %100 : tensor<2xui32> loc(#loc1125) + %102 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %103 = stablehlo.add %95, %102 : tensor<2xui32> loc(#loc1130) + %104 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %105 = stablehlo.add %101, %104 : tensor<2xui32> loc(#loc1130) + %106 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %107 = stablehlo.add %105, %106 : tensor<2xui32> loc(#loc1130) + %108 = stablehlo.add %103, %107 : tensor<2xui32> loc(#loc1130) + %109 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %110 = stablehlo.shift_left %107, %109 : tensor<2xui32> loc(#loc1137) + %111 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %112 = stablehlo.shift_right_logical %107, %111 : tensor<2xui32> loc(#loc1138) + %113 = stablehlo.or %110, %112 : tensor<2xui32> loc(#loc1139) + %114 = stablehlo.xor %108, %113 : tensor<2xui32> loc(#loc1125) + %115 = stablehlo.add %108, %114 : tensor<2xui32> loc(#loc1130) + %116 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %117 = stablehlo.shift_left %114, %116 : tensor<2xui32> loc(#loc1137) + %118 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %119 = stablehlo.shift_right_logical %114, %118 : tensor<2xui32> loc(#loc1138) + %120 = stablehlo.or %117, %119 : tensor<2xui32> loc(#loc1139) + %121 = stablehlo.xor %115, %120 : tensor<2xui32> loc(#loc1125) + %122 = stablehlo.add %115, %121 : tensor<2xui32> loc(#loc1130) + %123 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %124 = stablehlo.shift_left %121, %123 : tensor<2xui32> loc(#loc1137) + %125 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %126 = stablehlo.shift_right_logical %121, %125 : tensor<2xui32> loc(#loc1138) + %127 = stablehlo.or %124, %126 : tensor<2xui32> loc(#loc1139) + %128 = stablehlo.xor %122, %127 : tensor<2xui32> loc(#loc1125) + %129 = stablehlo.add %122, %128 : tensor<2xui32> loc(#loc1130) + %130 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %131 = stablehlo.shift_left %128, %130 : tensor<2xui32> loc(#loc1137) + %132 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %133 = stablehlo.shift_right_logical %128, %132 : tensor<2xui32> loc(#loc1138) + %134 = stablehlo.or %131, %133 : tensor<2xui32> loc(#loc1139) + %135 = stablehlo.xor %129, %134 : tensor<2xui32> loc(#loc1125) + %136 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %137 = stablehlo.add %129, %136 : tensor<2xui32> loc(#loc1130) + %138 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %139 = stablehlo.add %135, %138 : tensor<2xui32> loc(#loc1130) + %140 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %141 = stablehlo.add %139, %140 : tensor<2xui32> loc(#loc1130) + %142 = stablehlo.add %137, %141 : tensor<2xui32> loc(#loc1130) + %143 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %144 = stablehlo.shift_left %141, %143 : tensor<2xui32> loc(#loc1137) + %145 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %146 = stablehlo.shift_right_logical %141, %145 : tensor<2xui32> loc(#loc1138) + %147 = stablehlo.or %144, %146 : tensor<2xui32> loc(#loc1139) + %148 = stablehlo.xor %142, %147 : tensor<2xui32> loc(#loc1125) + %149 = stablehlo.add %142, %148 : tensor<2xui32> loc(#loc1130) + %150 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %151 = stablehlo.shift_left %148, %150 : tensor<2xui32> loc(#loc1137) + %152 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %153 = stablehlo.shift_right_logical %148, %152 : tensor<2xui32> loc(#loc1138) + %154 = stablehlo.or %151, %153 : tensor<2xui32> loc(#loc1139) + %155 = stablehlo.xor %149, %154 : tensor<2xui32> loc(#loc1125) + %156 = stablehlo.add %149, %155 : tensor<2xui32> loc(#loc1130) + %157 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %158 = stablehlo.shift_left %155, %157 : tensor<2xui32> loc(#loc1137) + %159 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %160 = stablehlo.shift_right_logical %155, %159 : tensor<2xui32> loc(#loc1138) + %161 = stablehlo.or %158, %160 : tensor<2xui32> loc(#loc1139) + %162 = stablehlo.xor %156, %161 : tensor<2xui32> loc(#loc1125) + %163 = stablehlo.add %156, %162 : tensor<2xui32> loc(#loc1130) + %164 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1137) + %165 = stablehlo.shift_left %162, %164 : tensor<2xui32> loc(#loc1137) + %166 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1138) + %167 = stablehlo.shift_right_logical %162, %166 : tensor<2xui32> loc(#loc1138) + %168 = stablehlo.or %165, %167 : tensor<2xui32> loc(#loc1139) + %169 = stablehlo.xor %163, %168 : tensor<2xui32> loc(#loc1125) + %170 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %171 = stablehlo.add %163, %170 : tensor<2xui32> loc(#loc1130) + %172 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %173 = stablehlo.add %169, %172 : tensor<2xui32> loc(#loc1130) + %174 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<2xui32> loc(#loc1130) + %175 = stablehlo.add %173, %174 : tensor<2xui32> loc(#loc1130) + return %171, %175 : tensor<2xui32>, tensor<2xui32> loc(#loc786) + } loc(#loc786) + func.func private @threefry2x32_27(%arg0: tensor loc(callsite(#loc199 at #loc592)), %arg1: tensor loc(callsite(#loc199 at #loc592)), %arg2: tensor<1x4xui32> loc(callsite(#loc199 at #loc592)), %arg3: tensor<1x4xui32> loc(callsite(#loc199 at #loc592))) -> (tensor<1x4xui32>, tensor<1x4xui32>) { + %c = stablehlo.constant dense<5> : tensor loc(#loc172) + %c_0 = stablehlo.constant dense<4> : tensor loc(#loc172) + %c_1 = stablehlo.constant dense<2> : tensor loc(#loc172) + %c_2 = stablehlo.constant dense<8> : tensor loc(#loc) + %c_3 = stablehlo.constant dense<24> : tensor loc(#loc) + %c_4 = stablehlo.constant dense<16> : tensor loc(#loc) + %c_5 = stablehlo.constant dense<3> : tensor loc(#loc) + %c_6 = stablehlo.constant dense<29> : tensor loc(#loc) + %c_7 = stablehlo.constant dense<1> : tensor loc(#loc172) + %c_8 = stablehlo.constant dense<6> : tensor loc(#loc) + %c_9 = stablehlo.constant dense<26> : tensor loc(#loc) + %c_10 = stablehlo.constant dense<17> : tensor loc(#loc) + %c_11 = stablehlo.constant dense<15> : tensor loc(#loc) + %c_12 = stablehlo.constant dense<19> : tensor loc(#loc) + %c_13 = stablehlo.constant dense<13> : tensor loc(#loc) + %c_14 = stablehlo.constant dense<466688986> : tensor loc(#loc172) + %0 = stablehlo.xor %arg0, %arg1 : tensor loc(#loc1125) + %1 = stablehlo.xor %0, %c_14 : tensor loc(#loc1125) + %2 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %3 = stablehlo.add %arg2, %2 : tensor<1x4xui32> loc(#loc1130) + %4 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %5 = stablehlo.add %arg3, %4 : tensor<1x4xui32> loc(#loc1130) + %6 = stablehlo.add %3, %5 : tensor<1x4xui32> loc(#loc1130) + %7 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %8 = stablehlo.shift_left %5, %7 : tensor<1x4xui32> loc(#loc1137) + %9 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %10 = stablehlo.shift_right_logical %5, %9 : tensor<1x4xui32> loc(#loc1138) + %11 = stablehlo.or %8, %10 : tensor<1x4xui32> loc(#loc1139) + %12 = stablehlo.xor %6, %11 : tensor<1x4xui32> loc(#loc1125) + %13 = stablehlo.add %6, %12 : tensor<1x4xui32> loc(#loc1130) + %14 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %15 = stablehlo.shift_left %12, %14 : tensor<1x4xui32> loc(#loc1137) + %16 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %17 = stablehlo.shift_right_logical %12, %16 : tensor<1x4xui32> loc(#loc1138) + %18 = stablehlo.or %15, %17 : tensor<1x4xui32> loc(#loc1139) + %19 = stablehlo.xor %13, %18 : tensor<1x4xui32> loc(#loc1125) + %20 = stablehlo.add %13, %19 : tensor<1x4xui32> loc(#loc1130) + %21 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %22 = stablehlo.shift_left %19, %21 : tensor<1x4xui32> loc(#loc1137) + %23 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %24 = stablehlo.shift_right_logical %19, %23 : tensor<1x4xui32> loc(#loc1138) + %25 = stablehlo.or %22, %24 : tensor<1x4xui32> loc(#loc1139) + %26 = stablehlo.xor %20, %25 : tensor<1x4xui32> loc(#loc1125) + %27 = stablehlo.add %20, %26 : tensor<1x4xui32> loc(#loc1130) + %28 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %29 = stablehlo.shift_left %26, %28 : tensor<1x4xui32> loc(#loc1137) + %30 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %31 = stablehlo.shift_right_logical %26, %30 : tensor<1x4xui32> loc(#loc1138) + %32 = stablehlo.or %29, %31 : tensor<1x4xui32> loc(#loc1139) + %33 = stablehlo.xor %27, %32 : tensor<1x4xui32> loc(#loc1125) + %34 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %35 = stablehlo.add %27, %34 : tensor<1x4xui32> loc(#loc1130) + %36 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %37 = stablehlo.add %33, %36 : tensor<1x4xui32> loc(#loc1130) + %38 = stablehlo.broadcast_in_dim %c_7, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %39 = stablehlo.add %37, %38 : tensor<1x4xui32> loc(#loc1130) + %40 = stablehlo.add %35, %39 : tensor<1x4xui32> loc(#loc1130) + %41 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %42 = stablehlo.shift_left %39, %41 : tensor<1x4xui32> loc(#loc1137) + %43 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %44 = stablehlo.shift_right_logical %39, %43 : tensor<1x4xui32> loc(#loc1138) + %45 = stablehlo.or %42, %44 : tensor<1x4xui32> loc(#loc1139) + %46 = stablehlo.xor %40, %45 : tensor<1x4xui32> loc(#loc1125) + %47 = stablehlo.add %40, %46 : tensor<1x4xui32> loc(#loc1130) + %48 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %49 = stablehlo.shift_left %46, %48 : tensor<1x4xui32> loc(#loc1137) + %50 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %51 = stablehlo.shift_right_logical %46, %50 : tensor<1x4xui32> loc(#loc1138) + %52 = stablehlo.or %49, %51 : tensor<1x4xui32> loc(#loc1139) + %53 = stablehlo.xor %47, %52 : tensor<1x4xui32> loc(#loc1125) + %54 = stablehlo.add %47, %53 : tensor<1x4xui32> loc(#loc1130) + %55 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %56 = stablehlo.shift_left %53, %55 : tensor<1x4xui32> loc(#loc1137) + %57 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %58 = stablehlo.shift_right_logical %53, %57 : tensor<1x4xui32> loc(#loc1138) + %59 = stablehlo.or %56, %58 : tensor<1x4xui32> loc(#loc1139) + %60 = stablehlo.xor %54, %59 : tensor<1x4xui32> loc(#loc1125) + %61 = stablehlo.add %54, %60 : tensor<1x4xui32> loc(#loc1130) + %62 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %63 = stablehlo.shift_left %60, %62 : tensor<1x4xui32> loc(#loc1137) + %64 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %65 = stablehlo.shift_right_logical %60, %64 : tensor<1x4xui32> loc(#loc1138) + %66 = stablehlo.or %63, %65 : tensor<1x4xui32> loc(#loc1139) + %67 = stablehlo.xor %61, %66 : tensor<1x4xui32> loc(#loc1125) + %68 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %69 = stablehlo.add %61, %68 : tensor<1x4xui32> loc(#loc1130) + %70 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %71 = stablehlo.add %67, %70 : tensor<1x4xui32> loc(#loc1130) + %72 = stablehlo.broadcast_in_dim %c_1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %73 = stablehlo.add %71, %72 : tensor<1x4xui32> loc(#loc1130) + %74 = stablehlo.add %69, %73 : tensor<1x4xui32> loc(#loc1130) + %75 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %76 = stablehlo.shift_left %73, %75 : tensor<1x4xui32> loc(#loc1137) + %77 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %78 = stablehlo.shift_right_logical %73, %77 : tensor<1x4xui32> loc(#loc1138) + %79 = stablehlo.or %76, %78 : tensor<1x4xui32> loc(#loc1139) + %80 = stablehlo.xor %74, %79 : tensor<1x4xui32> loc(#loc1125) + %81 = stablehlo.add %74, %80 : tensor<1x4xui32> loc(#loc1130) + %82 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %83 = stablehlo.shift_left %80, %82 : tensor<1x4xui32> loc(#loc1137) + %84 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %85 = stablehlo.shift_right_logical %80, %84 : tensor<1x4xui32> loc(#loc1138) + %86 = stablehlo.or %83, %85 : tensor<1x4xui32> loc(#loc1139) + %87 = stablehlo.xor %81, %86 : tensor<1x4xui32> loc(#loc1125) + %88 = stablehlo.add %81, %87 : tensor<1x4xui32> loc(#loc1130) + %89 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %90 = stablehlo.shift_left %87, %89 : tensor<1x4xui32> loc(#loc1137) + %91 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %92 = stablehlo.shift_right_logical %87, %91 : tensor<1x4xui32> loc(#loc1138) + %93 = stablehlo.or %90, %92 : tensor<1x4xui32> loc(#loc1139) + %94 = stablehlo.xor %88, %93 : tensor<1x4xui32> loc(#loc1125) + %95 = stablehlo.add %88, %94 : tensor<1x4xui32> loc(#loc1130) + %96 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %97 = stablehlo.shift_left %94, %96 : tensor<1x4xui32> loc(#loc1137) + %98 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %99 = stablehlo.shift_right_logical %94, %98 : tensor<1x4xui32> loc(#loc1138) + %100 = stablehlo.or %97, %99 : tensor<1x4xui32> loc(#loc1139) + %101 = stablehlo.xor %95, %100 : tensor<1x4xui32> loc(#loc1125) + %102 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %103 = stablehlo.add %95, %102 : tensor<1x4xui32> loc(#loc1130) + %104 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %105 = stablehlo.add %101, %104 : tensor<1x4xui32> loc(#loc1130) + %106 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %107 = stablehlo.add %105, %106 : tensor<1x4xui32> loc(#loc1130) + %108 = stablehlo.add %103, %107 : tensor<1x4xui32> loc(#loc1130) + %109 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %110 = stablehlo.shift_left %107, %109 : tensor<1x4xui32> loc(#loc1137) + %111 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %112 = stablehlo.shift_right_logical %107, %111 : tensor<1x4xui32> loc(#loc1138) + %113 = stablehlo.or %110, %112 : tensor<1x4xui32> loc(#loc1139) + %114 = stablehlo.xor %108, %113 : tensor<1x4xui32> loc(#loc1125) + %115 = stablehlo.add %108, %114 : tensor<1x4xui32> loc(#loc1130) + %116 = stablehlo.broadcast_in_dim %c_6, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %117 = stablehlo.shift_left %114, %116 : tensor<1x4xui32> loc(#loc1137) + %118 = stablehlo.broadcast_in_dim %c_5, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %119 = stablehlo.shift_right_logical %114, %118 : tensor<1x4xui32> loc(#loc1138) + %120 = stablehlo.or %117, %119 : tensor<1x4xui32> loc(#loc1139) + %121 = stablehlo.xor %115, %120 : tensor<1x4xui32> loc(#loc1125) + %122 = stablehlo.add %115, %121 : tensor<1x4xui32> loc(#loc1130) + %123 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %124 = stablehlo.shift_left %121, %123 : tensor<1x4xui32> loc(#loc1137) + %125 = stablehlo.broadcast_in_dim %c_4, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %126 = stablehlo.shift_right_logical %121, %125 : tensor<1x4xui32> loc(#loc1138) + %127 = stablehlo.or %124, %126 : tensor<1x4xui32> loc(#loc1139) + %128 = stablehlo.xor %122, %127 : tensor<1x4xui32> loc(#loc1125) + %129 = stablehlo.add %122, %128 : tensor<1x4xui32> loc(#loc1130) + %130 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %131 = stablehlo.shift_left %128, %130 : tensor<1x4xui32> loc(#loc1137) + %132 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %133 = stablehlo.shift_right_logical %128, %132 : tensor<1x4xui32> loc(#loc1138) + %134 = stablehlo.or %131, %133 : tensor<1x4xui32> loc(#loc1139) + %135 = stablehlo.xor %129, %134 : tensor<1x4xui32> loc(#loc1125) + %136 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %137 = stablehlo.add %129, %136 : tensor<1x4xui32> loc(#loc1130) + %138 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %139 = stablehlo.add %135, %138 : tensor<1x4xui32> loc(#loc1130) + %140 = stablehlo.broadcast_in_dim %c_0, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %141 = stablehlo.add %139, %140 : tensor<1x4xui32> loc(#loc1130) + %142 = stablehlo.add %137, %141 : tensor<1x4xui32> loc(#loc1130) + %143 = stablehlo.broadcast_in_dim %c_13, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %144 = stablehlo.shift_left %141, %143 : tensor<1x4xui32> loc(#loc1137) + %145 = stablehlo.broadcast_in_dim %c_12, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %146 = stablehlo.shift_right_logical %141, %145 : tensor<1x4xui32> loc(#loc1138) + %147 = stablehlo.or %144, %146 : tensor<1x4xui32> loc(#loc1139) + %148 = stablehlo.xor %142, %147 : tensor<1x4xui32> loc(#loc1125) + %149 = stablehlo.add %142, %148 : tensor<1x4xui32> loc(#loc1130) + %150 = stablehlo.broadcast_in_dim %c_11, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %151 = stablehlo.shift_left %148, %150 : tensor<1x4xui32> loc(#loc1137) + %152 = stablehlo.broadcast_in_dim %c_10, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %153 = stablehlo.shift_right_logical %148, %152 : tensor<1x4xui32> loc(#loc1138) + %154 = stablehlo.or %151, %153 : tensor<1x4xui32> loc(#loc1139) + %155 = stablehlo.xor %149, %154 : tensor<1x4xui32> loc(#loc1125) + %156 = stablehlo.add %149, %155 : tensor<1x4xui32> loc(#loc1130) + %157 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %158 = stablehlo.shift_left %155, %157 : tensor<1x4xui32> loc(#loc1137) + %159 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %160 = stablehlo.shift_right_logical %155, %159 : tensor<1x4xui32> loc(#loc1138) + %161 = stablehlo.or %158, %160 : tensor<1x4xui32> loc(#loc1139) + %162 = stablehlo.xor %156, %161 : tensor<1x4xui32> loc(#loc1125) + %163 = stablehlo.add %156, %162 : tensor<1x4xui32> loc(#loc1130) + %164 = stablehlo.broadcast_in_dim %c_8, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1137) + %165 = stablehlo.shift_left %162, %164 : tensor<1x4xui32> loc(#loc1137) + %166 = stablehlo.broadcast_in_dim %c_9, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1138) + %167 = stablehlo.shift_right_logical %162, %166 : tensor<1x4xui32> loc(#loc1138) + %168 = stablehlo.or %165, %167 : tensor<1x4xui32> loc(#loc1139) + %169 = stablehlo.xor %163, %168 : tensor<1x4xui32> loc(#loc1125) + %170 = stablehlo.broadcast_in_dim %1, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %171 = stablehlo.add %163, %170 : tensor<1x4xui32> loc(#loc1130) + %172 = stablehlo.broadcast_in_dim %arg0, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %173 = stablehlo.add %169, %172 : tensor<1x4xui32> loc(#loc1130) + %174 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1x4xui32> loc(#loc1130) + %175 = stablehlo.add %173, %174 : tensor<1x4xui32> loc(#loc1130) + return %171, %175 : tensor<1x4xui32>, tensor<1x4xui32> loc(#loc786) + } loc(#loc786) + func.func private @_diag(%arg0: tensor<16xi32> loc(unknown)) -> tensor<16x16xi32> { + %c = stablehlo.constant dense<0> : tensor loc(#loc) + %0 = stablehlo.pad %arg0, %c, low = [0], high = [0], interior = [0] : (tensor<16xi32>, tensor) -> tensor<16xi32> loc(#loc1141) + %1 = stablehlo.iota dim = 0 : tensor<16x16xi32> loc(#loc1142) + %2 = stablehlo.iota dim = 1 : tensor<16x16xi32> loc(#loc1142) + %3 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<16x16xi32> loc(#loc1143) + %4 = stablehlo.add %1, %3 : tensor<16x16xi32> loc(#loc1143) + %5 = stablehlo.compare EQ, %4, %2, SIGNED : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi1> loc(#loc1144) + %6 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<16xi32> loc(#loc1145) + %7 = call @_where(%5, %0, %6) : (tensor<16x16xi1>, tensor<16xi32>, tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1146) + return %7 : tensor<16x16xi32> loc(#loc1199) + } loc(#loc1199) + func.func private @_where(%arg0: tensor<16x16xi1> loc(unknown), %arg1: tensor<16xi32> loc(unknown), %arg2: tensor<16xi32> loc(unknown)) -> tensor<16x16xi32> { + %0 = stablehlo.broadcast_in_dim %arg1, dims = [1] : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1145) + %1 = stablehlo.broadcast_in_dim %arg2, dims = [1] : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1145) + %2 = stablehlo.select %arg0, %0, %1 : tensor<16x16xi1>, tensor<16x16xi32> loc(#loc1148) + return %2 : tensor<16x16xi32> loc(#loc1200) + } loc(#loc1200) + func.func private @_where_79(%arg0: tensor<1xi1> loc(unknown), %arg1: tensor loc(unknown), %arg2: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %0 = stablehlo.convert %arg1 : tensor loc(#loc1150) + %1 = stablehlo.broadcast_in_dim %0, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1151) + %2 = stablehlo.select %arg0, %1, %arg2 : tensor<1xi1>, tensor<1xi32> loc(#loc1152) + return %2 : tensor<1xi32> loc(#loc1201) + } loc(#loc1201) + func.func private @_cumsum_with_promotion(%arg0: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %0 = call @cumsum(%arg0) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc642) + return %0 : tensor<1xi32> loc(#loc1201) + } loc(#loc1201) + func.func private @cumsum(%arg0: tensor<1xi32> loc(callsite(#loc259 at #loc545))) -> tensor<1xi32> { + %c = stablehlo.constant dense<0> : tensor loc(#loc1153) + %0 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor loc(#loc1153) + %1 = "stablehlo.reduce_window"(%arg0, %0) <{window_dimensions = array}> ({ + ^bb0(%arg1: tensor loc("reduce_window_sum"), %arg2: tensor loc("reduce_window_sum")): + %2 = stablehlo.add %arg1, %arg2 : tensor loc(#loc1153) + stablehlo.return %2 : tensor loc(#loc1153) + }) : (tensor<1xi32>, tensor) -> tensor<1xi32> loc(#loc1153) + return %1 : tensor<1xi32> loc(#loc642) + } loc(#loc642) + func.func private @_diag_156(%arg0: tensor<16xi32> loc(unknown)) -> tensor<16x16xi32> { + %c = stablehlo.constant dense<0> : tensor loc(#loc) + %0 = stablehlo.pad %arg0, %c, low = [0], high = [0], interior = [0] : (tensor<16xi32>, tensor) -> tensor<16xi32> loc(#loc1155) + %1 = stablehlo.iota dim = 0 : tensor<16x16xi32> loc(#loc1156) + %2 = stablehlo.iota dim = 1 : tensor<16x16xi32> loc(#loc1156) + %3 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<16x16xi32> loc(#loc1157) + %4 = stablehlo.add %1, %3 : tensor<16x16xi32> loc(#loc1157) + %5 = stablehlo.compare EQ, %4, %2, SIGNED : (tensor<16x16xi32>, tensor<16x16xi32>) -> tensor<16x16xi1> loc(#loc1158) + %6 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<16xi32> loc(#loc1159) + %7 = call @_where_157(%5, %0, %6) : (tensor<16x16xi1>, tensor<16xi32>, tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1160) + return %7 : tensor<16x16xi32> loc(#loc1202) + } loc(#loc1202) + func.func private @_where_157(%arg0: tensor<16x16xi1> loc(unknown), %arg1: tensor<16xi32> loc(unknown), %arg2: tensor<16xi32> loc(unknown)) -> tensor<16x16xi32> { + %0 = stablehlo.broadcast_in_dim %arg1, dims = [1] : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1159) + %1 = stablehlo.broadcast_in_dim %arg2, dims = [1] : (tensor<16xi32>) -> tensor<16x16xi32> loc(#loc1159) + %2 = stablehlo.select %arg0, %0, %1 : tensor<16x16xi1>, tensor<16x16xi32> loc(#loc1162) + return %2 : tensor<16x16xi32> loc(#loc1203) + } loc(#loc1203) + func.func private @_roll_static(%arg0: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %0 = stablehlo.slice %arg0 [0:1] : (tensor<1xi32>) -> tensor<1xi32> loc(#loc1164) + %1 = stablehlo.slice %arg0 [0:0] : (tensor<1xi32>) -> tensor<0xi32> loc(#loc1164) + %2 = stablehlo.concatenate %0, %1, dim = 0 : (tensor<1xi32>, tensor<0xi32>) -> tensor<1xi32> loc(#loc1165) + return %2 : tensor<1xi32> loc(#loc1204) + } loc(#loc1204) + func.func private @cumsum_213(%arg0: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %0 = call @cumsum(%arg0) : (tensor<1xi32>) -> tensor<1xi32> loc(#loc750) + return %0 : tensor<1xi32> loc(#loc1204) + } loc(#loc1204) + func.func private @cumsum_218(%arg0: tensor<128xi32> loc(unknown)) -> tensor<128xi32> { + %0 = call @cumsum_219(%arg0) : (tensor<128xi32>) -> tensor<128xi32> loc(#loc750) + return %0 : tensor<128xi32> loc(#loc1204) + } loc(#loc1204) + func.func private @cumsum_219(%arg0: tensor<128xi32> loc(callsite(#loc309 at #loc587))) -> tensor<128xi32> { + %c = stablehlo.constant dense<0> : tensor loc(#loc1166) + %0 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor loc(#loc1166) + %1 = "stablehlo.reduce_window"(%arg0, %0) <{padding = dense<[[127, 0]]> : tensor<1x2xi64>, window_dimensions = array}> ({ + ^bb0(%arg1: tensor loc("reduce_window_sum"), %arg2: tensor loc("reduce_window_sum")): + %2 = stablehlo.add %arg1, %arg2 : tensor loc(#loc1166) + stablehlo.return %2 : tensor loc(#loc1166) + }) : (tensor<128xi32>, tensor) -> tensor<128xi32> loc(#loc1166) + return %1 : tensor<128xi32> loc(#loc750) + } loc(#loc750) + func.func private @_take(%arg0: tensor<1xi32> loc(unknown), %arg1: tensor<128xi32> loc(unknown)) -> tensor<128xi32> { + %c = stablehlo.constant dense<-2147483648> : tensor loc(#loc1167) + %c_0 = stablehlo.constant dense : tensor loc(#loc1168) + %c_1 = stablehlo.constant dense<0> : tensor<1xi32> loc(#loc1167) + %c_2 = stablehlo.constant dense<1> : tensor loc(#loc1204) + %c_3 = stablehlo.constant dense<0> : tensor loc(#loc) + %0 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<128xi32> loc(#loc1169) + %1 = stablehlo.compare LT, %arg1, %0, SIGNED : (tensor<128xi32>, tensor<128xi32>) -> tensor<128xi1> loc(#loc1169) + %2 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<128xi32> loc(#loc1170) + %3 = stablehlo.add %arg1, %2 : tensor<128xi32> loc(#loc1170) + %4 = call @_where_224(%1, %3, %arg1) : (tensor<128xi1>, tensor<128xi32>, tensor<128xi32>) -> tensor<128xi32> loc(#loc1171) + %5 = stablehlo.broadcast_in_dim %4, dims = [0] : (tensor<128xi32>) -> tensor<128x1xi32> loc(#loc1172) + %6 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<128x1xi32> loc(#loc1173) + %7 = stablehlo.compare GE, %5, %6, SIGNED : (tensor<128x1xi32>, tensor<128x1xi32>) -> tensor<128x1xi1> loc(#loc1173) + %8 = stablehlo.broadcast_in_dim %c_1, dims = [1] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1172) + %9 = stablehlo.broadcast_in_dim %8, dims = [0, 1] : (tensor<1x1xi32>) -> tensor<128x1xi32> loc(#loc1174) + %10 = stablehlo.compare LE, %5, %9, SIGNED : (tensor<128x1xi32>, tensor<128x1xi32>) -> tensor<128x1xi1> loc(#loc1174) + %11 = stablehlo.and %7, %10 : tensor<128x1xi1> loc(#loc1175) + %12 = stablehlo.reduce(%11 init: %c_0) applies stablehlo.and across dimensions = [1] : (tensor<128x1xi1>, tensor) -> tensor<128xi1> loc(#loc1168) + %13 = "stablehlo.gather"(%arg0, %5) <{dimension_numbers = #stablehlo.gather, slice_sizes = array}> : (tensor<1xi32>, tensor<128x1xi32>) -> tensor<128xi32> loc(#loc1167) + %14 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<128xi32> loc(#loc1172) + %15 = stablehlo.select %12, %13, %14 : tensor<128xi1>, tensor<128xi32> loc(#loc1176) + return %15 : tensor<128xi32> loc(#loc1204) + } loc(#loc1204) + func.func private @_where_224(%arg0: tensor<128xi1> loc(unknown), %arg1: tensor<128xi32> loc(unknown), %arg2: tensor<128xi32> loc(unknown)) -> tensor<128xi32> { + %0 = stablehlo.select %arg0, %arg1, %arg2 : tensor<128xi1>, tensor<128xi32> loc(#loc1176) + return %0 : tensor<128xi32> loc(#loc1205) + } loc(#loc1205) + func.func private @_where_237(%arg0: tensor loc(unknown), %arg1: tensor loc(unknown), %arg2: tensor<1xf32> loc(unknown)) -> tensor<1xf32> { + %0 = stablehlo.broadcast_in_dim %arg1, dims = [] : (tensor) -> tensor<1xf32> loc(#loc1179) + %1 = stablehlo.select %arg0, %0, %arg2 : tensor, tensor<1xf32> loc(#loc1180) + return %1 : tensor<1xf32> loc(#loc1206) + } loc(#loc1206) + func.func private @_where_239(%arg0: tensor loc(unknown), %arg1: tensor<1xf32> loc(unknown), %arg2: tensor<1xf32> loc(unknown)) -> tensor<1xf32> { + %0 = stablehlo.select %arg0, %arg1, %arg2 : tensor, tensor<1xf32> loc(#loc1182) + return %0 : tensor<1xf32> loc(#loc1207) + } loc(#loc1207) + func.func private @_take_248(%arg0: tensor<1xi32> loc(unknown), %arg1: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %c = stablehlo.constant dense<-2147483648> : tensor loc(#loc1184) + %c_0 = stablehlo.constant dense : tensor loc(#loc1185) + %c_1 = stablehlo.constant dense<0> : tensor<1xi32> loc(#loc1184) + %c_2 = stablehlo.constant dense<1> : tensor loc(#loc1208) + %c_3 = stablehlo.constant dense<0> : tensor loc(#loc) + %0 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1186) + %1 = stablehlo.compare LT, %arg1, %0, SIGNED : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1> loc(#loc1186) + %2 = stablehlo.broadcast_in_dim %c_2, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1187) + %3 = stablehlo.add %arg1, %2 : tensor<1xi32> loc(#loc1187) + %4 = call @_where_249(%1, %3, %arg1) : (tensor<1xi1>, tensor<1xi32>, tensor<1xi32>) -> tensor<1xi32> loc(#loc1188) + %5 = stablehlo.broadcast_in_dim %4, dims = [0] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1189) + %6 = stablehlo.broadcast_in_dim %c_3, dims = [] : (tensor) -> tensor<1x1xi32> loc(#loc1190) + %7 = stablehlo.compare GE, %5, %6, SIGNED : (tensor<1x1xi32>, tensor<1x1xi32>) -> tensor<1x1xi1> loc(#loc1190) + %8 = stablehlo.broadcast_in_dim %c_1, dims = [1] : (tensor<1xi32>) -> tensor<1x1xi32> loc(#loc1189) + %9 = stablehlo.compare LE, %5, %8, SIGNED : (tensor<1x1xi32>, tensor<1x1xi32>) -> tensor<1x1xi1> loc(#loc1191) + %10 = stablehlo.and %7, %9 : tensor<1x1xi1> loc(#loc1192) + %11 = stablehlo.reduce(%10 init: %c_0) applies stablehlo.and across dimensions = [1] : (tensor<1x1xi1>, tensor) -> tensor<1xi1> loc(#loc1185) + %12 = "stablehlo.gather"(%arg0, %5) <{dimension_numbers = #stablehlo.gather, slice_sizes = array}> : (tensor<1xi32>, tensor<1x1xi32>) -> tensor<1xi32> loc(#loc1184) + %13 = stablehlo.broadcast_in_dim %c, dims = [] : (tensor) -> tensor<1xi32> loc(#loc1189) + %14 = stablehlo.select %11, %12, %13 : tensor<1xi1>, tensor<1xi32> loc(#loc1193) + return %14 : tensor<1xi32> loc(#loc1208) + } loc(#loc1208) + func.func private @_where_249(%arg0: tensor<1xi1> loc(unknown), %arg1: tensor<1xi32> loc(unknown), %arg2: tensor<1xi32> loc(unknown)) -> tensor<1xi32> { + %0 = stablehlo.select %arg0, %arg1, %arg2 : tensor<1xi1>, tensor<1xi32> loc(#loc1193) + return %0 : tensor<1xi32> loc(#loc1209) + } loc(#loc1209) +} loc(#loc) +#loc13 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1985:64 to :85) +#loc14 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1991:30 to :67) +#loc20 = loc("/usr/local/lib/python3.12/dist-packages/_pytest/python.py":1720:8 to :54) +#loc21 = loc("/usr/local/lib/python3.12/dist-packages/_pytest/runner.py":179:8 to :22) +#loc22 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1700:21 to :81) +#loc23 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":2065:28 to :97) +#loc26 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1701:38 to :79) +#loc27 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1701:28 to :98) +#loc28 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1703:17 to 1705:9) +#loc29 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1706:15 to :39) +#loc30 = loc("/opt/transformerengine/transformer_engine/jax/quantize/helper.py":630:21 to :46) +#loc38 = loc("/opt/transformerengine/transformer_engine/jax/quantize/helper.py":643:10 to :26) +#loc39 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1007:13 to :46) +#loc40 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1062:26 to 1072:9) +#loc41 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1497:29 to 1506:9) +#loc42 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_dense.py":197:32 to 207:5) +#loc43 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_dense.py":79:13 to 93:5) +#loc44 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":800:16 to 813:13) +#loc45 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1009:11 to :45) +#loc46 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1011:52 to 1028:9) +#loc47 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":801:12 to :40) +#loc48 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":942:13 to 948:5) +#loc49 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/normalization.py":1073:14 to 1078:9) +#loc50 = loc("/opt/transformerengine/transformer_engine/jax/quantize/hadamard.py":44:8 to :19) +#loc51 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":808:21 to :37) +#loc52 = loc("/opt/transformerengine/transformer_engine/jax/quantize/hadamard.py":44:8 to :23) +#loc53 = loc("/opt/transformerengine/transformer_engine/jax/quantize/hadamard.py":46:16 to :36) +#loc54 = loc("/opt/transformerengine/transformer_engine/jax/quantize/hadamard.py":46:12 to :36) +#loc55 = loc("/opt/transformerengine/transformer_engine/jax/quantize/hadamard.py":46:11 to :58) +#loc56 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/amax.py":409:26 to 416:5) +#loc57 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":810:34 to 816:9) +#loc58 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":858:19 to :46) +#loc59 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":871:8 to 887:5) +#loc60 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_dense.py":212:20 to 218:5) +#loc61 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":803:17 to :48) +#loc62 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/amax.py":382:11 to 386:5) +#loc63 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":836:19 to 840:13) +#loc64 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":876:56 to :84) +#loc65 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1322:11 to :39) +#loc66 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1857:14 to 1866:5) +#loc67 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_dense.py":224:13 to 231:5) +#loc68 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":187:11 to :52) +#loc69 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1325:31 to :68) +#loc70 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1326:31 to :68) +#loc71 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1327:16 to :59) +#loc72 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1332:15 to :44) +#loc73 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1334:21 to :50) +#loc74 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1336:11 to 1356:5) +#loc75 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1544:23 to 1546:13) +#loc76 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":911:21 to :70) +#loc81 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":912:22 to :72) +#loc82 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":674:28 to :40) +#loc83 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":762:15 to :51) +#loc84 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":927:8 to :66) +#loc85 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":674:42 to :54) +#loc86 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":675:28 to :40) +#loc87 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":675:42 to :54) +#loc88 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":676:28 to :40) +#loc89 = loc("/opt/transformerengine/transformer_engine/jax/attention.py":676:42 to :54) +#loc90 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":233:19 to :59) +#loc91 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":3361:11 to :89) +#loc94 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":234:19 to :53) +#loc95 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":3363:16 to :48) +#loc96 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/attention.py":3401:29 to :60) +#loc98 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1586:12 to :72) +#loc99 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":522:24 to 524:9) +#loc100 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":1591:14 to 1604:12) +#loc101 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":195:15 to 201:5) +#loc102 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":101:13 to 112:5) +#loc103 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":540:12 to 548:9) +#loc104 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":204:20 to 209:5) +#loc105 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":214:13 to 222:5) +#loc106 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":2154:12 to :24) +#loc107 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":1060:29 to 1062:9) +#loc108 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":2224:20 to 2253:49) +#loc109 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":1063:29 to 1065:9) +#loc110 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":295:32 to 305:5) +#loc111 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":121:13 to 144:5) +#loc112 = loc("/opt/transformerengine/transformer_engine/jax/flax/module.py":1197:18 to 1217:13) +#loc113 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":308:22 to 314:5) +#loc114 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":318:19 to 326:5) +#loc115 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/activation.py":1302:12 to :40) +#loc116 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/activation.py":1336:14 to 1344:9) +#loc117 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":343:21 to 354:5) +#loc118 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/activation.py":1304:11 to :39) +#loc119 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/activation.py":1307:37 to 1322:9) +#loc120 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/activation.py":1348:14 to 1353:9) +#loc121 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":358:22 to 363:5) +#loc122 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_mlp.py":367:19 to 375:5) +#loc123 = loc("/opt/transformerengine/transformer_engine/jax/flax/transformer.py":2272:12 to :24) +#loc124 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1968:24 to :48) +#loc125 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/softmax.py":880:11 to :92) +#loc126 = loc("/opt/transformerengine/transformer_engine/jax/softmax.py":50:17 to :61) +#loc127 = loc("/opt/transformerengine/transformer_engine/jax/softmax.py":32:13 to :70) +#loc128 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1969:24 to :96) +#loc129 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1970:28 to :177) +#loc130 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/softmax.py":892:11 to 894:5) +#loc131 = loc("/opt/transformerengine/transformer_engine/jax/softmax.py":46:17 to :74) +#loc132 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1971:24 to :115) +#loc133 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1973:24 to :50) +#loc134 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/softmax.py":925:11 to 927:5) +#loc135 = loc("/opt/transformerengine/transformer_engine/jax/softmax.py":48:17 to :81) +#loc136 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1974:24 to :128) +#loc137 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1986:20 to :31) +#loc138 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":1994:20 to :48) +#loc139 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1216:12 to :47) +#loc144 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1226:31 to :41) +#loc145 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1226:23 to :78) +#loc146 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1228:12 to :32) +#loc150 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1230:23 to :88) +#loc152 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1232:48 to :63) +#loc153 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":66:16 to :30) +#loc154 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1232:24 to :95) +#loc155 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":67:10 to :24) +#loc156 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":67:9 to :39) +#loc157 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":68:19 to :29) +#loc158 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":68:9 to :41) +#loc159 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":69:19 to :37) +#loc160 = loc("/opt/transformerengine/transformer_engine/jax/quantize/quantizer.py":69:9 to :49) +#loc161 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1233:36 to :48) +#loc162 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1233:20 to :49) +#loc163 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1249:8 to 1259:5) +#loc164 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/quantization.py":1205:22 to :68) +#loc165 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":464:24 to 466:9) +#loc166 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":1933:35 to :61) +#loc167 = loc("/opt/transformerengine/transformer_engine/jax/dense.py":519:13 to 528:5) +#loc168 = loc("/opt/transformerengine/transformer_engine/jax/cpp_extensions/gemm.py":2077:13 to 2095:5) +#loc169 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":2002:24 to :35) +#loc170 = loc("/opt/transformerengine/tests/jax/test_custom_call_compute.py":2002:16 to :35) +#loc172 = loc("threefry2x32") +#loc173 = loc("/opt/transformerengine/transformer_engine/jax/layernorm_dense.py":147:16 to 161:5) +#loc175 = loc("TestFFICompatibility.test_generate_hlo..f"(#loc13)) +#loc176 = loc("TestFFICompatibility.test_generate_hlo..train_step"(#loc14)) +#loc182 = loc("Function.runtest"(#loc20)) +#loc183 = loc("pytest_runtest_call"(#loc21)) +#loc184 = loc("RelativePositionBiases.__call__"(#loc22)) +#loc185 = loc("TransformerLayer.__call__"(#loc23)) +#loc188 = loc("RelativePositionBiases.__call__"(#loc26)) +#loc189 = loc("RelativePositionBiases.__call__"(#loc27)) +#loc190 = loc("RelativePositionBiases.__call__"(#loc28)) +#loc191 = loc("RelativePositionBiases.__call__"(#loc29)) +#loc192 = loc("NVFP4ScalingQuantizeConfig._make_stochastic_rounding_rng_state"(#loc30)) +#loc200 = loc("NVFP4ScalingQuantizeConfig._make_stochastic_rounding_rng_state"(#loc38)) +#loc201 = loc("layernorm_fwd"(#loc39)) +#loc202 = loc("layernorm_fwd"(#loc40)) +#loc203 = loc("normalization_fwd"(#loc41)) +#loc204 = loc("_layernorm_dense_fwd_rule"(#loc42)) +#loc205 = loc("layernorm_dense"(#loc43)) +#loc206 = loc("LayerNormDenseGeneral.__call__"(#loc44)) +#loc207 = loc("layernorm_fwd"(#loc45)) +#loc208 = loc("layernorm_fwd"(#loc46)) +#loc209 = loc("_quantize_dbias_impl"(#loc47)) +#loc210 = loc("quantize"(#loc48)) +#loc211 = loc("layernorm_fwd"(#loc49)) +#loc212 = loc("get_rht_matrix"(#loc50)) +#loc213 = loc("_quantize_dbias_impl"(#loc51)) +#loc214 = loc("get_rht_matrix"(#loc52)) +#loc215 = loc("get_rht_matrix"(#loc53)) +#loc216 = loc("get_rht_matrix"(#loc54)) +#loc217 = loc("get_rht_matrix"(#loc55)) +#loc218 = loc("calculate_post_rht_amax"(#loc56)) +#loc219 = loc("_quantize_dbias_impl"(#loc57)) +#loc220 = loc("_quantize_dbias_impl"(#loc58)) +#loc221 = loc("_quantize_dbias_impl"(#loc59)) +#loc222 = loc("_layernorm_dense_fwd_rule"(#loc60)) +#loc223 = loc("_quantize_dbias_impl"(#loc61)) +#loc224 = loc("calculate_amax"(#loc62)) +#loc225 = loc("_quantize_dbias_impl"(#loc63)) +#loc226 = loc("_quantize_dbias_impl"(#loc64)) +#loc227 = loc("_te_gemm"(#loc65)) +#loc228 = loc("gemm"(#loc66)) +#loc229 = loc("_layernorm_dense_fwd_rule"(#loc67)) +#loc230 = loc("_get_nvfp4_tensor_scale_inv"(#loc68)) +#loc231 = loc("_te_gemm"(#loc69)) +#loc232 = loc("_te_gemm"(#loc70)) +#loc233 = loc("_te_gemm"(#loc71)) +#loc234 = loc("_te_gemm"(#loc72)) +#loc235 = loc("_te_gemm"(#loc73)) +#loc236 = loc("_te_gemm"(#loc74)) +#loc237 = loc("MultiHeadAttention.__call__"(#loc75)) +#loc238 = loc("_legacy_fused_attn"(#loc76)) +#loc243 = loc("_legacy_fused_attn"(#loc81)) +#loc244 = loc("SequenceDescriptor.__init__"(#loc82)) +#loc245 = loc("SequenceDescriptor.from_seqlens"(#loc83)) +#loc246 = loc("_legacy_fused_attn"(#loc84)) +#loc247 = loc("SequenceDescriptor.__init__"(#loc85)) +#loc248 = loc("SequenceDescriptor.__init__"(#loc86)) +#loc249 = loc("SequenceDescriptor.__init__"(#loc87)) +#loc250 = loc("SequenceDescriptor.__init__"(#loc88)) +#loc251 = loc("SequenceDescriptor.__init__"(#loc89)) +#loc252 = loc("_FusedAttnRNGStateChecker.check_seed"(#loc90)) +#loc253 = loc("fused_attn_fwd"(#loc91)) +#loc256 = loc("_FusedAttnRNGStateChecker.check_seed"(#loc94)) +#loc257 = loc("fused_attn_fwd"(#loc95)) +#loc258 = loc("fused_attn_fwd"(#loc96)) +#loc260 = loc("MultiHeadAttention.__call__"(#loc98)) +#loc261 = loc("DenseGeneral.__call__"(#loc99)) +#loc262 = loc("MultiHeadAttention.__call__"(#loc100)) +#loc263 = loc("_dense_fwd_rule"(#loc101)) +#loc264 = loc("dense"(#loc102)) +#loc265 = loc("DenseGeneral.__call__"(#loc103)) +#loc266 = loc("_dense_fwd_rule"(#loc104)) +#loc267 = loc("_dense_fwd_rule"(#loc105)) +#loc268 = loc("TransformerLayer.__call__"(#loc106)) +#loc269 = loc("LayerNormMLP.__call__"(#loc107)) +#loc270 = loc("TransformerLayer.__call__"(#loc108)) +#loc271 = loc("LayerNormMLP.__call__"(#loc109)) +#loc272 = loc("_layernorm_mlp_fwd_rule"(#loc110)) +#loc273 = loc("layernorm_mlp"(#loc111)) +#loc274 = loc("LayerNormMLP.__call__"(#loc112)) +#loc275 = loc("_layernorm_mlp_fwd_rule"(#loc113)) +#loc276 = loc("_layernorm_mlp_fwd_rule"(#loc114)) +#loc277 = loc("act_lu"(#loc115)) +#loc278 = loc("act_lu"(#loc116)) +#loc279 = loc("_layernorm_mlp_fwd_rule"(#loc117)) +#loc280 = loc("act_lu"(#loc118)) +#loc281 = loc("act_lu"(#loc119)) +#loc282 = loc("act_lu"(#loc120)) +#loc283 = loc("_layernorm_mlp_fwd_rule"(#loc121)) +#loc284 = loc("_layernorm_mlp_fwd_rule"(#loc122)) +#loc285 = loc("TransformerLayer.__call__"(#loc123)) +#loc286 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc124)) +#loc287 = loc("scaled_softmax_fwd"(#loc125)) +#loc288 = loc("_softmax_fwd_rule"(#loc126)) +#loc289 = loc("softmax"(#loc127)) +#loc290 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc128)) +#loc291 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc129)) +#loc292 = loc("scaled_masked_softmax_fwd"(#loc130)) +#loc293 = loc("_softmax_fwd_rule"(#loc131)) +#loc294 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc132)) +#loc295 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc133)) +#loc296 = loc("scaled_upper_triang_masked_softmax_fwd"(#loc134)) +#loc297 = loc("_softmax_fwd_rule"(#loc135)) +#loc298 = loc("TestFFICompatibility.test_generate_hlo..Model.__call__"(#loc136)) +#loc299 = loc("TestFFICompatibility.test_generate_hlo..f"(#loc137)) +#loc300 = loc("TestFFICompatibility.test_generate_hlo..train_step"(#loc138)) +#loc301 = loc("grouped_quantize"(#loc139)) +#loc306 = loc("grouped_quantize"(#loc144)) +#loc307 = loc("grouped_quantize"(#loc145)) +#loc308 = loc("grouped_quantize"(#loc146)) +#loc310 = loc("grouped_quantize"(#loc150)) +#loc311 = loc("grouped_quantize"(#loc152)) +#loc312 = loc("compute_scale_from_amax"(#loc153)) +#loc313 = loc("grouped_quantize"(#loc154)) +#loc314 = loc("compute_scale_from_amax"(#loc155)) +#loc315 = loc("compute_scale_from_amax"(#loc156)) +#loc316 = loc("compute_scale_from_amax"(#loc157)) +#loc317 = loc("compute_scale_from_amax"(#loc158)) +#loc318 = loc("compute_scale_from_amax"(#loc159)) +#loc319 = loc("compute_scale_from_amax"(#loc160)) +#loc320 = loc("grouped_quantize"(#loc161)) +#loc321 = loc("grouped_quantize"(#loc162)) +#loc322 = loc("grouped_quantize"(#loc163)) +#loc323 = loc("grouped_quantize"(#loc164)) +#loc324 = loc("_grouped_dense_fwd_rule"(#loc165)) +#loc325 = loc("grouped_gemm"(#loc166)) +#loc326 = loc("_grouped_dense_fwd_rule"(#loc167)) +#loc327 = loc("grouped_gemm"(#loc168)) +#loc328 = loc("TestFFICompatibility.test_generate_hlo..train_step"(#loc169)) +#loc329 = loc("TestFFICompatibility.test_generate_hlo..train_step"(#loc170)) +#loc331 = loc("_layernorm_dense"(#loc173)) +#loc332 = loc(callsite(#loc183 at #loc179)) +#loc334 = loc(callsite(#loc176 at #loc177)) +#loc336 = loc(callsite(#loc197 at #loc186)) +#loc337 = loc(callsite(#loc196 at #loc197)) +#loc338 = loc(callsite(#loc187 at #loc176)) +#loc339 = loc(callsite(#loc181 at #loc182)) +#loc340 = loc(callsite(#loc177 at #loc178)) +#loc341 = loc(callsite(#loc270 at #loc186)) +#loc342 = loc(callsite(#loc182 at #loc183)) +#loc343 = loc(callsite(#loc179 at #loc180)) +#loc345 = loc(callsite(#loc206 at #loc196)) +#loc346 = loc(callsite(#loc182 at #loc332)) +#loc348 = loc(callsite(#loc187 at #loc334)) +#loc350 = loc(callsite(#loc196 at #loc336)) +#loc351 = loc(callsite(#loc206 at #loc337)) +#loc352 = loc(callsite(#loc186 at #loc338)) +#loc353 = loc(callsite(#loc242 at #loc336)) +#loc354 = loc(callsite(#loc180 at #loc339)) +#loc355 = loc(callsite(#loc176 at #loc340)) +#loc356 = loc(callsite(#loc270 at #loc335)) +#loc357 = loc(callsite(#loc274 at #loc341)) +#loc358 = loc(callsite(#loc181 at #loc342)) +#loc359 = loc(callsite(#loc178 at #loc343)) +#loc360 = loc(callsite(#loc177 at #loc344)) +#loc361 = loc(callsite(#loc183 at #loc343)) +#loc363 = loc(callsite(#loc205 at #loc345)) +#loc364 = loc(callsite(#loc181 at #loc346)) +#loc366 = loc(callsite(#loc186 at #loc348)) +#loc367 = loc(callsite(#loc196 at #loc349)) +#loc368 = loc(callsite(#loc206 at #loc350)) +#loc369 = loc(callsite(#loc205 at #loc351)) +#loc370 = loc(callsite(#loc197 at #loc352)) +#loc372 = loc(callsite(#loc241 at #loc353)) +#loc373 = loc(callsite(#loc262 at #loc349)) +#loc374 = loc(callsite(#loc179 at #loc354)) +#loc375 = loc(callsite(#loc187 at #loc355)) +#loc376 = loc(callsite(#loc270 at #loc352)) +#loc377 = loc(callsite(#loc274 at #loc356)) +#loc378 = loc(callsite(#loc273 at #loc357)) +#loc379 = loc(callsite(#loc180 at #loc358)) +#loc380 = loc(callsite(#loc177 at #loc359)) +#loc381 = loc(callsite(#loc176 at #loc360)) +#loc382 = loc(callsite(#loc182 at #loc361)) +#loc384 = loc(callsite(#loc331 at #loc363)) +#loc385 = loc(callsite(#loc180 at #loc364)) +#loc387 = loc(callsite(#loc197 at #loc366)) +#loc388 = loc(callsite(#loc206 at #loc367)) +#loc389 = loc(callsite(#loc205 at #loc368)) +#loc390 = loc(callsite(#loc204 at #loc369)) +#loc391 = loc(callsite(#loc196 at #loc370)) +#loc393 = loc(callsite(#loc240 at #loc372)) +#loc394 = loc(callsite(#loc262 at #loc370)) +#loc395 = loc(callsite(#loc265 at #loc373)) +#loc396 = loc(callsite(#loc178 at #loc374)) +#loc397 = loc(callsite(#loc186 at #loc375)) +#loc398 = loc(callsite(#loc274 at #loc376)) +#loc399 = loc(callsite(#loc273 at #loc377)) +#loc400 = loc(callsite(#loc272 at #loc378)) +#loc401 = loc(callsite(#loc270 at #loc366)) +#loc402 = loc(callsite(#loc179 at #loc379)) +#loc403 = loc(callsite(#loc176 at #loc380)) +#loc404 = loc(callsite(#loc187 at #loc381)) +#loc405 = loc(callsite(#loc181 at #loc382)) +#loc406 = loc(callsite(#loc305 at #loc380)) +#loc408 = loc(callsite(#loc204 at #loc384)) +#loc409 = loc(callsite(#loc179 at #loc385)) +#loc410 = loc(callsite(#loc176 at #loc386)) +#loc411 = loc(callsite(#loc196 at #loc387)) +#loc412 = loc(callsite(#loc205 at #loc388)) +#loc413 = loc(callsite(#loc204 at #loc389)) +#loc414 = loc(callsite(#loc203 at #loc390)) +#loc415 = loc(callsite(#loc206 at #loc391)) +#loc416 = loc(callsite(#loc242 at #loc387)) +#loc418 = loc(callsite(#loc239 at #loc393)) +#loc419 = loc(callsite(#loc262 at #loc387)) +#loc420 = loc(callsite(#loc265 at #loc394)) +#loc421 = loc(callsite(#loc264 at #loc395)) +#loc422 = loc(callsite(#loc177 at #loc396)) +#loc423 = loc(callsite(#loc270 at #loc397)) +#loc424 = loc(callsite(#loc273 at #loc398)) +#loc425 = loc(callsite(#loc272 at #loc399)) +#loc426 = loc(callsite(#loc203 at #loc400)) +#loc427 = loc(callsite(#loc274 at #loc401)) +#loc428 = loc(callsite(#loc279 at #loc399)) +#loc429 = loc(callsite(#loc178 at #loc402)) +#loc430 = loc(callsite(#loc187 at #loc403)) +#loc431 = loc(callsite(#loc186 at #loc404)) +#loc432 = loc(callsite(#loc180 at #loc405)) +#loc434 = loc(callsite(#loc304 at #loc406)) +#loc436 = loc(callsite(#loc203 at #loc408)) +#loc437 = loc(callsite(#loc178 at #loc409)) +#loc438 = loc(callsite(#loc187 at #loc410)) +#loc439 = loc(callsite(#loc195 at #loc411)) +#loc440 = loc(callsite(#loc204 at #loc412)) +#loc441 = loc(callsite(#loc203 at #loc413)) +#loc442 = loc(callsite(#loc211 at #loc414)) +#loc443 = loc(callsite(#loc205 at #loc415)) +#loc444 = loc(callsite(#loc222 at #loc412)) +#loc445 = loc(callsite(#loc229 at #loc412)) +#loc446 = loc(callsite(#loc241 at #loc416)) +#loc448 = loc(callsite(#loc255 at #loc418)) +#loc449 = loc(callsite(#loc261 at #loc419)) +#loc450 = loc(callsite(#loc264 at #loc420)) +#loc451 = loc(callsite(#loc263 at #loc421)) +#loc452 = loc(callsite(#loc266 at #loc421)) +#loc453 = loc(callsite(#loc267 at #loc421)) +#loc454 = loc(callsite(#loc176 at #loc422)) +#loc455 = loc(callsite(#loc269 at #loc423)) +#loc456 = loc(callsite(#loc271 at #loc423)) +#loc457 = loc(callsite(#loc272 at #loc424)) +#loc458 = loc(callsite(#loc203 at #loc425)) +#loc459 = loc(callsite(#loc211 at #loc426)) +#loc460 = loc(callsite(#loc273 at #loc427)) +#loc461 = loc(callsite(#loc275 at #loc424)) +#loc462 = loc(callsite(#loc276 at #loc424)) +#loc463 = loc(callsite(#loc279 at #loc424)) +#loc464 = loc(callsite(#loc282 at #loc428)) +#loc465 = loc(callsite(#loc283 at #loc424)) +#loc466 = loc(callsite(#loc284 at #loc424)) +#loc467 = loc(callsite(#loc177 at #loc429)) +#loc468 = loc(callsite(#loc290 at #loc430)) +#loc469 = loc(callsite(#loc294 at #loc430)) +#loc470 = loc(callsite(#loc298 at #loc430)) +#loc471 = loc(callsite(#loc186 at #loc430)) +#loc472 = loc(callsite(#loc197 at #loc431)) +#loc473 = loc(callsite(#loc179 at #loc432)) +#loc475 = loc(callsite(#loc303 at #loc434)) +#loc477 = loc(callsite(#loc211 at #loc436)) +#loc478 = loc(callsite(#loc177 at #loc437)) +#loc479 = loc(callsite(#loc186 at #loc438)) +#loc480 = loc(callsite(#loc194 at #loc439)) +#loc481 = loc(callsite(#loc203 at #loc440)) +#loc482 = loc(callsite(#loc211 at #loc441)) +#loc483 = loc(callsite(#loc210 at #loc442)) +#loc484 = loc(callsite(#loc222 at #loc443)) +#loc485 = loc(callsite(#loc210 at #loc444)) +#loc486 = loc(callsite(#loc229 at #loc443)) +#loc487 = loc(callsite(#loc228 at #loc445)) +#loc488 = loc(callsite(#loc240 at #loc446)) +#loc489 = loc(callsite(#loc246 at #loc447)) +#loc490 = loc(callsite(#loc254 at #loc448)) +#loc492 = loc(callsite(#loc194 at #loc449)) +#loc493 = loc(callsite(#loc263 at #loc450)) +#loc494 = loc(callsite(#loc210 at #loc451)) +#loc495 = loc(callsite(#loc266 at #loc450)) +#loc496 = loc(callsite(#loc210 at #loc452)) +#loc497 = loc(callsite(#loc267 at #loc450)) +#loc498 = loc(callsite(#loc228 at #loc453)) +#loc499 = loc(callsite(#loc187 at #loc454)) +#loc500 = loc(callsite(#loc194 at #loc455)) +#loc501 = loc(callsite(#loc194 at #loc456)) +#loc502 = loc(callsite(#loc203 at #loc457)) +#loc503 = loc(callsite(#loc211 at #loc458)) +#loc504 = loc(callsite(#loc210 at #loc459)) +#loc505 = loc(callsite(#loc275 at #loc460)) +#loc506 = loc(callsite(#loc210 at #loc461)) +#loc507 = loc(callsite(#loc276 at #loc460)) +#loc508 = loc(callsite(#loc228 at #loc462)) +#loc509 = loc(callsite(#loc279 at #loc460)) +#loc510 = loc(callsite(#loc282 at #loc463)) +#loc511 = loc(callsite(#loc210 at #loc464)) +#loc512 = loc(callsite(#loc283 at #loc460)) +#loc513 = loc(callsite(#loc210 at #loc465)) +#loc514 = loc(callsite(#loc284 at #loc460)) +#loc515 = loc(callsite(#loc228 at #loc466)) +#loc516 = loc(callsite(#loc176 at #loc467)) +#loc517 = loc(callsite(#loc289 at #loc468)) +#loc518 = loc(callsite(#loc289 at #loc469)) +#loc519 = loc(callsite(#loc289 at #loc470)) +#loc520 = loc(callsite(#loc270 at #loc471)) +#loc521 = loc(callsite(#loc262 at #loc472)) +#loc522 = loc(callsite(#loc196 at #loc472)) +#loc523 = loc(callsite(#loc178 at #loc473)) +#loc525 = loc(callsite(#loc302 at #loc475)) +#loc526 = loc(callsite(#loc324 at #loc475)) +#loc528 = loc(callsite(#loc210 at #loc477)) +#loc529 = loc(callsite(#loc176 at #loc478)) +#loc530 = loc(callsite(#loc185 at #loc479)) +#loc531 = loc(callsite(#loc193 at #loc480)) +#loc532 = loc(callsite(#loc202 at #loc481)) +#loc533 = loc(callsite(#loc210 at #loc482)) +#loc534 = loc(callsite(#loc213 at #loc483)) +#loc535 = loc(callsite(#loc219 at #loc483)) +#loc536 = loc(callsite(#loc210 at #loc484)) +#loc537 = loc(callsite(#loc225 at #loc485)) +#loc538 = loc(callsite(#loc228 at #loc486)) +#loc539 = loc(callsite(#loc231 at #loc487)) +#loc540 = loc(callsite(#loc232 at #loc487)) +#loc541 = loc(callsite(#loc197 at #loc479)) +#loc542 = loc(callsite(#loc239 at #loc488)) +#loc543 = loc(callsite(#loc245 at #loc489)) +#loc544 = loc(callsite(#loc253 at #loc490)) +#loc546 = loc(callsite(#loc193 at #loc492)) +#loc547 = loc(callsite(#loc210 at #loc493)) +#loc548 = loc(callsite(#loc213 at #loc494)) +#loc549 = loc(callsite(#loc219 at #loc494)) +#loc550 = loc(callsite(#loc210 at #loc495)) +#loc551 = loc(callsite(#loc225 at #loc496)) +#loc552 = loc(callsite(#loc228 at #loc497)) +#loc553 = loc(callsite(#loc231 at #loc498)) +#loc554 = loc(callsite(#loc232 at #loc498)) +#loc555 = loc(callsite(#loc186 at #loc499)) +#loc556 = loc(callsite(#loc193 at #loc500)) +#loc557 = loc(callsite(#loc193 at #loc501)) +#loc558 = loc(callsite(#loc202 at #loc502)) +#loc559 = loc(callsite(#loc210 at #loc503)) +#loc560 = loc(callsite(#loc213 at #loc504)) +#loc561 = loc(callsite(#loc219 at #loc504)) +#loc562 = loc(callsite(#loc210 at #loc505)) +#loc563 = loc(callsite(#loc225 at #loc506)) +#loc564 = loc(callsite(#loc228 at #loc507)) +#loc565 = loc(callsite(#loc231 at #loc508)) +#loc566 = loc(callsite(#loc232 at #loc508)) +#loc567 = loc(callsite(#loc278 at #loc509)) +#loc568 = loc(callsite(#loc210 at #loc510)) +#loc569 = loc(callsite(#loc213 at #loc511)) +#loc570 = loc(callsite(#loc219 at #loc511)) +#loc571 = loc(callsite(#loc210 at #loc512)) +#loc572 = loc(callsite(#loc225 at #loc513)) +#loc573 = loc(callsite(#loc228 at #loc514)) +#loc574 = loc(callsite(#loc231 at #loc515)) +#loc575 = loc(callsite(#loc232 at #loc515)) +#loc576 = loc(callsite(#loc187 at #loc516)) +#loc577 = loc(callsite(#loc288 at #loc517)) +#loc578 = loc(callsite(#loc293 at #loc518)) +#loc579 = loc(callsite(#loc297 at #loc519)) +#loc580 = loc(callsite(#loc298 at #loc499)) +#loc581 = loc(callsite(#loc294 at #loc499)) +#loc582 = loc(callsite(#loc290 at #loc499)) +#loc583 = loc(callsite(#loc274 at #loc520)) +#loc584 = loc(callsite(#loc265 at #loc521)) +#loc585 = loc(callsite(#loc206 at #loc522)) +#loc586 = loc(callsite(#loc177 at #loc523)) +#loc588 = loc(callsite(#loc313 at #loc525)) +#loc589 = loc(callsite(#loc324 at #loc524)) +#loc590 = loc(callsite(#loc313 at #loc526)) +#loc591 = loc(callsite(#loc326 at #loc524)) +#loc593 = loc(callsite(#loc213 at #loc528)) +#loc594 = loc(callsite(#loc175 at #loc529)) +#loc595 = loc(callsite(#loc184 at #loc530)) +#loc596 = loc(callsite(#loc188 at #loc530)) +#loc597 = loc(callsite(#loc189 at #loc530)) +#loc598 = loc(callsite(#loc190 at #loc530)) +#loc599 = loc(callsite(#loc191 at #loc530)) +#loc600 = loc(callsite(#loc192 at #loc531)) +#loc601 = loc(callsite(#loc198 at #loc531)) +#loc602 = loc(callsite(#loc199 at #loc531)) +#loc603 = loc(callsite(#loc200 at #loc531)) +#loc604 = loc(callsite(#loc201 at #loc532)) +#loc605 = loc(callsite(#loc207 at #loc532)) +#loc606 = loc(callsite(#loc208 at #loc532)) +#loc607 = loc(callsite(#loc209 at #loc533)) +#loc608 = loc(callsite(#loc212 at #loc534)) +#loc609 = loc(callsite(#loc214 at #loc534)) +#loc610 = loc(callsite(#loc215 at #loc534)) +#loc611 = loc(callsite(#loc216 at #loc534)) +#loc612 = loc(callsite(#loc217 at #loc534)) +#loc613 = loc(callsite(#loc218 at #loc535)) +#loc614 = loc(callsite(#loc220 at #loc533)) +#loc615 = loc(callsite(#loc221 at #loc533)) +#loc616 = loc(callsite(#loc209 at #loc536)) +#loc617 = loc(callsite(#loc223 at #loc536)) +#loc618 = loc(callsite(#loc224 at #loc537)) +#loc619 = loc(callsite(#loc220 at #loc536)) +#loc620 = loc(callsite(#loc226 at #loc536)) +#loc621 = loc(callsite(#loc221 at #loc536)) +#loc622 = loc(callsite(#loc227 at #loc538)) +#loc623 = loc(callsite(#loc230 at #loc539)) +#loc624 = loc(callsite(#loc230 at #loc540)) +#loc625 = loc(callsite(#loc233 at #loc538)) +#loc626 = loc(callsite(#loc234 at #loc538)) +#loc627 = loc(callsite(#loc235 at #loc538)) +#loc628 = loc(callsite(#loc236 at #loc538)) +#loc629 = loc(callsite(#loc237 at #loc541)) +#loc630 = loc(callsite(#loc238 at #loc542)) +#loc631 = loc(callsite(#loc243 at #loc542)) +#loc632 = loc(callsite(#loc244 at #loc543)) +#loc633 = loc(callsite(#loc247 at #loc543)) +#loc634 = loc(callsite(#loc248 at #loc543)) +#loc635 = loc(callsite(#loc249 at #loc543)) +#loc636 = loc(callsite(#loc250 at #loc543)) +#loc637 = loc(callsite(#loc251 at #loc543)) +#loc638 = loc(callsite(#loc252 at #loc544)) +#loc639 = loc(callsite(#loc256 at #loc544)) +#loc640 = loc(callsite(#loc257 at #loc545)) +#loc641 = loc(callsite(#loc258 at #loc545)) +#loc643 = loc(callsite(#loc260 at #loc541)) +#loc644 = loc(callsite(#loc192 at #loc546)) +#loc645 = loc(callsite(#loc198 at #loc546)) +#loc646 = loc(callsite(#loc199 at #loc546)) +#loc647 = loc(callsite(#loc200 at #loc546)) +#loc648 = loc(callsite(#loc209 at #loc547)) +#loc649 = loc(callsite(#loc212 at #loc548)) +#loc650 = loc(callsite(#loc214 at #loc548)) +#loc651 = loc(callsite(#loc215 at #loc548)) +#loc652 = loc(callsite(#loc216 at #loc548)) +#loc653 = loc(callsite(#loc217 at #loc548)) +#loc654 = loc(callsite(#loc218 at #loc549)) +#loc655 = loc(callsite(#loc220 at #loc547)) +#loc656 = loc(callsite(#loc221 at #loc547)) +#loc657 = loc(callsite(#loc209 at #loc550)) +#loc658 = loc(callsite(#loc223 at #loc550)) +#loc659 = loc(callsite(#loc224 at #loc551)) +#loc660 = loc(callsite(#loc220 at #loc550)) +#loc661 = loc(callsite(#loc226 at #loc550)) +#loc662 = loc(callsite(#loc221 at #loc550)) +#loc663 = loc(callsite(#loc227 at #loc552)) +#loc664 = loc(callsite(#loc230 at #loc553)) +#loc665 = loc(callsite(#loc230 at #loc554)) +#loc666 = loc(callsite(#loc233 at #loc552)) +#loc667 = loc(callsite(#loc234 at #loc552)) +#loc668 = loc(callsite(#loc235 at #loc552)) +#loc669 = loc(callsite(#loc236 at #loc552)) +#loc670 = loc(callsite(#loc268 at #loc555)) +#loc671 = loc(callsite(#loc192 at #loc556)) +#loc672 = loc(callsite(#loc198 at #loc556)) +#loc673 = loc(callsite(#loc199 at #loc556)) +#loc674 = loc(callsite(#loc200 at #loc556)) +#loc675 = loc(callsite(#loc192 at #loc557)) +#loc676 = loc(callsite(#loc198 at #loc557)) +#loc677 = loc(callsite(#loc199 at #loc557)) +#loc678 = loc(callsite(#loc200 at #loc557)) +#loc679 = loc(callsite(#loc201 at #loc558)) +#loc680 = loc(callsite(#loc207 at #loc558)) +#loc681 = loc(callsite(#loc208 at #loc558)) +#loc682 = loc(callsite(#loc209 at #loc559)) +#loc683 = loc(callsite(#loc212 at #loc560)) +#loc684 = loc(callsite(#loc214 at #loc560)) +#loc685 = loc(callsite(#loc215 at #loc560)) +#loc686 = loc(callsite(#loc216 at #loc560)) +#loc687 = loc(callsite(#loc217 at #loc560)) +#loc688 = loc(callsite(#loc218 at #loc561)) +#loc689 = loc(callsite(#loc220 at #loc559)) +#loc690 = loc(callsite(#loc221 at #loc559)) +#loc691 = loc(callsite(#loc209 at #loc562)) +#loc692 = loc(callsite(#loc223 at #loc562)) +#loc693 = loc(callsite(#loc224 at #loc563)) +#loc694 = loc(callsite(#loc220 at #loc562)) +#loc695 = loc(callsite(#loc226 at #loc562)) +#loc696 = loc(callsite(#loc221 at #loc562)) +#loc697 = loc(callsite(#loc227 at #loc564)) +#loc698 = loc(callsite(#loc230 at #loc565)) +#loc699 = loc(callsite(#loc230 at #loc566)) +#loc700 = loc(callsite(#loc233 at #loc564)) +#loc701 = loc(callsite(#loc234 at #loc564)) +#loc702 = loc(callsite(#loc235 at #loc564)) +#loc703 = loc(callsite(#loc236 at #loc564)) +#loc704 = loc(callsite(#loc277 at #loc567)) +#loc705 = loc(callsite(#loc280 at #loc567)) +#loc706 = loc(callsite(#loc281 at #loc567)) +#loc707 = loc(callsite(#loc209 at #loc568)) +#loc708 = loc(callsite(#loc212 at #loc569)) +#loc709 = loc(callsite(#loc214 at #loc569)) +#loc710 = loc(callsite(#loc215 at #loc569)) +#loc711 = loc(callsite(#loc216 at #loc569)) +#loc712 = loc(callsite(#loc217 at #loc569)) +#loc713 = loc(callsite(#loc218 at #loc570)) +#loc714 = loc(callsite(#loc220 at #loc568)) +#loc715 = loc(callsite(#loc221 at #loc568)) +#loc716 = loc(callsite(#loc209 at #loc571)) +#loc717 = loc(callsite(#loc223 at #loc571)) +#loc718 = loc(callsite(#loc224 at #loc572)) +#loc719 = loc(callsite(#loc220 at #loc571)) +#loc720 = loc(callsite(#loc226 at #loc571)) +#loc721 = loc(callsite(#loc221 at #loc571)) +#loc722 = loc(callsite(#loc227 at #loc573)) +#loc723 = loc(callsite(#loc230 at #loc574)) +#loc724 = loc(callsite(#loc230 at #loc575)) +#loc725 = loc(callsite(#loc233 at #loc573)) +#loc726 = loc(callsite(#loc234 at #loc573)) +#loc727 = loc(callsite(#loc235 at #loc573)) +#loc728 = loc(callsite(#loc236 at #loc573)) +#loc729 = loc(callsite(#loc285 at #loc555)) +#loc730 = loc(callsite(#loc286 at #loc576)) +#loc731 = loc(callsite(#loc287 at #loc577)) +#loc732 = loc(callsite(#loc291 at #loc576)) +#loc733 = loc(callsite(#loc292 at #loc578)) +#loc734 = loc(callsite(#loc295 at #loc576)) +#loc735 = loc(callsite(#loc296 at #loc579)) +#loc736 = loc(callsite(#loc299 at #loc529)) +#loc737 = loc(callsite(#loc289 at #loc580)) +#loc738 = loc(callsite(#loc289 at #loc581)) +#loc739 = loc(callsite(#loc289 at #loc582)) +#loc740 = loc(callsite(#loc273 at #loc583)) +#loc741 = loc(callsite(#loc264 at #loc584)) +#loc742 = loc(callsite(#loc255 at #loc542)) +#loc743 = loc(callsite(#loc205 at #loc585)) +#loc744 = loc(callsite(#loc176 at #loc586)) +#loc745 = loc(callsite(#loc300 at #loc586)) +#loc746 = loc(callsite(#loc301 at #loc587)) +#loc747 = loc(callsite(#loc306 at #loc587)) +#loc748 = loc(callsite(#loc307 at #loc587)) +#loc749 = loc(callsite(#loc308 at #loc587)) +#loc751 = loc(callsite(#loc310 at #loc587)) +#loc752 = loc(callsite(#loc311 at #loc587)) +#loc753 = loc(callsite(#loc312 at #loc588)) +#loc754 = loc(callsite(#loc314 at #loc588)) +#loc755 = loc(callsite(#loc315 at #loc588)) +#loc756 = loc(callsite(#loc316 at #loc588)) +#loc757 = loc(callsite(#loc317 at #loc588)) +#loc758 = loc(callsite(#loc318 at #loc588)) +#loc759 = loc(callsite(#loc319 at #loc588)) +#loc760 = loc(callsite(#loc320 at #loc587)) +#loc761 = loc(callsite(#loc321 at #loc587)) +#loc762 = loc(callsite(#loc322 at #loc587)) +#loc763 = loc(callsite(#loc323 at #loc589)) +#loc764 = loc(callsite(#loc301 at #loc589)) +#loc765 = loc(callsite(#loc306 at #loc589)) +#loc766 = loc(callsite(#loc307 at #loc589)) +#loc767 = loc(callsite(#loc308 at #loc589)) +#loc768 = loc(callsite(#loc309 at #loc589)) +#loc769 = loc(callsite(#loc310 at #loc589)) +#loc770 = loc(callsite(#loc311 at #loc589)) +#loc771 = loc(callsite(#loc312 at #loc590)) +#loc772 = loc(callsite(#loc314 at #loc590)) +#loc773 = loc(callsite(#loc315 at #loc590)) +#loc774 = loc(callsite(#loc316 at #loc590)) +#loc775 = loc(callsite(#loc317 at #loc590)) +#loc776 = loc(callsite(#loc318 at #loc590)) +#loc777 = loc(callsite(#loc319 at #loc590)) +#loc778 = loc(callsite(#loc320 at #loc589)) +#loc779 = loc(callsite(#loc321 at #loc589)) +#loc780 = loc(callsite(#loc322 at #loc589)) +#loc781 = loc(callsite(#loc325 at #loc591)) +#loc782 = loc(callsite(#loc327 at #loc591)) +#loc783 = loc(callsite(#loc328 at #loc586)) +#loc784 = loc(callsite(#loc329 at #loc586)) +#loc787 = loc(callsite(#loc212 at #loc593)) +#loc788 = loc("jit(train_step)/jvp()/random_seed"(#loc594)) +#loc789 = loc("jit(train_step)/jvp()/shift_right_logical"(#loc594)) +#loc790 = loc("jit(train_step)/jvp()/convert_element_type"(#loc594)) +#loc791 = loc("jit(train_step)/jvp()/broadcast_in_dim"(#loc594)) +#loc792 = loc("jit(train_step)/jvp()/and"(#loc594)) +#loc793 = loc("jit(train_step)/jvp()/concatenate"(#loc594)) +#loc794 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/relpos_bias/iota"(#loc595)) +#loc795 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/relpos_bias/eq"(#loc596)) +#loc796 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/relpos_bias/convert_element_type"(#loc597)) +#loc797 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/relpos_bias/dot_general"(#loc598)) +#loc798 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/relpos_bias/broadcast_in_dim"(#loc599)) +#loc799 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/jit(_threefry_fold_in)"(#loc600)) +#loc800 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/jit(fold_in)"(#loc601)) +#loc801 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/jit(_randint)"(#loc602)) +#loc802 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/bitcast_convert_type"(#loc603)) +#loc803 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc604)) +#loc804 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc605)) +#loc805 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/te_norm_forward_ffi"(#loc606)) +#loc806 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc607)) +#loc807 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/jit(_diag)"(#loc608)) +#loc808 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/dot_general"(#loc609)) +#loc809 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/sqrt"(#loc610)) +#loc810 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/convert_element_type"(#loc611)) +#loc811 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/div"(#loc611)) +#loc812 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/convert_element_type"(#loc612)) +#loc813 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/te_rht_amax_ffi"(#loc613)) +#loc814 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc614)) +#loc815 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/te_dbias_quantize_ffi"(#loc615)) +#loc816 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/slice"(#loc615)) +#loc817 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc616)) +#loc818 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc617)) +#loc819 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/abs"(#loc618)) +#loc820 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/reduce_max"(#loc618)) +#loc821 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc618)) +#loc822 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/convert_element_type"(#loc618)) +#loc823 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/reshape"(#loc618)) +#loc824 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc619)) +#loc825 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc620)) +#loc826 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/te_dbias_quantize_ffi"(#loc621)) +#loc827 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/slice"(#loc621)) +#loc828 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc622)) +#loc829 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/div"(#loc623)) +#loc830 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/div"(#loc624)) +#loc831 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/mul"(#loc625)) +#loc832 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc626)) +#loc833 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc627)) +#loc834 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/te_gemm_ffi"(#loc628)) +#loc835 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/reshape"(#loc629)) +#loc836 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc630)) +#loc837 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc631)) +#loc838 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc632)) +#loc839 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc633)) +#loc840 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc634)) +#loc841 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc635)) +#loc842 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc636)) +#loc843 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc637)) +#loc844 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc638)) +#loc845 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc639)) +#loc846 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/reshape"(#loc639)) +#loc847 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc640)) +#loc848 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc641)) +#loc849 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/lt"(#loc642)) +#loc850 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/jit(_where)"(#loc642)) +#loc851 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/jit(_cumsum_with_promotion)"(#loc642)) +#loc852 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc642)) +#loc853 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/concatenate"(#loc642)) +#loc854 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/te_fused_attn_forward_ffi"(#loc642)) +#loc855 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/reshape"(#loc643)) +#loc856 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/out.generate_quantizer_set/jit(_threefry_fold_in)"(#loc644)) +#loc857 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/out.generate_quantizer_set/jit(fold_in)"(#loc645)) +#loc858 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/out.generate_quantizer_set/jit(_randint)"(#loc646)) +#loc859 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/out.generate_quantizer_set/bitcast_convert_type"(#loc647)) +#loc860 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc648)) +#loc861 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/jit(_diag)"(#loc649)) +#loc862 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/dot_general"(#loc650)) +#loc863 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/sqrt"(#loc651)) +#loc864 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/convert_element_type"(#loc652)) +#loc865 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/div"(#loc652)) +#loc866 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/convert_element_type"(#loc653)) +#loc867 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/te_rht_amax_ffi"(#loc654)) +#loc868 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc655)) +#loc869 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/te_dbias_quantize_ffi"(#loc656)) +#loc870 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/slice"(#loc656)) +#loc871 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc657)) +#loc872 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc658)) +#loc873 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/abs"(#loc659)) +#loc874 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/reduce_max"(#loc659)) +#loc875 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc659)) +#loc876 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/convert_element_type"(#loc659)) +#loc877 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/reshape"(#loc659)) +#loc878 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc660)) +#loc879 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc661)) +#loc880 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/te_dbias_quantize_ffi"(#loc662)) +#loc881 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/slice"(#loc662)) +#loc882 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc663)) +#loc883 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/div"(#loc664)) +#loc884 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/div"(#loc665)) +#loc885 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/mul"(#loc666)) +#loc886 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc667)) +#loc887 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc668)) +#loc888 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/out/te_gemm_ffi"(#loc669)) +#loc889 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/add"(#loc670)) +#loc890 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(_threefry_fold_in)"(#loc671)) +#loc891 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(fold_in)"(#loc672)) +#loc892 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(_randint)"(#loc673)) +#loc893 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/bitcast_convert_type"(#loc674)) +#loc894 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(_threefry_fold_in)"(#loc675)) +#loc895 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(fold_in)"(#loc676)) +#loc896 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/jit(_randint)"(#loc677)) +#loc897 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mlp.generate_quantizer_set/bitcast_convert_type"(#loc678)) +#loc898 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc679)) +#loc899 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc680)) +#loc900 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_norm_forward_ffi"(#loc681)) +#loc901 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc682)) +#loc902 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/jit(_diag)"(#loc683)) +#loc903 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/dot_general"(#loc684)) +#loc904 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/sqrt"(#loc685)) +#loc905 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc686)) +#loc906 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc686)) +#loc907 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc687)) +#loc908 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_rht_amax_ffi"(#loc688)) +#loc909 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc689)) +#loc910 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_dbias_quantize_ffi"(#loc690)) +#loc911 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/slice"(#loc690)) +#loc912 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc691)) +#loc913 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc692)) +#loc914 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/abs"(#loc693)) +#loc915 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/reduce_max"(#loc693)) +#loc916 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc693)) +#loc917 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc693)) +#loc918 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/reshape"(#loc693)) +#loc919 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc694)) +#loc920 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc695)) +#loc921 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_dbias_quantize_ffi"(#loc696)) +#loc922 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/slice"(#loc696)) +#loc923 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc697)) +#loc924 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc698)) +#loc925 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc699)) +#loc926 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mul"(#loc700)) +#loc927 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc701)) +#loc928 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc702)) +#loc929 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_gemm_ffi"(#loc703)) +#loc930 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc704)) +#loc931 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc705)) +#loc932 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_act_lu_ffi"(#loc706)) +#loc933 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc707)) +#loc934 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/jit(_diag)"(#loc708)) +#loc935 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/dot_general"(#loc709)) +#loc936 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/sqrt"(#loc710)) +#loc937 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc711)) +#loc938 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc711)) +#loc939 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc712)) +#loc940 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_rht_amax_ffi"(#loc713)) +#loc941 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc714)) +#loc942 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_dbias_quantize_ffi"(#loc715)) +#loc943 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/slice"(#loc715)) +#loc944 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc716)) +#loc945 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc717)) +#loc946 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/abs"(#loc718)) +#loc947 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/reduce_max"(#loc718)) +#loc948 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc718)) +#loc949 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/convert_element_type"(#loc718)) +#loc950 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/reshape"(#loc718)) +#loc951 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc719)) +#loc952 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc720)) +#loc953 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_dbias_quantize_ffi"(#loc721)) +#loc954 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/slice"(#loc721)) +#loc955 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc722)) +#loc956 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc723)) +#loc957 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/div"(#loc724)) +#loc958 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/mul"(#loc725)) +#loc959 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc726)) +#loc960 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/broadcast_in_dim"(#loc727)) +#loc961 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/mlp/te_gemm_ffi"(#loc728)) +#loc962 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/add"(#loc729)) +#loc963 = loc("jit(train_step)/jvp(Model)/reshape"(#loc730)) +#loc964 = loc("jit(train_step)/jvp(Model)/te_scaled_softmax_forward_ffi"(#loc731)) +#loc965 = loc("jit(train_step)/jvp(Model)/convert_element_type"(#loc732)) +#loc966 = loc("jit(train_step)/jvp(Model)/te_scaled_masked_softmax_forward_ffi"(#loc733)) +#loc967 = loc("jit(train_step)/jvp(Model)/reshape"(#loc734)) +#loc968 = loc("jit(train_step)/jvp(Model)/te_scaled_upper_triang_masked_softmax_forward_ffi"(#loc735)) +#loc969 = loc("jit(train_step)/jvp()/convert_element_type"(#loc736)) +#loc970 = loc("jit(train_step)/jvp()/reduce_sum"(#loc736)) +#loc971 = loc("jit(train_step)/jvp()/div"(#loc736)) +#loc972 = loc("jit(train_step)/transpose(jvp())/div"(#loc736)) +#loc973 = loc("jit(train_step)/transpose(jvp())/broadcast_in_dim"(#loc736)) +#loc974 = loc("jit(train_step)/transpose(jvp())/convert_element_type"(#loc736)) +#loc975 = loc("jit(train_step)/transpose(jvp(Model))/te_scaled_upper_triang_masked_softmax_backward_ffi"(#loc737)) +#loc976 = loc("jit(train_step)/transpose(jvp(Model))/reshape"(#loc734)) +#loc977 = loc("jit(train_step)/transpose(jvp(Model))/te_scaled_masked_softmax_backward_ffi"(#loc738)) +#loc978 = loc("jit(train_step)/transpose(jvp(Model))/te_scaled_softmax_backward_ffi"(#loc739)) +#loc979 = loc("jit(train_step)/transpose(jvp(Model))/reshape"(#loc730)) +#loc980 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/broadcast_in_dim"(#loc740)) +#loc981 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/jit(_diag)"(#loc740)) +#loc982 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/dot_general"(#loc740)) +#loc983 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/sqrt"(#loc740)) +#loc984 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/convert_element_type"(#loc740)) +#loc985 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/div"(#loc740)) +#loc986 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/te_rht_amax_ffi"(#loc740)) +#loc987 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/te_dbias_quantize_ffi"(#loc740)) +#loc988 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/slice"(#loc740)) +#loc989 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/mul"(#loc740)) +#loc990 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/te_gemm_ffi"(#loc740)) +#loc991 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/te_dact_dbias_quantize_ffi"(#loc740)) +#loc992 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/te_norm_backward_ffi"(#loc740)) +#loc993 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/add_any"(#loc740)) +#loc994 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/broadcast_in_dim"(#loc741)) +#loc995 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/jit(_diag)"(#loc741)) +#loc996 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/dot_general"(#loc741)) +#loc997 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/sqrt"(#loc741)) +#loc998 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/convert_element_type"(#loc741)) +#loc999 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/div"(#loc741)) +#loc1000 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/te_rht_amax_ffi"(#loc741)) +#loc1001 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/te_dbias_quantize_ffi"(#loc741)) +#loc1002 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/slice"(#loc741)) +#loc1003 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/mul"(#loc741)) +#loc1004 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/out/te_gemm_ffi"(#loc741)) +#loc1005 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/reshape"(#loc643)) +#loc1006 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/broadcast_in_dim"(#loc742)) +#loc1007 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/lt"(#loc742)) +#loc1008 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/jit(_where)"(#loc742)) +#loc1009 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/jit(_cumsum_with_promotion)"(#loc742)) +#loc1010 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/concatenate"(#loc742)) +#loc1011 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/DotProductAttention_0/_FusedDotProductAttention_0/te_fused_attn_backward_ffi"(#loc742)) +#loc1012 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/reshape"(#loc629)) +#loc1013 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/broadcast_in_dim"(#loc743)) +#loc1014 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/jit(_diag)"(#loc743)) +#loc1015 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/dot_general"(#loc743)) +#loc1016 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/sqrt"(#loc743)) +#loc1017 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/convert_element_type"(#loc743)) +#loc1018 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/div"(#loc743)) +#loc1019 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/te_rht_amax_ffi"(#loc743)) +#loc1020 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/te_dbias_quantize_ffi"(#loc743)) +#loc1021 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/slice"(#loc743)) +#loc1022 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/mul"(#loc743)) +#loc1023 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/te_gemm_ffi"(#loc743)) +#loc1024 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/attention/qkv/te_norm_backward_ffi"(#loc743)) +#loc1025 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/relpos_bias/reduce_sum"(#loc599)) +#loc1026 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/relpos_bias/dot_general"(#loc598)) +#loc1027 = loc("jit(train_step)/broadcast_in_dim"(#loc744)) +#loc1028 = loc("jit(train_step)/reshape"(#loc745)) +#loc1029 = loc("jit(train_step)/broadcast_in_dim"(#loc746)) +#loc1030 = loc("jit(train_step)/abs"(#loc747)) +#loc1031 = loc("jit(train_step)/reduce_max"(#loc748)) +#loc1032 = loc("jit(train_step)/iota"(#loc749)) +#loc1033 = loc("jit(train_step)/jit(_roll_static)"(#loc750)) +#loc1034 = loc("jit(train_step)/broadcast_in_dim"(#loc750)) +#loc1035 = loc("jit(train_step)/scatter"(#loc750)) +#loc1036 = loc("jit(train_step)/jit(cumsum)"(#loc750)) +#loc1037 = loc("jit(train_step)/lt"(#loc750)) +#loc1038 = loc("jit(train_step)/add"(#loc750)) +#loc1039 = loc("jit(train_step)/select_n"(#loc750)) +#loc1040 = loc("jit(train_step)/scatter-add"(#loc750)) +#loc1041 = loc("jit(train_step)/sub"(#loc750)) +#loc1042 = loc("jit(train_step)/jit(_take)"(#loc750)) +#loc1043 = loc("jit(train_step)/broadcast_in_dim"(#loc751)) +#loc1044 = loc("jit(train_step)/scatter-max"(#loc751)) +#loc1045 = loc("jit(train_step)/max"(#loc751)) +#loc1046 = loc("jit(train_step)/slice"(#loc752)) +#loc1047 = loc("jit(train_step)/squeeze"(#loc752)) +#loc1048 = loc("jit(train_step)/broadcast_in_dim"(#loc753)) +#loc1049 = loc("jit(train_step)/convert_element_type"(#loc754)) +#loc1050 = loc("jit(train_step)/div"(#loc754)) +#loc1051 = loc("jit(train_step)/div"(#loc755)) +#loc1052 = loc("jit(train_step)/gt"(#loc756)) +#loc1053 = loc("jit(train_step)/jit(_where)"(#loc757)) +#loc1054 = loc("jit(train_step)/is_finite"(#loc758)) +#loc1055 = loc("jit(train_step)/jit(_where)"(#loc759)) +#loc1056 = loc("jit(train_step)/slice"(#loc760)) +#loc1057 = loc("jit(train_step)/squeeze"(#loc760)) +#loc1058 = loc("jit(train_step)/broadcast_in_dim"(#loc761)) +#loc1059 = loc("jit(train_step)/scatter"(#loc761)) +#loc1060 = loc("jit(train_step)/te_grouped_quantize_ffi"(#loc762)) +#loc1061 = loc("jit(train_step)/broadcast_in_dim"(#loc763)) +#loc1062 = loc("jit(train_step)/broadcast_in_dim"(#loc764)) +#loc1063 = loc("jit(train_step)/abs"(#loc765)) +#loc1064 = loc("jit(train_step)/reduce_max"(#loc766)) +#loc1065 = loc("jit(train_step)/iota"(#loc767)) +#loc1066 = loc("jit(train_step)/jit(_roll_static)"(#loc768)) +#loc1067 = loc("jit(train_step)/broadcast_in_dim"(#loc768)) +#loc1068 = loc("jit(train_step)/scatter"(#loc768)) +#loc1069 = loc("jit(train_step)/jit(cumsum)"(#loc768)) +#loc1070 = loc("jit(train_step)/lt"(#loc768)) +#loc1071 = loc("jit(train_step)/add"(#loc768)) +#loc1072 = loc("jit(train_step)/select_n"(#loc768)) +#loc1073 = loc("jit(train_step)/scatter-add"(#loc768)) +#loc1074 = loc("jit(train_step)/sub"(#loc768)) +#loc1075 = loc("jit(train_step)/jit(_take)"(#loc768)) +#loc1076 = loc("jit(train_step)/broadcast_in_dim"(#loc769)) +#loc1077 = loc("jit(train_step)/scatter-max"(#loc769)) +#loc1078 = loc("jit(train_step)/max"(#loc769)) +#loc1079 = loc("jit(train_step)/slice"(#loc770)) +#loc1080 = loc("jit(train_step)/squeeze"(#loc770)) +#loc1081 = loc("jit(train_step)/broadcast_in_dim"(#loc771)) +#loc1082 = loc("jit(train_step)/convert_element_type"(#loc772)) +#loc1083 = loc("jit(train_step)/div"(#loc772)) +#loc1084 = loc("jit(train_step)/div"(#loc773)) +#loc1085 = loc("jit(train_step)/gt"(#loc774)) +#loc1086 = loc("jit(train_step)/jit(_where)"(#loc775)) +#loc1087 = loc("jit(train_step)/is_finite"(#loc776)) +#loc1088 = loc("jit(train_step)/jit(_where)"(#loc777)) +#loc1089 = loc("jit(train_step)/slice"(#loc778)) +#loc1090 = loc("jit(train_step)/squeeze"(#loc778)) +#loc1091 = loc("jit(train_step)/broadcast_in_dim"(#loc779)) +#loc1092 = loc("jit(train_step)/scatter"(#loc779)) +#loc1093 = loc("jit(train_step)/te_grouped_quantize_ffi"(#loc780)) +#loc1094 = loc("jit(train_step)/broadcast_in_dim"(#loc781)) +#loc1095 = loc("jit(train_step)/te_grouped_gemm_ffi"(#loc782)) +#loc1096 = loc("jit(train_step)/convert_element_type"(#loc783)) +#loc1097 = loc("jit(train_step)/reduce_sum"(#loc783)) +#loc1098 = loc("jit(train_step)/div"(#loc783)) +#loc1099 = loc("jit(train_step)/add"(#loc784)) +#loc1100 = loc("jit"(#loc600)) +#loc1101 = loc("shift_right_logical"(#loc785)) +#loc1102 = loc("broadcast_in_dim"(#loc785)) +#loc1103 = loc("and"(#loc785)) +#loc1104 = loc("concatenate"(#loc785)) +#loc1105 = loc("slice"(#loc785)) +#loc1106 = loc("squeeze"(#loc785)) +#loc1107 = loc("split"(#loc785)) +#loc1108 = loc("xor"(#loc785)) +#loc1109 = loc("add"(#loc785)) +#loc1110 = loc("shift_left"(#loc785)) +#loc1111 = loc("or"(#loc785)) +#loc1112 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/jit"(#loc601)) +#loc1113 = loc("convert_element_type"(#loc785)) +#loc1114 = loc("jit(_threefry_fold_in)"(#loc785)) +#loc1115 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/qkv.generate_quantizer_set/jit"(#loc602)) +#loc1116 = loc("jit(clip)"(#loc786)) +#loc1117 = loc("gt"(#loc786)) +#loc1118 = loc("convert_element_type"(#loc786)) +#loc1119 = loc("broadcast_in_dim"(#loc786)) +#loc1120 = loc("jit(_threefry_split)"(#loc786)) +#loc1121 = loc("slice"(#loc786)) +#loc1122 = loc("squeeze"(#loc786)) +#loc1123 = loc("iota_2x32_shape"(#loc786)) +#loc1124 = loc(""(#loc786)) +#loc1125 = loc("xor"(#loc786)) +#loc1126 = loc("sub"(#loc786)) +#loc1127 = loc("le"(#loc786)) +#loc1128 = loc("select_n"(#loc786)) +#loc1129 = loc("and"(#loc786)) +#loc1130 = loc("add"(#loc786)) +#loc1131 = loc("rem"(#loc786)) +#loc1132 = loc("mul"(#loc786)) +#loc1133 = loc("jit"(#loc786)) +#loc1134 = loc("max"(#loc786)) +#loc1135 = loc("min"(#loc786)) +#loc1136 = loc("concatenate"(#loc786)) +#loc1137 = loc("shift_left"(#loc786)) +#loc1138 = loc("shift_right_logical"(#loc786)) +#loc1139 = loc("or"(#loc786)) +#loc1140 = loc("jit(train_step)/jvp(Model)/TransformerLayer_0/attention/qkv/jit"(#loc608)) +#loc1141 = loc("pad"(#loc787)) +#loc1142 = loc("iota"(#loc787)) +#loc1143 = loc("add"(#loc787)) +#loc1144 = loc("eq"(#loc787)) +#loc1145 = loc("broadcast_in_dim"(#loc787)) +#loc1146 = loc("jit(_where)"(#loc787)) +#loc1147 = loc("jit"(#loc787)) +#loc1148 = loc("select_n"(#loc787)) +#loc1149 = loc("jit"(#loc642)) +#loc1150 = loc("convert_element_type"(#loc642)) +#loc1151 = loc("broadcast_in_dim"(#loc642)) +#loc1152 = loc("select_n"(#loc642)) +#loc1153 = loc("reduce_window_sum"(#loc642)) +#loc1154 = loc("jit(train_step)/transpose(jvp(Model))/TransformerLayer_0/mlp/jit"(#loc740)) +#loc1155 = loc("pad"(#loc740)) +#loc1156 = loc("iota"(#loc740)) +#loc1157 = loc("add"(#loc740)) +#loc1158 = loc("eq"(#loc740)) +#loc1159 = loc("broadcast_in_dim"(#loc740)) +#loc1160 = loc("jit(_where)"(#loc740)) +#loc1161 = loc("jit"(#loc740)) +#loc1162 = loc("select_n"(#loc740)) +#loc1163 = loc("jit(train_step)/jit"(#loc750)) +#loc1164 = loc("slice"(#loc750)) +#loc1165 = loc("concatenate"(#loc750)) +#loc1166 = loc("reduce_window_sum"(#loc750)) +#loc1167 = loc("gather"(#loc750)) +#loc1168 = loc("reduce_and"(#loc750)) +#loc1169 = loc("lt"(#loc750)) +#loc1170 = loc("add"(#loc750)) +#loc1171 = loc("jit(_where)"(#loc750)) +#loc1172 = loc("broadcast_in_dim"(#loc750)) +#loc1173 = loc("ge"(#loc750)) +#loc1174 = loc("le"(#loc750)) +#loc1175 = loc("and"(#loc750)) +#loc1176 = loc("select_n"(#loc750)) +#loc1177 = loc("jit"(#loc750)) +#loc1178 = loc("jit(train_step)/jit"(#loc757)) +#loc1179 = loc("broadcast_in_dim"(#loc757)) +#loc1180 = loc("select_n"(#loc757)) +#loc1181 = loc("jit(train_step)/jit"(#loc759)) +#loc1182 = loc("select_n"(#loc759)) +#loc1183 = loc("jit(train_step)/jit"(#loc768)) +#loc1184 = loc("gather"(#loc768)) +#loc1185 = loc("reduce_and"(#loc768)) +#loc1186 = loc("lt"(#loc768)) +#loc1187 = loc("add"(#loc768)) +#loc1188 = loc("jit(_where)"(#loc768)) +#loc1189 = loc("broadcast_in_dim"(#loc768)) +#loc1190 = loc("ge"(#loc768)) +#loc1191 = loc("le"(#loc768)) +#loc1192 = loc("and"(#loc768)) +#loc1193 = loc("select_n"(#loc768)) +#loc1194 = loc("jit"(#loc768)) +#loc1195 = loc("jit:"(#loc1100)) +#loc1196 = loc("jit:"(#loc1112)) +#loc1197 = loc("jit:"(#loc1115)) +#loc1198 = loc("jit:"(#loc1133)) +#loc1199 = loc("jit:"(#loc1140)) +#loc1200 = loc("jit:"(#loc1147)) +#loc1201 = loc("jit:"(#loc1149)) +#loc1202 = loc("jit:"(#loc1154)) +#loc1203 = loc("jit:"(#loc1161)) +#loc1204 = loc("jit:"(#loc1163)) +#loc1205 = loc("jit:"(#loc1177)) +#loc1206 = loc("jit:"(#loc1178)) +#loc1207 = loc("jit:"(#loc1181)) +#loc1208 = loc("jit:"(#loc1183)) +#loc1209 = loc("jit:"(#loc1194)) diff --git a/tests/jax/test_custom_call_compute.py b/tests/jax/test_custom_call_compute.py index c8bd9d47c31..65b25c69478 100644 --- a/tests/jax/test_custom_call_compute.py +++ b/tests/jax/test_custom_call_compute.py @@ -2,6 +2,7 @@ # # See LICENSE for license information. +from io import StringIO import jax import jax.numpy as jnp import pytest @@ -9,6 +10,8 @@ from functools import reduce from typing import Union import operator +import os +import re from utils import ( assert_allclose, @@ -1921,3 +1924,179 @@ def test_grouped_dense_grad_fp8(self, fwd_bwd_dtype, scaling_mode, input_shape): assert_allclose(prim_dgrad, ref_dgrad, dtype=bwd_dtype) assert_allclose(prim_wgrad, ref_wgrad, dtype=bwd_dtype) assert_allclose(prim_dbias, ref_dbias, dtype=dtype) + +class TestFFICompatibility: + + HLO_DIR = os.path.join(os.path.dirname(__file__), 'ffi_hlo') + + @pytest.fixture(name="ffi_hlo_name") + def hlo_fixture(shape): + for file in os.listdir(TestFFICompatibility.HLO_DIR): + file_path = os.path.join(TestFFICompatibility.HLO_DIR, file) + if os.path.isfile(file_path): + yield file.split('.')[0] + + @pytest.mark.skipif(os.getenv("NVTE_JAX_FFI_HLO_GENERATE", "0") != "1", reason="HLO generation not enabled") + def test_generate_hlo(self): + """ Run this test with NVTE_JAX_FFI_HLO_GENERATE=1 to generate StableHLO text files for FFI compatibility tests. Use this when intentionally changing FFI bindings and breaking compatibility changes are required. + + Instructions: + 1. `CUDA_VISIBLE_DEVICES=0 XLA_FLAGS="$XLA_FLAGS --xla_dump_to=./tests/jax/ffi_hlo_dump" NVTE_JAX_FFI_HLO_GENERATE=1 pytest tests/jax/test_custom_call_compute.py::TestFFICompatibility::test_generate_hlo -s` + 2. Find `tests/jax/ffi_hlo_dump/jit_train_step_/module.mlir` and copy it to the `tests/jax/ffi_hlo/` directory named transformer_stablehlo.txt + """ + import math + from transformer_engine.common.recipe import NVFP4BlockScaling, Float8CurrentScaling + from transformer_engine.jax import autocast, MeshResource, softmax + from transformer_engine.jax.flax import TransformerLayer + import flax.linen as nn + + with autocast(enabled=True, recipe=NVFP4BlockScaling(), mesh_resource=MeshResource()): + class Model(nn.Module): + """ This module does not represent any meaningful model, it is just to cover all FFI calls. """ + + @nn.compact + def __call__(self, x): + # Covers most of the FFI calls + x = TransformerLayer( + hidden_dropout=0.0, + attention_dropout=0.0, + intermediate_dropout=0.0, + dtype=jnp.bfloat16, + )(x) + + # Arbitrarily call softmax multiple times to cover all softmax FFI calls + x = x.reshape((1, *x.shape)) + x = softmax.softmax(x, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED) + mask1 = self.variable('collection', 'mask1', lambda: jax.random.bernoulli(jax.random.PRNGKey(0), shape=x.shape).astype(jnp.bfloat16)).value.astype(jnp.uint8) + x = softmax.softmax(x, mask=mask1, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_MASKED) + mask2 = self.variable('collection', 'mask2', lambda: (1.0 - jnp.tril(jnp.ones_like(x))).astype(jnp.bfloat16)).value.astype(jnp.uint8) + x = x.reshape((-1, 1, 32, 32)) + x = softmax.softmax(x, mask=mask2, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_UPPER_TRIANG_MASKED) + return x + + + model = Model() + input_shape = (1, 128, 512) + x = jnp.ones(input_shape, dtype=jnp.bfloat16) + + var_collect = model.init(jax.random.PRNGKey(0), x) + + def f(var_collect, x): + x = model.apply(var_collect, x, rngs={'sr_rng': jax.random.PRNGKey(0)}) + x = jnp.mean(x) # fake loss function for value_and_grad + return x + + @jax.jit + def train_step(var_collect, x, grouped_kernel): + loss, grads = jax.value_and_grad(f)(var_collect, x) + + # Arbitrarily call grouped quantize and GEMM to cover remaining FFI calls + x = x.reshape((-1, x.shape[-1])) + x = grouped_dense( + x, + grouped_kernel, + contracting_dims=((1,), (1,)), + group_sizes=jnp.array([x.shape[0]], dtype=jnp.int32), + quantizer_set=QuantizerFactory.create_set(n_groups=1, fp8_recipe=Float8CurrentScaling(), quantize_meta_set=QuantizeMetaSet(QuantizeMeta(), QuantizeMeta(), QuantizeMeta())), + ) + loss += jnp.mean(x) + + return loss, grads + + grouped_kernel = jnp.zeros((1, x.shape[-1], x.shape[-1]), dtype=jnp.bfloat16) + train_step(var_collect, x, grouped_kernel) + + def _get_hlo_text_from_file(self, hlo_name: str) -> str: + """ Reads the StableHLO text from a file given its name. """ + hlo_file_path = os.path.join(self.HLO_DIR, f"{hlo_name}.txt") + with open(hlo_file_path, 'r') as f: + hlo_text = f.read() + return hlo_text + + def _make_args_based_on_input_tensor_shape_and_dtype(self, stablehlo_text: str): + """ Parses the StableHLO text to extract input tensor shapes and dtypes, and creates dummy JAX arrays accordingly. """ + # Parse function signature to extract argument information + # Pattern matches: @main(%arg0: tensor<32x32xbf16>, %arg1: tensor<64xf32>, ...) + pattern = r'@main\((.*?)\{' + match = re.search(pattern, stablehlo_text) + + if not match: + raise ValueError("Could not find @main function signature in StableHLO text") + + args_str = match.group(1) + + # Parse individual arguments + # Pattern matches: %arg0: tensor<32x32xbf16> + arg_pattern = r'%arg(\d+):\s*tensor<([^>]+)>' + arg_matches = re.findall(arg_pattern, args_str) + + parsed_args = [] + for arg_num, shape_and_dtype_str in arg_matches: + print(f"Parsing argument {arg_num} with shape and dtype: {shape_and_dtype_str}") + # Parse shape: "32x32xbf16" -> [32, 32] + dtype_str = shape_and_dtype_str.split('x')[-1] + shape = [int(dim) for dim in shape_and_dtype_str.split('x')[:-1]] + + # Map StableHLO dtype to JAX dtype + dtype_map = { + 'bf16': jnp.bfloat16, + 'f32': jnp.float32, + 'f16': jnp.float16, + 'f8E4M3FN': jnp.float8_e4m3fn, + 'f8E5M2': jnp.float8_e5m2, + 'i32': jnp.int32, + 'ui32': jnp.uint32, + } + dtype = dtype_map.get(dtype_str) + + parsed_args.append(jnp.ones( + shape, + dtype=dtype + )) + return parsed_args + + def test_ffi_compatibility(self, ffi_hlo_name): + """ Tests that the current FFI bindings are compatible with the provided HLO and there are no API mismatches. """ + from jax.extend.backend import get_backend + + stablehlo_text = self._get_hlo_text_from_file(ffi_hlo_name) + args = self._make_args_based_on_input_tensor_shape_and_dtype(stablehlo_text) + + client = get_backend('cuda') + executable = client.compile_and_load(stablehlo_text.encode('utf-8'), executable_devices=jax.devices()[:1]) + results = executable.execute(args) + print(results) # No need to assert anything here, just ensure it runs without error + + def test_all_primitive_ffi_tested(self): + """ Ensures that all our TE primitives with FFI bindings are included in the FFI HLO compatibility tests. """ + # Open all HLO files and extract primitive FFI names + tested_hlos = set() + for file in os.listdir(self.HLO_DIR): + file_path = os.path.join(self.HLO_DIR, file) + if os.path.isfile(file_path) and file.endswith('.txt'): + with open(file_path, 'r') as f: + hlo_text = f.read() + # Extract primitive name from HLO text + pattern = r'stablehlo.custom_call @(.+?)\(' + matches = re.findall(pattern, hlo_text) + if matches: + for match in matches: + primitive_name = match + tested_hlos.add(primitive_name) + + # Assert that all registered primitives have corresponding FFI tests + import transformer_engine_jax + + KNOWN_MISSING_FFI_TESTS = { + # dequantize does not have a JAX primitive currently + 'te_dequantize_ffi', + # needs testing + 'te_grouped_gemm_d2h_group_sizes_ffi', + } + + unmatched_primitives = set() + for primitive_ffi_name, _ in transformer_engine_jax.registrations().items(): + if primitive_ffi_name not in tested_hlos and primitive_ffi_name not in KNOWN_MISSING_FFI_TESTS: + unmatched_primitives.add(primitive_ffi_name) + + assert len(unmatched_primitives) == 0, f"The following primitives do not have FFI tests: {unmatched_primitives}" \ No newline at end of file From 45519796d9b8ca547b406cb4cdb43981bdbf1e8b Mon Sep 17 00:00:00 2001 From: Jeremy Berchtold Date: Tue, 16 Dec 2025 08:13:58 -0800 Subject: [PATCH 2/5] Skip nvfp4 test on unsupported arch Signed-off-by: Jeremy Berchtold --- tests/jax/test_custom_call_compute.py | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/jax/test_custom_call_compute.py b/tests/jax/test_custom_call_compute.py index 65b25c69478..2b7039068d8 100644 --- a/tests/jax/test_custom_call_compute.py +++ b/tests/jax/test_custom_call_compute.py @@ -2055,6 +2055,7 @@ def _make_args_based_on_input_tensor_shape_and_dtype(self, stablehlo_text: str): )) return parsed_args + @pytest.mark.skipif(is_fp4_supported, reason=fp4_unsupported_reason) def test_ffi_compatibility(self, ffi_hlo_name): """ Tests that the current FFI bindings are compatible with the provided HLO and there are no API mismatches. """ from jax.extend.backend import get_backend From 6d5d210241d485f7a2d09294414f4e4d5f50296b Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Tue, 16 Dec 2025 16:23:08 +0000 Subject: [PATCH 3/5] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- tests/jax/test_custom_call_compute.py | 127 ++++++++++++++++---------- 1 file changed, 78 insertions(+), 49 deletions(-) diff --git a/tests/jax/test_custom_call_compute.py b/tests/jax/test_custom_call_compute.py index 2b7039068d8..b2c5775cc0d 100644 --- a/tests/jax/test_custom_call_compute.py +++ b/tests/jax/test_custom_call_compute.py @@ -1925,21 +1925,24 @@ def test_grouped_dense_grad_fp8(self, fwd_bwd_dtype, scaling_mode, input_shape): assert_allclose(prim_wgrad, ref_wgrad, dtype=bwd_dtype) assert_allclose(prim_dbias, ref_dbias, dtype=dtype) + class TestFFICompatibility: - HLO_DIR = os.path.join(os.path.dirname(__file__), 'ffi_hlo') + HLO_DIR = os.path.join(os.path.dirname(__file__), "ffi_hlo") @pytest.fixture(name="ffi_hlo_name") def hlo_fixture(shape): for file in os.listdir(TestFFICompatibility.HLO_DIR): file_path = os.path.join(TestFFICompatibility.HLO_DIR, file) if os.path.isfile(file_path): - yield file.split('.')[0] + yield file.split(".")[0] - @pytest.mark.skipif(os.getenv("NVTE_JAX_FFI_HLO_GENERATE", "0") != "1", reason="HLO generation not enabled") + @pytest.mark.skipif( + os.getenv("NVTE_JAX_FFI_HLO_GENERATE", "0") != "1", reason="HLO generation not enabled" + ) def test_generate_hlo(self): - """ Run this test with NVTE_JAX_FFI_HLO_GENERATE=1 to generate StableHLO text files for FFI compatibility tests. Use this when intentionally changing FFI bindings and breaking compatibility changes are required. - + """Run this test with NVTE_JAX_FFI_HLO_GENERATE=1 to generate StableHLO text files for FFI compatibility tests. Use this when intentionally changing FFI bindings and breaking compatibility changes are required. + Instructions: 1. `CUDA_VISIBLE_DEVICES=0 XLA_FLAGS="$XLA_FLAGS --xla_dump_to=./tests/jax/ffi_hlo_dump" NVTE_JAX_FFI_HLO_GENERATE=1 pytest tests/jax/test_custom_call_compute.py::TestFFICompatibility::test_generate_hlo -s` 2. Find `tests/jax/ffi_hlo_dump/jit_train_step_/module.mlir` and copy it to the `tests/jax/ffi_hlo/` directory named transformer_stablehlo.txt @@ -1949,10 +1952,11 @@ def test_generate_hlo(self): from transformer_engine.jax import autocast, MeshResource, softmax from transformer_engine.jax.flax import TransformerLayer import flax.linen as nn - + with autocast(enabled=True, recipe=NVFP4BlockScaling(), mesh_resource=MeshResource()): + class Model(nn.Module): - """ This module does not represent any meaningful model, it is just to cover all FFI calls. """ + """This module does not represent any meaningful model, it is just to cover all FFI calls.""" @nn.compact def __call__(self, x): @@ -1967,14 +1971,29 @@ def __call__(self, x): # Arbitrarily call softmax multiple times to cover all softmax FFI calls x = x.reshape((1, *x.shape)) x = softmax.softmax(x, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED) - mask1 = self.variable('collection', 'mask1', lambda: jax.random.bernoulli(jax.random.PRNGKey(0), shape=x.shape).astype(jnp.bfloat16)).value.astype(jnp.uint8) - x = softmax.softmax(x, mask=mask1, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_MASKED) - mask2 = self.variable('collection', 'mask2', lambda: (1.0 - jnp.tril(jnp.ones_like(x))).astype(jnp.bfloat16)).value.astype(jnp.uint8) + mask1 = self.variable( + "collection", + "mask1", + lambda: jax.random.bernoulli(jax.random.PRNGKey(0), shape=x.shape).astype( + jnp.bfloat16 + ), + ).value.astype(jnp.uint8) + x = softmax.softmax( + x, mask=mask1, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_MASKED + ) + mask2 = self.variable( + "collection", + "mask2", + lambda: (1.0 - jnp.tril(jnp.ones_like(x))).astype(jnp.bfloat16), + ).value.astype(jnp.uint8) x = x.reshape((-1, 1, 32, 32)) - x = softmax.softmax(x, mask=mask2, softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_UPPER_TRIANG_MASKED) + x = softmax.softmax( + x, + mask=mask2, + softmax_fusion_type=softmax.SoftmaxFusionType.SCALED_UPPER_TRIANG_MASKED, + ) return x - model = Model() input_shape = (1, 128, 512) x = jnp.ones(input_shape, dtype=jnp.bfloat16) @@ -1982,8 +2001,8 @@ def __call__(self, x): var_collect = model.init(jax.random.PRNGKey(0), x) def f(var_collect, x): - x = model.apply(var_collect, x, rngs={'sr_rng': jax.random.PRNGKey(0)}) - x = jnp.mean(x) # fake loss function for value_and_grad + x = model.apply(var_collect, x, rngs={"sr_rng": jax.random.PRNGKey(0)}) + x = jnp.mean(x) # fake loss function for value_and_grad return x @jax.jit @@ -1997,27 +2016,33 @@ def train_step(var_collect, x, grouped_kernel): grouped_kernel, contracting_dims=((1,), (1,)), group_sizes=jnp.array([x.shape[0]], dtype=jnp.int32), - quantizer_set=QuantizerFactory.create_set(n_groups=1, fp8_recipe=Float8CurrentScaling(), quantize_meta_set=QuantizeMetaSet(QuantizeMeta(), QuantizeMeta(), QuantizeMeta())), + quantizer_set=QuantizerFactory.create_set( + n_groups=1, + fp8_recipe=Float8CurrentScaling(), + quantize_meta_set=QuantizeMetaSet( + QuantizeMeta(), QuantizeMeta(), QuantizeMeta() + ), + ), ) loss += jnp.mean(x) return loss, grads - + grouped_kernel = jnp.zeros((1, x.shape[-1], x.shape[-1]), dtype=jnp.bfloat16) train_step(var_collect, x, grouped_kernel) def _get_hlo_text_from_file(self, hlo_name: str) -> str: - """ Reads the StableHLO text from a file given its name. """ + """Reads the StableHLO text from a file given its name.""" hlo_file_path = os.path.join(self.HLO_DIR, f"{hlo_name}.txt") - with open(hlo_file_path, 'r') as f: + with open(hlo_file_path, "r") as f: hlo_text = f.read() return hlo_text def _make_args_based_on_input_tensor_shape_and_dtype(self, stablehlo_text: str): - """ Parses the StableHLO text to extract input tensor shapes and dtypes, and creates dummy JAX arrays accordingly. """ + """Parses the StableHLO text to extract input tensor shapes and dtypes, and creates dummy JAX arrays accordingly.""" # Parse function signature to extract argument information # Pattern matches: @main(%arg0: tensor<32x32xbf16>, %arg1: tensor<64xf32>, ...) - pattern = r'@main\((.*?)\{' + pattern = r"@main\((.*?)\{" match = re.search(pattern, stablehlo_text) if not match: @@ -2027,77 +2052,81 @@ def _make_args_based_on_input_tensor_shape_and_dtype(self, stablehlo_text: str): # Parse individual arguments # Pattern matches: %arg0: tensor<32x32xbf16> - arg_pattern = r'%arg(\d+):\s*tensor<([^>]+)>' + arg_pattern = r"%arg(\d+):\s*tensor<([^>]+)>" arg_matches = re.findall(arg_pattern, args_str) parsed_args = [] for arg_num, shape_and_dtype_str in arg_matches: print(f"Parsing argument {arg_num} with shape and dtype: {shape_and_dtype_str}") # Parse shape: "32x32xbf16" -> [32, 32] - dtype_str = shape_and_dtype_str.split('x')[-1] - shape = [int(dim) for dim in shape_and_dtype_str.split('x')[:-1]] - + dtype_str = shape_and_dtype_str.split("x")[-1] + shape = [int(dim) for dim in shape_and_dtype_str.split("x")[:-1]] + # Map StableHLO dtype to JAX dtype dtype_map = { - 'bf16': jnp.bfloat16, - 'f32': jnp.float32, - 'f16': jnp.float16, - 'f8E4M3FN': jnp.float8_e4m3fn, - 'f8E5M2': jnp.float8_e5m2, - 'i32': jnp.int32, - 'ui32': jnp.uint32, + "bf16": jnp.bfloat16, + "f32": jnp.float32, + "f16": jnp.float16, + "f8E4M3FN": jnp.float8_e4m3fn, + "f8E5M2": jnp.float8_e5m2, + "i32": jnp.int32, + "ui32": jnp.uint32, } dtype = dtype_map.get(dtype_str) - - parsed_args.append(jnp.ones( - shape, - dtype=dtype - )) + + parsed_args.append(jnp.ones(shape, dtype=dtype)) return parsed_args @pytest.mark.skipif(is_fp4_supported, reason=fp4_unsupported_reason) def test_ffi_compatibility(self, ffi_hlo_name): - """ Tests that the current FFI bindings are compatible with the provided HLO and there are no API mismatches. """ + """Tests that the current FFI bindings are compatible with the provided HLO and there are no API mismatches.""" from jax.extend.backend import get_backend stablehlo_text = self._get_hlo_text_from_file(ffi_hlo_name) args = self._make_args_based_on_input_tensor_shape_and_dtype(stablehlo_text) - client = get_backend('cuda') - executable = client.compile_and_load(stablehlo_text.encode('utf-8'), executable_devices=jax.devices()[:1]) + client = get_backend("cuda") + executable = client.compile_and_load( + stablehlo_text.encode("utf-8"), executable_devices=jax.devices()[:1] + ) results = executable.execute(args) - print(results) # No need to assert anything here, just ensure it runs without error + print(results) # No need to assert anything here, just ensure it runs without error def test_all_primitive_ffi_tested(self): - """ Ensures that all our TE primitives with FFI bindings are included in the FFI HLO compatibility tests. """ + """Ensures that all our TE primitives with FFI bindings are included in the FFI HLO compatibility tests.""" # Open all HLO files and extract primitive FFI names tested_hlos = set() for file in os.listdir(self.HLO_DIR): file_path = os.path.join(self.HLO_DIR, file) - if os.path.isfile(file_path) and file.endswith('.txt'): - with open(file_path, 'r') as f: + if os.path.isfile(file_path) and file.endswith(".txt"): + with open(file_path, "r") as f: hlo_text = f.read() # Extract primitive name from HLO text - pattern = r'stablehlo.custom_call @(.+?)\(' + pattern = r"stablehlo.custom_call @(.+?)\(" matches = re.findall(pattern, hlo_text) if matches: for match in matches: primitive_name = match tested_hlos.add(primitive_name) - + # Assert that all registered primitives have corresponding FFI tests import transformer_engine_jax KNOWN_MISSING_FFI_TESTS = { # dequantize does not have a JAX primitive currently - 'te_dequantize_ffi', + "te_dequantize_ffi", # needs testing - 'te_grouped_gemm_d2h_group_sizes_ffi', + "te_grouped_gemm_d2h_group_sizes_ffi", } unmatched_primitives = set() for primitive_ffi_name, _ in transformer_engine_jax.registrations().items(): - if primitive_ffi_name not in tested_hlos and primitive_ffi_name not in KNOWN_MISSING_FFI_TESTS: + if ( + primitive_ffi_name not in tested_hlos + and primitive_ffi_name not in KNOWN_MISSING_FFI_TESTS + ): unmatched_primitives.add(primitive_ffi_name) - assert len(unmatched_primitives) == 0, f"The following primitives do not have FFI tests: {unmatched_primitives}" \ No newline at end of file + assert ( + len(unmatched_primitives) == 0 + ), f"The following primitives do not have FFI tests: {unmatched_primitives}" From 28a07c16350b657612b522284b71b0db65b03149 Mon Sep 17 00:00:00 2001 From: Jeremy Berchtold Date: Fri, 19 Dec 2025 08:46:00 -0800 Subject: [PATCH 4/5] fix test skipping Signed-off-by: Jeremy Berchtold --- tests/jax/test_custom_call_compute.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/jax/test_custom_call_compute.py b/tests/jax/test_custom_call_compute.py index b2c5775cc0d..2bc2c46ae08 100644 --- a/tests/jax/test_custom_call_compute.py +++ b/tests/jax/test_custom_call_compute.py @@ -1925,7 +1925,7 @@ def test_grouped_dense_grad_fp8(self, fwd_bwd_dtype, scaling_mode, input_shape): assert_allclose(prim_wgrad, ref_wgrad, dtype=bwd_dtype) assert_allclose(prim_dbias, ref_dbias, dtype=dtype) - +@pytest.mark.skipif(not is_fp4_supported, reason=fp4_unsupported_reason) class TestFFICompatibility: HLO_DIR = os.path.join(os.path.dirname(__file__), "ffi_hlo") @@ -2077,7 +2077,6 @@ def _make_args_based_on_input_tensor_shape_and_dtype(self, stablehlo_text: str): parsed_args.append(jnp.ones(shape, dtype=dtype)) return parsed_args - @pytest.mark.skipif(is_fp4_supported, reason=fp4_unsupported_reason) def test_ffi_compatibility(self, ffi_hlo_name): """Tests that the current FFI bindings are compatible with the provided HLO and there are no API mismatches.""" from jax.extend.backend import get_backend From 069f7511b407c0acdddbd8830ae13059725a2378 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Fri, 19 Dec 2025 16:46:58 +0000 Subject: [PATCH 5/5] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- tests/jax/test_custom_call_compute.py | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/jax/test_custom_call_compute.py b/tests/jax/test_custom_call_compute.py index 2bc2c46ae08..3c7a82544a0 100644 --- a/tests/jax/test_custom_call_compute.py +++ b/tests/jax/test_custom_call_compute.py @@ -1925,6 +1925,7 @@ def test_grouped_dense_grad_fp8(self, fwd_bwd_dtype, scaling_mode, input_shape): assert_allclose(prim_wgrad, ref_wgrad, dtype=bwd_dtype) assert_allclose(prim_dbias, ref_dbias, dtype=dtype) + @pytest.mark.skipif(not is_fp4_supported, reason=fp4_unsupported_reason) class TestFFICompatibility: