facebookresearch / facebookresearch/blt
Triton compilation error occurred during inference
- Dominant language
- Python
- Stars
- 2.1k
- Forks
- 196
- PR merge metrics
- No merged PRs in 30d
Description
Please kindly ask the author to help solve the problems encountered during the reasoning process
1.Wrong core issue in brief
Triton compilation error:
NameError('s10 is not defined')
2.After configuring and running eval using a pre trained model, I received this error message:
[rank0]:[W605 21:07:13.811336253 ProcessGroupNCCL.cpp:4115] [PG ID 0 PG GUID 0 Rank 0] using GPU 0 to perform barrier as devices used by this process are currently unknown. This can potentially cause a hang if this rank to GPU mapping is incorrect.Specify device_ids in barrier() to force use of a particular device,or call init_process_group() with a device_id.
WARNING:root:Token embedding dimensions may cause issues - input dim 512 is not divisible by output dim 1280 or vice versa
[rank0]: Traceback (most recent call last):
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/output_graph.py", line 1446, in _call_user_compiler
[rank0]: compiled_fn = compiler_fn(gm, self.example_inputs())
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/repro/after_dynamo.py", line 129, in __call__
[rank0]: compiled_gm = compiler_fn(gm, example_inputs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/__init__.py", line 2235, in __call__
[rank0]: return compile_fx(model_, inputs_, config_patches=self.config)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 1521, in compile_fx
[rank0]: return aot_autograd(
[rank0]: ^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/backends/common.py", line 72, in __call__
[rank0]: cg = aot_module_simplified(gm, example_inputs, **self.kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_functorch/aot_autograd.py", line 1071, in aot_module_simplified
[rank0]: compiled_fn = dispatch_and_compile()
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_functorch/aot_autograd.py", line 1056, in dispatch_and_compile
[rank0]: compiled_fn, _ = create_aot_dispatcher_function(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_functorch/aot_autograd.py", line 522, in create_aot_dispatcher_function
[rank0]: return _create_aot_dispatcher_function(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_functorch/aot_autograd.py", line 759, in _create_aot_dispatcher_function
[rank0]: compiled_fn, fw_metadata = compiler_fn(
[rank0]: ^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_functorch/_aot_autograd/jit_compile_runtime_wrappers.py", line 179, in aot_dispatch_base
[rank0]: compiled_fw = compiler(fw_module, updated_flat_args)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 1350, in fw_compiler_base
[rank0]: return _fw_compiler_base(model, example_inputs, is_inference)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 1421, in _fw_compiler_base
[rank0]: return inner_compile(
[rank0]: ^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 475, in compile_fx_inner
[rank0]: return wrap_compiler_debug(_compile_fx_inner, compiler_name="inductor")(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/repro/after_aot.py", line 85, in debug_wrapper
[rank0]: inner_compiled_fn = compiler_fn(gm, example_inputs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 661, in _compile_fx_inner
[rank0]: compiled_graph = FxGraphCache.load(
[rank0]: ^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/codecache.py", line 1370, in load
[rank0]: compiled_graph = compile_fx_fn(
[rank0]: ^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 570, in codegen_and_compile
[rank0]: compiled_graph = fx_codegen_and_compile(gm, example_inputs, **fx_kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_fx.py", line 878, in fx_codegen_and_compile
[rank0]: compiled_fn = graph.compile_to_fn()
[rank0]: ^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/graph.py", line 1913, in compile_to_fn
[rank0]: return self.compile_to_module().call
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/graph.py", line 1839, in compile_to_module
[rank0]: return self._compile_to_module()
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/graph.py", line 1867, in _compile_to_module
[rank0]: mod = PyCodeCache.load_by_key_path(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/codecache.py", line 2876, in load_by_key_path
[rank0]: mod = _reload_python_module(key, path)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/compile_tasks.py", line 45, in _reload_python_module
[rank0]: exec(code, mod.__dict__, mod.__dict__)
[rank0]: File "/tmp/torchinductor_xxx/yi/cyi5dcbwxi23ie2hk2rhcnt3zczuwx4ozn7s2xsu5nmx25t72zsd.py", line 474, in
[rank0]: async_compile.wait(globals())
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/async_compile.py", line 276, in wait
[rank0]: scope[key] = result.result()
[rank0]: ^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/codecache.py", line 3341, in result
[rank0]: result = self.future.result()
[rank0]: ^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/concurrent/futures/_base.py", line 456, in result
[rank0]: return self.__get_result()
[rank0]: ^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/concurrent/futures/_base.py", line 401, in __get_result
[rank0]: raise self._exception
[rank0]: torch._inductor.compile_worker.subproc_pool.SubprocException: An exception occurred in a subprocess:
[rank0]: triton.compiler.errors.CompilationError: at 56:48:
[rank0]: m = offs_m
[rank0]: n = offs_n
[rank0]: post_mod_scores = (qk)
[rank0]: if CHECK_BLOCK_BOUNDARY:
[rank0]: # Mask out the elements that are out of the KV_LEN for non divisible seqlen.
[rank0]: post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf"))
[rank0]: if not IS_FULL_BLOCKS:
[rank0]: mask_mod_output = tl.load(in_ptr8 + (m)*s10 + (n) + (off_z)*s10*s9)
[rank0]: ^
[rank0]: NameError('s10 is not defined')
[rank0]: The above exception was the direct cause of the following exception:
[rank0]: triton.compiler.errors.CompilationError: at 57:28:
[rank0]: acc, l_i, m_i,
[rank0]: # Offsets
[rank0]: off_z, off_h, offs_m, offs_n,
[rank0]: MATMUL_PRECISION, RCP_LN2,
[rank0]: IS_FULL_BLOCKS,
[rank0]: )
[rank0]: else:
[rank0]: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen,
[rank0]: # it's on par or slightly faster than only applying to the last block in fwd.
[rank0]: # However, we choose different strategy for bwd, where we only apply mod & mask
[rank0]: # to the last block because it's faster a lot.
[rank0]: acc, l_i, m_i = forward_block_mn(
[rank0]: ^
[rank0]: The above exception was the direct cause of the following exception:
[rank0]: Traceback (most recent call last):
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_worker/subproc_pool.py", line 270, in do_job
[rank0]: result = job()
[rank0]: ^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/compile_tasks.py", line 68, in _worker_compile_triton
[rank0]: load_kernel().precompile(warm_cache_only=True)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/triton_heuristics.py", line 244, in precompile
[rank0]: compiled_binary, launcher = self._precompile_config(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/triton_heuristics.py", line 428, in _precompile_config
[rank0]: triton.compile(*compile_args, **compile_kwargs),
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 276, in compile
[rank0]: module = src.make_ir(options, codegen_fns, context)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 113, in make_ir
[rank0]: return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: triton.compiler.errors.CompilationError: at 148:20:
[rank0]: )
[rank0]: V_block_ptr = tl.make_block_ptr(
[rank0]: base=V,
[rank0]: shape=(KV_LEN, V_HEAD_DIM),
[rank0]: strides=(stride_vn, stride_vk),
[rank0]: offsets=(kv_start, 0),
[rank0]: block_shape=(BLOCK_N, V_HEAD_DIM),
[rank0]: order=(1, 0)
[rank0]: )
[rank0]: offs_n = kv_start + tl.arange(0, BLOCK_N)
[rank0]: acc, l_i, m_i = forward_inner(
[rank0]: ^
[rank0]: The above exception was the direct cause of the following exception:
[rank0]: Traceback (most recent call last):
[rank0]: File "", line 198, in _run_module_as_main
[rank0]: File "", line 88, in _run_code
[rank0]: File "/home/xxx/blt-main/bytelatent/eval.py", line 417, in
[rank0]: main()
[rank0]: File "/home/xxx/blt-main/bytelatent/eval.py", line 413, in main
[rank0]: launch_eval(eval_args)
[rank0]: File "/home/xxx/blt-main/bytelatent/eval.py", line 342, in launch_eval
[rank0]: ppl_results[source] = eval_ppl_on_path(
[rank0]: ^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 116, in decorate_context
[rank0]: return func(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/blt-main/bytelatent/eval.py", line 243, in eval_ppl_on_path
[rank0]: pred = model(x, patch_lengths=patch_lengths)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1736, in _wrapped_call_impl
[rank0]: return self._call_impl(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1747, in _call_impl
[rank0]: return forward_call(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/blt-main/bytelatent/model/blt.py", line 1127, in forward
[rank0]: output, _ = self.local_decoder(
[rank0]: ^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1736, in _wrapped_call_impl
[rank0]: return self._call_impl(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1747, in _call_impl
[rank0]: return forward_call(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/blt-main/bytelatent/model/local_models.py", line 465, in forward
[rank0]: h_cross = self.cross_attn_layers[i](
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1736, in _wrapped_call_impl
[rank0]: return self._call_impl(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1747, in _call_impl
[rank0]: return forward_call(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/blt-main/bytelatent/model/latent_transformer.py", line 139, in forward
[rank0]: output = flex_attention_comp(xq, xk, xv, block_mask=mask)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/eval_frame.py", line 465, in _fn
[rank0]: return fn(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 1269, in __call__
[rank0]: return self._torchdynamo_orig_callable(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 1064, in __call__
[rank0]: result = self._inner_convert(
[rank0]: ^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 526, in __call__
[rank0]: return _compile(
[rank0]: ^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 924, in _compile
[rank0]: guarded_code = compile_inner(code, one_graph, hooks, transform)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 666, in compile_inner
[rank0]: return _compile_inner(code, one_graph, hooks, transform)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_utils_internal.py", line 87, in wrapper_function
[rank0]: return function(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 699, in _compile_inner
[rank0]: out_code = transform_code_object(code, transform)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/bytecode_transformation.py", line 1322, in transform_code_object
[rank0]: transformations(instructions, code_options)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 219, in _fn
[rank0]: return fn(*args, **kwargs)
[rank0]: ^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/convert_frame.py", line 634, in transform
[rank0]: tracer.run()
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/symbolic_convert.py", line 2796, in run
[rank0]: super().run()
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/symbolic_convert.py", line 983, in run
[rank0]: while self.step():
[rank0]: ^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/symbolic_convert.py", line 895, in step
[rank0]: self.dispatch_table[inst.opcode](self, inst)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/symbolic_convert.py", line 2987, in RETURN_VALUE
[rank0]: self._return(inst)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/symbolic_convert.py", line 2972, in _return
[rank0]: self.output.compile_subgraph(
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/output_graph.py", line 1117, in compile_subgraph
[rank0]: self.compile_and_call_fx_graph(tx, list(reversed(stack_values)), root)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/output_graph.py", line 1369, in compile_and_call_fx_graph
[rank0]: compiled_fn = self.call_user_compiler(gm)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/output_graph.py", line 1416, in call_user_compiler
[rank0]: return self._call_user_compiler(gm)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_dynamo/output_graph.py", line 1465, in _call_user_compiler
[rank0]: raise BackendCompilerFailed(self.compiler_fn, e) from e
[rank0]: torch._dynamo.exc.BackendCompilerFailed: backend='inductor' raised:
[rank0]: SubprocException: An exception occurred in a subprocess:
[rank0]: triton.compiler.errors.CompilationError: at 56:48:
[rank0]: m = offs_m
[rank0]: n = offs_n
[rank0]: post_mod_scores = (qk)
[rank0]: if CHECK_BLOCK_BOUNDARY:
[rank0]: # Mask out the elements that are out of the KV_LEN for non divisible seqlen.
[rank0]: post_mod_scores = tl.where(offs_n < KV_LEN, post_mod_scores, float("-inf"))
[rank0]: if not IS_FULL_BLOCKS:
[rank0]: mask_mod_output = tl.load(in_ptr8 + (m)*s10 + (n) + (off_z)*s10*s9)
[rank0]: ^
[rank0]: NameError('s10 is not defined')
[rank0]: The above exception was the direct cause of the following exception:
[rank0]: triton.compiler.errors.CompilationError: at 57:28:
[rank0]: acc, l_i, m_i,
[rank0]: # Offsets
[rank0]: off_z, off_h, offs_m, offs_n,
[rank0]: MATMUL_PRECISION, RCP_LN2,
[rank0]: IS_FULL_BLOCKS,
[rank0]: )
[rank0]: else:
[rank0]: # Benchmark shows even we applied mod & mask to each block for non divisible seqlen,
[rank0]: # it's on par or slightly faster than only applying to the last block in fwd.
[rank0]: # However, we choose different strategy for bwd, where we only apply mod & mask
[rank0]: # to the last block because it's faster a lot.
[rank0]: acc, l_i, m_i = forward_block_mn(
[rank0]: ^
[rank0]: The above exception was the direct cause of the following exception:
[rank0]: Traceback (most recent call last):
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/compile_worker/subproc_pool.py", line 270, in do_job
[rank0]: result = job()
[rank0]: ^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/compile_tasks.py", line 68, in _worker_compile_triton
[rank0]: load_kernel().precompile(warm_cache_only=True)
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/triton_heuristics.py", line 244, in precompile
[rank0]: compiled_binary, launcher = self._precompile_config(
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/torch/_inductor/runtime/triton_heuristics.py", line 428, in _precompile_config
[rank0]: triton.compile(*compile_args, **compile_kwargs),
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 276, in compile
[rank0]: module = src.make_ir(options, codegen_fns, context)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: File "/home/xxx/anaconda3/envs/blt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 113, in make_ir
[rank0]: return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns)
[rank0]: ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
[rank0]: triton.compiler.errors.CompilationError: at 148:20:
[rank0]: )
[rank0]: V_block_ptr = tl.make_block_ptr(
[rank0]: base=V,
[rank0]: shape=(KV_LEN, V_HEAD_DIM),
[rank0]: strides=(stride_vn, stride_vk),
[rank0]: offsets=(kv_start, 0),
[rank0]: block_shape=(BLOCK_N, V_HEAD_DIM),
[rank0]: order=(1, 0)
[rank0]: )
[rank0]: offs_n = kv_start + tl.arange(0, BLOCK_N)
[rank0]: acc, l_i, m_i = forward_inner(
[rank0]: ^
[rank0]: Set TORCH_LOGS="+dynamo" and TORCHDYNAMO_VERBOSE=1 for more information
[rank0]: You can suppress this exception and fall back to eager by setting:
[rank0]: import torch._dynamo
[rank0]: torch._dynamo.config.suppress_errors = True
world size {1}
[rank0]:[W605 21:07:38.118181372 ProcessGroupNCCL.cpp:1250] Warning: WARNING: process group has NOT been destroyed before we destruct ProcessGroupNCCL. On normal program exit, the application should call destroy_process_group to ensure that any pending NCCL operations have finished in this process. In rare cases this process can exit before this point and block the progress of another member of the process group. This constraint has always been present, but this warning has only been added since PyTorch 2.4 (function operator())
Contributor guide
Assessment
This issue has not been assessed yet.