V1003 10:10:50.629000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/torch/_dynamo/convert_frame.py", 0]} V1003 10:10:50.630000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/test/inductor/test_codecache.py", 1]} V1003 10:10:50.630000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/torch/_inductor/test_case.py", 2]} V1003 10:10:50.630000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/torch/_dynamo/test_case.py", 3]} V1003 10:10:50.631000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/torch/testing/_internal/common_utils.py", 4]} V1003 10:10:50.631000 2235078 torch/_logging/structured.py:22] {"str": ["/home/oulgen/.conda/envs/py311/lib/python3.11/unittest/main.py", 5]} V1003 10:10:50.631000 2235078 torch/_logging/structured.py:22] {"str": ["/home/oulgen/.conda/envs/py311/lib/python3.11/unittest/runner.py", 6]} V1003 10:10:50.631000 2235078 torch/_logging/structured.py:22] {"str": ["/home/oulgen/.conda/envs/py311/lib/python3.11/unittest/suite.py", 7]} V1003 10:10:50.631000 2235078 torch/_logging/structured.py:22] {"str": ["/home/oulgen/.conda/envs/py311/lib/python3.11/unittest/case.py", 8]} V1003 10:10:50.632000 2235078 torch/_logging/structured.py:22] {"str": ["/home/oulgen/.conda/envs/py311/lib/python3.11/contextlib.py", 9]} V1003 10:10:50.632000 2235078 torch/_logging/structured.py:22] {"str": ["/data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py", 10]} V1003 10:10:50.632000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 395, "name": "test_flex_attention_caching", "filename": 1}, {"line": 380, "name": "fn", "filename": 1}, {"line": 1062, "name": "flex_attention", "filename": 10}, {"line": 1049, "name": "_flex_attention_hop_wrapper", "filename": 10}]}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.632000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "d442394c2c9a973ef4b64ef5deddb37f"} { "name": "_compile.compile_inner", "ts": 1727975450632606.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.633000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "686bcefec76c75096eb61274391c2867"} { "name": "entire_frame_compile", "ts": 1727975450632606.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.639000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 0, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.640000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.640000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 0, "source": "L['args'][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.835000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 0, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.836000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 4, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.836000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 4, "source": "L['args'][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.837000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 0, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.837000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 5, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.838000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 5, "source": "L['args'][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.839000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 0, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.839000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 6, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.839000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 6, "source": "L['args'][4][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.840000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 0, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.841000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.841000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 7, "source": "L['args'][4][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.842000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 0, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.842000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.843000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 8, "source": "L['args'][4][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.844000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 0, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.844000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.844000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 9, "source": "L['args'][4][3]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.845000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 0, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.846000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.846000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 10, "source": "L['args'][4][4]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.847000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 0, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.847000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.847000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 11, "source": "L['args'][4][5]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.848000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 0, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.849000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 12, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.849000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 12, "source": "L['args'][4][6]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.850000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 0, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.850000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 13, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 0}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.850000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 0, "id": 13, "source": "L['args'][4][7]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:50.864000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_args_0_": [1, 4, 512, 64], "l_args_1_": [1, 4, 512, 64], "l_args_2_": [1, 4, 512, 64], "l_args_4_0_": [1, 1, 16], "l_args_4_1_": [1, 1, 16, 16], "l_args_4_2_": [1, 1, 16], "l_args_4_3_": [1, 1, 16, 16], "l_args_4_4_": [1, 1, 16], "l_args_4_5_": [1, 1, 16, 16], "l_args_4_6_": [1, 1, 16], "l_args_4_7_": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "getitem": [1, 4, 512, 64], "getitem_1": [1, 4, 512]}}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "36de5ad6eb1efc648a27dc62c107a2ca"} class GraphModule(torch.nn.Module): def forward(self, L_args_0_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_1_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_2_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_4_0_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_1_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_2_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_3_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_4_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_5_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_6_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_7_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_args_0_ = L_args_0_ l_args_1_ = L_args_1_ l_args_2_ = L_args_2_ l_args_4_0_ = L_args_4_0_ l_args_4_1_ = L_args_4_1_ l_args_4_2_ = L_args_4_2_ l_args_4_3_ = L_args_4_3_ l_args_4_4_ = L_args_4_4_ l_args_4_5_ = L_args_4_5_ l_args_4_6_ = L_args_4_6_ l_args_4_7_ = L_args_4_7_ # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1050 in _flex_attention_hop_wrapper, code: return flex_attention_hop(*args, **kwargs) child_1: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_args_0_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_args_0_, l_args_1_, l_args_2_, score_mod_0, (l_args_4_0_, l_args_4_1_, l_args_4_2_, l_args_4_3_, l_args_4_4_, l_args_4_5_, l_args_4_6_, l_args_4_7_, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_args_0_ = l_args_1_ = l_args_2_ = score_mod_0 = l_args_4_0_ = l_args_4_1_ = l_args_4_2_ = l_args_4_3_ = l_args_4_4_ = l_args_4_5_ = l_args_4_6_ = l_args_4_7_ = mask_fn_0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0] getitem_1: "f32[1, 4, 512][2048, 512, 1]cuda:0" = flex_attention[1]; flex_attention = None return (getitem, getitem_1) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:377 in score_mod, code: return score + (q - kv) sub: "i32[][]cuda:0" = child_3 - child_4; child_3 = child_4 = None add: "f32[][]cuda:0" = child + sub; child = sub = None return add class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:10:50.865000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "78b3ca52339404a8a8b475af3393e9be"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975450865801.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.866000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "f7bc91b808bbb22050236e69d68d305e"} { "name": "backend_compile", "ts": 1727975450865801.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.866000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "4a09b2e379da345bd3263251d88c5014"} { "name": "backend_compile", "ts": 1727975450866663.5, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 0, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.867000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "53dff51299e2eac55b28bb862cc5cb64"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975450867057.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 0, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.893000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "e693a6590122b61ee5c84da6a497ef23"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['args'], accessed_by=DictGetItemGuardAccessor(args) | | +- TYPE_MATCH: ___check_type_id(L['args'], 8815232) | | +- LENGTH_CHECK: len(L['args']) == 7 | | +- GuardManager: source=L['args'][0], accessed_by=TupleGetItemGuardAccessor(0) | | | +- TENSOR_MATCH: check_tensor(L['args'][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][0], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['args'][0], L['args'][1], L['args'][2], L['args'][4][0], L['args'][4][1], L['args'][4][2], L['args'][4][3], L['args'][4][4], L['args'][4][5], L['args'][4][6], L['args'][4][7]) | | +- GuardManager: source=L['args'][1], accessed_by=TupleGetItemGuardAccessor(1) | | | +- TENSOR_MATCH: check_tensor(L['args'][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][1], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][2], accessed_by=TupleGetItemGuardAccessor(2) | | | +- TENSOR_MATCH: check_tensor(L['args'][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][2], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][3], accessed_by=TupleGetItemGuardAccessor(3) | | | +- GuardManager: source=L['args'][3].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['args'][3].__code__, 140413271879296) | | +- GuardManager: source=L['args'][4], accessed_by=TupleGetItemGuardAccessor(4) | | | +- TYPE_MATCH: ___check_type_id(L['args'][4], 8815232) | | | +- LENGTH_CHECK: len(L['args'][4]) == 11 | | | +- GuardManager: source=L['args'][4][0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][0], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][1], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][2], accessed_by=TupleGetItemGuardAccessor(2) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][2], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][3], accessed_by=TupleGetItemGuardAccessor(3) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][3], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][3], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][4], accessed_by=TupleGetItemGuardAccessor(4) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][4], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][4], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][5], accessed_by=TupleGetItemGuardAccessor(5) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][5], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][5], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][6], accessed_by=TupleGetItemGuardAccessor(6) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][6], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][6], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][7], accessed_by=TupleGetItemGuardAccessor(7) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][7], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][7], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][8], accessed_by=TupleGetItemGuardAccessor(8) | | | | +- EQUALS_MATCH: L['args'][4][8] == 128 | | | +- GuardManager: source=L['args'][4][9], accessed_by=TupleGetItemGuardAccessor(9) | | | | +- EQUALS_MATCH: L['args'][4][9] == 128 | | | +- GuardManager: source=L['args'][4][10], accessed_by=TupleGetItemGuardAccessor(10) | | | | +- GuardManager: source=L['args'][4][10].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(L['args'][4][10].__code__, 140413271880128) | | +- GuardManager: source=L['args'][5], accessed_by=TupleGetItemGuardAccessor(5) | | | +- EQUALS_MATCH: L['args'][5] == 0.125 | | +- GuardManager: source=L['args'][6], accessed_by=TupleGetItemGuardAccessor(6) | | | +- DICT_LENGTH: len(L['args'][6]) == 3 | | | +- GuardManager: source=L['args'][6]['ROWS_GUARANTEED_SAFE'], accessed_by=DictGetItemGuardAccessor(ROWS_GUARANTEED_SAFE) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['ROWS_GUARANTEED_SAFE'], 8910592) | | | +- GuardManager: source=L['args'][6]['PRESCALE_QK'], accessed_by=DictGetItemGuardAccessor(PRESCALE_QK) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['PRESCALE_QK'], 8910592) | | | +- GuardManager: source=L['args'][6]['OUTPUT_LOGSUMEXP'], accessed_by=DictGetItemGuardAccessor(OUTPUT_LOGSUMEXP) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['OUTPUT_LOGSUMEXP'], 8910592) | +- GuardManager: source=L['kwargs'], accessed_by=DictGetItemGuardAccessor(kwargs) | | +- DICT_LENGTH: not L['kwargs'] | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['flex_attention_hop'], accessed_by=DictGetItemGuardAccessor(flex_attention_hop) | | | +- TYPE_MATCH: ___check_type_id(G['flex_attention_hop'], 96992544) | | | +- GuardManager: source=G['flex_attention_hop'].__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | +- EQUALS_MATCH: G['flex_attention_hop'].__name__ == 'flex_attention' | | +- GuardManager: source=G['__builtins_dict___0'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___0) | | | +- GuardManager: source=G['__builtins_dict___0']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___0']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___0']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___0']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___0']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___0']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___0']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___0']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) V1003 10:10:50.894000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "20353735c197adbac78fed570e5a0d91"} { "name": "entire_frame_compile", "ts": 1727975450894280.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 0, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.894000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "6dc3af6842c03572276aaba0318e48d3"} { "name": "_compile.compile_inner", "ts": 1727975450894591.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 0, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:50.895000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "0/0", "frame_key": "1", "co_name": "_flex_attention_hop_wrapper", "co_filename": "/data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py", "co_firstlineno": 1049, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 55, "shape_env_guard_count": 0, "graph_op_count": 12, "graph_node_count": 26, "graph_input_count": 11, "start_time": 1727975450.6325822, "entire_frame_compile_time_s": 0.26139068603515625, "backend_compile_time_s": 0.0006465911865234375, "inductor_compile_time_s": null, "code_gen_time_s": null, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 0, "structured_logging_overhead_s": 0.019217516000000004, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.499000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 395, "name": "test_flex_attention_caching", "filename": 1}, {"line": 379, "name": "fn", "filename": 1}]}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.499000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "e10001bc2249e6200cda3bf42b80be57"} { "name": "_compile.compile_inner", "ts": 1727975453499593.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:53.499000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "ff0561bc609e54d7e5775bdf8f81ff48"} { "name": "entire_frame_compile", "ts": 1727975453499593.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:53.502000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 9, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.503000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.503000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 0, "source": "L['q']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.511000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 9, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.512000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 1, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.512000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 1, "source": "L['k']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.513000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 9, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.514000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 2, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.514000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 2, "source": "L['v']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.532000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 9, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.532000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 3, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.532000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 3, "source": "L['block_mask'].kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.676000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 9, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.677000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.677000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 7, "source": "L['block_mask'].kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.679000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 9, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.679000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.679000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 8, "source": "L['block_mask'].full_kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.680000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 9, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.681000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.681000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 9, "source": "L['block_mask'].full_kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.682000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 9, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.683000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.683000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 10, "source": "L['block_mask'].q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.684000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 9, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.685000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.685000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 11, "source": "L['block_mask'].q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.686000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 9, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.686000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 12, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.687000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 12, "source": "L['block_mask'].full_q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.688000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 9, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.688000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 13, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 9}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.689000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 9, "id": 13, "source": "L['block_mask'].full_q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:10:53.697000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_q_": [1, 4, 512, 64], "l_k_": [1, 4, 512, 64], "l_v_": [1, 4, 512, 64], "l_block_mask_kv_num_blocks": [1, 1, 16], "l_block_mask_kv_indices": [1, 1, 16, 16], "l_block_mask_full_kv_num_blocks": [1, 1, 16], "l_block_mask_full_kv_indices": [1, 1, 16, 16], "l_block_mask_q_num_blocks": [1, 1, 16], "l_block_mask_q_indices": [1, 1, 16, 16], "l_block_mask_full_q_num_blocks": [1, 1, 16], "l_block_mask_full_q_indices": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "out": [1, 4, 512, 64]}}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0365bd76b3474fb6d96e4c3c42585fb7"} class GraphModule(torch.nn.Module): def forward(self, L_q_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_k_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_v_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_block_mask_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_q_ = L_q_ l_k_ = L_k_ l_v_ = L_v_ l_block_mask_kv_num_blocks = L_block_mask_kv_num_blocks l_block_mask_kv_indices = L_block_mask_kv_indices l_block_mask_full_kv_num_blocks = L_block_mask_full_kv_num_blocks l_block_mask_full_kv_indices = L_block_mask_full_kv_indices l_block_mask_q_num_blocks = L_block_mask_q_num_blocks l_block_mask_q_indices = L_block_mask_q_indices l_block_mask_full_q_num_blocks = L_block_mask_full_q_num_blocks l_block_mask_full_q_indices = L_block_mask_full_q_indices # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( child_1: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_q_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_q_, l_k_, l_v_, score_mod_0, (l_block_mask_kv_num_blocks, l_block_mask_kv_indices, l_block_mask_full_kv_num_blocks, l_block_mask_full_kv_indices, l_block_mask_q_num_blocks, l_block_mask_q_indices, l_block_mask_full_q_num_blocks, l_block_mask_full_q_indices, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_q_ = l_k_ = l_v_ = score_mod_0 = l_block_mask_kv_num_blocks = l_block_mask_kv_indices = l_block_mask_full_kv_num_blocks = l_block_mask_full_kv_indices = l_block_mask_q_num_blocks = l_block_mask_q_indices = l_block_mask_full_q_num_blocks = l_block_mask_full_q_indices = mask_fn_0 = None out: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (out,) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:377 in score_mod, code: return score + (q - kv) sub: "i32[][]cuda:0" = child_3 - child_4; child_3 = child_4 = None add: "f32[][]cuda:0" = child + sub; child = sub = None return add class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:10:53.698000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "f3daea7b3445ed3e94cb46786b2c5d08"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975453698084.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:53.698000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "d8ed1bc2cb1632abc21c22f112ed0fec"} { "name": "backend_compile", "ts": 1727975453698084.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:53.708000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "18faf90ed7d7bb7f081083da150205b3"} { "name": "create_aot_dispatcher_function", "ts": 1727975453708777.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:53.842000 2235078 torch/_functorch/_aot_autograd/dispatch_and_compile_graph.py:215] {"aot_forward_graph": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "f6cec159cb1cee52405afcfe40b62c03"} class (torch.nn.Module): def forward(self, arg0_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg1_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg2_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg3_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg4_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg5_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg6_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg7_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg8_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg9_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg10_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (getitem,) class sdpa_score0(torch.nn.Module): def forward(self, arg0_1: "f32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0", arg4_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sub: "i32[][]cuda:0" = torch.ops.aten.sub.Tensor(arg3_1, arg4_1); arg3_1 = arg4_1 = None add: "f32[][]cuda:0" = torch.ops.aten.add.Tensor(arg0_1, sub); arg0_1 = sub = None return add class sdpa_mask0(torch.nn.Module): def forward(self, arg0_1: "i32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = torch.ops.aten.ge.Tensor(arg2_1, arg3_1); arg2_1 = arg3_1 = None return ge V1003 10:10:53.843000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "5dd81e8de7669fa4f8da5f234591e1dc"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975453842696.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:54.072000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "df64f9080a1fb9082793b2c995d4ba13"} { "name": "compile_fx_inner", "ts": 1727975454072213.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:54.072000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "6fe6927d5157dc6d0578283e86ead8a3"} { "name": "inductor_compile", "ts": 1727975454072213.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.234000 2235078 torch/_inductor/compile_fx.py:731] {"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c507028eda974d39e4ca6593c6320460"} import torch from torch import tensor, device import torch.fx as fx from torch._dynamo.testing import rand_strided from math import inf import torch._inductor.inductor_prims import torch._dynamo.config import torch._inductor.config import torch._functorch.config import torch.fx.experimental._config torch._dynamo.config.log_compilation_metrics = False torch._dynamo.config.fake_tensor_cache_crosscheck_enabled = True torch._inductor.config.fx_graph_remote_cache = False torch._inductor.config.autotune_local_cache = False torch._inductor.config.autotune_remote_cache = False torch._functorch.config.unlift_effect_tokens = True isolate_fails_code_str = None # torch version: 2.5.0a0+git7647c39 # torch cuda version: 12.0 # torch git version: 7647c398ff87daf70260854cf0a7f7993b3abc76 # CUDA Info: # nvcc: NVIDIA (R) Cuda compiler driver # Copyright (c) 2005-2023 NVIDIA Corporation # Built on Fri_Jan__6_16:45:21_PST_2023 # Cuda compilation tools, release 12.0, V12.0.140 # Build cuda_12.0.r12.0/compiler.32267302_0 # GPU Hardware Info: # NVIDIA PG509-210 : 8 from torch.nn import * class Repro(torch.nn.Module): def __init__(self) -> None: super().__init__() self.sdpa_score0 = () self.sdpa_mask0 = () def forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1): sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem = flex_attention[0]; flex_attention = None return (getitem,) def load_args(reader): buf0 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf0, (1, 4, 512, 64), is_leaf=True) # arg0_1 buf1 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf1, (1, 4, 512, 64), is_leaf=True) # arg1_1 buf2 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf2, (1, 4, 512, 64), is_leaf=True) # arg2_1 buf3 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf3, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg3_1 buf4 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf4, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg4_1 buf5 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf5, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg5_1 buf6 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf6, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg6_1 buf7 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf7, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg7_1 buf8 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf8, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg8_1 buf9 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf9, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg9_1 buf10 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf10, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg10_1 load_args._version = 0 mod = Repro() if __name__ == '__main__': from torch._dynamo.repro.after_aot import run_repro with torch.no_grad(): run_repro(mod, load_args, accuracy=False, command='run', save_dir=None, tracing_mode='real', check_str=None) # To run it separately, do # mod, args = run_repro(mod, load_args, accuracy=False, command='get_args', save_dir=None, tracing_mode='real', check_str=None) # mod(*args) V1003 10:10:56.463000 2235078 torch/_inductor/compile_fx.py:795] {"inductor_post_grad_graph": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "f6cec159cb1cee52405afcfe40b62c03"} class (torch.nn.Module): def forward(self, arg0_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg1_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg2_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg3_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg4_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg5_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg6_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg7_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg8_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg9_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg10_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (getitem,) class sdpa_score0(torch.nn.Module): def forward(self, arg0_1: "f32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0", arg4_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sub: "i32[][]cuda:0" = torch.ops.aten.sub.Tensor(arg3_1, arg4_1); arg3_1 = arg4_1 = None add: "f32[][]cuda:0" = torch.ops.aten.add.Tensor(arg0_1, sub); arg0_1 = sub = None return add class sdpa_mask0(torch.nn.Module): def forward(self, arg0_1: "i32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = torch.ops.aten.ge.Tensor(arg2_1, arg3_1); arg2_1 = arg3_1 = None return ge V1003 10:10:56.496000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "fc5eec6dd74463eae278f61fbe01c218"} { "name": "GraphLowering.run", "ts": 1727975456496145.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.814000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "ceafe5eb3d5703b5ca9729fc880cd47b"} { "name": "GraphLowering.run", "ts": 1727975456813991.2, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.814000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "5103da0c4e3f17344951cb8876a0ef08"} { "name": "GraphLowering.compile_to_module", "ts": 1727975456814764.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.815000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0cc1abc97b9fd26bc3c8bcfa0fb94399"} { "name": "code_gen", "ts": 1727975456814764.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.820000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "b680f8a68ec3cd182eef8fd76116df1f"} { "name": "Scheduler.__init__", "ts": 1727975456820242.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.826000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "3ea7d3d1a284720131038fa646c6b219"} { "name": "Scheduler.__init__", "ts": 1727975456826429.2, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.826000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "6328137516750744d43158caca6502ff"} { "name": "Scheduler.codegen", "ts": 1727975456826813.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.837000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "ec3d70ea4b3cc1330ead5e46c08a16ac"} { "name": "Scheduler.codegen", "ts": 1727975456837107.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.837000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "bdcac968c06e063fd73d86a38c67b967"} { "name": "PythonWrapperCodegen.generate", "ts": 1727975456837466.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.840000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "4f0fc9e676bebc5b3ec366bb1e137c82"} { "name": "PythonWrapperCodegen.generate", "ts": 1727975456840311.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:10:56.841000 2235078 torch/_inductor/graph.py:1899] {"inductor_output_code": {"filename": "/tmp/oulgen/tmp4z1i5ywe/kp/ckpysuucou6gm55terbvpynnfevubjpwkfm3ubzxyauw2bgggpvi.py"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "187381168c614ba1181a424c67fdb77c"} # AOT ID: ['0_inference'] from ctypes import c_void_p, c_long, c_int import torch import math import random import os import tempfile from math import inf, nan from torch._inductor.hooks import run_intermediate_hooks from torch._inductor.utils import maybe_profile from torch._inductor.codegen.memory_planning import _align as align from torch import device, empty_strided from torch._inductor.async_compile import AsyncCompile from torch._inductor.select_algorithm import extern_kernels from torch._inductor.codegen.multi_kernel import MultiKernelCall import torch._inductor.kernel.flex_attention from torch._C import _cuda_getCurrentRawStream as get_raw_stream import triton import triton.language as tl from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph aten = torch.ops.aten inductor_ops = torch.ops.inductor _quantized = torch.ops._quantized assert_size_stride = torch._C._dynamo.guards.assert_size_stride empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor alloc_from_pool = torch.ops.inductor._alloc_from_pool async_compile = AsyncCompile() # kernel path: /tmp/oulgen/tmp4z1i5ywe/ke/ckedjitfju7kpaxnpatsaz3gzkdz6znc4u5v2e22w2xzbjdvecy5.py # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] # Source node to ATen node mapping: # flex_attention => flex_attention # Graph fragment: # %flex_attention : [num_users=1] = call_function[target=torch.ops.higher_order.flex_attention](args = (%arg0_1, %arg1_1, %arg2_1, %sdpa_score0, (%arg3_1, %arg4_1, %arg5_1, %arg6_1, %arg7_1, %arg8_1, %arg9_1, %arg10_1, 128, 128, %sdpa_mask0), 0.125, {ROWS_GUARANTEED_SAFE: False, PRESCALE_QK: False, OUTPUT_LOGSUMEXP: False}, (), ()), kwargs = {}) triton_tem_fused_0 = async_compile.triton('triton_tem_fused_0', ''' import triton import triton.language as tl from triton.compiler.compiler import AttrsDescriptor from torch._inductor.runtime import triton_helpers, triton_heuristics from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties @triton_heuristics.template( num_stages=3, num_warps=4, triton_meta={'signature': {'arg_Q': '*fp32', 'arg_K': '*fp32', 'arg_V': '*fp32', 'arg_LSE': '*fp32', 'arg_KV_NUM_BLKS': '*i32', 'arg_KV_IDX': '*i32', 'arg_FULL_KV_NUM_BLKS': '*i32', 'arg_FULL_KV_IDX': '*i32', 'out_ptr0': '*fp32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108, warp_size=32), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7, 8), equal_to_1=())]}, inductor_meta={'kernel_name': 'triton_tem_fused_0', 'backend_hash': 'FB2CA426CF35F271C56C0D69873498391AC248E25890F2B631CA8B52D56952BD', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': False, 'autotune_pointwise': True, 'autotune_remote_cache': False, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}, ) @triton.jit def triton_tem_fused_0(arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0): ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 Q = arg_Q K = arg_K V = arg_V LSE = arg_LSE KV_NUM_BLKS = arg_KV_NUM_BLKS KV_IDX = arg_KV_IDX FULL_KV_NUM_BLKS = arg_FULL_KV_NUM_BLKS FULL_KV_IDX = arg_FULL_KV_IDX # Sub notation for this kernel: # # Q: Query, K: Key, V: Value # M: Number of queries, N: Number of keys/values, D: Model dimension # QK_HEAD_DIM: The dimension of the query and key embeddings # V_HEAD_DIM: The dimension of the value embeddings # z: Batch size, h: Number of heads, m: Number of queries per head, k: Number of keys per head # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups. # # The following FULL_* and PARTIAL_* is defined in the block sparse mask grid, rather than the thread block grid. # KV_NUM_BLKS: The number of KV blocks (that may or may not require masking) for each query. # KV_IDX: The indices of KV blocks (that may or may not require masking) for each query. # FULL_KV_NUM_BLKS: The number of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_KV_IDX: The indices of fully unmasked KV blocks (so we don't need masking) for each query. # # OUTPUT_LOGSUMEXP: We only need to store the logsumexp if we require grad # # (Modifiable) Performance tuning options # BLOCK_M: The thread block size across the seqlen dim of Q. # BLOCK_N: Iterate over BLOCK_N across the seqlen dim of K/V in each thread block. # The below are kernel options that can be applied for certain score_mods, # or involve a numerics vs. perf tradeoff # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base. Has # about 20% more numerical error, but slightly faster. # ROWS_GUARANTEED_SAFE: Is it guaranteed that at least one value in each row # is not masked out? If so, we can skip an extra safety check tl.static_assert(SPARSE_Q_BLOCK_SIZE >= BLOCK_M and SPARSE_Q_BLOCK_SIZE % BLOCK_M == 0) tl.static_assert(SPARSE_KV_BLOCK_SIZE >= BLOCK_N and SPARSE_KV_BLOCK_SIZE % BLOCK_N == 0) # Define strides of inputs stride_qz, stride_qh, stride_qm, stride_qk = 131072, 32768, 64, 1 stride_kz, stride_kh, stride_kn, stride_kk = 131072, 32768, 64, 1 stride_vz, stride_vh, stride_vn, stride_vk = 131072, 32768, 64, 1 ZQ = 1 HQ = 4 Q_LEN = 512 ZKV = 1 KV_LEN = 512 MATMUL_PRECISION = Q.dtype.element_ty q_start = tl.program_id(0) off_zq = tl.program_id(1) // HQ off_hq = tl.program_id(1) % HQ # We support two cases for batch dimension. a) (ZKV == ZQ) where off_zkv = off_zq. # b) (ZKV == 1 and ZQ > 1) where KV is broadcasted along the batch dimension and off_zkv=0. off_zkv = off_zq % ZKV off_hkv = off_hq // GQA_SHARED_HEADS off_g = off_hq % GQA_SHARED_HEADS q_offset = off_zq * stride_qz + off_hq * stride_qh k_offset = off_zkv * stride_kz + off_hkv * stride_kh v_offset = off_zkv * stride_vz + off_hkv * stride_vh Q = Q + q_offset K = K + k_offset V = V + v_offset SPARSE_Z = 1 SPARSE_HQ = 1 sparse_idx_z = off_zq % SPARSE_Z sparse_idx_hq = off_hq % SPARSE_HQ SPARSE_Q_MULTIPLE: tl.constexpr = (SPARSE_Q_BLOCK_SIZE // BLOCK_M) SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) stride_kv_num_blks_h = 16 stride_kv_idx_h = 256 stride_kv_idx_m = 16 # initialize pointer to m and l m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf") l_i = tl.zeros([BLOCK_M], dtype=tl.float32) acc = tl.zeros([BLOCK_M, V_HEAD_DIM], dtype=tl.float32) offs_m = q_start * BLOCK_M + tl.arange(0, BLOCK_M) # KV_IDX and KV_NUM_BLKS are always contiguous. sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq sparse_kv_num_blks_offset = sparse_hz_offset * stride_kv_num_blks_h + q_start // SPARSE_Q_MULTIPLE sparse_kv_idx_offset = sparse_hz_offset * stride_kv_idx_h + (q_start // SPARSE_Q_MULTIPLE) * stride_kv_idx_m # noqa: B950 Q_block_ptr = tl.make_block_ptr( base=Q, shape=(Q_LEN, QK_HEAD_DIM), strides=(stride_qm, stride_qk), offsets=(q_start * BLOCK_M, 0), block_shape=(BLOCK_M, QK_HEAD_DIM), order=(1, 0) ) # load q: it stays in SRAM throughout the inner loop. if IS_DIVISIBLE: q = tl.load(Q_block_ptr) else: # boundary check is not free, so we only do it when necessary. q = tl.load(Q_block_ptr, boundary_check=(0,), padding_option = "zero") # ~~~~~~~~~~~~~~ normal blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We don't know anything "special" about these blocks, so we need to apply # both score_mod and mask_mod to it kv_indices = KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=False, ) # ~~~~~~~~~~~~~~ "full" blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We know these blocks are guaranteed to be "full", so we don't need to # apply mask_mod to them - only score_mod if HAS_FULL_BLOCKS: # FULL_KV_IDX and FULL_KV_NUM_BLKS are always contiguous. kv_indices = FULL_KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # [Note] Handle fully masked out rows: # Li will be the sum(e^(-inf)) == 0.0 for masked out rows, mi will be -inf. # We set Li to 1.0 which will result in lse/out = 0.0 | after the log(li) + mi(0.0) step l_i = tl.where(l_i == 0.0, 1, l_i) acc = acc / l_i[:, None] idx_zq = tl.program_id(1) // HQ idx_hq = tl.program_id(1) % HQ idx_m = offs_m[:, None] idx_d = tl.arange(0, V_HEAD_DIM)[None, :] mask = idx_m < Q_LEN # TODO generalize and add proper mask support xindex = idx_d + (64*idx_m) + (32768*idx_hq) + (131072*idx_zq) tl.store(out_ptr0 + (tl.broadcast_to(idx_d + (64*idx_m) + (32768*idx_hq), acc.shape)), acc, mask) # TODO dont want to write this if we dont require grad if OUTPUT_LOGSUMEXP: off_hz = tl.program_id(1) l_ptrs = LSE + off_hz * Q_LEN + offs_m lse = m_i + tl.math.log2(l_i) if IS_DIVISIBLE: tl.store(l_ptrs, lse) else: tl.store(l_ptrs, lse, mask=offs_m < Q_LEN) @triton.jit def forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets used as inputs to score_mod & mask_mod # of size [BLOCK_M, BLOCK_N] or scalar. off_z, off_h, offs_m, offs_n, # blocksparse data kv_indices, kv_num_blocks, # start kv and end kv block block_n_start, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) RCP_LN2: tl.constexpr = 1.44269504 if PRESCALE_QK: q = (q * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) # loop over k, v and update accumulator until block_n_end for start_n in range(block_n_start, block_n_end): if IS_DIVISIBLE: acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) else: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen, # it's on par or slightly faster than only applying to the last block in fwd. # However, we choose different strategy for bwd, where we only apply mod & mask # to the last block because it's faster a lot. acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) # update pointers offset = get_offset_for_next_block( start_n, kv_indices, kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N ) V_block_ptr = tl.advance(V_block_ptr, (offset, 0)) K_block_ptr = tl.advance(K_block_ptr, (0, offset)) offs_n = offs_n + offset return acc, l_i, m_i @triton.jit def get_offset_for_next_block(loop_iter, col_indices, total_blocks, SPARSE_BLOCK, SPARSE_BLOCK_MULTIPLE, BLOCK): cur_block_idx = loop_iter // SPARSE_BLOCK_MULTIPLE cur_block = tl.load(col_indices + cur_block_idx, eviction_policy="evict_last") next_block = tl.load(col_indices + cur_block_idx + 1, eviction_policy="evict_last", mask=cur_block_idx + 1 < total_blocks) needs_jump = (loop_iter + 1) % SPARSE_BLOCK_MULTIPLE == 0 jump_to_block = (next_block - cur_block ) * SPARSE_BLOCK - (SPARSE_BLOCK_MULTIPLE - 1) * BLOCK offset = jump_to_block * needs_jump + (1 - needs_jump) * BLOCK return offset @triton.jit def forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 # -- load k -- if IS_DIVISIBLE: k = tl.load(K_block_ptr) else: k = tl.load(K_block_ptr, boundary_check=(1,), padding_option = "zero") # -- compute qk --- qk = tl.dot(q, k, input_precision=FLOAT32_PRECISION) # TODO: use cuda matmul when q_len <= 2. if not PRESCALE_QK: qk *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ if CHECK_BLOCK_BOUNDARY: # If this is the last block of a non divisible seqlen, we still need to load [BLOCK_M, BLOCK_N] elements, # which is larger than the actual number of elements. To avoid access memory out of bound, # we need to mask out the elements that are out of Q_LEN & KV_LEN. m = offs_m % Q_LEN n = offs_n % KV_LEN else: m = offs_m n = offs_n tmp0 = (m) - (n) tmp1 = tmp0.to(tl.float32) tmp2 = (qk) + tmp1 post_mod_scores = tmp2 if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: tmp3 = (m) >= (n) mask_mod_output = tmp3 if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n < KV_LEN, mask_mod_output, float("-inf")) # apply mask for partially unmasked blocks post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) # TODO: In the case that score_mod is linear, this can be LICMed if not PRESCALE_QK: post_mod_scores *= RCP_LN2 # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # -- compute scaling constant --- m_ij = tl.maximum(m_i, tl.max(post_mod_scores, 1)) if not ROWS_GUARANTEED_SAFE: masked_out_rows = (m_ij == float("-inf")) m_ij_masked = tl.where(masked_out_rows, 0, m_ij) else: m_ij_masked = m_ij alpha = tl.math.exp2(m_i - m_ij_masked) p = tl.math.exp2(post_mod_scores - m_ij_masked[:, None]) # NB: l_i update is pulled up here since it's a bit faster # NB: For headdim=256, it's faster to move it back down to after m_i = # m_ij l_i = l_i * alpha + tl.sum(p, 1) # # -- scale and update acc -- acc = acc * alpha[:, None] if IS_DIVISIBLE: v = tl.load(V_block_ptr) else: v = tl.load(V_block_ptr, boundary_check=(0,), padding_option = "zero") acc = tl.dot(p.to(MATMUL_PRECISION), v, acc, input_precision=FLOAT32_PRECISION) # -- update m_i m_i = m_ij return acc, l_i, m_i ''', device_str='cuda') meta0 = {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False, 'FLOAT32_PRECISION': "'ieee'", 'IS_DIVISIBLE': True, 'SM_SCALE': 0.125, 'GQA_SHARED_HEADS': 1, 'HAS_FULL_BLOCKS': True, 'QK_HEAD_DIM': 64, 'V_HEAD_DIM': 64, 'BLOCK_M': 128, 'BLOCK_N': 32, 'SPARSE_Q_BLOCK_SIZE': 128, 'SPARSE_KV_BLOCK_SIZE': 128} async_compile.wait(globals()) del async_compile def call(args): arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1 = args args.clear() assert_size_stride(arg0_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg1_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg2_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg3_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg4_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg5_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg6_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg7_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg8_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg9_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg10_1, (1, 1, 16, 16), (256, 256, 16, 1)) buf0 = empty_strided_cuda((1, 4, 512), (2048, 512, 1), torch.float32) with torch.cuda._DeviceGuard(0): torch.cuda.set_device(0) buf1 = empty_strided_cuda((1, 4, 512, 64), (131072, 32768, 64, 1), torch.float32) # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] stream0 = get_raw_stream(0) triton_tem_fused_0.run(arg0_1, arg1_1, arg2_1, buf0, arg3_1, arg4_1, arg5_1, arg6_1, buf1, grid=torch._inductor.kernel.flex_attention.flex_attention_grid(1, 4, 512, 64, meta0), stream=stream0) del arg0_1 del arg1_1 del arg2_1 del arg3_1 del arg4_1 del arg5_1 del arg6_1 del buf0 return (buf1, ) def benchmark_compiled_module(times=10, repeat=10): from torch._dynamo.testing import rand_strided from torch._inductor.utils import print_performance arg0_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg1_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg2_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg3_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg4_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg5_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg6_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg7_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg8_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg9_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg10_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) fn = lambda: call([arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1]) return print_performance(fn, times=times, repeat=repeat) if __name__ == "__main__": from torch._inductor.wrapper_benchmark import compiled_module_main compiled_module_main('None', benchmark_compiled_module) V1003 10:11:02.564000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "a51341d292b7eede808d8a68c1fd524e"} { "name": "code_gen", "ts": 1727975462564030.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.564000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0723c01321e1ba5c17eafbd9000d9e2d"} { "name": "GraphLowering.compile_to_module", "ts": 1727975462564643.5, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.603000 2235078 torch/_dynamo/utils.py:1020] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "489dd15ea036ef0a9a0677625f691d7c"} { "name": "fx_graph_cache_miss", "ts": 1727975456229537.0, "args": { "key": "f4lkea5y7lzhlshohvr3aqpd7bchdflfs7j5wn7mrurponawoutk", "components": [ "[n7x23yy6fih6vdcjzlzbhy3d6vx3ilu7ylp3zz6wkultu4yvnzn] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False" ], "cache_event_time": 1727975456229537044, "cache_state": "miss", "time_taken_ns": 6372999079 }, "ph": "i", "cat": "dynamo_timed", "tid": 0, "pid": 0, "s": "p" } V1003 10:11:02.604000 2235078 torch/_inductor/codecache.py:1463] {"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "785133baf5b9f8dad73506774a2a34dc"} {"key": "f4lkea5y7lzhlshohvr3aqpd7bchdflfs7j5wn7mrurponawoutk", "components": ["[n7x23yy6fih6vdcjzlzbhy3d6vx3ilu7ylp3zz6wkultu4yvnzn] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False"], "cache_event_time": 1727975456229537044, "cache_state": "miss", "time_taken_ns": 6372999079} V1003 10:11:02.605000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "20079c82f956785a2eeb02be390ae599"} { "name": "inductor_compile", "ts": 1727975462605207.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.605000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "194073b2ec42e1f2dbb4aff2a005e416"} { "name": "compile_fx_inner", "ts": 1727975462605545.2, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.606000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "ae70f86dc28b7efeea5cae51d87a7327"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975462606301.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.609000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "7d733902132e302c6ed3910c31586fe0"} { "name": "create_aot_dispatcher_function", "ts": 1727975462609845.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.610000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c2cdcd6e0e2a687934149efb13d13ca8"} { "name": "backend_compile", "ts": 1727975462610479.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.610000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "d20e4a17a90e3e33d30bd927bed546fc"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975462610762.5, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.642000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "35a49d1f2da1768efdba707d805b4b97"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['k'], accessed_by=DictGetItemGuardAccessor(k) | | +- TYPE_MATCH: ___check_type_id(L['k'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['k'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['k'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['k'], L['q'], L['v'], L['block_mask'].q_indices, L['block_mask'].kv_indices, L['block_mask'].q_num_blocks, L['block_mask'].kv_num_blocks, L['block_mask'].full_q_indices, L['block_mask'].full_kv_indices, L['block_mask'].full_q_num_blocks, L['block_mask'].full_kv_num_blocks) | +- GuardManager: source=L['q'], accessed_by=DictGetItemGuardAccessor(q) | | +- TYPE_MATCH: ___check_type_id(L['q'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['q'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['q'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['v'], accessed_by=DictGetItemGuardAccessor(v) | | +- TYPE_MATCH: ___check_type_id(L['v'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['v'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['v'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['score_mod'], accessed_by=DictGetItemGuardAccessor(score_mod) | | +- GuardManager: source=L['score_mod'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['score_mod'].__code__, 140413271879296) | +- GuardManager: source=L['block_mask'], accessed_by=DictGetItemGuardAccessor(block_mask) | | +- TYPE_MATCH: ___check_type_id(L['block_mask'], 387600320) | | +- GuardManager: source=L['block_mask'].mask_mod, accessed_by=GetAttrGuardAccessor(mask_mod) | | | +- GuardManager: source=L['block_mask'].mask_mod.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].mask_mod.__code__, 140413271880128) | | +- GuardManager: source=L['block_mask'].q_indices, accessed_by=GetAttrGuardAccessor(q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE, accessed_by=GetAttrGuardAccessor(BLOCK_SIZE) | | | +- TYPE_MATCH: ___check_type_id(L['block_mask'].BLOCK_SIZE, 8815232) | | | +- LENGTH_CHECK: len(L['block_mask'].BLOCK_SIZE) == 2 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[0] == 128 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[1] == 128 | | +- GuardManager: source=L['block_mask'].kv_indices, accessed_by=GetAttrGuardAccessor(kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].q_num_blocks, accessed_by=GetAttrGuardAccessor(q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].kv_num_blocks, accessed_by=GetAttrGuardAccessor(kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_indices, accessed_by=GetAttrGuardAccessor(full_q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_indices, accessed_by=GetAttrGuardAccessor(full_kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_num_blocks, accessed_by=GetAttrGuardAccessor(full_q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_num_blocks, accessed_by=GetAttrGuardAccessor(full_kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=GetAttrGuardAccessor(as_tuple) | | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=FuncDefaultsGuardAccessor | | | | +- GuardManager: source=L['block_mask'].as_tuple.__defaults__[0], accessed_by=GetItemGuardAccessor(0) | | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].as_tuple.__defaults__[0], 8911040) | +- GuardManager: source=L['flex_attention'], accessed_by=DictGetItemGuardAccessor(flex_attention) | | +- GuardManager: source=L['flex_attention'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__code__, 387082992) | | +- GuardManager: source=L['flex_attention'], accessed_by=FuncDefaultsGuardAccessor | | | +- GuardManager: source=L['flex_attention'].__defaults__[2], accessed_by=GetItemGuardAccessor(2) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[2], 8825760) | | | +- GuardManager: source=L['flex_attention'].__defaults__[3], accessed_by=GetItemGuardAccessor(3) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[3], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[4], accessed_by=GetItemGuardAccessor(4) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[4], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[5], accessed_by=GetItemGuardAccessor(5) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[5], 8825760) | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['__builtins_dict___2'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___2) | | | +- GuardManager: source=G['__builtins_dict___2']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___2']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___2']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___2']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___2']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___2']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___2']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___2']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_comptime'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_comptime) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_comptime'], 140410226912176) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_decorators) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'], 140410226910096) | | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, accessed_by=GetAttrGuardAccessor(is_compiling) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, 140410376252096) | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot__utils) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot__utils'], 140409673896784) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, accessed_by=GetAttrGuardAccessor(_SUPPORTED_HEAD_DIMS) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, 8844320) | | | | +- LENGTH_CHECK: len(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS) == 10 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0], accessed_by=ListGetItemGuardAccessor(0) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0] == 2 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1], accessed_by=ListGetItemGuardAccessor(1) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1] == 4 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2], accessed_by=ListGetItemGuardAccessor(2) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2] == 8 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3], accessed_by=ListGetItemGuardAccessor(3) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3] == 16 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4], accessed_by=ListGetItemGuardAccessor(4) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4] == 32 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5], accessed_by=ListGetItemGuardAccessor(5) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5] == 64 | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot_flex_attention) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], 140409673895824) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, accessed_by=GetAttrGuardAccessor(math) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, 140413266939392) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, accessed_by=GetAttrGuardAccessor(sqrt) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, 140413266943072) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, accessed_by=GetAttrGuardAccessor(torch) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, 140413267918368) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, accessed_by=GetAttrGuardAccessor(_dynamo) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, 140413260098400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static, accessed_by=GetAttrGuardAccessor(mark_static) | | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, 123166432) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, accessed_by=GetAttrGuardAccessor(compiler) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, 140410826010400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, accessed_by=GetAttrGuardAccessor(is_dynamo_compiling) | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, 140410826132992) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device, accessed_by=GetAttrGuardAccessor(_validate_device) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, 140409673611088) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, accessed_by=GetAttrGuardAccessor(flex_attention_hop) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, 96992544) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__ == 'flex_attention' | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim, accessed_by=GetAttrGuardAccessor(_supported_head_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, 140409673231376) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim, accessed_by=GetAttrGuardAccessor(_validate_embed_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, 388086512) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input, accessed_by=GetAttrGuardAccessor(_validate_sdpa_input) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, 387915104) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options, accessed_by=GetAttrGuardAccessor(_apply_kernel_options) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, 140409683680752) V1003 10:11:02.643000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "3937c61e6c98ab6b52d8e00fc2327f9d"} { "name": "entire_frame_compile", "ts": 1727975462642948.2, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.643000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "af7756663bbea703738c5d774f3ca273"} { "name": "_compile.compile_inner", "ts": 1727975462643261.5, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.643000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "1/0", "frame_key": "2", "co_name": "fn", "co_filename": "/data/users/oulgen/pytorch/test/inductor/test_codecache.py", "co_firstlineno": 379, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 80, "shape_env_guard_count": 0, "graph_op_count": 11, "graph_node_count": 25, "graph_input_count": 11, "start_time": 1727975453.4995801, "entire_frame_compile_time_s": 9.143243551254272, "backend_compile_time_s": 8.912311315536499, "inductor_compile_time_s": 8.53272271156311, "code_gen_time_s": 5.749065160751343, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 0, "structured_logging_overhead_s": 0.034261202, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.649000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 405, "name": "test_flex_attention_caching", "filename": 1}, {"line": 380, "name": "fn", "filename": 1}, {"line": 1062, "name": "flex_attention", "filename": 10}, {"line": 1049, "name": "_flex_attention_hop_wrapper", "filename": 10}]}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.649000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "a2ef97fb78450a26e1e87b56945ab331"} { "name": "_compile.compile_inner", "ts": 1727975462649375.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.649000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "3a33794e999255e5aaa28e163fcdc260"} { "name": "entire_frame_compile", "ts": 1727975462649375.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.652000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 240, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.652000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.653000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 0, "source": "L['args'][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.768000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 240, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.769000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 4, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.769000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 4, "source": "L['args'][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.770000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 240, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.770000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 5, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.771000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 5, "source": "L['args'][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.772000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 240, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.772000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 6, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.772000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 6, "source": "L['args'][4][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.773000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 240, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.774000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.774000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 7, "source": "L['args'][4][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.775000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 240, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.775000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.775000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 8, "source": "L['args'][4][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.776000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 240, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.777000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.777000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 9, "source": "L['args'][4][3]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.778000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 240, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.778000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.778000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 10, "source": "L['args'][4][4]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.779000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 240, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.780000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.780000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 11, "source": "L['args'][4][5]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.781000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 240, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.781000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 12, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.781000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 12, "source": "L['args'][4][6]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.782000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 240, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.783000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 13, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 240}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.783000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 240, "id": 13, "source": "L['args'][4][7]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.791000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_args_0_": [1, 4, 512, 64], "l_args_1_": [1, 4, 512, 64], "l_args_2_": [1, 4, 512, 64], "l_args_4_0_": [1, 1, 16], "l_args_4_1_": [1, 1, 16, 16], "l_args_4_2_": [1, 1, 16], "l_args_4_3_": [1, 1, 16, 16], "l_args_4_4_": [1, 1, 16], "l_args_4_5_": [1, 1, 16, 16], "l_args_4_6_": [1, 1, 16], "l_args_4_7_": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "getitem": [1, 4, 512, 64], "getitem_1": [1, 4, 512]}}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "36de5ad6eb1efc648a27dc62c107a2ca"} class GraphModule(torch.nn.Module): def forward(self, L_args_0_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_1_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_2_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_4_0_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_1_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_2_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_3_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_4_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_5_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_6_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_7_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_args_0_ = L_args_0_ l_args_1_ = L_args_1_ l_args_2_ = L_args_2_ l_args_4_0_ = L_args_4_0_ l_args_4_1_ = L_args_4_1_ l_args_4_2_ = L_args_4_2_ l_args_4_3_ = L_args_4_3_ l_args_4_4_ = L_args_4_4_ l_args_4_5_ = L_args_4_5_ l_args_4_6_ = L_args_4_6_ l_args_4_7_ = L_args_4_7_ # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1050 in _flex_attention_hop_wrapper, code: return flex_attention_hop(*args, **kwargs) child_1: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_args_0_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_args_0_, l_args_1_, l_args_2_, score_mod_0, (l_args_4_0_, l_args_4_1_, l_args_4_2_, l_args_4_3_, l_args_4_4_, l_args_4_5_, l_args_4_6_, l_args_4_7_, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_args_0_ = l_args_1_ = l_args_2_ = score_mod_0 = l_args_4_0_ = l_args_4_1_ = l_args_4_2_ = l_args_4_3_ = l_args_4_4_ = l_args_4_5_ = l_args_4_6_ = l_args_4_7_ = mask_fn_0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0] getitem_1: "f32[1, 4, 512][2048, 512, 1]cuda:0" = flex_attention[1]; flex_attention = None return (getitem, getitem_1) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:377 in score_mod, code: return score + (q - kv) sub: "i32[][]cuda:0" = child_3 - child_4; child_3 = child_4 = None add: "f32[][]cuda:0" = child + sub; child = sub = None return add class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:11:02.792000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "bbe4f56afd5b46c1bcfb8b89b3682d1c"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975462792636.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.793000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "42f674bdb6c8ba79dc796eb9e177d39e"} { "name": "backend_compile", "ts": 1727975462792636.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.793000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "70d8c05c0dc42c6de392be3605cefc5d"} { "name": "backend_compile", "ts": 1727975462793358.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.793000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "51a3e0876977448806622fe8baf64561"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975462793614.8, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.818000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "2b6b7ba9e9bb8df50f6d571c50ac67a8"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['args'], accessed_by=DictGetItemGuardAccessor(args) | | +- TYPE_MATCH: ___check_type_id(L['args'], 8815232) | | +- LENGTH_CHECK: len(L['args']) == 7 | | +- GuardManager: source=L['args'][0], accessed_by=TupleGetItemGuardAccessor(0) | | | +- TENSOR_MATCH: check_tensor(L['args'][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][0], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['args'][0], L['args'][1], L['args'][2], L['args'][4][0], L['args'][4][1], L['args'][4][2], L['args'][4][3], L['args'][4][4], L['args'][4][5], L['args'][4][6], L['args'][4][7]) | | +- GuardManager: source=L['args'][1], accessed_by=TupleGetItemGuardAccessor(1) | | | +- TENSOR_MATCH: check_tensor(L['args'][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][1], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][2], accessed_by=TupleGetItemGuardAccessor(2) | | | +- TENSOR_MATCH: check_tensor(L['args'][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][2], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][3], accessed_by=TupleGetItemGuardAccessor(3) | | | +- GuardManager: source=L['args'][3].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['args'][3].__code__, 140413271879296) | | +- GuardManager: source=L['args'][4], accessed_by=TupleGetItemGuardAccessor(4) | | | +- TYPE_MATCH: ___check_type_id(L['args'][4], 8815232) | | | +- LENGTH_CHECK: len(L['args'][4]) == 11 | | | +- GuardManager: source=L['args'][4][0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][0], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][1], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][2], accessed_by=TupleGetItemGuardAccessor(2) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][2], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][3], accessed_by=TupleGetItemGuardAccessor(3) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][3], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][3], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][4], accessed_by=TupleGetItemGuardAccessor(4) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][4], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][4], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][5], accessed_by=TupleGetItemGuardAccessor(5) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][5], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][5], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][6], accessed_by=TupleGetItemGuardAccessor(6) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][6], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][6], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][7], accessed_by=TupleGetItemGuardAccessor(7) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][7], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][7], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][8], accessed_by=TupleGetItemGuardAccessor(8) | | | | +- EQUALS_MATCH: L['args'][4][8] == 128 | | | +- GuardManager: source=L['args'][4][9], accessed_by=TupleGetItemGuardAccessor(9) | | | | +- EQUALS_MATCH: L['args'][4][9] == 128 | | | +- GuardManager: source=L['args'][4][10], accessed_by=TupleGetItemGuardAccessor(10) | | | | +- GuardManager: source=L['args'][4][10].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(L['args'][4][10].__code__, 140413271880128) | | +- GuardManager: source=L['args'][5], accessed_by=TupleGetItemGuardAccessor(5) | | | +- EQUALS_MATCH: L['args'][5] == 0.125 | | +- GuardManager: source=L['args'][6], accessed_by=TupleGetItemGuardAccessor(6) | | | +- DICT_LENGTH: len(L['args'][6]) == 3 | | | +- GuardManager: source=L['args'][6]['ROWS_GUARANTEED_SAFE'], accessed_by=DictGetItemGuardAccessor(ROWS_GUARANTEED_SAFE) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['ROWS_GUARANTEED_SAFE'], 8910592) | | | +- GuardManager: source=L['args'][6]['PRESCALE_QK'], accessed_by=DictGetItemGuardAccessor(PRESCALE_QK) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['PRESCALE_QK'], 8910592) | | | +- GuardManager: source=L['args'][6]['OUTPUT_LOGSUMEXP'], accessed_by=DictGetItemGuardAccessor(OUTPUT_LOGSUMEXP) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['OUTPUT_LOGSUMEXP'], 8910592) | +- GuardManager: source=L['kwargs'], accessed_by=DictGetItemGuardAccessor(kwargs) | | +- DICT_LENGTH: not L['kwargs'] | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['flex_attention_hop'], accessed_by=DictGetItemGuardAccessor(flex_attention_hop) | | | +- TYPE_MATCH: ___check_type_id(G['flex_attention_hop'], 96992544) | | | +- GuardManager: source=G['flex_attention_hop'].__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | +- EQUALS_MATCH: G['flex_attention_hop'].__name__ == 'flex_attention' | | +- GuardManager: source=G['__builtins_dict___4'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___4) | | | +- GuardManager: source=G['__builtins_dict___4']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___4']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___4']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___4']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___4']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___4']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___4']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___4']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) V1003 10:11:02.818000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "0a198eacfa70f28e771164e9f21c0377"} { "name": "entire_frame_compile", "ts": 1727975462818771.2, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.819000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "fe22d36e59a389d50249ec2f0d81445a"} { "name": "_compile.compile_inner", "ts": 1727975462819091.0, "args": { "cache_stats": { "fxgraph_cache_hit": 0, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.819000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "0/0", "frame_key": "1", "co_name": "_flex_attention_hop_wrapper", "co_filename": "/data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py", "co_firstlineno": 1049, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 55, "shape_env_guard_count": 0, "graph_op_count": 12, "graph_node_count": 26, "graph_input_count": 11, "start_time": 1727975462.6493652, "entire_frame_compile_time_s": 0.16928768157958984, "backend_compile_time_s": 0.0006258487701416016, "inductor_compile_time_s": null, "code_gen_time_s": null, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 0, "structured_logging_overhead_s": 0.032939939, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.823000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 405, "name": "test_flex_attention_caching", "filename": 1}, {"line": 379, "name": "fn", "filename": 1}]}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.824000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "5c328b3ed733b4526b0979bdb2929653"} { "name": "_compile.compile_inner", "ts": 1727975462824011.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.824000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "7333e4a8a9e70773a0188ced97e391ef"} { "name": "entire_frame_compile", "ts": 1727975462824011.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:02.826000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 256, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.827000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.827000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 0, "source": "L['q']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.836000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 256, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.836000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 1, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.836000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 1, "source": "L['k']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.838000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 256, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.838000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 2, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.838000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 2, "source": "L['v']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.847000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 256, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.847000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 3, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.847000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 3, "source": "L['block_mask'].kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.985000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 256, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.985000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.985000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 7, "source": "L['block_mask'].kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.987000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 256, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.987000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.987000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 8, "source": "L['block_mask'].full_kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.989000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 256, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.989000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.989000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 9, "source": "L['block_mask'].full_kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.991000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 256, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.991000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.991000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 10, "source": "L['block_mask'].q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.992000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 256, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.993000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.993000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 11, "source": "L['block_mask'].q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.994000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 256, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.995000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 12, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.995000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 12, "source": "L['block_mask'].full_q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.996000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 256, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.997000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 13, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 256}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:02.997000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 256, "id": 13, "source": "L['block_mask'].full_q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.005000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_q_": [1, 4, 512, 64], "l_k_": [1, 4, 512, 64], "l_v_": [1, 4, 512, 64], "l_block_mask_kv_num_blocks": [1, 1, 16], "l_block_mask_kv_indices": [1, 1, 16, 16], "l_block_mask_full_kv_num_blocks": [1, 1, 16], "l_block_mask_full_kv_indices": [1, 1, 16, 16], "l_block_mask_q_num_blocks": [1, 1, 16], "l_block_mask_q_indices": [1, 1, 16, 16], "l_block_mask_full_q_num_blocks": [1, 1, 16], "l_block_mask_full_q_indices": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "out": [1, 4, 512, 64]}}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0365bd76b3474fb6d96e4c3c42585fb7"} class GraphModule(torch.nn.Module): def forward(self, L_q_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_k_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_v_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_block_mask_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_q_ = L_q_ l_k_ = L_k_ l_v_ = L_v_ l_block_mask_kv_num_blocks = L_block_mask_kv_num_blocks l_block_mask_kv_indices = L_block_mask_kv_indices l_block_mask_full_kv_num_blocks = L_block_mask_full_kv_num_blocks l_block_mask_full_kv_indices = L_block_mask_full_kv_indices l_block_mask_q_num_blocks = L_block_mask_q_num_blocks l_block_mask_q_indices = L_block_mask_q_indices l_block_mask_full_q_num_blocks = L_block_mask_full_q_num_blocks l_block_mask_full_q_indices = L_block_mask_full_q_indices # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( child_1: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_q_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_q_, l_k_, l_v_, score_mod_0, (l_block_mask_kv_num_blocks, l_block_mask_kv_indices, l_block_mask_full_kv_num_blocks, l_block_mask_full_kv_indices, l_block_mask_q_num_blocks, l_block_mask_q_indices, l_block_mask_full_q_num_blocks, l_block_mask_full_q_indices, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_q_ = l_k_ = l_v_ = score_mod_0 = l_block_mask_kv_num_blocks = l_block_mask_kv_indices = l_block_mask_full_kv_num_blocks = l_block_mask_full_kv_indices = l_block_mask_q_num_blocks = l_block_mask_q_indices = l_block_mask_full_q_num_blocks = l_block_mask_full_q_indices = mask_fn_0 = None out: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (out,) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:377 in score_mod, code: return score + (q - kv) sub: "i32[][]cuda:0" = child_3 - child_4; child_3 = child_4 = None add: "f32[][]cuda:0" = child + sub; child = sub = None return add class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:11:03.006000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "d0fdf0bbe0d39c70da0305c36b970c1c"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975463006504.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.006000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "2c12bb411b8d14ec0cf2ffdf2c4427c3"} { "name": "backend_compile", "ts": 1727975463006504.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.012000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "17a34747abc18cf69e8e2a0141a3ab09"} { "name": "create_aot_dispatcher_function", "ts": 1727975463012509.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.138000 2235078 torch/_functorch/_aot_autograd/dispatch_and_compile_graph.py:215] {"aot_forward_graph": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "f6cec159cb1cee52405afcfe40b62c03"} class (torch.nn.Module): def forward(self, arg0_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg1_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg2_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg3_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg4_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg5_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg6_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg7_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg8_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg9_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg10_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (getitem,) class sdpa_score0(torch.nn.Module): def forward(self, arg0_1: "f32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0", arg4_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sub: "i32[][]cuda:0" = torch.ops.aten.sub.Tensor(arg3_1, arg4_1); arg3_1 = arg4_1 = None add: "f32[][]cuda:0" = torch.ops.aten.add.Tensor(arg0_1, sub); arg0_1 = sub = None return add class sdpa_mask0(torch.nn.Module): def forward(self, arg0_1: "i32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = torch.ops.aten.ge.Tensor(arg2_1, arg3_1); arg2_1 = arg3_1 = None return ge V1003 10:11:03.139000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "96a6ba3562cc8fdc9267e57d6099f6dd"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975463139343.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.141000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "4405e9a49d13d3764e0174322946bd85"} { "name": "compile_fx_inner", "ts": 1727975463141246.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.141000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "98ce9c89d4855d24e864b2a507237295"} { "name": "inductor_compile", "ts": 1727975463141246.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.282000 2235078 torch/_inductor/codecache.py:1138] {"inductor_output_code": {"filename": "/tmp/oulgen/tmp4z1i5ywe/kp/ckpysuucou6gm55terbvpynnfevubjpwkfm3ubzxyauw2bgggpvi.py"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "187381168c614ba1181a424c67fdb77c"} # AOT ID: ['0_inference'] from ctypes import c_void_p, c_long, c_int import torch import math import random import os import tempfile from math import inf, nan from torch._inductor.hooks import run_intermediate_hooks from torch._inductor.utils import maybe_profile from torch._inductor.codegen.memory_planning import _align as align from torch import device, empty_strided from torch._inductor.async_compile import AsyncCompile from torch._inductor.select_algorithm import extern_kernels from torch._inductor.codegen.multi_kernel import MultiKernelCall import torch._inductor.kernel.flex_attention from torch._C import _cuda_getCurrentRawStream as get_raw_stream import triton import triton.language as tl from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph aten = torch.ops.aten inductor_ops = torch.ops.inductor _quantized = torch.ops._quantized assert_size_stride = torch._C._dynamo.guards.assert_size_stride empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor alloc_from_pool = torch.ops.inductor._alloc_from_pool async_compile = AsyncCompile() # kernel path: /tmp/oulgen/tmp4z1i5ywe/ke/ckedjitfju7kpaxnpatsaz3gzkdz6znc4u5v2e22w2xzbjdvecy5.py # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] # Source node to ATen node mapping: # flex_attention => flex_attention # Graph fragment: # %flex_attention : [num_users=1] = call_function[target=torch.ops.higher_order.flex_attention](args = (%arg0_1, %arg1_1, %arg2_1, %sdpa_score0, (%arg3_1, %arg4_1, %arg5_1, %arg6_1, %arg7_1, %arg8_1, %arg9_1, %arg10_1, 128, 128, %sdpa_mask0), 0.125, {ROWS_GUARANTEED_SAFE: False, PRESCALE_QK: False, OUTPUT_LOGSUMEXP: False}, (), ()), kwargs = {}) triton_tem_fused_0 = async_compile.triton('triton_tem_fused_0', ''' import triton import triton.language as tl from triton.compiler.compiler import AttrsDescriptor from torch._inductor.runtime import triton_helpers, triton_heuristics from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties @triton_heuristics.template( num_stages=3, num_warps=4, triton_meta={'signature': {'arg_Q': '*fp32', 'arg_K': '*fp32', 'arg_V': '*fp32', 'arg_LSE': '*fp32', 'arg_KV_NUM_BLKS': '*i32', 'arg_KV_IDX': '*i32', 'arg_FULL_KV_NUM_BLKS': '*i32', 'arg_FULL_KV_IDX': '*i32', 'out_ptr0': '*fp32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108, warp_size=32), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7, 8), equal_to_1=())]}, inductor_meta={'kernel_name': 'triton_tem_fused_0', 'backend_hash': 'FB2CA426CF35F271C56C0D69873498391AC248E25890F2B631CA8B52D56952BD', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': False, 'autotune_pointwise': True, 'autotune_remote_cache': False, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}, ) @triton.jit def triton_tem_fused_0(arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0): ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 Q = arg_Q K = arg_K V = arg_V LSE = arg_LSE KV_NUM_BLKS = arg_KV_NUM_BLKS KV_IDX = arg_KV_IDX FULL_KV_NUM_BLKS = arg_FULL_KV_NUM_BLKS FULL_KV_IDX = arg_FULL_KV_IDX # Sub notation for this kernel: # # Q: Query, K: Key, V: Value # M: Number of queries, N: Number of keys/values, D: Model dimension # QK_HEAD_DIM: The dimension of the query and key embeddings # V_HEAD_DIM: The dimension of the value embeddings # z: Batch size, h: Number of heads, m: Number of queries per head, k: Number of keys per head # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups. # # The following FULL_* and PARTIAL_* is defined in the block sparse mask grid, rather than the thread block grid. # KV_NUM_BLKS: The number of KV blocks (that may or may not require masking) for each query. # KV_IDX: The indices of KV blocks (that may or may not require masking) for each query. # FULL_KV_NUM_BLKS: The number of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_KV_IDX: The indices of fully unmasked KV blocks (so we don't need masking) for each query. # # OUTPUT_LOGSUMEXP: We only need to store the logsumexp if we require grad # # (Modifiable) Performance tuning options # BLOCK_M: The thread block size across the seqlen dim of Q. # BLOCK_N: Iterate over BLOCK_N across the seqlen dim of K/V in each thread block. # The below are kernel options that can be applied for certain score_mods, # or involve a numerics vs. perf tradeoff # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base. Has # about 20% more numerical error, but slightly faster. # ROWS_GUARANTEED_SAFE: Is it guaranteed that at least one value in each row # is not masked out? If so, we can skip an extra safety check tl.static_assert(SPARSE_Q_BLOCK_SIZE >= BLOCK_M and SPARSE_Q_BLOCK_SIZE % BLOCK_M == 0) tl.static_assert(SPARSE_KV_BLOCK_SIZE >= BLOCK_N and SPARSE_KV_BLOCK_SIZE % BLOCK_N == 0) # Define strides of inputs stride_qz, stride_qh, stride_qm, stride_qk = 131072, 32768, 64, 1 stride_kz, stride_kh, stride_kn, stride_kk = 131072, 32768, 64, 1 stride_vz, stride_vh, stride_vn, stride_vk = 131072, 32768, 64, 1 ZQ = 1 HQ = 4 Q_LEN = 512 ZKV = 1 KV_LEN = 512 MATMUL_PRECISION = Q.dtype.element_ty q_start = tl.program_id(0) off_zq = tl.program_id(1) // HQ off_hq = tl.program_id(1) % HQ # We support two cases for batch dimension. a) (ZKV == ZQ) where off_zkv = off_zq. # b) (ZKV == 1 and ZQ > 1) where KV is broadcasted along the batch dimension and off_zkv=0. off_zkv = off_zq % ZKV off_hkv = off_hq // GQA_SHARED_HEADS off_g = off_hq % GQA_SHARED_HEADS q_offset = off_zq * stride_qz + off_hq * stride_qh k_offset = off_zkv * stride_kz + off_hkv * stride_kh v_offset = off_zkv * stride_vz + off_hkv * stride_vh Q = Q + q_offset K = K + k_offset V = V + v_offset SPARSE_Z = 1 SPARSE_HQ = 1 sparse_idx_z = off_zq % SPARSE_Z sparse_idx_hq = off_hq % SPARSE_HQ SPARSE_Q_MULTIPLE: tl.constexpr = (SPARSE_Q_BLOCK_SIZE // BLOCK_M) SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) stride_kv_num_blks_h = 16 stride_kv_idx_h = 256 stride_kv_idx_m = 16 # initialize pointer to m and l m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf") l_i = tl.zeros([BLOCK_M], dtype=tl.float32) acc = tl.zeros([BLOCK_M, V_HEAD_DIM], dtype=tl.float32) offs_m = q_start * BLOCK_M + tl.arange(0, BLOCK_M) # KV_IDX and KV_NUM_BLKS are always contiguous. sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq sparse_kv_num_blks_offset = sparse_hz_offset * stride_kv_num_blks_h + q_start // SPARSE_Q_MULTIPLE sparse_kv_idx_offset = sparse_hz_offset * stride_kv_idx_h + (q_start // SPARSE_Q_MULTIPLE) * stride_kv_idx_m # noqa: B950 Q_block_ptr = tl.make_block_ptr( base=Q, shape=(Q_LEN, QK_HEAD_DIM), strides=(stride_qm, stride_qk), offsets=(q_start * BLOCK_M, 0), block_shape=(BLOCK_M, QK_HEAD_DIM), order=(1, 0) ) # load q: it stays in SRAM throughout the inner loop. if IS_DIVISIBLE: q = tl.load(Q_block_ptr) else: # boundary check is not free, so we only do it when necessary. q = tl.load(Q_block_ptr, boundary_check=(0,), padding_option = "zero") # ~~~~~~~~~~~~~~ normal blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We don't know anything "special" about these blocks, so we need to apply # both score_mod and mask_mod to it kv_indices = KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=False, ) # ~~~~~~~~~~~~~~ "full" blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We know these blocks are guaranteed to be "full", so we don't need to # apply mask_mod to them - only score_mod if HAS_FULL_BLOCKS: # FULL_KV_IDX and FULL_KV_NUM_BLKS are always contiguous. kv_indices = FULL_KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # [Note] Handle fully masked out rows: # Li will be the sum(e^(-inf)) == 0.0 for masked out rows, mi will be -inf. # We set Li to 1.0 which will result in lse/out = 0.0 | after the log(li) + mi(0.0) step l_i = tl.where(l_i == 0.0, 1, l_i) acc = acc / l_i[:, None] idx_zq = tl.program_id(1) // HQ idx_hq = tl.program_id(1) % HQ idx_m = offs_m[:, None] idx_d = tl.arange(0, V_HEAD_DIM)[None, :] mask = idx_m < Q_LEN # TODO generalize and add proper mask support xindex = idx_d + (64*idx_m) + (32768*idx_hq) + (131072*idx_zq) tl.store(out_ptr0 + (tl.broadcast_to(idx_d + (64*idx_m) + (32768*idx_hq), acc.shape)), acc, mask) # TODO dont want to write this if we dont require grad if OUTPUT_LOGSUMEXP: off_hz = tl.program_id(1) l_ptrs = LSE + off_hz * Q_LEN + offs_m lse = m_i + tl.math.log2(l_i) if IS_DIVISIBLE: tl.store(l_ptrs, lse) else: tl.store(l_ptrs, lse, mask=offs_m < Q_LEN) @triton.jit def forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets used as inputs to score_mod & mask_mod # of size [BLOCK_M, BLOCK_N] or scalar. off_z, off_h, offs_m, offs_n, # blocksparse data kv_indices, kv_num_blocks, # start kv and end kv block block_n_start, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) RCP_LN2: tl.constexpr = 1.44269504 if PRESCALE_QK: q = (q * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) # loop over k, v and update accumulator until block_n_end for start_n in range(block_n_start, block_n_end): if IS_DIVISIBLE: acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) else: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen, # it's on par or slightly faster than only applying to the last block in fwd. # However, we choose different strategy for bwd, where we only apply mod & mask # to the last block because it's faster a lot. acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) # update pointers offset = get_offset_for_next_block( start_n, kv_indices, kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N ) V_block_ptr = tl.advance(V_block_ptr, (offset, 0)) K_block_ptr = tl.advance(K_block_ptr, (0, offset)) offs_n = offs_n + offset return acc, l_i, m_i @triton.jit def get_offset_for_next_block(loop_iter, col_indices, total_blocks, SPARSE_BLOCK, SPARSE_BLOCK_MULTIPLE, BLOCK): cur_block_idx = loop_iter // SPARSE_BLOCK_MULTIPLE cur_block = tl.load(col_indices + cur_block_idx, eviction_policy="evict_last") next_block = tl.load(col_indices + cur_block_idx + 1, eviction_policy="evict_last", mask=cur_block_idx + 1 < total_blocks) needs_jump = (loop_iter + 1) % SPARSE_BLOCK_MULTIPLE == 0 jump_to_block = (next_block - cur_block ) * SPARSE_BLOCK - (SPARSE_BLOCK_MULTIPLE - 1) * BLOCK offset = jump_to_block * needs_jump + (1 - needs_jump) * BLOCK return offset @triton.jit def forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 # -- load k -- if IS_DIVISIBLE: k = tl.load(K_block_ptr) else: k = tl.load(K_block_ptr, boundary_check=(1,), padding_option = "zero") # -- compute qk --- qk = tl.dot(q, k, input_precision=FLOAT32_PRECISION) # TODO: use cuda matmul when q_len <= 2. if not PRESCALE_QK: qk *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ if CHECK_BLOCK_BOUNDARY: # If this is the last block of a non divisible seqlen, we still need to load [BLOCK_M, BLOCK_N] elements, # which is larger than the actual number of elements. To avoid access memory out of bound, # we need to mask out the elements that are out of Q_LEN & KV_LEN. m = offs_m % Q_LEN n = offs_n % KV_LEN else: m = offs_m n = offs_n tmp0 = (m) - (n) tmp1 = tmp0.to(tl.float32) tmp2 = (qk) + tmp1 post_mod_scores = tmp2 if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: tmp3 = (m) >= (n) mask_mod_output = tmp3 if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n < KV_LEN, mask_mod_output, float("-inf")) # apply mask for partially unmasked blocks post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) # TODO: In the case that score_mod is linear, this can be LICMed if not PRESCALE_QK: post_mod_scores *= RCP_LN2 # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # -- compute scaling constant --- m_ij = tl.maximum(m_i, tl.max(post_mod_scores, 1)) if not ROWS_GUARANTEED_SAFE: masked_out_rows = (m_ij == float("-inf")) m_ij_masked = tl.where(masked_out_rows, 0, m_ij) else: m_ij_masked = m_ij alpha = tl.math.exp2(m_i - m_ij_masked) p = tl.math.exp2(post_mod_scores - m_ij_masked[:, None]) # NB: l_i update is pulled up here since it's a bit faster # NB: For headdim=256, it's faster to move it back down to after m_i = # m_ij l_i = l_i * alpha + tl.sum(p, 1) # # -- scale and update acc -- acc = acc * alpha[:, None] if IS_DIVISIBLE: v = tl.load(V_block_ptr) else: v = tl.load(V_block_ptr, boundary_check=(0,), padding_option = "zero") acc = tl.dot(p.to(MATMUL_PRECISION), v, acc, input_precision=FLOAT32_PRECISION) # -- update m_i m_i = m_ij return acc, l_i, m_i ''', device_str='cuda') meta0 = {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False, 'FLOAT32_PRECISION': "'ieee'", 'IS_DIVISIBLE': True, 'SM_SCALE': 0.125, 'GQA_SHARED_HEADS': 1, 'HAS_FULL_BLOCKS': True, 'QK_HEAD_DIM': 64, 'V_HEAD_DIM': 64, 'BLOCK_M': 128, 'BLOCK_N': 32, 'SPARSE_Q_BLOCK_SIZE': 128, 'SPARSE_KV_BLOCK_SIZE': 128} async_compile.wait(globals()) del async_compile def call(args): arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1 = args args.clear() assert_size_stride(arg0_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg1_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg2_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg3_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg4_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg5_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg6_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg7_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg8_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg9_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg10_1, (1, 1, 16, 16), (256, 256, 16, 1)) buf0 = empty_strided_cuda((1, 4, 512), (2048, 512, 1), torch.float32) with torch.cuda._DeviceGuard(0): torch.cuda.set_device(0) buf1 = empty_strided_cuda((1, 4, 512, 64), (131072, 32768, 64, 1), torch.float32) # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] stream0 = get_raw_stream(0) triton_tem_fused_0.run(arg0_1, arg1_1, arg2_1, buf0, arg3_1, arg4_1, arg5_1, arg6_1, buf1, grid=torch._inductor.kernel.flex_attention.flex_attention_grid(1, 4, 512, 64, meta0), stream=stream0) del arg0_1 del arg1_1 del arg2_1 del arg3_1 del arg4_1 del arg5_1 del arg6_1 del buf0 return (buf1, ) def benchmark_compiled_module(times=10, repeat=10): from torch._dynamo.testing import rand_strided from torch._inductor.utils import print_performance arg0_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg1_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg2_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg3_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg4_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg5_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg6_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg7_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg8_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg9_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg10_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) fn = lambda: call([arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1]) return print_performance(fn, times=times, repeat=repeat) if __name__ == "__main__": from torch._inductor.wrapper_benchmark import compiled_module_main compiled_module_main('None', benchmark_compiled_module) V1003 10:11:03.283000 2235078 torch/_dynamo/utils.py:1020] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "55abfbf07c349adf63e8871178bf974b"} { "name": "fx_graph_cache_hit", "ts": 1727975463283186.5, "args": { "key": "f4lkea5y7lzhlshohvr3aqpd7bchdflfs7j5wn7mrurponawoutk", "components": [ "[n7x23yy6fih6vdcjzlzbhy3d6vx3ilu7ylp3zz6wkultu4yvnzn] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False" ], "cache_event_time": 1727975463283186408, "cache_state": "hit", "time_saved_ns": 6372999079 }, "ph": "i", "cat": "dynamo_timed", "tid": 0, "pid": 0, "s": "p" } V1003 10:11:03.284000 2235078 torch/_inductor/codecache.py:1463] {"artifact": {"name": "fx_graph_cache_hit", "encoding": "json"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "b87f8df3afd49918e4211191e3fc10a5"} {"key": "f4lkea5y7lzhlshohvr3aqpd7bchdflfs7j5wn7mrurponawoutk", "components": ["[n7x23yy6fih6vdcjzlzbhy3d6vx3ilu7ylp3zz6wkultu4yvnzn] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False"], "cache_event_time": 1727975463283186408, "cache_state": "hit", "time_saved_ns": 6372999079} V1003 10:11:03.284000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "d86af378fd420082506385ca5acfaccc"} { "name": "inductor_compile", "ts": 1727975463284819.8, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.285000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "45657f0a65e5f7e3f3fcfb73f437d842"} { "name": "compile_fx_inner", "ts": 1727975463285151.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.285000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "12df42838dd8ed7060ce2ffd52760f3c"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975463285792.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.289000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "1bd305e520f60c583bf3f70d53bf1988"} { "name": "create_aot_dispatcher_function", "ts": 1727975463289199.8, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.289000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "02f3ea38f17491dd957dd1864f5411e6"} { "name": "backend_compile", "ts": 1727975463289789.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.290000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "573fd1c41404d17c4da25ab7cfc13f4b"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975463290072.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.320000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "6627aec2e890cf84ba5ee985a77950cc"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['k'], accessed_by=DictGetItemGuardAccessor(k) | | +- TYPE_MATCH: ___check_type_id(L['k'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['k'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['k'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['k'], L['q'], L['v'], L['block_mask'].q_indices, L['block_mask'].kv_indices, L['block_mask'].q_num_blocks, L['block_mask'].kv_num_blocks, L['block_mask'].full_q_indices, L['block_mask'].full_kv_indices, L['block_mask'].full_q_num_blocks, L['block_mask'].full_kv_num_blocks) | +- GuardManager: source=L['q'], accessed_by=DictGetItemGuardAccessor(q) | | +- TYPE_MATCH: ___check_type_id(L['q'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['q'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['q'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['v'], accessed_by=DictGetItemGuardAccessor(v) | | +- TYPE_MATCH: ___check_type_id(L['v'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['v'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['v'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['score_mod'], accessed_by=DictGetItemGuardAccessor(score_mod) | | +- GuardManager: source=L['score_mod'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['score_mod'].__code__, 140413271879296) | +- GuardManager: source=L['block_mask'], accessed_by=DictGetItemGuardAccessor(block_mask) | | +- TYPE_MATCH: ___check_type_id(L['block_mask'], 387600320) | | +- GuardManager: source=L['block_mask'].mask_mod, accessed_by=GetAttrGuardAccessor(mask_mod) | | | +- GuardManager: source=L['block_mask'].mask_mod.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].mask_mod.__code__, 140413271880128) | | +- GuardManager: source=L['block_mask'].q_indices, accessed_by=GetAttrGuardAccessor(q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE, accessed_by=GetAttrGuardAccessor(BLOCK_SIZE) | | | +- TYPE_MATCH: ___check_type_id(L['block_mask'].BLOCK_SIZE, 8815232) | | | +- LENGTH_CHECK: len(L['block_mask'].BLOCK_SIZE) == 2 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[0] == 128 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[1] == 128 | | +- GuardManager: source=L['block_mask'].kv_indices, accessed_by=GetAttrGuardAccessor(kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].q_num_blocks, accessed_by=GetAttrGuardAccessor(q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].kv_num_blocks, accessed_by=GetAttrGuardAccessor(kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_indices, accessed_by=GetAttrGuardAccessor(full_q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_indices, accessed_by=GetAttrGuardAccessor(full_kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_num_blocks, accessed_by=GetAttrGuardAccessor(full_q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_num_blocks, accessed_by=GetAttrGuardAccessor(full_kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=GetAttrGuardAccessor(as_tuple) | | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=FuncDefaultsGuardAccessor | | | | +- GuardManager: source=L['block_mask'].as_tuple.__defaults__[0], accessed_by=GetItemGuardAccessor(0) | | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].as_tuple.__defaults__[0], 8911040) | +- GuardManager: source=L['flex_attention'], accessed_by=DictGetItemGuardAccessor(flex_attention) | | +- GuardManager: source=L['flex_attention'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__code__, 387082992) | | +- GuardManager: source=L['flex_attention'], accessed_by=FuncDefaultsGuardAccessor | | | +- GuardManager: source=L['flex_attention'].__defaults__[2], accessed_by=GetItemGuardAccessor(2) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[2], 8825760) | | | +- GuardManager: source=L['flex_attention'].__defaults__[3], accessed_by=GetItemGuardAccessor(3) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[3], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[4], accessed_by=GetItemGuardAccessor(4) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[4], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[5], accessed_by=GetItemGuardAccessor(5) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[5], 8825760) | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['__builtins_dict___6'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___6) | | | +- GuardManager: source=G['__builtins_dict___6']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___6']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___6']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___6']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___6']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___6']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___6']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___6']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_comptime'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_comptime) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_comptime'], 140410226912176) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_decorators) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'], 140410226910096) | | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, accessed_by=GetAttrGuardAccessor(is_compiling) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, 140410376252096) | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot__utils) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot__utils'], 140409673896784) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, accessed_by=GetAttrGuardAccessor(_SUPPORTED_HEAD_DIMS) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, 8844320) | | | | +- LENGTH_CHECK: len(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS) == 10 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0], accessed_by=ListGetItemGuardAccessor(0) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0] == 2 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1], accessed_by=ListGetItemGuardAccessor(1) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1] == 4 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2], accessed_by=ListGetItemGuardAccessor(2) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2] == 8 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3], accessed_by=ListGetItemGuardAccessor(3) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3] == 16 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4], accessed_by=ListGetItemGuardAccessor(4) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4] == 32 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5], accessed_by=ListGetItemGuardAccessor(5) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5] == 64 | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot_flex_attention) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], 140409673895824) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, accessed_by=GetAttrGuardAccessor(math) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, 140413266939392) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, accessed_by=GetAttrGuardAccessor(sqrt) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, 140413266943072) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, accessed_by=GetAttrGuardAccessor(torch) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, 140413267918368) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, accessed_by=GetAttrGuardAccessor(_dynamo) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, 140413260098400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static, accessed_by=GetAttrGuardAccessor(mark_static) | | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, 123166432) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, accessed_by=GetAttrGuardAccessor(compiler) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, 140410826010400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, accessed_by=GetAttrGuardAccessor(is_dynamo_compiling) | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, 140410826132992) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device, accessed_by=GetAttrGuardAccessor(_validate_device) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, 140409673611088) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, accessed_by=GetAttrGuardAccessor(flex_attention_hop) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, 96992544) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__ == 'flex_attention' | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim, accessed_by=GetAttrGuardAccessor(_supported_head_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, 140409673231376) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim, accessed_by=GetAttrGuardAccessor(_validate_embed_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, 388086512) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input, accessed_by=GetAttrGuardAccessor(_validate_sdpa_input) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, 387915104) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options, accessed_by=GetAttrGuardAccessor(_apply_kernel_options) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, 140409683680752) V1003 10:11:03.321000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "341d7b7d938fce7537065544860a0efd"} { "name": "entire_frame_compile", "ts": 1727975463321242.8, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.321000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "02dcbee7c15ab2085e7ee42a1679beb8"} { "name": "_compile.compile_inner", "ts": 1727975463321536.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.321000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "1/0", "frame_key": "2", "co_name": "fn", "co_filename": "/data/users/oulgen/pytorch/test/inductor/test_codecache.py", "co_firstlineno": 379, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 80, "shape_env_guard_count": 0, "graph_op_count": 11, "graph_node_count": 25, "graph_input_count": 11, "start_time": 1727975462.8240004, "entire_frame_compile_time_s": 0.4971275329589844, "backend_compile_time_s": 0.2832028865814209, "inductor_compile_time_s": 0.1434309482574463, "code_gen_time_s": null, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 6.372999079, "structured_logging_overhead_s": 0.055921114, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.325000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 414, "name": "test_flex_attention_caching", "filename": 1}, {"line": 386, "name": "fn2", "filename": 1}, {"line": 1062, "name": "flex_attention", "filename": 10}, {"line": 1049, "name": "_flex_attention_hop_wrapper", "filename": 10}]}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.325000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "af0d73822951ce32af4202533e2a8290"} { "name": "_compile.compile_inner", "ts": 1727975463325378.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.325000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "dc5cd9b8a50a9005377afd5d28ba36fb"} { "name": "entire_frame_compile", "ts": 1727975463325378.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.328000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 348, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.328000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.328000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 0, "source": "L['args'][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.441000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 348, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.442000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 2, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.442000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 2, "source": "L['args'][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.443000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 348, "size": 524288}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.444000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 3, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.444000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 3, "source": "L['args'][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.445000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 348, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.445000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 4, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.445000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 4, "source": "L['args'][4][0]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.446000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 348, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.447000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 5, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.447000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 5, "source": "L['args'][4][1]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.448000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 348, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.448000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 6, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.448000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 6, "source": "L['args'][4][2]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.449000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 348, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.450000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.450000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 7, "source": "L['args'][4][3]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.451000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 348, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.451000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.451000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 8, "source": "L['args'][4][4]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.452000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 348, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.453000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.453000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 9, "source": "L['args'][4][5]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.454000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 348, "size": 64}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.454000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.455000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 10, "source": "L['args'][4][6]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.456000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 348, "size": 1024}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.456000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 348}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.456000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 348, "id": 11, "source": "L['args'][4][7]"}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.464000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_args_0_": [1, 4, 512, 64], "l_args_1_": [1, 4, 512, 64], "l_args_2_": [1, 4, 512, 64], "l_args_4_0_": [1, 1, 16], "l_args_4_1_": [1, 1, 16, 16], "l_args_4_2_": [1, 1, 16], "l_args_4_3_": [1, 1, 16, 16], "l_args_4_4_": [1, 1, 16], "l_args_4_5_": [1, 1, 16, 16], "l_args_4_6_": [1, 1, 16], "l_args_4_7_": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "getitem": [1, 4, 512, 64], "getitem_1": [1, 4, 512]}}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "c759d677da3d86354d83c3018e999089"} class GraphModule(torch.nn.Module): def forward(self, L_args_0_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_1_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_2_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_args_4_0_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_1_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_2_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_3_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_4_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_5_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_args_4_6_: "i32[1, 1, 16][16, 16, 1]cuda:0", L_args_4_7_: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_args_0_ = L_args_0_ l_args_1_ = L_args_1_ l_args_2_ = L_args_2_ l_args_4_0_ = L_args_4_0_ l_args_4_1_ = L_args_4_1_ l_args_4_2_ = L_args_4_2_ l_args_4_3_ = L_args_4_3_ l_args_4_4_ = L_args_4_4_ l_args_4_5_ = L_args_4_5_ l_args_4_6_ = L_args_4_6_ l_args_4_7_ = L_args_4_7_ # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1050 in _flex_attention_hop_wrapper, code: return flex_attention_hop(*args, **kwargs) child_1: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_args_0_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_args_0_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_args_0_, l_args_1_, l_args_2_, score_mod_0, (l_args_4_0_, l_args_4_1_, l_args_4_2_, l_args_4_3_, l_args_4_4_, l_args_4_5_, l_args_4_6_, l_args_4_7_, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_args_0_ = l_args_1_ = l_args_2_ = score_mod_0 = l_args_4_0_ = l_args_4_1_ = l_args_4_2_ = l_args_4_3_ = l_args_4_4_ = l_args_4_5_ = l_args_4_6_ = l_args_4_7_ = mask_fn_0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0] getitem_1: "f32[1, 4, 512][2048, 512, 1]cuda:0" = flex_attention[1]; flex_attention = None return (getitem, getitem_1) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): return child class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:11:03.465000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "3f2a0e3577c204c69d66ef7550317ca5"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975463465456.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.465000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "b48e24df1eadf9ac86a2058fefc95dfb"} { "name": "backend_compile", "ts": 1727975463465456.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.466000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "9ed8f1c18e744cba5c7e86e34c900750"} { "name": "backend_compile", "ts": 1727975463466215.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.466000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "c19895e106748f6b125e2df40adb8b9e"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975463466480.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.490000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "6c94a4177e3b74bb7d4cee6cdd33898c"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['args'], accessed_by=DictGetItemGuardAccessor(args) | | +- TYPE_MATCH: ___check_type_id(L['args'], 8815232) | | +- LENGTH_CHECK: len(L['args']) == 7 | | +- GuardManager: source=L['args'][0], accessed_by=TupleGetItemGuardAccessor(0) | | | +- TENSOR_MATCH: check_tensor(L['args'][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][0], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['args'][0], L['args'][1], L['args'][2], L['args'][4][0], L['args'][4][1], L['args'][4][2], L['args'][4][3], L['args'][4][4], L['args'][4][5], L['args'][4][6], L['args'][4][7]) | | +- GuardManager: source=L['args'][1], accessed_by=TupleGetItemGuardAccessor(1) | | | +- TENSOR_MATCH: check_tensor(L['args'][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][1], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][2], accessed_by=TupleGetItemGuardAccessor(2) | | | +- TENSOR_MATCH: check_tensor(L['args'][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | | +- NO_HASATTR: hasattr(L['args'][2], '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['args'][3], accessed_by=TupleGetItemGuardAccessor(3) | | | +- GuardManager: source=L['args'][3].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['args'][3].__code__, 140413275216112) | | +- GuardManager: source=L['args'][4], accessed_by=TupleGetItemGuardAccessor(4) | | | +- TYPE_MATCH: ___check_type_id(L['args'][4], 8815232) | | | +- LENGTH_CHECK: len(L['args'][4]) == 11 | | | +- GuardManager: source=L['args'][4][0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][0], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][0], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][1], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][1], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][2], accessed_by=TupleGetItemGuardAccessor(2) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][2], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][2], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][3], accessed_by=TupleGetItemGuardAccessor(3) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][3], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][3], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][4], accessed_by=TupleGetItemGuardAccessor(4) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][4], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][4], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][5], accessed_by=TupleGetItemGuardAccessor(5) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][5], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][5], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][6], accessed_by=TupleGetItemGuardAccessor(6) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][6], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][6], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][7], accessed_by=TupleGetItemGuardAccessor(7) | | | | +- TENSOR_MATCH: check_tensor(L['args'][4][7], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | | +- NO_HASATTR: hasattr(L['args'][4][7], '_dynamo_dynamic_indices') == False | | | | +- NO_TENSOR_ALIASING | | | +- GuardManager: source=L['args'][4][8], accessed_by=TupleGetItemGuardAccessor(8) | | | | +- EQUALS_MATCH: L['args'][4][8] == 128 | | | +- GuardManager: source=L['args'][4][9], accessed_by=TupleGetItemGuardAccessor(9) | | | | +- EQUALS_MATCH: L['args'][4][9] == 128 | | | +- GuardManager: source=L['args'][4][10], accessed_by=TupleGetItemGuardAccessor(10) | | | | +- GuardManager: source=L['args'][4][10].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(L['args'][4][10].__code__, 140413271880128) | | +- GuardManager: source=L['args'][5], accessed_by=TupleGetItemGuardAccessor(5) | | | +- EQUALS_MATCH: L['args'][5] == 0.125 | | +- GuardManager: source=L['args'][6], accessed_by=TupleGetItemGuardAccessor(6) | | | +- DICT_LENGTH: len(L['args'][6]) == 3 | | | +- GuardManager: source=L['args'][6]['ROWS_GUARANTEED_SAFE'], accessed_by=DictGetItemGuardAccessor(ROWS_GUARANTEED_SAFE) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['ROWS_GUARANTEED_SAFE'], 8910592) | | | +- GuardManager: source=L['args'][6]['PRESCALE_QK'], accessed_by=DictGetItemGuardAccessor(PRESCALE_QK) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['PRESCALE_QK'], 8910592) | | | +- GuardManager: source=L['args'][6]['OUTPUT_LOGSUMEXP'], accessed_by=DictGetItemGuardAccessor(OUTPUT_LOGSUMEXP) | | | | +- ID_MATCH: ___check_obj_id(L['args'][6]['OUTPUT_LOGSUMEXP'], 8910592) | +- GuardManager: source=L['kwargs'], accessed_by=DictGetItemGuardAccessor(kwargs) | | +- DICT_LENGTH: not L['kwargs'] | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['flex_attention_hop'], accessed_by=DictGetItemGuardAccessor(flex_attention_hop) | | | +- TYPE_MATCH: ___check_type_id(G['flex_attention_hop'], 96992544) | | | +- GuardManager: source=G['flex_attention_hop'].__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | +- EQUALS_MATCH: G['flex_attention_hop'].__name__ == 'flex_attention' | | +- GuardManager: source=G['__builtins_dict___8'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___8) | | | +- GuardManager: source=G['__builtins_dict___8']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___8']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___8']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___8']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___8']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___8']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___8']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___8']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) V1003 10:11:03.491000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "eda00b5e5981dad4abbde14a07a38b9c"} { "name": "entire_frame_compile", "ts": 1727975463491399.2, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.491000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "dc3c58a766d43af7e167ce88702b4e5c"} { "name": "_compile.compile_inner", "ts": 1727975463491701.2, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 1, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.492000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "0/0", "frame_key": "1", "co_name": "_flex_attention_hop_wrapper", "co_filename": "/data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py", "co_firstlineno": 1049, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 55, "shape_env_guard_count": 0, "graph_op_count": 12, "graph_node_count": 26, "graph_input_count": 11, "start_time": 1727975463.3253677, "entire_frame_compile_time_s": 0.16592144966125488, "backend_compile_time_s": 0.0006587505340576172, "inductor_compile_time_s": null, "code_gen_time_s": null, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 0, "structured_logging_overhead_s": 0.04651781599999999, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.496000 2235078 torch/_dynamo/convert_frame.py:915] {"dynamo_start": {"stack": [{"line": 916, "name": "", "filename": 1}, {"line": 14, "name": "run_tests", "filename": 2}, {"line": 38, "name": "run_tests", "filename": 3}, {"line": 1273, "name": "run_tests", "filename": 4}, {"line": 102, "name": "__init__", "filename": 5}, {"line": 274, "name": "runTests", "filename": 5}, {"line": 217, "name": "run", "filename": 6}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 84, "name": "__call__", "filename": 7}, {"line": 122, "name": "run", "filename": 7}, {"line": 678, "name": "__call__", "filename": 8}, {"line": 3116, "name": "run", "filename": 4}, {"line": 3088, "name": "_run_custom", "filename": 4}, {"line": 623, "name": "run", "filename": 8}, {"line": 579, "name": "_callTestMethod", "filename": 8}, {"line": 2983, "name": "wrapper", "filename": 4}, {"line": 81, "name": "inner", "filename": 9}, {"line": 81, "name": "inner", "filename": 9}, {"line": 414, "name": "test_flex_attention_caching", "filename": 1}, {"line": 385, "name": "fn2", "filename": 1}]}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.496000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "68d9557f97487c1ecd28ffaf4020aff9"} { "name": "_compile.compile_inner", "ts": 1727975463496392.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.496000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "eab0483a7d9a0ee867d4596b4367f9e6"} { "name": "entire_frame_compile", "ts": 1727975463496392.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.499000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 0, "describer_id": 362, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.499000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 0, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 0, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.499000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 0, "source": "L['q']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.508000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 1, "describer_id": 362, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.508000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 1, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 1, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.509000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 1, "source": "L['k']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.510000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 2, "describer_id": 362, "size": 524288}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.510000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 2, "ndim": 4, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1, 4, 512, 64], "is_leaf": true, "stride": [131072, 32768, 64, 1], "storage": 2, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.511000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 2, "source": "L['v']"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.520000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 3, "describer_id": 362, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.520000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 3, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 3, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.521000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 3, "source": "L['block_mask'].kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.653000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 4, "describer_id": 362, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.654000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 5, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 4, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.654000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 5, "source": "L['block_mask'].kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.655000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 5, "describer_id": 362, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.656000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 6, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 5, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.656000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 6, "source": "L['block_mask'].full_kv_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.657000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 6, "describer_id": 362, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.658000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 7, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 6, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.658000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 7, "source": "L['block_mask'].full_kv_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.659000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 7, "describer_id": 362, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.659000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 8, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 7, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.660000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 8, "source": "L['block_mask'].q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.661000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 8, "describer_id": 362, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.661000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 9, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 8, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.661000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 9, "source": "L['block_mask'].q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.663000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 9, "describer_id": 362, "size": 64}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.663000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 10, "ndim": 3, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16], "is_leaf": true, "stride": [16, 16, 1], "storage": 9, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.663000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 10, "source": "L['block_mask'].full_q_num_blocks"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.665000 2235078 torch/_subclasses/meta_utils.py:204] {"describe_storage": {"id": 10, "describer_id": 362, "size": 1024}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.665000 2235078 torch/_subclasses/meta_utils.py:417] {"describe_tensor": {"id": 11, "ndim": 4, "dtype": "torch.int32", "device": "device(type='cuda', index=0)", "size": [1, 1, 16, 16], "is_leaf": true, "stride": [256, 256, 16, 1], "storage": 10, "view_func": "", "describer_id": 362}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.665000 2235078 torch/_subclasses/meta_utils.py:1640] {"describe_source": {"describer_id": 362, "id": 11, "source": "L['block_mask'].full_q_indices"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} V1003 10:11:03.673000 2235078 torch/_dynamo/output_graph.py:1347] {"dynamo_output_graph": {"sizes": {"l_q_": [1, 4, 512, 64], "l_k_": [1, 4, 512, 64], "l_v_": [1, 4, 512, 64], "l_block_mask_kv_num_blocks": [1, 1, 16], "l_block_mask_kv_indices": [1, 1, 16, 16], "l_block_mask_full_kv_num_blocks": [1, 1, 16], "l_block_mask_full_kv_indices": [1, 1, 16, 16], "l_block_mask_q_num_blocks": [1, 1, 16], "l_block_mask_q_indices": [1, 1, 16, 16], "l_block_mask_full_q_num_blocks": [1, 1, 16], "l_block_mask_full_q_indices": [1, 1, 16, 16], "child_1": [], "child_2": [], "child_3": [], "child_4": [], "child": [], "child_5": [], "child_6": [], "child_7": [], "child_8": [], "out": [1, 4, 512, 64]}}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "487c69baa1e2981125efdf989466d11f"} class GraphModule(torch.nn.Module): def forward(self, L_q_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_k_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_v_: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", L_block_mask_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_kv_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_kv_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", L_block_mask_full_q_num_blocks: "i32[1, 1, 16][16, 16, 1]cuda:0", L_block_mask_full_q_indices: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): l_q_ = L_q_ l_k_ = L_k_ l_v_ = L_v_ l_block_mask_kv_num_blocks = L_block_mask_kv_num_blocks l_block_mask_kv_indices = L_block_mask_kv_indices l_block_mask_full_kv_num_blocks = L_block_mask_full_kv_num_blocks l_block_mask_full_kv_indices = L_block_mask_full_kv_indices l_block_mask_q_num_blocks = L_block_mask_q_num_blocks l_block_mask_q_indices = L_block_mask_q_indices l_block_mask_full_q_num_blocks = L_block_mask_full_q_num_blocks l_block_mask_full_q_indices = L_block_mask_full_q_indices # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( child_1: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_1 = None child_2: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_2 = None child_3: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_3 = None child_4: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_4 = None child: "f32[][]cuda:0" = l_q_.new_empty([], requires_grad = False); child = None score_mod_0 = self.score_mod_0 child_5: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_5 = None child_6: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_6 = None child_7: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_7 = None child_8: "i32[][]cuda:0" = l_q_.new_empty([], dtype = torch.int32); child_8 = None mask_fn_0 = self.mask_fn_0 flex_attention = torch.ops.higher_order.flex_attention(l_q_, l_k_, l_v_, score_mod_0, (l_block_mask_kv_num_blocks, l_block_mask_kv_indices, l_block_mask_full_kv_num_blocks, l_block_mask_full_kv_indices, l_block_mask_q_num_blocks, l_block_mask_q_indices, l_block_mask_full_q_num_blocks, l_block_mask_full_q_indices, 128, 128, mask_fn_0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); l_q_ = l_k_ = l_v_ = score_mod_0 = l_block_mask_kv_num_blocks = l_block_mask_kv_indices = l_block_mask_full_kv_num_blocks = l_block_mask_full_kv_indices = l_block_mask_q_num_blocks = l_block_mask_q_indices = l_block_mask_full_q_num_blocks = l_block_mask_full_q_indices = mask_fn_0 = None out: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (out,) class score_mod_0(torch.nn.Module): def forward(self, child: "f32[][]cuda:0", child_1: "i32[][]cuda:0", child_2: "i32[][]cuda:0", child_3: "i32[][]cuda:0", child_4: "i32[][]cuda:0"): return child class mask_fn_0(torch.nn.Module): def forward(self, child_5: "i32[][]cuda:0", child_6: "i32[][]cuda:0", child_7: "i32[][]cuda:0", child_8: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = child_7 >= child_8; child_7 = child_8 = None return ge V1003 10:11:03.674000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "94724b556ec4c03944e43538685317ee"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975463674357.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.674000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "726b102bddec88824ed82f86c8c02d39"} { "name": "backend_compile", "ts": 1727975463674357.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.680000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c103e7f8badb6398fb5a927c8aeef126"} { "name": "create_aot_dispatcher_function", "ts": 1727975463680196.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.798000 2235078 torch/_functorch/_aot_autograd/dispatch_and_compile_graph.py:215] {"aot_forward_graph": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "fe8d53ae08cd360e1378a8a527d186f0"} class (torch.nn.Module): def forward(self, arg0_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg1_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg2_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg3_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg4_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg5_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg6_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg7_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg8_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg9_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg10_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (getitem,) class sdpa_score0(torch.nn.Module): def forward(self, arg0_1: "f32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0", arg4_1: "i32[][]cuda:0"): return arg0_1 class sdpa_mask0(torch.nn.Module): def forward(self, arg0_1: "i32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = torch.ops.aten.ge.Tensor(arg2_1, arg3_1); arg2_1 = arg3_1 = None return ge V1003 10:11:03.799000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "19af874521da2f24c9f873fb757d7dc6"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975463799121.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.801000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "bb6048e2282795589bd2400ba093dfc8"} { "name": "compile_fx_inner", "ts": 1727975463801024.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.801000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "b7b2a66f7ec1d6b6227990c8cd3b1d21"} { "name": "inductor_compile", "ts": 1727975463801024.5, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.816000 2235078 torch/_inductor/compile_fx.py:731] {"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c507028eda974d39e4ca6593c6320460"} import torch from torch import tensor, device import torch.fx as fx from torch._dynamo.testing import rand_strided from math import inf import torch._inductor.inductor_prims import torch._dynamo.config import torch._inductor.config import torch._functorch.config import torch.fx.experimental._config torch._dynamo.config.log_compilation_metrics = False torch._dynamo.config.fake_tensor_cache_crosscheck_enabled = True torch._inductor.config.fx_graph_remote_cache = False torch._inductor.config.autotune_local_cache = False torch._inductor.config.autotune_remote_cache = False torch._functorch.config.unlift_effect_tokens = True isolate_fails_code_str = None # torch version: 2.5.0a0+git7647c39 # torch cuda version: 12.0 # torch git version: 7647c398ff87daf70260854cf0a7f7993b3abc76 # CUDA Info: # nvcc: NVIDIA (R) Cuda compiler driver # Copyright (c) 2005-2023 NVIDIA Corporation # Built on Fri_Jan__6_16:45:21_PST_2023 # Cuda compilation tools, release 12.0, V12.0.140 # Build cuda_12.0.r12.0/compiler.32267302_0 # GPU Hardware Info: # NVIDIA PG509-210 : 8 from torch.nn import * class Repro(torch.nn.Module): def __init__(self) -> None: super().__init__() self.sdpa_score0 = () self.sdpa_mask0 = () def forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1): sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem = flex_attention[0]; flex_attention = None return (getitem,) def load_args(reader): buf0 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf0, (1, 4, 512, 64), is_leaf=True) # arg0_1 buf1 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf1, (1, 4, 512, 64), is_leaf=True) # arg1_1 buf2 = reader.storage(None, 524288, device=device(type='cuda', index=0)) reader.tensor(buf2, (1, 4, 512, 64), is_leaf=True) # arg2_1 buf3 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf3, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg3_1 buf4 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf4, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg4_1 buf5 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf5, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg5_1 buf6 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf6, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg6_1 buf7 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf7, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg7_1 buf8 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf8, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg8_1 buf9 = reader.storage(None, 64, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf9, (1, 1, 16), dtype=torch.int32, is_leaf=True) # arg9_1 buf10 = reader.storage(None, 1024, device=device(type='cuda', index=0), dtype_hint=torch.int32) reader.tensor(buf10, (1, 1, 16, 16), dtype=torch.int32, is_leaf=True) # arg10_1 load_args._version = 0 mod = Repro() if __name__ == '__main__': from torch._dynamo.repro.after_aot import run_repro with torch.no_grad(): run_repro(mod, load_args, accuracy=False, command='run', save_dir=None, tracing_mode='real', check_str=None) # To run it separately, do # mod, args = run_repro(mod, load_args, accuracy=False, command='get_args', save_dir=None, tracing_mode='real', check_str=None) # mod(*args) V1003 10:11:03.827000 2235078 torch/_inductor/compile_fx.py:795] {"inductor_post_grad_graph": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "fe8d53ae08cd360e1378a8a527d186f0"} class (torch.nn.Module): def forward(self, arg0_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg1_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg2_1: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0", arg3_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg4_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg5_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg6_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg7_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg8_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0", arg9_1: "i32[1, 1, 16][16, 16, 1]cuda:0", arg10_1: "i32[1, 1, 16, 16][256, 256, 16, 1]cuda:0"): # File: /data/users/oulgen/pytorch/torch/nn/attention/flex_attention.py:1032 in flex_attention, code: out, lse = flex_attention_hop( sdpa_score0 = self.sdpa_score0 sdpa_mask0 = self.sdpa_mask0 flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None getitem: "f32[1, 4, 512, 64][131072, 32768, 64, 1]cuda:0" = flex_attention[0]; flex_attention = None return (getitem,) class sdpa_score0(torch.nn.Module): def forward(self, arg0_1: "f32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0", arg4_1: "i32[][]cuda:0"): return arg0_1 class sdpa_mask0(torch.nn.Module): def forward(self, arg0_1: "i32[][]cuda:0", arg1_1: "i32[][]cuda:0", arg2_1: "i32[][]cuda:0", arg3_1: "i32[][]cuda:0"): # File: /data/users/oulgen/pytorch/test/inductor/test_codecache.py:373 in , code: lambda b, h, q, kv: q >= kv, None, None, 2048, 2048 ge: "b8[][]cuda:0" = torch.ops.aten.ge.Tensor(arg2_1, arg3_1); arg2_1 = arg3_1 = None return ge V1003 10:11:03.828000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "e24475885265d7add113de685b5b4785"} { "name": "GraphLowering.run", "ts": 1727975463828251.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.863000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "df613d910e7ef57b6d235421c34afbec"} { "name": "GraphLowering.run", "ts": 1727975463863685.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.864000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "5fd8784b67f1838294fd144c5df7385a"} { "name": "GraphLowering.compile_to_module", "ts": 1727975463864454.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.864000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "2614bec07c124d967f10a9f150623f5a"} { "name": "code_gen", "ts": 1727975463864454.0, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.868000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c941c5142cd262302a4ea27ba7d7eda9"} { "name": "Scheduler.__init__", "ts": 1727975463868177.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.871000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "19fd40416e912c56022b05ddd45c6f91"} { "name": "Scheduler.__init__", "ts": 1727975463871282.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.871000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "1ae33988f4240ac670d91f8fe8bc3613"} { "name": "Scheduler.codegen", "ts": 1727975463871646.2, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.881000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "4ea7f96a10ffbc186f91b8fbf392a78d"} { "name": "Scheduler.codegen", "ts": 1727975463881456.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.881000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "72cc59f88d66ead0086c546b18c184d7"} { "name": "PythonWrapperCodegen.generate", "ts": 1727975463881839.8, "args": null, "ph": "B", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.884000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "20792f3cea34e70b42f7c6d998aaa598"} { "name": "PythonWrapperCodegen.generate", "ts": 1727975463884612.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:03.885000 2235078 torch/_inductor/graph.py:1899] {"inductor_output_code": {"filename": "/tmp/oulgen/tmp4z1i5ywe/si/csitvhfwicmtxv44ng5kkavccd5rcpalvwhjof7rk2hputpzelxm.py"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "c453185bb704fad4bd52bdb872abe96c"} # AOT ID: ['2_inference'] from ctypes import c_void_p, c_long, c_int import torch import math import random import os import tempfile from math import inf, nan from torch._inductor.hooks import run_intermediate_hooks from torch._inductor.utils import maybe_profile from torch._inductor.codegen.memory_planning import _align as align from torch import device, empty_strided from torch._inductor.async_compile import AsyncCompile from torch._inductor.select_algorithm import extern_kernels from torch._inductor.codegen.multi_kernel import MultiKernelCall import torch._inductor.kernel.flex_attention from torch._C import _cuda_getCurrentRawStream as get_raw_stream import triton import triton.language as tl from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph aten = torch.ops.aten inductor_ops = torch.ops.inductor _quantized = torch.ops._quantized assert_size_stride = torch._C._dynamo.guards.assert_size_stride empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor alloc_from_pool = torch.ops.inductor._alloc_from_pool async_compile = AsyncCompile() # kernel path: /tmp/oulgen/tmp4z1i5ywe/vr/cvrnbiynkuy34l7cucuaotjk5vdulwxhwmvmwz7gy5423p3sv57h.py # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] # Source node to ATen node mapping: # flex_attention => flex_attention # Graph fragment: # %flex_attention : [num_users=1] = call_function[target=torch.ops.higher_order.flex_attention](args = (%arg0_1, %arg1_1, %arg2_1, %sdpa_score0, (%arg3_1, %arg4_1, %arg5_1, %arg6_1, %arg7_1, %arg8_1, %arg9_1, %arg10_1, 128, 128, %sdpa_mask0), 0.125, {ROWS_GUARANTEED_SAFE: False, PRESCALE_QK: False, OUTPUT_LOGSUMEXP: False}, (), ()), kwargs = {}) triton_tem_fused_0 = async_compile.triton('triton_tem_fused_0', ''' import triton import triton.language as tl from triton.compiler.compiler import AttrsDescriptor from torch._inductor.runtime import triton_helpers, triton_heuristics from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties @triton_heuristics.template( num_stages=3, num_warps=4, triton_meta={'signature': {'arg_Q': '*fp32', 'arg_K': '*fp32', 'arg_V': '*fp32', 'arg_LSE': '*fp32', 'arg_KV_NUM_BLKS': '*i32', 'arg_KV_IDX': '*i32', 'arg_FULL_KV_NUM_BLKS': '*i32', 'arg_FULL_KV_IDX': '*i32', 'out_ptr0': '*fp32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108, warp_size=32), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7, 8), equal_to_1=())]}, inductor_meta={'kernel_name': 'triton_tem_fused_0', 'backend_hash': 'FB2CA426CF35F271C56C0D69873498391AC248E25890F2B631CA8B52D56952BD', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': False, 'autotune_pointwise': True, 'autotune_remote_cache': False, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}, ) @triton.jit def triton_tem_fused_0(arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0): ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 Q = arg_Q K = arg_K V = arg_V LSE = arg_LSE KV_NUM_BLKS = arg_KV_NUM_BLKS KV_IDX = arg_KV_IDX FULL_KV_NUM_BLKS = arg_FULL_KV_NUM_BLKS FULL_KV_IDX = arg_FULL_KV_IDX # Sub notation for this kernel: # # Q: Query, K: Key, V: Value # M: Number of queries, N: Number of keys/values, D: Model dimension # QK_HEAD_DIM: The dimension of the query and key embeddings # V_HEAD_DIM: The dimension of the value embeddings # z: Batch size, h: Number of heads, m: Number of queries per head, k: Number of keys per head # GQA_SHARED_HEADS: number of query heads sharing one kv head in GQA setups. # # The following FULL_* and PARTIAL_* is defined in the block sparse mask grid, rather than the thread block grid. # KV_NUM_BLKS: The number of KV blocks (that may or may not require masking) for each query. # KV_IDX: The indices of KV blocks (that may or may not require masking) for each query. # FULL_KV_NUM_BLKS: The number of fully unmasked KV blocks (so we don't need masking) for each query. # FULL_KV_IDX: The indices of fully unmasked KV blocks (so we don't need masking) for each query. # # OUTPUT_LOGSUMEXP: We only need to store the logsumexp if we require grad # # (Modifiable) Performance tuning options # BLOCK_M: The thread block size across the seqlen dim of Q. # BLOCK_N: Iterate over BLOCK_N across the seqlen dim of K/V in each thread block. # The below are kernel options that can be applied for certain score_mods, # or involve a numerics vs. perf tradeoff # PRESCALE_QK: Whether to pre-scale QK by 1/sqrt(d) and change of base. Has # about 20% more numerical error, but slightly faster. # ROWS_GUARANTEED_SAFE: Is it guaranteed that at least one value in each row # is not masked out? If so, we can skip an extra safety check tl.static_assert(SPARSE_Q_BLOCK_SIZE >= BLOCK_M and SPARSE_Q_BLOCK_SIZE % BLOCK_M == 0) tl.static_assert(SPARSE_KV_BLOCK_SIZE >= BLOCK_N and SPARSE_KV_BLOCK_SIZE % BLOCK_N == 0) # Define strides of inputs stride_qz, stride_qh, stride_qm, stride_qk = 131072, 32768, 64, 1 stride_kz, stride_kh, stride_kn, stride_kk = 131072, 32768, 64, 1 stride_vz, stride_vh, stride_vn, stride_vk = 131072, 32768, 64, 1 ZQ = 1 HQ = 4 Q_LEN = 512 ZKV = 1 KV_LEN = 512 MATMUL_PRECISION = Q.dtype.element_ty q_start = tl.program_id(0) off_zq = tl.program_id(1) // HQ off_hq = tl.program_id(1) % HQ # We support two cases for batch dimension. a) (ZKV == ZQ) where off_zkv = off_zq. # b) (ZKV == 1 and ZQ > 1) where KV is broadcasted along the batch dimension and off_zkv=0. off_zkv = off_zq % ZKV off_hkv = off_hq // GQA_SHARED_HEADS off_g = off_hq % GQA_SHARED_HEADS q_offset = off_zq * stride_qz + off_hq * stride_qh k_offset = off_zkv * stride_kz + off_hkv * stride_kh v_offset = off_zkv * stride_vz + off_hkv * stride_vh Q = Q + q_offset K = K + k_offset V = V + v_offset SPARSE_Z = 1 SPARSE_HQ = 1 sparse_idx_z = off_zq % SPARSE_Z sparse_idx_hq = off_hq % SPARSE_HQ SPARSE_Q_MULTIPLE: tl.constexpr = (SPARSE_Q_BLOCK_SIZE // BLOCK_M) SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) stride_kv_num_blks_h = 16 stride_kv_idx_h = 256 stride_kv_idx_m = 16 # initialize pointer to m and l m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf") l_i = tl.zeros([BLOCK_M], dtype=tl.float32) acc = tl.zeros([BLOCK_M, V_HEAD_DIM], dtype=tl.float32) offs_m = q_start * BLOCK_M + tl.arange(0, BLOCK_M) # KV_IDX and KV_NUM_BLKS are always contiguous. sparse_hz_offset = sparse_idx_z * SPARSE_HQ + sparse_idx_hq sparse_kv_num_blks_offset = sparse_hz_offset * stride_kv_num_blks_h + q_start // SPARSE_Q_MULTIPLE sparse_kv_idx_offset = sparse_hz_offset * stride_kv_idx_h + (q_start // SPARSE_Q_MULTIPLE) * stride_kv_idx_m # noqa: B950 Q_block_ptr = tl.make_block_ptr( base=Q, shape=(Q_LEN, QK_HEAD_DIM), strides=(stride_qm, stride_qk), offsets=(q_start * BLOCK_M, 0), block_shape=(BLOCK_M, QK_HEAD_DIM), order=(1, 0) ) # load q: it stays in SRAM throughout the inner loop. if IS_DIVISIBLE: q = tl.load(Q_block_ptr) else: # boundary check is not free, so we only do it when necessary. q = tl.load(Q_block_ptr, boundary_check=(0,), padding_option = "zero") # ~~~~~~~~~~~~~~ normal blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We don't know anything "special" about these blocks, so we need to apply # both score_mod and mask_mod to it kv_indices = KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=False, ) # ~~~~~~~~~~~~~~ "full" blocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # We know these blocks are guaranteed to be "full", so we don't need to # apply mask_mod to them - only score_mod if HAS_FULL_BLOCKS: # FULL_KV_IDX and FULL_KV_NUM_BLKS are always contiguous. kv_indices = FULL_KV_IDX + sparse_kv_idx_offset kv_start = tl.load(kv_indices) * SPARSE_KV_BLOCK_SIZE # first kv block we're loading kv_num_blocks = tl.load(FULL_KV_NUM_BLKS + sparse_kv_num_blks_offset) block_n_end = tl.minimum(kv_num_blocks * SPARSE_KV_MULTIPLE, tl.maximum(tl.cdiv(KV_LEN, BLOCK_N), 1)) K_block_ptr = tl.make_block_ptr( base=K, shape=(QK_HEAD_DIM, KV_LEN), strides=(stride_kk, stride_kn), offsets=(0, kv_start), block_shape=(QK_HEAD_DIM, BLOCK_N), order=(0, 1) ) V_block_ptr = tl.make_block_ptr( base=V, shape=(KV_LEN, V_HEAD_DIM), strides=(stride_vn, stride_vk), offsets=(kv_start, 0), block_shape=(BLOCK_N, V_HEAD_DIM), order=(1, 0) ) offs_n = kv_start + tl.arange(0, BLOCK_N) acc, l_i, m_i = forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, acc, l_i, m_i, off_zq, off_hq, offs_m[:, None], offs_n[None, :], kv_indices, kv_num_blocks, 0, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS=True, ) # [Note] Handle fully masked out rows: # Li will be the sum(e^(-inf)) == 0.0 for masked out rows, mi will be -inf. # We set Li to 1.0 which will result in lse/out = 0.0 | after the log(li) + mi(0.0) step l_i = tl.where(l_i == 0.0, 1, l_i) acc = acc / l_i[:, None] idx_zq = tl.program_id(1) // HQ idx_hq = tl.program_id(1) % HQ idx_m = offs_m[:, None] idx_d = tl.arange(0, V_HEAD_DIM)[None, :] mask = idx_m < Q_LEN # TODO generalize and add proper mask support xindex = idx_d + (64*idx_m) + (32768*idx_hq) + (131072*idx_zq) tl.store(out_ptr0 + (tl.broadcast_to(idx_d + (64*idx_m) + (32768*idx_hq), acc.shape)), acc, mask) # TODO dont want to write this if we dont require grad if OUTPUT_LOGSUMEXP: off_hz = tl.program_id(1) l_ptrs = LSE + off_hz * Q_LEN + offs_m lse = m_i + tl.math.log2(l_i) if IS_DIVISIBLE: tl.store(l_ptrs, lse) else: tl.store(l_ptrs, lse, mask=offs_m < Q_LEN) @triton.jit def forward_inner( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets used as inputs to score_mod & mask_mod # of size [BLOCK_M, BLOCK_N] or scalar. off_z, off_h, offs_m, offs_n, # blocksparse data kv_indices, kv_num_blocks, # start kv and end kv block block_n_start, block_n_end, MATMUL_PRECISION, IS_FULL_BLOCKS, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_MULTIPLE: tl.constexpr = (SPARSE_KV_BLOCK_SIZE // BLOCK_N) RCP_LN2: tl.constexpr = 1.44269504 if PRESCALE_QK: q = (q * SM_SCALE * RCP_LN2).to(MATMUL_PRECISION) # loop over k, v and update accumulator until block_n_end for start_n in range(block_n_start, block_n_end): if IS_DIVISIBLE: acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, ) else: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen, # it's on par or slightly faster than only applying to the last block in fwd. # However, we choose different strategy for bwd, where we only apply mod & mask # to the last block because it's faster a lot. acc, l_i, m_i = forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=True, ) # update pointers offset = get_offset_for_next_block( start_n, kv_indices, kv_num_blocks, SPARSE_KV_BLOCK_SIZE, SPARSE_KV_MULTIPLE, BLOCK_N ) V_block_ptr = tl.advance(V_block_ptr, (offset, 0)) K_block_ptr = tl.advance(K_block_ptr, (0, offset)) offs_n = offs_n + offset return acc, l_i, m_i @triton.jit def get_offset_for_next_block(loop_iter, col_indices, total_blocks, SPARSE_BLOCK, SPARSE_BLOCK_MULTIPLE, BLOCK): cur_block_idx = loop_iter // SPARSE_BLOCK_MULTIPLE cur_block = tl.load(col_indices + cur_block_idx, eviction_policy="evict_last") next_block = tl.load(col_indices + cur_block_idx + 1, eviction_policy="evict_last", mask=cur_block_idx + 1 < total_blocks) needs_jump = (loop_iter + 1) % SPARSE_BLOCK_MULTIPLE == 0 jump_to_block = (next_block - cur_block ) * SPARSE_BLOCK - (SPARSE_BLOCK_MULTIPLE - 1) * BLOCK offset = jump_to_block * needs_jump + (1 - needs_jump) * BLOCK return offset @triton.jit def forward_block_mn( arg_Q, arg_K, arg_V, arg_LSE, arg_KV_NUM_BLKS, arg_KV_IDX, arg_FULL_KV_NUM_BLKS, arg_FULL_KV_IDX, out_ptr0, q, K_block_ptr, V_block_ptr, Q_LEN, KV_LEN, # accumulated values acc, l_i, m_i, # Offsets off_z, off_h, offs_m, offs_n, MATMUL_PRECISION, RCP_LN2, IS_FULL_BLOCKS, CHECK_BLOCK_BOUNDARY=False, ): # Redefines all kernel parameters (BLOCK_M, etc.) so we don't need to plumb them all through ROWS_GUARANTEED_SAFE : tl.constexpr = False PRESCALE_QK : tl.constexpr = False OUTPUT_LOGSUMEXP : tl.constexpr = False FLOAT32_PRECISION : tl.constexpr = 'ieee' IS_DIVISIBLE : tl.constexpr = True SM_SCALE : tl.constexpr = 0.125 GQA_SHARED_HEADS : tl.constexpr = 1 HAS_FULL_BLOCKS : tl.constexpr = True QK_HEAD_DIM : tl.constexpr = 64 V_HEAD_DIM : tl.constexpr = 64 BLOCK_M : tl.constexpr = 128 BLOCK_N : tl.constexpr = 32 SPARSE_Q_BLOCK_SIZE : tl.constexpr = 128 SPARSE_KV_BLOCK_SIZE : tl.constexpr = 128 # -- load k -- if IS_DIVISIBLE: k = tl.load(K_block_ptr) else: k = tl.load(K_block_ptr, boundary_check=(1,), padding_option = "zero") # -- compute qk --- qk = tl.dot(q, k, input_precision=FLOAT32_PRECISION) # TODO: use cuda matmul when q_len <= 2. if not PRESCALE_QK: qk *= SM_SCALE # ~~~~~~~~~~~~~~~~~~~ Apply score modification ~~~~~~~~~~~~~~~~~~~ if CHECK_BLOCK_BOUNDARY: # If this is the last block of a non divisible seqlen, we still need to load [BLOCK_M, BLOCK_N] elements, # which is larger than the actual number of elements. To avoid access memory out of bound, # we need to mask out the elements that are out of Q_LEN & KV_LEN. m = offs_m % Q_LEN n = offs_n % KV_LEN else: m = offs_m n = offs_n post_mod_scores = (qk) if CHECK_BLOCK_BOUNDARY: # Mask out the elements that are out of the KV_LEN for non divisible seqlen. post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf")) if not IS_FULL_BLOCKS: tmp0 = (m) >= (n) mask_mod_output = tmp0 if CHECK_BLOCK_BOUNDARY: mask_mod_output = tl.where(offs_n < KV_LEN, mask_mod_output, float("-inf")) # apply mask for partially unmasked blocks post_mod_scores = tl.where(mask_mod_output, post_mod_scores, float("-inf")) # TODO: In the case that score_mod is linear, this can be LICMed if not PRESCALE_QK: post_mod_scores *= RCP_LN2 # ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ # -- compute scaling constant --- m_ij = tl.maximum(m_i, tl.max(post_mod_scores, 1)) if not ROWS_GUARANTEED_SAFE: masked_out_rows = (m_ij == float("-inf")) m_ij_masked = tl.where(masked_out_rows, 0, m_ij) else: m_ij_masked = m_ij alpha = tl.math.exp2(m_i - m_ij_masked) p = tl.math.exp2(post_mod_scores - m_ij_masked[:, None]) # NB: l_i update is pulled up here since it's a bit faster # NB: For headdim=256, it's faster to move it back down to after m_i = # m_ij l_i = l_i * alpha + tl.sum(p, 1) # # -- scale and update acc -- acc = acc * alpha[:, None] if IS_DIVISIBLE: v = tl.load(V_block_ptr) else: v = tl.load(V_block_ptr, boundary_check=(0,), padding_option = "zero") acc = tl.dot(p.to(MATMUL_PRECISION), v, acc, input_precision=FLOAT32_PRECISION) # -- update m_i m_i = m_ij return acc, l_i, m_i ''', device_str='cuda') meta0 = {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False, 'FLOAT32_PRECISION': "'ieee'", 'IS_DIVISIBLE': True, 'SM_SCALE': 0.125, 'GQA_SHARED_HEADS': 1, 'HAS_FULL_BLOCKS': True, 'QK_HEAD_DIM': 64, 'V_HEAD_DIM': 64, 'BLOCK_M': 128, 'BLOCK_N': 32, 'SPARSE_Q_BLOCK_SIZE': 128, 'SPARSE_KV_BLOCK_SIZE': 128} async_compile.wait(globals()) del async_compile def call(args): arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1 = args args.clear() assert_size_stride(arg0_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg1_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg2_1, (1, 4, 512, 64), (131072, 32768, 64, 1)) assert_size_stride(arg3_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg4_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg5_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg6_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg7_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg8_1, (1, 1, 16, 16), (256, 256, 16, 1)) assert_size_stride(arg9_1, (1, 1, 16), (16, 16, 1)) assert_size_stride(arg10_1, (1, 1, 16, 16), (256, 256, 16, 1)) buf0 = empty_strided_cuda((1, 4, 512), (2048, 512, 1), torch.float32) with torch.cuda._DeviceGuard(0): torch.cuda.set_device(0) buf1 = empty_strided_cuda((1, 4, 512, 64), (131072, 32768, 64, 1), torch.float32) # Topologically Sorted Source Nodes: [flex_attention], Original ATen: [] stream0 = get_raw_stream(0) triton_tem_fused_0.run(arg0_1, arg1_1, arg2_1, buf0, arg3_1, arg4_1, arg5_1, arg6_1, buf1, grid=torch._inductor.kernel.flex_attention.flex_attention_grid(1, 4, 512, 64, meta0), stream=stream0) del arg0_1 del arg1_1 del arg2_1 del arg3_1 del arg4_1 del arg5_1 del arg6_1 del buf0 return (buf1, ) def benchmark_compiled_module(times=10, repeat=10): from torch._dynamo.testing import rand_strided from torch._inductor.utils import print_performance arg0_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg1_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg2_1 = rand_strided((1, 4, 512, 64), (131072, 32768, 64, 1), device='cuda:0', dtype=torch.float32) arg3_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg4_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg5_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg6_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg7_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg8_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) arg9_1 = rand_strided((1, 1, 16), (16, 16, 1), device='cuda:0', dtype=torch.int32) arg10_1 = rand_strided((1, 1, 16, 16), (256, 256, 16, 1), device='cuda:0', dtype=torch.int32) fn = lambda: call([arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1]) return print_performance(fn, times=times, repeat=repeat) if __name__ == "__main__": from torch._inductor.wrapper_benchmark import compiled_module_main compiled_module_main('None', benchmark_compiled_module) V1003 10:11:10.314000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0c6c8a1109a650b2f61d8b8a52500c9f"} { "name": "code_gen", "ts": 1727975470313939.2, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.314000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "3d6e3c947033eed0a9cadd2013455900"} { "name": "GraphLowering.compile_to_module", "ts": 1727975470314540.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.316000 2235078 torch/_dynamo/utils.py:1020] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "17b87920e3cd4803f6a7235280e07423"} { "name": "fx_graph_cache_miss", "ts": 1727975463815134.0, "args": { "key": "flfpwf422lzhb6yszystfjdqdgiiqyds6rdbqc57o3a6n4uehjgp", "components": [ "[nhtxa6qinb75ty5rqjx3bdkokm2s3bz3nrhke3zmvm73qypm2md] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False" ], "cache_event_time": 1727975463815134085, "cache_state": "miss", "time_taken_ns": 6500727742 }, "ph": "i", "cat": "dynamo_timed", "tid": 0, "pid": 0, "s": "p" } V1003 10:11:10.317000 2235078 torch/_inductor/codecache.py:1463] {"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "2b8d4494cf5a9a8e000f402106cf2ba5"} {"key": "flfpwf422lzhb6yszystfjdqdgiiqyds6rdbqc57o3a6n4uehjgp", "components": ["[nhtxa6qinb75ty5rqjx3bdkokm2s3bz3nrhke3zmvm73qypm2md] gm: (\n (sdpa_score0): ()\n (sdpa_mask0): ()\n)\n\n\n\ndef forward(self, arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1):\n sdpa_score0 = self.sdpa_score0\n sdpa_mask0 = self.sdpa_mask0\n flex_attention = torch.ops.higher_order.flex_attention(arg0_1, arg1_1, arg2_1, sdpa_score0, (arg3_1, arg4_1, arg5_1, arg6_1, arg7_1, arg8_1, arg9_1, arg10_1, 128, 128, sdpa_mask0), 0.125, {'ROWS_GUARANTEED_SAFE': False, 'PRESCALE_QK': False, 'OUTPUT_LOGSUMEXP': False}, (), ()); arg0_1 = arg1_1 = arg2_1 = sdpa_score0 = arg3_1 = arg4_1 = arg5_1 = arg6_1 = arg7_1 = arg8_1 = arg9_1 = arg10_1 = sdpa_mask0 = None\n getitem = flex_attention[0]; flex_attention = None\n return (getitem,)\n \n# To see more debug info, please use `graph_module.print_readable()`", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[0]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[1]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[avf2u3luxvyabchjhbddapcjn5gev47wfdtkrprayuhv6lf2z6u] example_inputs[2]: TensorMetadata(dtype=torch.float32, shape=torch.Size([1, 4, 512, 64]), stride=(131072, 32768, 64, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[3]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[4]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[5]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[6]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[7]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[8]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[zsk3gejenkcvvwhiyk36u5zdnlrcs6wgy3pina3csuierfd2zri] example_inputs[9]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16]), stride=(16, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[hnbjjzmb63q27mbr22eubaelyb423burv27meouma6ccysmwu6g] example_inputs[10]: TensorMetadata(dtype=torch.int32, shape=torch.Size([1, 1, 16, 16]), stride=(256, 256, 16, 1), device=device(type='cuda', index=0), layout=torch.strided, memory_format=torch.contiguous_format, storage_offset=0, storage_bytes=None, requires_grad=False, is_quantized=False, is_conj=False, is_neg=False, is_inference=False, is_sparse=False, is_coalesced=None, dense_dim=None, sparse_dim=None)", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[aot_mode]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[cpp_wrapper]: False", "[xq2hdkbfkbcuye6rgtypayrkhqf4cntij2dsd24rei3lsknakkf] fx_kwargs[cudagraphs]: BoxedBool(value=False)", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[extern_node_serializer]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] fx_kwargs[is_backward]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] fx_kwargs[is_inference]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] fx_kwargs[layout_opt]: None", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] fx_kwargs[static_input_idxs]: []", "[f44ag5aflby2bkxl7a4k6whljrk7jat7bmreuxklei4p3czhk7p] fx_kwargs[user_visible_outputs]: {'getitem': None}", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inputs_to_check[0]: 0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inputs_to_check[1]: 1", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inputs_to_check[2]: 2", "[kcuxe2zwm3mzv2uk6adm6iskoy35bqfv725twacrdewod2dbl5d] inputs_to_check[3]: 3", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inputs_to_check[4]: 4", "[qs5hilycp4ew4ivtc7m5jaxp7q4pm5slioxw3fi3ur6ei65ybz4] inputs_to_check[5]: 5", "[agkvbkaha53nbz3aeeuhvxjvvc4glhfjofzkg6g2qjoo2e5otcx] inputs_to_check[6]: 6", "[j3s5elu6itwgjafc7rzhy4whrbufl6kfmlufjhh25grt643bk5f] inputs_to_check[7]: 7", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inputs_to_check[8]: 8", "[qlgfiyqewrmkgqth2qm6wkq2ja5lzkapg3ypgnvoyfqqnidaoj3] inputs_to_check[9]: 9", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inputs_to_check[10]: 10", "[du4vyrfyozrfxcf6kk6ma7oqwatapifazeelfsawmsiu6gjdtxp] deterministic_algorithms_settings: (False, False, True)", "[qiptf2633zubseuei4bkisoq3not35l6lud6p23p4qmcsxiw2uq] cuda_matmul_settings: (False, True, True)", "[7uhqwjfn75ek3woo3k7em2mluon5hx2ojvzlevlvjvz6xfxjhzl] torch_version: ", "[c3z7bmoxyo6gl5hi47v6dc7jwsl55b3asd75nr25uyengi5ah3p] system_info[device]: {'name': 'NVIDIA PG509-210'}", "[3fb7kae6ogkdd4zcm3fkjoipdpybxhn4aoxzv7z7xsfwq233e4l] system_info[version]: {'triton': '3.1.0+5fe38ffd73dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-835d4fc33500e1accafc5c5e00f4f73d87432c114860c04b68849bf6f942b8e5-dc767c8fadcf23ea82d79e257c37d44077eae7f681cf967565fd43e9c017937b-23d635e690d670bf61798e1259674b78c0ed5ba222ab6a455f329f27a758fc2d-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-20b017e9c4d858ab05e783f77df50b86c6d6eee5d79f3f4b158562b4a54f8443-f44338a31e0534290b08653050804c3fabbde403a6d3004ae04f0c28495f0802-e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855-a979896b9c0acfd41dd953b90bdc4b10968f7c0b45a286eae3f829aaddb2bb55-da771298f7bc45d24a61f35ef51742304421df1ab49d50bf1fc510dd5a46ea4b-a8fb7be728d460b7ec64ab62edb8af1bbca8994fd718cde7178d46bad64530a1-71330f394e584b0df29595d49f6ac8ac0c5503db9147090dc58ad888cebac7be-f24adfd52383f7866791ebaa5d45a5d2cc826e56ee2fd285f438e85d201fe643-a34be0d3ae4b3ac9aede195cfda42f8a0a097b2bc9642fb59673ce6b3b607f10-36130a37af1b19a0dec569aa08d30b00c74c8f02b6b632999d86dea169146792-36d42f0429aae027cb985b53b9abc616fae4dad9e0ea03953e1e9fb46d0fb9a0-e5d2cb724c08d0ef4130f3ba858d22cf21f834bfd970a5388aa6ad2a6bab91f9', 'cuda': '12.0'}", "[z5x2bdhir5lzlbti73vdbfulnuu5vinzpwgmmgf4rjb775tzl3h] system_info[hash]: 9698c97edde4a99a2f3b54bbd0db5291bbcdb75c83acb376ccff61fb0bf0ac1a", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[TYPE_CHECKING]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[abi_compatible]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aggressive_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[allow_buffer_reuse]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[allow_stack_allocation]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[always_keep_tensor_constants]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_compile]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.debug_dump_consts_bin]: False", "[ngkkx5e6z7erl6da23zb2cmsctz4yvaqyameyg5hbqln4wrhh7x] inductor_config[aot_inductor.debug_intermediate_value_printer]: 0", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[aot_inductor.filtered_kernel_names]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.force_mmap_weights]: False", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[aot_inductor.metadata]: {}", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.output_path]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.package_cpp_only]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_in_spec]: ", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[aot_inductor.serialized_out_spec]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[aot_inductor.use_runtime_constant_folding]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[assert_indirect_indexing]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[assume_aligned_inputs]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[autoheuristic_collect]: ", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[autoheuristic_log_path]: DEFAULT", "[jwbrgxes7vjqumngs5hyj6gn5nytv2whnppnzngvaagfmawhkkd] inductor_config[autoheuristic_use]: mixed_mm", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[autotune_fallback_to_aten]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_in_subproc]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_local_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_multi_device]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[autotune_remote_cache]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[b2b_gemm_pass]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[batch_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_combo_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[benchmark_harness]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[benchmark_kernel]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[bw_outputs_user_visible]: True", "[b4ha3ravs3qv237q65hpfqegbnoww7tf2ahcbu2i7xo6te5spqs] inductor_config[c_shim_version]: 2", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[check_stack_no_cycles_TESTING_ONLY]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernel_allow_mixed_sizes]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernel_foreach_dynamic_shapes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[combo_kernels]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[combo_kernels_autotune]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[comment_origin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[comprehensive_padding]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[compute_all_bounds]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[constant_and_index_propagation]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[conv_1x1_as_mm]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_check_all_directions]: False", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[coordinate_descent_search_radius]: 1", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[coordinate_descent_tuning]: False", "[c7zj4qytmety6keurs3hsh5wn7foxp3dqx4kym2ucszzcb2ngrf] inductor_config[cpp.cxx]: (None, 'g++')", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[cpp.descriptive_names]: original_aten", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.dynamic_threads]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_floating_point_contract_flag]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_kernel_profile]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_loop_tail_vec]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.enable_tiling_heuristics]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp.enable_unsafe_math_opt_flag]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.fallback_scatter_reduce_sum]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_cache_blocking]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cpp.gemm_max_k_slices]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.gemm_thread_factors]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_log1p_bug_TESTING_ONLY]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.inject_relu_bug_TESTING_ONLY]: None", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[cpp.max_horizontal_fusion_size]: 16", "[g7rrnbg5yonzux3cfj5ovre5lob3ayda7qcfpxjvtwmiz4uicii] inductor_config[cpp.min_chunk_size]: 4096", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.no_redundant_loops]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.simdlen]: None", "[sz3im5ogc6asp7g4uqocnovype63tkdexzfrniv6hn2oank3biu] inductor_config[cpp.threads]: -1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cpp.vec_isa_ok]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cpp.weight_prepack]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cpp_wrapper]: False", "[bsvfcwwoczx2rlkdz2eta6doujsymyihmi46hhwk6clrrvwcb6m] inductor_config[cpu_backend]: cpp", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.arch]: None", "[tvyftmtdmezlejo2xllu7awzv4pzc4vm4fub4b3gpl5jptjkosi] inductor_config[cuda.compile_opt_level]: -O1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cuda_cxx]: None", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[cuda.cutlass_backend_min_gemm_size]: 1", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_max_profiling_configs]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.cutlass_op_allowlist_regex]: None", "[lwkz5chtpji756gurqw4foijfi7zfgljtnn5nmnvdi2skpt4mgh] inductor_config[cuda.cutlass_op_denylist_regex]: pingpong", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_cuda_lto]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_debug_info]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.enable_ptxas_info]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[cuda.generate_test_runner]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[cuda.use_fast_math]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[cuda.version]: None", "[caw4ly2z672k6kjfahoxwpajp5idhhtrpgf3ma2clylcp7c7aid] inductor_config[cuda_backend]: triton", "[pikr7bbcoixfzftsazp5ggufhdklj24babfry77bl4nuvyrrcp4] inductor_config[custom_op_default_layout_constraint]: needs_fixed_stride_order", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[dce]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_index_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[debug_ir_traceback]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[decompose_mem_bound_mm]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[developer_warnings]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[disable_cpp_codegen]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_padding_cpu]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[disable_progress]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[dynamic_scale_rblock]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[efficient_conv_bn_eval_fx_passes]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[emulate_precision_casts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[enable_auto_functionalized_v2]: False", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[enabled_metric_tables]: ", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[epilogue_fusion]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[epilogue_fusion_first]: False", "[lxxtoqhcoepwfokeiibd575gnxo3uzwiv4hmpomlwkpzqz3qzsh] inductor_config[estimate_op_runtime]: default", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[external_matmul]: []", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fallback_random]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_disable_caches]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_fuse_int_mm_with_mul]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_layout_optimization]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_same_precision]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[force_shape_pad]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[freezing_discard_parameters]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[fx_graph_cache]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[fx_graph_remote_cache]: False", "[62lrdx35b7hnumwb7mp5oc5y5csm2abylvtdzfloct3noaqov3n] inductor_config[fx_passes_numeric_check]: {'pre_grad': False, 'post_grad': False, 'precision': 0.0001, 'num_iterations': 1, 'requires_optimizer': True}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[generate_intermediate_hooks]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[global_cache_dir]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[group_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.asserts]: False", "[ljhgflgihidopsfsdcbqynv27nceykby3nutyd5jlcpq7n6e7l4] inductor_config[halide.cpu_target]: host", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.debug]: False", "[wx7vmsmrdpk5ue2txlywp3lj3faqmdjphs5fgg2ehzsyno7uovg] inductor_config[halide.gpu_target]: host-cuda", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[halide.scan_kernels]: False", "[k5ogk6345jvklsnu7g2njqstiz2g6pm5wmqpgg3kasrmuqwjvl6] inductor_config[halide.scheduler_cpu]: Adams2019", "[svgytlua5wcyeia7wq7e6zgh5tsueikrnzchmdmouvmkpfsc2zq] inductor_config[halide.scheduler_cuda]: Anderson2021", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[implicit_fallbacks]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[inplace_buffers]: True", "[5fxczt3ciyxitdhizb7sfsgn7fhpczcqsngttnt5ot2wyctk7co] inductor_config[inter_node_bw]: 25", "[yezuzjtg4h3jjur4jwtwiehbyixa7eonq4tqsqmwqve2lvvmrem] inductor_config[intra_node_bw]: 300", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[is_nightly_or_source]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[is_predispatch]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[joint_custom_pre_pass]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[joint_graph_constant_folding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[keep_output_stride]: True", "[j6c55jha5r2sdys2rwq7uqhtleea5dgjcye7nicfgft36v7xfvp] inductor_config[kernel_name_max_ops]: 10", "[4p2fdjlvxrcw7c7fvzm5huhtqxnro4kvkx56f7p5zyrxqkwooov] inductor_config[layout_opt_default]: 1", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[layout_optimization]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[loop_ordering_after_fusion]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune]: False", "[uqlsbif4zxd75vt522p52txyuguieipi2lwz5g5awt56lccqk7s] inductor_config[max_autotune_conv_backends]: ATEN,TRITON", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_gemm]: False", "[2y7luesktjrque3nr7qtxnum2mkbeegzdrsvkm3rvdlhqboajhx] inductor_config[max_autotune_gemm_backends]: ATEN,TRITON,CPP", "[jvchmi66fvqzlemhr5fcqorz5trfdtdalzfagtj2aolmimwqhdq] inductor_config[max_autotune_gemm_search_space]: DEFAULT", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[max_autotune_pointwise]: False", "[bh33ranllcgilhgmgr3qvygzxjm6isq5iexnfm3zx6fnr2zwlp2] inductor_config[max_autotune_subproc_graceful_timeout_seconds]: 1.0", "[iglov24t7x5ruci344aer2tm6nqshi4veuw4wxlssxtu46cx76m] inductor_config[max_autotune_subproc_result_timeout_seconds]: 60.0", "[pwoh5aypf4fxbntdvwt67rppxorqos6xr3w7qzeun6kblbfg2ga] inductor_config[max_autotune_subproc_terminate_timeout_seconds]: 2.0", "[aghvyrrgwvxijco2pk5wzc3cgmmthrbmgxitiibxuuscxdwrjd3] inductor_config[max_epilogue_benchmarked_choices]: 1", "[jykiys6ynafs3zdylwa5ggq6j655mxeh42d6mtdi22gffkrmiac] inductor_config[max_fusion_size]: 64", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[max_pointwise_cat_inputs]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[memory_planning]: False", "[x75won4jmsgeb63pcvwr2y4eteyzzdhmf5rv6xhjppie4hx2yu5] inductor_config[memory_pool]: intermediates", "[v2td5s4lnsvyxvaevy4chx6kc5h3mm2axazbgwimqule5zrzao7] inductor_config[mixed_mm_choice]: heuristic", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[nan_asserts]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[optimize_scatter_upon_const_tensor]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_channels_last]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[pad_outputs]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[padding_alignment_bytes]: 128", "[dnnw5ks3yxrp7mwvihb2hh4tqx35ye637xt33x64kw4fvz2nyzg] inductor_config[padding_stride_threshold]: 1024", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pattern_matcher]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[permute_fusion]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[pick_loop_orders]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_post_pass]: None", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[post_grad_custom_pre_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[post_grad_fusion_options]: {}", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[pre_grad_custom_pass]: None", "[4bryyl4ahh5whyg3zwqebpwmjnx6w77nqgqbdjlowju6lkqtn7w] inductor_config[pre_grad_fusion_options]: {}", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[profile_bandwidth_output]: None", "[v3hzzlv4tjgvp3pyhmzagjd25orl6n7nynoa7svlhhwk73b7u3c] inductor_config[profile_bandwidth_regex]: ", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profile_bandwidth_with_do_bench_using_profiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[profiler_mark_wrapper_call]: False", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[realize_acc_reads_threshold]: 8", "[rr5m5hsocoyodldz7vcvaizdwvm2rt34evmqdxvng7wz3tufvo6] inductor_config[realize_opcount_threshold]: 30", "[lkkae3meylaixfif4thncru4hjqeaislawjoghffrbwuscaagei] inductor_config[realize_reads_threshold]: 4", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_compute_comm_overlap]: False", "[ssupi7bu3rrhdpg2jyegzncu3kg3nnhklyliqvutaxgs7y7k3dx] inductor_config[reorder_for_compute_comm_overlap_passes]: ['reorder_compute_for_overlap', 'sink_waits', 'raise_comms']", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[reorder_for_locality]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[reorder_for_peak_memory]: False", "[h25wqx6vliw4j5rtzzbv6latydxyei3deyg6v7wzvnzryfktuki] inductor_config[rocm.arch]: []", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.ck_dir]: None", "[oartxnko2l7d67tzwwm2otcumaut3n4wwcfgz3o377hmcveu5ft] inductor_config[rocm.ck_supported_arch]: ['gfx90a', 'gfx940', 'gfx941', 'gfx942']", "[klfqjprnpfhcdurgvuikvc4rpd5ynkpk77toousr5h3u5roty6p] inductor_config[rocm.compile_opt_level]: -O2", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.flush_denormals]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.is_debug]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.n_max_profiling_configs]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.print_kernel_resource_usage]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[rocm.rocm_home]: None", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.save_temps]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[rocm.use_fast_math]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[rocm.use_preselected_instances]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[save_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[search_autotune_cache]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[shape_padding]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[size_asserts]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[sleep_sec_TESTING_ONLY]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_cat_fx_passes]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[split_reductions]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[static_weight_shapes]: True", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.autotune_at_compile_time]: None", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_cublasLt]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.autotune_pointwise]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.codegen_upcast_to_fp32]: True", "[tuax46wac7rfv2trf5gcps6vleo3cq44lbnrdxtprvo3ljjaddj] inductor_config[triton.cudagraph_dynamic_shape_warn_limit]: 50", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_skip_dynamic_graphs]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_support_input_mutation]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.cudagraph_trees]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraph_trees_history_recording]: False", "[ljdqgtysl3vdf7j6attlz5gmjg2ncihnveojfyubosplmkrjgra] inductor_config[triton.cudagraph_unexpected_rerecord_limit]: 128", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.cudagraphs]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_graph]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.debug_sync_kernel]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.dense_indexing]: False", "[yrty22bseefglnysuoec4ji7j2rnaggdj3g33zzj7avogwfmgdw] inductor_config[triton.descriptive_names]: original_aten", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.divisible_by_16]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.fast_path_cudagraph_asserts]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraph_sync]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.force_cudagraphs_warmup]: False", "[tquy2we2efmowuj4wuqzcfcfdcrkzkzmwdae6hprj7fa64jpusq] inductor_config[triton.inject_relu_bug_TESTING_ONLY]: None", "[pr5nr4a7dthirgd2ljo3d2xakc63ywxugusu6mkmr6gmpeliyib] inductor_config[triton.max_tiles]: 2", "[fv6slhtedtydps5s5u2etitscliblzcidyitqf7krsv4e23fzk6] inductor_config[triton.min_split_scan_rblock]: 256", "[vrl5ktomgtzox5xucd3np6vug3vyj6hwwzahqijuwpmamlv7ohi] inductor_config[triton.multi_kernel]: 0", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.persistent_reductions]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.prefer_nd_tiling]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.skip_cudagraph_warmup]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.slow_path_cudagraph_asserts]: True", "[ebt2ncs4f5y7dn7btzi76mnouepvzad474tmp5iju4wiuumjl4s] inductor_config[triton.spill_threshold]: 16", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.store_cubin]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_pointwise_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.tiling_prevents_reduction_fusion]: True", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[triton.unique_kernel_names]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[triton.use_block_ptr]: False", "[vzzema5ityqj2wepdmkulue7q5pcevdr5h27oxxutf35d4tjume] inductor_config[triton_kernel_default_layout_constraint]: flexible_layout", "[wft6ljqsfr3x4m7fa5zuyb7cwknky4irrxz4bjr6uzr2yiopxqj] inductor_config[unbacked_symint_fallback]: 8192", "[yttmfmxblgcbsvbokguzowcorrcxz5uunxtcvsbe6nijgcx45he] inductor_config[unroll_reductions_threshold]: 8", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[unsafe_ignore_unsupported_triton_autotune_args]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[use_minimal_arrayref_interface]: False", "[cev5uo2jlwdhw2uyzcm7vr6cl23azjfw437f5r5lskm7spucos6] inductor_config[use_mixed_mm]: True", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[verbose_progress]: False", "[esstihe2nyydk4mhzpvox3qkajyu5y5t23hk3fi2me7jn75xi3o] inductor_config[warn_mix_layout]: False"], "cache_event_time": 1727975463815134085, "cache_state": "miss", "time_taken_ns": 6500727742} V1003 10:11:10.318000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "59401c365876e3190b828c9baef54d73"} { "name": "inductor_compile", "ts": 1727975470318146.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.318000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "75e1054f04f269f6411028767770c53c"} { "name": "compile_fx_inner", "ts": 1727975470318439.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.319000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "54543856a0e4177bd66c38146fda7120"} { "name": "compile_fx..fw_compiler_base", "ts": 1727975470319104.2, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.322000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0017891a0ab9c15e62e1b8125198b31d"} { "name": "create_aot_dispatcher_function", "ts": 1727975470322488.8, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.323000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "0927305357c25444fda9907c2d33a10f"} { "name": "backend_compile", "ts": 1727975470323094.0, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.323000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "f3906ae17fd7aa60928cc5980667652c"} { "name": "OutputGraph.call_user_compiler", "ts": 1727975470323375.2, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.354000 2235078 torch/_dynamo/guards.py:2311] {"dynamo_cpp_guards_str": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "1b9c89081ea4384ef6cc3f1c7563841f"} TREE_GUARD_MANAGER: +- RootGuardManager | +- DEFAULT_DEVICE: utils_device.CURRENT_DEVICE == None # _dynamo/output_graph.py:471 in init_ambient_guards | +- GLOBAL_STATE: ___check_global_state() | +- TORCH_FUNCTION_MODE_STACK: ___check_torch_function_mode_stack() | +- GuardManager: source=L['k'], accessed_by=DictGetItemGuardAccessor(k) | | +- TYPE_MATCH: ___check_type_id(L['k'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['k'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['k'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING: check_no_aliasing(L['k'], L['q'], L['v'], L['block_mask'].q_indices, L['block_mask'].kv_indices, L['block_mask'].q_num_blocks, L['block_mask'].kv_num_blocks, L['block_mask'].full_q_indices, L['block_mask'].full_kv_indices, L['block_mask'].full_q_num_blocks, L['block_mask'].full_kv_num_blocks) | +- GuardManager: source=L['q'], accessed_by=DictGetItemGuardAccessor(q) | | +- TYPE_MATCH: ___check_type_id(L['q'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['q'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['q'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['v'], accessed_by=DictGetItemGuardAccessor(v) | | +- TYPE_MATCH: ___check_type_id(L['v'], 82028112) | | +- TENSOR_MATCH: check_tensor(L['v'], Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.float32, device=0, requires_grad=False, size=[1, 4, 512, 64], stride=[131072, 32768, 64, 1]) | | +- NO_HASATTR: hasattr(L['v'], '_dynamo_dynamic_indices') == False | | +- NO_TENSOR_ALIASING | +- GuardManager: source=L['block_mask'], accessed_by=DictGetItemGuardAccessor(block_mask) | | +- TYPE_MATCH: ___check_type_id(L['block_mask'], 387600320) | | +- GuardManager: source=L['block_mask'].mask_mod, accessed_by=GetAttrGuardAccessor(mask_mod) | | | +- GuardManager: source=L['block_mask'].mask_mod.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].mask_mod.__code__, 140413271880128) | | +- GuardManager: source=L['block_mask'].q_indices, accessed_by=GetAttrGuardAccessor(q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE, accessed_by=GetAttrGuardAccessor(BLOCK_SIZE) | | | +- TYPE_MATCH: ___check_type_id(L['block_mask'].BLOCK_SIZE, 8815232) | | | +- LENGTH_CHECK: len(L['block_mask'].BLOCK_SIZE) == 2 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[0], accessed_by=TupleGetItemGuardAccessor(0) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[0] == 128 | | | +- GuardManager: source=L['block_mask'].BLOCK_SIZE[1], accessed_by=TupleGetItemGuardAccessor(1) | | | | +- EQUALS_MATCH: L['block_mask'].BLOCK_SIZE[1] == 128 | | +- GuardManager: source=L['block_mask'].kv_indices, accessed_by=GetAttrGuardAccessor(kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].q_num_blocks, accessed_by=GetAttrGuardAccessor(q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].kv_num_blocks, accessed_by=GetAttrGuardAccessor(kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_indices, accessed_by=GetAttrGuardAccessor(full_q_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_indices, accessed_by=GetAttrGuardAccessor(full_kv_indices) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_indices, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16, 16], stride=[256, 256, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_indices, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_q_num_blocks, accessed_by=GetAttrGuardAccessor(full_q_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_q_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_q_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].full_kv_num_blocks, accessed_by=GetAttrGuardAccessor(full_kv_num_blocks) | | | +- TENSOR_MATCH: check_tensor(L['block_mask'].full_kv_num_blocks, Tensor, DispatchKeySet(CUDA, BackendSelect, ADInplaceOrView, AutogradCUDA), torch.int32, device=0, requires_grad=False, size=[1, 1, 16], stride=[16, 16, 1]) | | | +- NO_HASATTR: hasattr(L['block_mask'].full_kv_num_blocks, '_dynamo_dynamic_indices') == False | | | +- NO_TENSOR_ALIASING | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=GetAttrGuardAccessor(as_tuple) | | | +- GuardManager: source=L['block_mask'].as_tuple, accessed_by=FuncDefaultsGuardAccessor | | | | +- GuardManager: source=L['block_mask'].as_tuple.__defaults__[0], accessed_by=GetItemGuardAccessor(0) | | | | | +- ID_MATCH: ___check_obj_id(L['block_mask'].as_tuple.__defaults__[0], 8911040) | +- GuardManager: source=L['score_mod2'], accessed_by=DictGetItemGuardAccessor(score_mod2) | | +- GuardManager: source=L['score_mod2'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['score_mod2'].__code__, 140413275216112) | +- GuardManager: source=L['flex_attention'], accessed_by=DictGetItemGuardAccessor(flex_attention) | | +- GuardManager: source=L['flex_attention'].__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__code__, 387082992) | | +- GuardManager: source=L['flex_attention'], accessed_by=FuncDefaultsGuardAccessor | | | +- GuardManager: source=L['flex_attention'].__defaults__[2], accessed_by=GetItemGuardAccessor(2) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[2], 8825760) | | | +- GuardManager: source=L['flex_attention'].__defaults__[3], accessed_by=GetItemGuardAccessor(3) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[3], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[4], accessed_by=GetItemGuardAccessor(4) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[4], 8910592) | | | +- GuardManager: source=L['flex_attention'].__defaults__[5], accessed_by=GetItemGuardAccessor(5) | | | | +- ID_MATCH: ___check_obj_id(L['flex_attention'].__defaults__[5], 8825760) | +- GuardManager: source=G, accessed_by=GlobalsGuardAccessor | | +- GuardManager: source=G['__builtins_dict___10'], accessed_by=DictGetItemGuardAccessor(__builtins_dict___10) | | | +- GuardManager: source=G['__builtins_dict___10']['len'], accessed_by=DictGetItemGuardAccessor(len) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['len'], 140413275558816) | | | +- GuardManager: source=G['__builtins_dict___10']['sum'], accessed_by=DictGetItemGuardAccessor(sum) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['sum'], 140413275559936) | | | +- GuardManager: source=G['__builtins_dict___10']['list'], accessed_by=DictGetItemGuardAccessor(list) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['list'], 8844320) | | | +- GuardManager: source=G['__builtins_dict___10']['type'], accessed_by=DictGetItemGuardAccessor(type) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['type'], 8813248) | | | +- GuardManager: source=G['__builtins_dict___10']['tuple'], accessed_by=DictGetItemGuardAccessor(tuple) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['tuple'], 8815232) | | | +- GuardManager: source=G['__builtins_dict___10']['object'], accessed_by=DictGetItemGuardAccessor(object) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['object'], 8813984) | | | +- GuardManager: source=G['__builtins_dict___10']['isinstance'], accessed_by=DictGetItemGuardAccessor(isinstance) | | | | +- ID_MATCH: ___check_obj_id(G['__builtins_dict___10']['isinstance'], 140413275558496) | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_utils_dot__pytree) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'], 140411217627952) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].TreeSpec, accessed_by=GetAttrGuardAccessor(TreeSpec) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].TreeSpec, 84866496) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf, accessed_by=GetAttrGuardAccessor(_is_leaf) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_leaf.__code__, 140411217262720) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, accessed_by=GetAttrGuardAccessor(_LEAF_SPEC) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC, 85171104) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, accessed_by=GetAttrGuardAccessor(type) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.type, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, accessed_by=GetAttrGuardAccessor(context) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.context, 8825760) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes, accessed_by=GetAttrGuardAccessor(num_nodes) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_nodes == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves, accessed_by=GetAttrGuardAccessor(num_leaves) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_leaves == 1 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children, accessed_by=GetAttrGuardAccessor(num_children) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.num_children == 0 | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, accessed_by=GetAttrGuardAccessor(children_specs) | | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs, 8844320) | | | | | +- LENGTH_CHECK: not G['__import_torch_dot_utils_dot__pytree']._LEAF_SPEC.children_specs | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type, accessed_by=GetAttrGuardAccessor(_get_node_type) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._get_node_type.__code__, 140411217262448) | | | +- DictGuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES, accessed_by=GetAttrGuardAccessor(SUPPORTED_NODES) | | | | +- DICT_VERSION: ___dict_version(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES) == 519596 | | | | +- KeyValueManager pair at index=1 | | | | | +- ValueManager: GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]] | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn, accessed_by=GetAttrGuardAccessor(flatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].flatten_fn.__code__, 140411196281984) | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn, accessed_by=GetAttrGuardAccessor(unflatten_fn) | | | | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES[list(G['__import_torch_dot_utils_dot__pytree'].SUPPORTED_NODES.keys())[1]].unflatten_fn.__code__, 140411217182288) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper, accessed_by=GetAttrGuardAccessor(_tree_flatten_helper) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._tree_flatten_helper.__code__, 140411217413040) | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance, accessed_by=GetAttrGuardAccessor(_is_namedtuple_instance) | | | | +- GuardManager: source=G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_utils_dot__pytree']._is_namedtuple_instance.__code__, 140411217412592) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_comptime'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_comptime) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_comptime'], 140410226912176) | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot__dynamo_dot_decorators) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'], 140410226910096) | | | +- GuardManager: source=G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, accessed_by=GetAttrGuardAccessor(is_compiling) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot__dynamo_dot_decorators'].is_compiling, 140410376252096) | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot__utils) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot__utils'], 140409673896784) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, accessed_by=GetAttrGuardAccessor(_SUPPORTED_HEAD_DIMS) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS, 8844320) | | | | +- LENGTH_CHECK: len(G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS) == 10 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0], accessed_by=ListGetItemGuardAccessor(0) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[0] == 2 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1], accessed_by=ListGetItemGuardAccessor(1) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[1] == 4 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2], accessed_by=ListGetItemGuardAccessor(2) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[2] == 8 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3], accessed_by=ListGetItemGuardAccessor(3) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[3] == 16 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4], accessed_by=ListGetItemGuardAccessor(4) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[4] == 32 | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5], accessed_by=ListGetItemGuardAccessor(5) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot__utils']._SUPPORTED_HEAD_DIMS[5] == 64 | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], accessed_by=DictGetItemGuardAccessor(__import_torch_dot_nn_dot_attention_dot_flex_attention) | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'], 140409673895824) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, accessed_by=GetAttrGuardAccessor(math) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math, 140413266939392) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, accessed_by=GetAttrGuardAccessor(sqrt) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].math.sqrt, 140413266943072) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, accessed_by=GetAttrGuardAccessor(torch) | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch, 140413267918368) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, accessed_by=GetAttrGuardAccessor(_dynamo) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo, 140413260098400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static, accessed_by=GetAttrGuardAccessor(mark_static) | | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch._dynamo.mark_static.__code__, 123166432) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, accessed_by=GetAttrGuardAccessor(compiler) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler, 140410826010400) | | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, accessed_by=GetAttrGuardAccessor(is_dynamo_compiling) | | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].torch.compiler.is_dynamo_compiling, 140410826132992) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device, accessed_by=GetAttrGuardAccessor(_validate_device) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_device.__code__, 140409673611088) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, accessed_by=GetAttrGuardAccessor(flex_attention_hop) | | | | +- TYPE_MATCH: ___check_type_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop, 96992544) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__, accessed_by=GetAttrGuardAccessor(__name__) | | | | | +- EQUALS_MATCH: G['__import_torch_dot_nn_dot_attention_dot_flex_attention'].flex_attention_hop.__name__ == 'flex_attention' | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim, accessed_by=GetAttrGuardAccessor(_supported_head_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._supported_head_dim.__code__, 140409673231376) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim, accessed_by=GetAttrGuardAccessor(_validate_embed_dim) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_embed_dim.__code__, 388086512) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input, accessed_by=GetAttrGuardAccessor(_validate_sdpa_input) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._validate_sdpa_input.__code__, 387915104) | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options, accessed_by=GetAttrGuardAccessor(_apply_kernel_options) | | | | +- GuardManager: source=G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, accessed_by=GetAttrGuardAccessor(__code__) | | | | | +- ID_MATCH: ___check_obj_id(G['__import_torch_dot_nn_dot_attention_dot_flex_attention']._apply_kernel_options.__code__, 140409683680752) V1003 10:11:10.355000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "68cb1c0ec8488a404219d0a05ec80df9"} { "name": "entire_frame_compile", "ts": 1727975470355549.8, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.355000 2235078 torch/_dynamo/utils.py:988] {"chromium_event": {}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "eb02aec5687f6fb5c0049bd4411b1b24"} { "name": "_compile.compile_inner", "ts": 1727975470355876.5, "args": { "cache_stats": { "fxgraph_cache_hit": 1, "fxgraph_cache_miss": 2, "fxgraph_cache_bypass": 0 } }, "ph": "E", "cat": "dynamo_timed", "tid": 0, "pid": 0 } V1003 10:11:10.356000 2235078 torch/_dynamo/utils.py:840] {"compilation_metrics": {"compile_id": "1/0", "frame_key": "2", "co_name": "fn2", "co_filename": "/data/users/oulgen/pytorch/test/inductor/test_codecache.py", "co_firstlineno": 385, "cache_size": 0, "accumulated_cache_size": 0, "guard_count": 80, "shape_env_guard_count": 0, "graph_op_count": 11, "graph_node_count": 25, "graph_input_count": 11, "start_time": 1727975463.4963815, "entire_frame_compile_time_s": 6.859049081802368, "backend_compile_time_s": 6.648646593093872, "inductor_compile_time_s": 6.517011404037476, "code_gen_time_s": 6.449295282363892, "fail_type": null, "fail_reason": null, "fail_user_frame_filename": null, "fail_user_frame_lineno": null, "non_compliant_ops": [], "compliant_custom_ops": [], "restart_reasons": [], "dynamo_time_before_restart_s": 0.0, "has_guarded_code": true, "possibly_missed_reinplacing_opportunities": 0, "remote_cache_time_saved_s": 0, "structured_logging_overhead_s": 0.085769178, "config_suppress_errors": false, "config_inline_inbuilt_nn_modules": true, "specialize_float": true}, "frame_id": 1, "frame_compile_id": 0, "attempt": 0}