GPT-OSS-20B not work in AMD GPUs

[tanreinama]

System Info

GPT-OSS-20B does not work on Radeon GPUs.
I tested it in both the native environment and the Docker container rocm/pytorch:rocm7.2.1_ubuntu24.04_py3.12_pytorch_release_2.9.1.
I tried updating Triton, but it still didn't work.
I tried those versions of Triton, triton-rocm 3.6.0, 3.5.1+rocm (included in rocm/pytorch), 3.6.0 nightly, and 3.7.0 nightly, but all resulted in errors.

@ivarflakstad

command log:

$ pip install -U transformers kernels accelerate
$ python
>>> from transformers import pipeline
>>> import torch
>>> model_id = "openai/gpt-oss-20b"
>>> pipe = pipeline(
...     "text-generation",
...     model=model_id,
...     torch_dtype="auto",
...     device_map="auto",
... )
`torch_dtype` is deprecated! Use `dtype` instead!
Fetching 42 files: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 42/42 [00:02<00:00, 16.72it/s]
Download complete: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 101/101 [00:02<00:00, 39.0B/s]
Loading weights: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 411/411 [01:16<00:00,  5.40it/s]
>>> messages = [
...     {"role": "user", "content": "Explain quantum mechanics clearly and concisely."},
... ]
>>> outputs = pipe(
...     messages,
...     max_new_tokens=256,
... )
Passing `generation_config` together with generation-related arguments=({'max_new_tokens'}) is deprecated and will be removed in future versions. Please pass either a `generation_config` object OR all generation parameters explicitly, but not both.
Both `max_new_tokens` (=256) and `max_length`(=20) seem to have been set. `max_new_tokens` will take precedence. Please refer to the documentation for more information. (https://huggingface.co/docs/transformers/main/en/main_classes/text_generation)
Traceback (most recent call last):
  File "<python-input-5>", line 1, in <module>
    outputs = pipe(
        messages,
        max_new_tokens=256,
    )
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/pipelines/text_generation.py", line 299, in __call__
    return super().__call__(text_inputs, **kwargs)
           ~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/pipelines/base.py", line 1264, in __call__
    return self.run_single(inputs, preprocess_params, forward_params, postprocess_params)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/pipelines/base.py", line 1271, in run_single
    model_outputs = self.forward(model_inputs, **forward_params)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/pipelines/base.py", line 1163, in forward
    model_outputs = self._forward(model_inputs, **forward_params)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/pipelines/text_generation.py", line 403, in _forward
    output = self.model.generate(input_ids=input_ids, attention_mask=attention_mask, **generate_kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/utils/_contextlib.py", line 124, in decorate_context
    return func(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/generation/utils.py", line 2543, in generate
    result = decoding_method(
        self,
    ...<5 lines>...
        **model_kwargs,
    )
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/generation/utils.py", line 2736, in _sample
    outputs = self._prefill(
        input_ids,
    ...<2 lines>...
        is_first_iteration=not generation_config.is_assistant,
    )
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/generation/utils.py", line 3768, in _prefill
    return self(**model_inputs, return_dict=True)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/utils/generic.py", line 876, in wrapper
    output = func(self, *args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 649, in forward
    outputs: MoeModelOutputWithPast = self.model(
                                      ~~~~~~~~~~^
        input_ids=input_ids,
        ^^^^^^^^^^^^^^^^^^^^
    ...<6 lines>...
        **kwargs,
        ^^^^^^^^^
    )
    ^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/utils/generic.py", line 952, in wrapper
    output = func(self, *args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/utils/output_capturing.py", line 248, in wrapper
    outputs = func(self, *args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 490, in forward
    hidden_states = decoder_layer(
        hidden_states,
    ...<5 lines>...
        **kwargs,
    )
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/modeling_layers.py", line 93, in __call__
    return super().__call__(*args, **kwargs)
           ~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 384, in forward
    hidden_states, _ = self.mlp(hidden_states)  # diff with llama: router scores
                       ~~~~~~~~^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/integrations/mxfp4.py", line 508, in mlp_forward
    routed_out = self.experts(hidden_states, routing_data, gather_idx, scatter_idx=scatter_idx)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1779, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/torch/nn/modules/module.py", line 1790, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/nama/.pyenv/versions/rocm/lib/python3.13/site-packages/transformers/integrations/mxfp4.py", line 411, in forward
    intermediate_cache3 = matmul_ogs(
        intermediate_cache1,
    ...<5 lines>...
        gammas=routing_data.gate_scal,
    )
  File "/home/nama/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 583, in matmul_ogs
    out = apply_postprocessing_features(scatter_indx, finalize_scatter_idxs, opt_flags, expt_token_offs_raw,
                                        num_indx, precision_config, routing_data,
                                        postprocessing_features, memory, fused_postprocess_activation, epilogue)
  File "/home/nama/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 252, in apply_postprocessing_features
    grid, (BLOCK_N, num_warps) = sorted([(compute_grid(*c), c) for c in candidates], key=lambda x: x[0][1])[0]
                                          ~~~~~~~~~~~~^^^^
  File "/home/nama/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 223, in compute_grid
    num_pid = target_info.num_sms() * (warps_per_sm // num_warps)
              ~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
TypeError: unsupported operand type(s) for *: 'NoneType' and 'int'

environment:

$ pip list
Package           Version
----------------- --------------
accelerate        1.13.0
annotated-doc     0.0.4
anyio             4.13.0
certifi           2026.2.25
click             8.3.2
filelock          3.25.2
fsspec            2026.2.0
h11               0.16.0
hf-xet            1.4.3
httpcore          1.0.9
httpx             0.28.1
huggingface_hub   1.9.0
idna              3.11
Jinja2            3.1.6
kernels           0.12.3
markdown-it-py    4.0.0
MarkupSafe        3.0.3
mdurl             0.1.2
mpmath            1.3.0
networkx          3.6.1
numpy             2.4.3
packaging         26.0
pillow            12.1.1
pip               25.3
psutil            7.2.2
Pygments          2.20.0
PyYAML            6.0.3
regex             2026.4.4
rich              14.3.3
safetensors       0.7.0
setuptools        70.2.0
shellingham       1.5.4
sympy             1.14.0
tokenizers        0.22.2
torch             2.11.0+rocm7.2
torchvision       0.26.0+rocm7.2
tqdm              4.67.3
transformers      5.5.0
triton-rocm       3.6.0
typer             0.24.1
typing_extensions 4.15.0

Tested triton version:

$ pip list|grep triton
triton             3.5.1+rocm7.2.1.gita272dfa8
$ pip install -U https://download-r2.pytorch.org/whl/nightly/triton_rocm-3.6.0%2Bgit6213a0e8-cp312-cp312-linux_x86_64.whl
$ pip install -U https://download-r2.pytorch.org/whl/nightly/triton_rocm-3.7.0%2Bgit9c288bc5-cp312-cp312-linux_x86_64.whl

Who can help?

No response

Information

• The official example scripts
• My own modified scripts

Tasks

• An officially supported task in the examples folder (such as GLUE/SQuAD, ...)
• My own task or dataset (give details below)

Reproduction

$ docker run -it --rm --network=host --device=/dev/kfd --device=/dev/dri --group-add=video --ipc=host --cap-add=SYS_PTRACE --security-opt seccomp=unconfined rocm/pytorch:rocm7.2.1_ubuntu24.04_py3.12_pytorch_release_2.9.1
# pip install -U transformers kernels accelerate
# python
>>> from transformers import pipeline
>>> import torch
>>> model_id = "openai/gpt-oss-20b"
>>> pipe = pipeline(
...     "text-generation",
...     model=model_id,
...     torch_dtype="auto",
...     device_map="auto",
... )
Warning: You are sending unauthenticated requests to the HF Hub. Please set a HF_TOKEN to enable higher rate limits and faster downloads.
config.json: 1.81kB [00:00, 1.49MB/s]
`torch_dtype` is deprecated! Use `dtype` instead!
model.safetensors.index.json: 36.4kB [00:00, 62.0MB/s]
Fetching 3 files: 100%|███████████████████████████████████████████████████████████████████████████████████████████████████████████| 3/3 [03:00<00:00, 60.12s/it]
Download complete: 100%|███████████████████████████████████████████████████████████████████████████████████████████████████| 13.8G/13.8G [03:00<00:00, 76.3MB/s]
Fetching 42 files: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 42/42 [00:01<00:00, 33.86it/s]
Download complete: : 249kB [00:01, 193kB/s]              ████████████████████████████████████████████████████████████████████▌  | 41/42 [00:01<00:00, 40.35it/s]
Loading weights: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 411/411 [00:14<00:00, 29.26it/s]
generation_config.json: 100%|██████████████████████████████████████████████████████████████████████████████████████████████████| 177/177 [00:00<00:00, 1.55MB/s]
tokenizer_config.json: 4.20kB [00:00, 9.92MB/s]
tokenizer.json: 100%|██████████████████████████████████████████████████████████████████████████████████████████████████████| 27.9M/27.9M [00:01<00:00, 19.6MB/s]
special_tokens_map.json: 100%|████████████████████████████████████████████████████████████████████████████████████████████████| 98.0/98.0 [00:00<00:00, 317kB/s]
chat_template.jinja: 16.7kB [00:00, 28.5MB/s]
>>> messages = [
...     {"role": "user", "content": "Explain quantum mechanics clearly and concisely."},
... ]
>>> outputs = pipe(
...     messages,
...     max_new_tokens=256,
... )
Passing `generation_config` together with generation-related arguments=({'max_new_tokens'}) is deprecated and will be removed in future versions. Please pass either a `generation_config` object OR all generation parameters explicitly, but not both.
Both `max_new_tokens` (=256) and `max_length`(=20) seem to have been set. `max_new_tokens` will take precedence. Please refer to the documentation for more information. (https://huggingface.co/docs/transformers/main/en/main_classes/text_generation)
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/text_generation.py", line 299, in __call__
    return super().__call__(text_inputs, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1264, in __call__
    return self.run_single(inputs, preprocess_params, forward_params, postprocess_params)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1271, in run_single
    model_outputs = self.forward(model_inputs, **forward_params)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1163, in forward
    model_outputs = self._forward(model_inputs, **forward_params)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/text_generation.py", line 403, in _forward
    output = self.model.generate(input_ids=input_ids, attention_mask=attention_mask, **generate_kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 120, in decorate_context
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 2543, in generate
    result = decoding_method(
             ^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 2736, in _sample
    outputs = self._prefill(
              ^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 3768, in _prefill
    return self(**model_inputs, return_dict=True)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/generic.py", line 876, in wrapper
    output = func(self, *args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 649, in forward
    outputs: MoeModelOutputWithPast = self.model(
                                      ^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/generic.py", line 952, in wrapper
    output = func(self, *args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/output_capturing.py", line 248, in wrapper
    outputs = func(self, *args, **kwargs)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 490, in forward
    hidden_states = decoder_layer(
                    ^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/modeling_layers.py", line 93, in __call__
    return super().__call__(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 384, in forward
    hidden_states, _ = self.mlp(hidden_states)  # diff with llama: router scores
                       ^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/integrations/mxfp4.py", line 508, in mlp_forward
    routed_out = self.experts(hidden_states, routing_data, gather_idx, scatter_idx=scatter_idx)
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/integrations/mxfp4.py", line 411, in forward
    intermediate_cache3 = matmul_ogs(
                          ^^^^^^^^^^^
  File "/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 583, in matmul_ogs
    out = apply_postprocessing_features(scatter_indx, finalize_scatter_idxs, opt_flags, expt_token_offs_raw,
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 252, in apply_postprocessing_features
    grid, (BLOCK_N, num_warps) = sorted([(compute_grid(*c), c) for c in candidates], key=lambda x: x[0][1])[0]
                                          ^^^^^^^^^^^^^^^^
  File "/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 223, in compute_grid
    num_pid = target_info.num_sms() * (warps_per_sm // num_warps)
              ~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
TypeError: unsupported operand type(s) for *: 'NoneType' and 'int'

Expected behavior

Execute without errors

[tanreinama]

with triton_rocm-3.4.0, those error found.

/# pip install -U https://download-r2.pytorch.org/whl/pytorch_triton_rocm-3.4.0-cp312-cp312-linux_x86_64.whl#sha256=7afe951b9fc38f1a5b3a7b98bebbaa092bf51e6192b699b4fade9b1ad6fc9c2c
Collecting pytorch-triton-rocm==3.4.0
  Downloading https://download-r2.pytorch.org/whl/pytorch_triton_rocm-3.4.0-cp312-cp312-linux_x86_64.whl (258.7 MB)
     ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ 258.7/258.7 MB 46.2 MB/s  0:00:05
Requirement already satisfied: setuptools>=40.8.0 in ./opt/venv/lib/python3.12/site-packages (from pytorch-triton-rocm==3.4.0) (82.0.1)
Installing collected packages: pytorch-triton-rocm
Successfully installed pytorch-triton-rocm-3.4.0
root@Z390:/# python
Python 3.12.3 (main, Mar  3 2026, 12:15:18) [GCC 13.3.0] on linux
Type "help", "copyright", "credits" or "license" for more information.
>>> from transformers import pipeline
>>> import torch
>>> model_id = "openai/gpt-oss-20b"
>>> pipe = pipeline(
...     "text-generation",
...     model=model_id,
...     torch_dtype="auto",
...     device_map="auto",
... )
`torch_dtype` is deprecated! Use `dtype` instead!
Fetching 42 files: 100%|█████████████████████████████████████████████████████████████████████████████████████████████████████| 42/42 [00:00<00:00, 10218.14it/s]
Download complete: : 0.00B [00:00, ?B/s]                                                                                                 | 0/42 [00:00<?, ?it/s]
Loading weights: 100%|████████████████████████████████████████████████████████████████████████████████████████████████████████| 411/411 [00:13<00:00, 30.55it/s]
>>> messages = [
...     {"role": "user", "content": "Explain quantum mechanics clearly and concisely."},
... ]
>>> outputs = pipe(
...     messages,
...     max_new_tokens=256,
... )
Passing `generation_config` together with generation-related arguments=({'max_new_tokens'}) is deprecated and will be removed in future versions. Please pass either a `generation_config` object OR all generation parameters explicitly, but not both.
Both `max_new_tokens` (=256) and `max_length`(=20) seem to have been set. `max_new_tokens` will take precedence. Please refer to the documentation for more information. (https://huggingface.co/docs/transformers/main/en/main_classes/text_generation)
python: /tmp/tmpbs5zbb0i/triton/third_party/amd/lib/TritonAMDGPUToLLVM/DotOpToLLVM/MFMA.cpp:775: LogicalResult mlir::triton::AMD::convertScaledMFMA(triton::DotScaledOp, triton::DotScaledOp::Adaptor, const LLVMTypeConverter *, ConversionPatternRewriter &): Assertion `isa<DotOperandEncodingAttr>(op.getA().getType().getEncoding()) && isa<DotOperandEncodingAttr>(op.getB().getType().getEncoding()) && "Both lhs and rhs should be in DotOperand layout."' failed.
#blocked = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [8, 4], warpsPerCTA = [8, 1], order = [1, 0]}>
#blocked1 = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [8], order = [0]}>
#blocked2 = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [1, 32], warpsPerCTA = [1, 8], order = [1, 0]}>
#blocked3 = #ttg.blocked<{sizePerThread = [16, 1], threadsPerWarp = [4, 8], warpsPerCTA = [1, 8], order = [0, 1]}>
#blocked4 = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [2, 16], warpsPerCTA = [8, 1], order = [1, 0]}>
#blocked5 = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [1, 32], warpsPerCTA = [2, 4], order = [1, 0]}>
#blocked6 = #ttg.blocked<{sizePerThread = [1, 1, 1], threadsPerWarp = [1, 16, 2], warpsPerCTA = [1, 8, 1], order = [2, 1, 0]}>
#blocked7 = #ttg.blocked<{sizePerThread = [1, 8, 2], threadsPerWarp = [2, 16, 1], warpsPerCTA = [8, 1, 1], order = [2, 1, 0]}>
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 8 : i32, ttg.target = "hip:gfx1201", "ttg.threads-per-warp" = 32 : i32} {
  tt.func public @_matmul_ogs_NNT_bf16xbf16xmxfp4_32x256x128x1_swiglu(%arg0: !tt.ptr<bf16> {tt.divisibility = 16 : i32}, %arg1: !tt.ptr<bf16> {tt.divisibility = 16 : i32}, %arg2: i32 {tt.divisibility = 16 : i32}, %arg3: i32 {tt.divisibility = 16 : i32}, %arg4: i32 {tt.divisibility = 16 : i32}, %arg5: !tt.ptr<bf16> {tt.divisibility = 16 : i32}, %arg6: !tt.ptr<bf16> {tt.divisibility = 16 : i32}, %arg7: i32 {tt.divisibility = 16 : i32}, %arg8: i32 {tt.divisibility = 16 : i32}, %arg9: !tt.ptr<i8> {tt.divisibility = 16 : i32}, %arg10: i32 {tt.divisibility = 16 : i32}, %arg11: i32 {tt.divisibility = 16 : i32}, %arg12: !tt.ptr<i8> {tt.divisibility = 16 : i32}, %arg13: i32 {tt.divisibility = 16 : i32}, %arg14: i32, %arg15: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %arg16: i32 {tt.divisibility = 16 : i32}, %arg17: i32, %arg18: i32 {tt.divisibility = 16 : i32}, %arg19: i32 {tt.divisibility = 16 : i32}, %arg20: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg21: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg22: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg23: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg24: !tt.ptr<i32> {tt.divisibility = 16 : i32}, %arg25: i32, %arg26: i32, %arg27: f32, %arg28: f32) attributes {noinline = false} {
    %c127_i32 = arith.constant 127 : i32
    %c1_i32 = arith.constant 1 : i32
    %c2_i32 = arith.constant 2 : i32
    %c0_i32 = arith.constant 0 : i32
    %c-1_i32 = arith.constant -1 : i32
    %c65535_i32 = arith.constant 65535 : i32
    %c16_i32 = arith.constant 16 : i32
    %c32_i32 = arith.constant 32 : i32
    %c128_i32 = arith.constant 128 : i32
    %c256_i32 = arith.constant 256 : i32
    %cst = arith.constant dense<4> : tensor<256x4xi32, #blocked>
    %cst_0 = arith.constant dense<0.000000e+00> : tensor<256xf32, #blocked1>
    %c4_i32 = arith.constant 4 : i32
    %c8_i32 = arith.constant 8 : i32
    %cst_1 = arith.constant dense<0.000000e+00> : tensor<32x256xf32, #blocked2>
    %cst_2 = arith.constant dense<0> : tensor<64x256xi8, #blocked3>
    %cst_3 = arith.constant dense<0.000000e+00> : tensor<32x128xbf16, #blocked4>
    %cst_4 = arith.constant 0.000000e+00 : f32
    %cst_5 = arith.constant dense<1.000000e+00> : tensor<32x128xf32, #blocked4>
    %cst_6 = arith.constant dense<64> : tensor<64x256xi32, #blocked3>
    %cst_7 = arith.constant dense<128> : tensor<32x128xi32, #blocked4>
    %cst_8 = arith.constant dense<32> : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
    %cst_9 = arith.constant dense<4> : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %0 = arith.divsi %arg18, %c2_i32 : i32
    %1 = tt.get_program_id x : i32
    %2 = tt.load %arg23 : !tt.ptr<i32>
    %3 = arith.subi %arg25, %2 : i32
    %4 = arith.subi %arg25, %3 : i32
    %5 = arith.muli %4, %arg26 : i32
    %6 = arith.cmpi sgt, %3, %c0_i32 : i32
    %7 = arith.cmpi sge, %1, %5 : i32
    %8 = arith.andi %6, %7 : i1
    cf.cond_br %8, ^bb1, ^bb2
  ^bb1:  // 2 preds: ^bb0, ^bb2
    tt.return
  ^bb2:  // pred: ^bb0
    %9 = arith.divsi %5, %c8_i32 : i32
    %10 = arith.remsi %5, %c8_i32 : i32
    %11 = arith.remsi %1, %c8_i32 : i32
    %12 = arith.divsi %1, %c8_i32 : i32
    %13 = arith.muli %11, %9 : i32
    %14 = arith.minsi %11, %10 : i32
    %15 = arith.addi %13, %14 : i32
    %16 = arith.addi %15, %12 : i32
    %17 = arith.remsi %16, %5 : i32
    %18 = arith.muli %arg26, %c4_i32 : i32
    %19 = arith.divsi %17, %18 : i32
    %20 = arith.muli %19, %c4_i32 : i32
    %21 = arith.subi %4, %20 : i32
    %22 = arith.minsi %21, %c4_i32 : i32
    %23 = arith.remsi %17, %22 : i32
    %24 = arith.addi %20, %23 : i32
    %25 = arith.remsi %17, %18 : i32
    %26 = arith.divsi %25, %22 : i32
    %27 = tt.addptr %arg24, %24 : !tt.ptr<i32>, i32
    %28 = tt.load %27 : !tt.ptr<i32>
    %29 = arith.cmpi eq, %28, %c-1_i32 : i32
    cf.cond_br %29, ^bb1, ^bb3
  ^bb3:  // pred: ^bb2
    %30 = arith.andi %28, %c65535_i32 : i32
    %31 = arith.shrsi %28, %c16_i32 : i32
    %32 = tt.addptr %arg21, %30 : !tt.ptr<i32>, i32
    %33 = tt.load %32 : !tt.ptr<i32>
    %34 = tt.addptr %arg22, %30 : !tt.ptr<i32>, i32
    %35 = tt.load %34 : !tt.ptr<i32>
    %36 = arith.muli %31, %c32_i32 : i32
    %37 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %38 = tt.splat %36 : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %39 = arith.addi %38, %37 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %40 = tt.splat %33 : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %41 = arith.remsi %39, %40 {tt.contiguity = dense<32> : tensor<1xi32>, tt.divisibility = dense<32> : tensor<1xi32>} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %42 = tt.addptr %arg20, %35 : !tt.ptr<i32>, i32
    %43 = tt.splat %42 : !tt.ptr<i32> -> tensor<32x!tt.ptr<i32>, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %44 = tt.addptr %43, %41 : tensor<32x!tt.ptr<i32>, #ttg.slice<{dim = 1, parent = #blocked4}>>, tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %45 = tt.load %44 : tensor<32x!tt.ptr<i32>, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %46 = arith.divsi %45, %cst_9 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %47 = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %48 = tt.expand_dims %46 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>> -> tensor<32x1xi32, #blocked4>
    %49 = tt.splat %arg8 : i32 -> tensor<32x1xi32, #blocked4>
    %50 = arith.muli %48, %49 : tensor<32x1xi32, #blocked4>
    %51 = tt.splat %arg5 : !tt.ptr<bf16> -> tensor<32x1x!tt.ptr<bf16>, #blocked4>
    %52 = tt.addptr %51, %50 : tensor<32x1x!tt.ptr<bf16>, #blocked4>, tensor<32x1xi32, #blocked4>
    %53 = tt.expand_dims %47 {axis = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>> -> tensor<1x128xi32, #blocked4>
    %54 = tt.broadcast %52 : tensor<32x1x!tt.ptr<bf16>, #blocked4> -> tensor<32x128x!tt.ptr<bf16>, #blocked4>
    %55 = tt.broadcast %53 : tensor<1x128xi32, #blocked4> -> tensor<32x128xi32, #blocked4>
    %56 = tt.addptr %54, %55 : tensor<32x128x!tt.ptr<bf16>, #blocked4>, tensor<32x128xi32, #blocked4>
    %57 = arith.cmpi sgt, %arg19, %c0_i32 : i32
    %58 = tt.splat %57 : i1 -> tensor<32x128xi1, #blocked4>
    %59 = tt.splat %arg19 : i32 -> tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %60 = arith.cmpi slt, %47, %59 : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %61 = tt.expand_dims %60 {axis = 0 : i32} : tensor<128xi1, #ttg.slice<{dim = 0, parent = #blocked4}>> -> tensor<1x128xi1, #blocked4>
    %62 = tt.broadcast %61 : tensor<1x128xi1, #blocked4> -> tensor<32x128xi1, #blocked4>
    %63 = arith.andi %58, %62 : tensor<32x128xi1, #blocked4>
    %64 = tt.load %56, %63, %cst_3 {amd.pipeliner_part = "prologue"} : tensor<32x128x!tt.ptr<bf16>, #blocked4>
    %65 = arith.muli %30, %arg13 : i32
    %66 = tt.addptr %arg12, %65 : !tt.ptr<i8>, i32
    %67 = arith.muli %26, %c256_i32 : i32
    %68 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
    %69 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>>
    %70 = tt.make_range {end = 256 : i32, start = 0 : i32} : tensor<256xi32, #blocked1>
    %71 = tt.splat %67 : i32 -> tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
    %72 = tt.splat %67 : i32 -> tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>>
    %73 = tt.splat %67 : i32 -> tensor<256xi32, #blocked1>
    %74 = arith.addi %71, %68 : tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
    %75 = arith.addi %72, %69 : tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>>
    %76 = arith.addi %73, %70 : tensor<256xi32, #blocked1>
    %77 = tt.splat %arg18 : i32 -> tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
    %78 = tt.splat %arg18 : i32 -> tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>>
    %79 = tt.splat %arg18 : i32 -> tensor<256xi32, #blocked1>
    %80 = arith.remsi %74, %77 {tt.contiguity = dense<256> : tensor<1xi32>, tt.divisibility = dense<256> : tensor<1xi32>} : tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
    %81 = arith.remsi %75, %78 {tt.contiguity = dense<256> : tensor<1xi32>, tt.divisibility = dense<256> : tensor<1xi32>} : tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>>
    %82 = tt.make_range {end = 4 : i32, start = 0 : i32} : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
    %83 = tt.expand_dims %82 {axis = 0 : i32} : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x4xi32, #blocked>
    %84 = tt.splat %66 : !tt.ptr<i8> -> tensor<1x4x!tt.ptr<i8>, #blocked>
    %85 = tt.addptr %84, %83 : tensor<1x4x!tt.ptr<i8>, #blocked>, tensor<1x4xi32, #blocked>
    %86 = tt.expand_dims %80 {axis = 1 : i32} : tensor<256xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<256x1xi32, #blocked>
    %87 = tt.splat %arg14 : i32 -> tensor<256x1xi32, #blocked>
    %88 = arith.muli %86, %87 : tensor<256x1xi32, #blocked>
    %89 = tt.broadcast %85 : tensor<1x4x!tt.ptr<i8>, #blocked> -> tensor<256x4x!tt.ptr<i8>, #blocked>
    %90 = tt.broadcast %88 : tensor<256x1xi32, #blocked> -> tensor<256x4xi32, #blocked>
    %91 = tt.addptr %89, %90 : tensor<256x4x!tt.ptr<i8>, #blocked>, tensor<256x4xi32, #blocked>
    %92 = tt.splat %57 : i1 -> tensor<256x4xi1, #blocked>
    %93 = arith.muli %82, %cst_8 : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
    %94 = tt.splat %arg19 : i32 -> tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
    %95 = arith.cmpi slt, %93, %94 : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
    %96 = tt.expand_dims %95 {axis = 0 : i32} : tensor<4xi1, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x4xi1, #blocked>
    %97 = tt.broadcast %96 : tensor<1x4xi1, #blocked> -> tensor<256x4xi1, #blocked>
    %98 = arith.andi %92, %97 : tensor<256x4xi1, #blocked>
    %99 = tt.load %91, %98 {amd.pipeliner_part = "prologue"} : tensor<256x4x!tt.ptr<i8>, #blocked>
    %100 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>>
    %101 = arith.muli %30, %arg10 : i32
    %102 = tt.addptr %arg9, %101 : !tt.ptr<i8>, i32
    %103 = tt.expand_dims %100 {axis = 1 : i32} : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>> -> tensor<64x1xi32, #blocked3>
    %104 = tt.expand_dims %81 {axis = 0 : i32} : tensor<256xi32, #ttg.slice<{dim = 0, parent = #blocked3}>> -> tensor<1x256xi32, #blocked3>
    %105 = tt.splat %arg11 : i32 -> tensor<1x256xi32, #blocked3>
    %106 = arith.muli %104, %105 : tensor<1x256xi32, #blocked3>
    %107 = tt.broadcast %103 : tensor<64x1xi32, #blocked3> -> tensor<64x256xi32, #blocked3>
    %108 = tt.broadcast %106 : tensor<1x256xi32, #blocked3> -> tensor<64x256xi32, #blocked3>
    %109 = arith.addi %107, %108 : tensor<64x256xi32, #blocked3>
    %110 = tt.splat %102 : !tt.ptr<i8> -> tensor<64x256x!tt.ptr<i8>, #blocked3>
    %111 = tt.addptr %110, %109 : tensor<64x256x!tt.ptr<i8>, #blocked3>, tensor<64x256xi32, #blocked3>
    %112 = tt.splat %57 : i1 -> tensor<64x256xi1, #blocked3>
    %113 = arith.divsi %arg19, %c2_i32 : i32
    %114 = tt.splat %113 : i32 -> tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>>
    %115 = arith.cmpi slt, %100, %114 : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>>
    %116 = tt.expand_dims %115 {axis = 1 : i32} : tensor<64xi1, #ttg.slice<{dim = 1, parent = #blocked3}>> -> tensor<64x1xi1, #blocked3>
    %117 = tt.broadcast %116 : tensor<64x1xi1, #blocked3> -> tensor<64x256xi1, #blocked3>
    %118 = arith.andi %112, %117 : tensor<64x256xi1, #blocked3>
    %119 = tt.load %111, %118, %cst_2 cacheModifier = cg {amd.pipeliner_part = "prologue"} : tensor<64x256x!tt.ptr<i8>, #blocked3>
    %120 = arith.subi %arg19, %c128_i32 : i32
    %121:7 = scf.for %arg29 = %c0_i32 to %120 step %c128_i32 iter_args(%arg30 = %cst_1, %arg31 = %91, %arg32 = %56, %arg33 = %111, %arg34 = %64, %arg35 = %119, %arg36 = %99) -> (tensor<32x256xf32, #blocked2>, tensor<256x4x!tt.ptr<i8>, #blocked>, tensor<32x128x!tt.ptr<bf16>, #blocked4>, tensor<64x256x!tt.ptr<i8>, #blocked3>, tensor<32x128xbf16, #blocked4>, tensor<64x256xi8, #blocked3>, tensor<256x4xi8, #blocked>)  : i32 {
      %177 = tt.addptr %arg31, %cst : tensor<256x4x!tt.ptr<i8>, #blocked>, tensor<256x4xi32, #blocked>
      %178 = tt.addptr %arg32, %cst_7 : tensor<32x128x!tt.ptr<bf16>, #blocked4>, tensor<32x128xi32, #blocked4>
      %179 = tt.addptr %arg33, %cst_6 : tensor<64x256x!tt.ptr<i8>, #blocked3>, tensor<64x256xi32, #blocked3>
      %180 = arith.addi %arg29, %c128_i32 : i32
      %181 = arith.subi %arg19, %180 : i32
      %182 = tt.splat %181 : i32 -> tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
      %183 = arith.cmpi slt, %47, %182 : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
      %184 = arith.divsi %181, %c2_i32 : i32
      %185 = tt.splat %184 : i32 -> tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>>
      %186 = arith.cmpi slt, %100, %185 : tensor<64xi32, #ttg.slice<{dim = 1, parent = #blocked3}>>
      %187 = tt.splat %181 : i32 -> tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
      %188 = arith.cmpi slt, %93, %187 : tensor<4xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
      %189 = tt.expand_dims %183 {axis = 0 : i32} : tensor<128xi1, #ttg.slice<{dim = 0, parent = #blocked4}>> -> tensor<1x128xi1, #blocked4>
      %190 = tt.broadcast %189 : tensor<1x128xi1, #blocked4> -> tensor<32x128xi1, #blocked4>
      %191 = tt.load %178, %190, %cst_3 : tensor<32x128x!tt.ptr<bf16>, #blocked4>
      %192 = ttg.convert_layout %arg34 : tensor<32x128xbf16, #blocked4> -> tensor<32x128xbf16, #blocked5>
      %193 = tt.expand_dims %186 {axis = 1 : i32} : tensor<64xi1, #ttg.slice<{dim = 1, parent = #blocked3}>> -> tensor<64x1xi1, #blocked3>
      %194 = tt.broadcast %193 : tensor<64x1xi1, #blocked3> -> tensor<64x256xi1, #blocked3>
      %195 = tt.load %179, %194, %cst_2 cacheModifier = cg : tensor<64x256x!tt.ptr<i8>, #blocked3>
      %196 = ttg.convert_layout %arg35 : tensor<64x256xi8, #blocked3> -> tensor<64x256xi8, #blocked2>
      %197 = tt.expand_dims %188 {axis = 0 : i32} : tensor<4xi1, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x4xi1, #blocked>
      %198 = tt.broadcast %197 : tensor<1x4xi1, #blocked> -> tensor<256x4xi1, #blocked>
      %199 = tt.load %177, %198 : tensor<256x4x!tt.ptr<i8>, #blocked>
      %200 = tt.dot_scaled %192, %196 scale %arg36, %arg30 lhs = bf16 rhs = e2m1 {fastMath = true} : tensor<32x128xbf16, #blocked5> * tensor<64x256xi8, #blocked2>, tensor<256x4xi8, #blocked> -> tensor<32x256xf32, #blocked2>
      scf.yield %200, %177, %178, %179, %191, %195, %199 : tensor<32x256xf32, #blocked2>, tensor<256x4x!tt.ptr<i8>, #blocked>, tensor<32x128x!tt.ptr<bf16>, #blocked4>, tensor<64x256x!tt.ptr<i8>, #blocked3>, tensor<32x128xbf16, #blocked4>, tensor<64x256xi8, #blocked3>, tensor<256x4xi8, #blocked>
    }
    %122 = arith.addi %arg19, %c127_i32 : i32
    %123 = arith.divsi %122, %c128_i32 : i32
    %124 = arith.cmpi sge, %123, %c1_i32 : i32
    %125 = ttg.convert_layout %121#4 : tensor<32x128xbf16, #blocked4> -> tensor<32x128xbf16, #blocked5>
    %126 = ttg.convert_layout %121#5 : tensor<64x256xi8, #blocked3> -> tensor<64x256xi8, #blocked2>
    %127 = scf.if %124 -> (tensor<32x256xf32, #blocked2>) {
      %177 = tt.dot_scaled %125, %126 scale %121#6, %121#0 lhs = bf16 rhs = e2m1 {fastMath = true} : tensor<32x128xbf16, #blocked5> * tensor<64x256xi8, #blocked2>, tensor<256x4xi8, #blocked> -> tensor<32x256xf32, #blocked2>
      scf.yield %177 : tensor<32x256xf32, #blocked2>
    } else {
      scf.yield %121#0 : tensor<32x256xf32, #blocked2>
    }
    %128 = arith.select %124, %127, %121#0 : tensor<32x256xf32, #blocked2>
    %129 = arith.cmpi slt, %39, %40 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>>
    %130 = arith.cmpi slt, %76, %79 : tensor<256xi32, #blocked1>
    %131 = arith.muli %30, %arg16 : i32
    %132 = tt.addptr %arg15, %131 : !tt.ptr<f32>, i32
    %133 = tt.splat %132 : !tt.ptr<f32> -> tensor<256x!tt.ptr<f32>, #blocked1>
    %134 = tt.addptr %133, %76 : tensor<256x!tt.ptr<f32>, #blocked1>, tensor<256xi32, #blocked1>
    %135 = tt.load %134, %130, %cst_0 : tensor<256x!tt.ptr<f32>, #blocked1>
    %136 = ttg.convert_layout %135 : tensor<256xf32, #blocked1> -> tensor<256xf32, #ttg.slice<{dim = 0, parent = #blocked2}>>
    %137 = tt.expand_dims %136 {axis = 0 : i32} : tensor<256xf32, #ttg.slice<{dim = 0, parent = #blocked2}>> -> tensor<1x256xf32, #blocked2>
    %138 = tt.broadcast %137 : tensor<1x256xf32, #blocked2> -> tensor<32x256xf32, #blocked2>
    %139 = arith.addf %128, %138 : tensor<32x256xf32, #blocked2>
    %140 = tt.reshape %139 : tensor<32x256xf32, #blocked2> -> tensor<32x128x2xf32, #blocked6>
    %141 = ttg.convert_layout %140 : tensor<32x128x2xf32, #blocked6> -> tensor<32x128x2xf32, #blocked7>
    %outLHS, %outRHS = tt.split %141 : tensor<32x128x2xf32, #blocked7> -> tensor<32x128xf32, #blocked4>
    %142 = tt.splat %arg28 : f32 -> tensor<32x128xf32, #blocked4>
    %143 = arith.minnumf %outLHS, %142 : tensor<32x128xf32, #blocked4>
    %144 = arith.minnumf %outRHS, %142 : tensor<32x128xf32, #blocked4>
    %145 = arith.subf %cst_4, %arg28 : f32
    %146 = tt.splat %145 : f32 -> tensor<32x128xf32, #blocked4>
    %147 = arith.maxnumf %146, %144 : tensor<32x128xf32, #blocked4>
    %148 = arith.subf %cst_4, %arg27 : f32
    %149 = tt.splat %148 : f32 -> tensor<32x128xf32, #blocked4>
    %150 = arith.mulf %149, %143 : tensor<32x128xf32, #blocked4>
    %151 = math.exp %150 : tensor<32x128xf32, #blocked4>
    %152 = arith.addf %151, %cst_5 : tensor<32x128xf32, #blocked4>
    %153 = arith.divf %143, %152 : tensor<32x128xf32, #blocked4>
    %154 = math.fma %153, %147, %153 : tensor<32x128xf32, #blocked4>
    %155 = arith.muli %26, %c128_i32 : i32
    %156 = tt.splat %155 : i32 -> tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %157 = arith.addi %156, %47 : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %158 = tt.splat %0 : i32 -> tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %159 = arith.cmpi slt, %157, %158 : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>>
    %160 = arith.muli %35, %arg4 : i32
    %161 = tt.addptr %arg1, %160 : !tt.ptr<bf16>, i32
    %162 = tt.expand_dims %39 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked4}>> -> tensor<32x1xi32, #blocked4>
    %163 = tt.splat %arg4 : i32 -> tensor<32x1xi32, #blocked4>
    %164 = arith.muli %162, %163 : tensor<32x1xi32, #blocked4>
    %165 = tt.splat %161 : !tt.ptr<bf16> -> tensor<32x1x!tt.ptr<bf16>, #blocked4>
    %166 = tt.addptr %165, %164 : tensor<32x1x!tt.ptr<bf16>, #blocked4>, tensor<32x1xi32, #blocked4>
    %167 = tt.expand_dims %157 {axis = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked4}>> -> tensor<1x128xi32, #blocked4>
    %168 = tt.broadcast %166 : tensor<32x1x!tt.ptr<bf16>, #blocked4> -> tensor<32x128x!tt.ptr<bf16>, #blocked4>
    %169 = tt.broadcast %167 : tensor<1x128xi32, #blocked4> -> tensor<32x128xi32, #blocked4>
    %170 = tt.addptr %168, %169 : tensor<32x128x!tt.ptr<bf16>, #blocked4>, tensor<32x128xi32, #blocked4>
    %171 = tt.expand_dims %129 {axis = 1 : i32} : tensor<32xi1, #ttg.slice<{dim = 1, parent = #blocked4}>> -> tensor<32x1xi1, #blocked4>
    %172 = tt.expand_dims %159 {axis = 0 : i32} : tensor<128xi1, #ttg.slice<{dim = 0, parent = #blocked4}>> -> tensor<1x128xi1, #blocked4>
    %173 = tt.broadcast %171 : tensor<32x1xi1, #blocked4> -> tensor<32x128xi1, #blocked4>
    %174 = tt.broadcast %172 : tensor<1x128xi1, #blocked4> -> tensor<32x128xi1, #blocked4>
    %175 = arith.andi %173, %174 : tensor<32x128xi1, #blocked4>
    %176 = arith.truncf %154 : tensor<32x128xf32, #blocked4> to tensor<32x128xbf16, #blocked4>
    tt.store %170, %176, %175 : tensor<32x128x!tt.ptr<bf16>, #blocked4>
    tt.return
  }
}

{-#
  external_resources: {
    mlir_reproducer: {
      pipeline: "builtin.module(optimize-amd-lds-usage{lds-limit=0 target-arch=gfx1201}, convert-scf-to-cf, convert-index-to-llvm{index-bitwidth=0}, allocate-shared-memory, convert-triton-amdgpu-to-llvm{arch=gfx1201 ftz=true}, canonicalize{  max-iterations=10 max-num-rewrites=-1 region-simplify=normal test-convergence=false top-down=true}, cse, convert-cf-to-llvm{index-bitwidth=0}, convert-arith-to-llvm{index-bitwidth=0}, canonicalize{  max-iterations=10 max-num-rewrites=-1 region-simplify=normal test-convergence=false top-down=true}, cse, symbol-dce, enable-line-info, convert-builtin-func-to-llvm{ftz=true})",
      disable_threading: false,
      verify_each: true
    }
  }
#-}
/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/specialize.py:33:0: error: Failures have been detected while processing an MLIR pass pipeline
/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/specialize.py:33:0: note: Pipeline failed while executing [`ConvertTritonAMDGPUToLLVM` on 'builtin.module' operation]: reproducer generated at `std::errs, please share the reproducer above with Triton project.`
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/text_generation.py", line 299, in __call__
    return super().__call__(text_inputs, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1264, in __call__
    return self.run_single(inputs, preprocess_params, forward_params, postprocess_params)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1271, in run_single
    model_outputs = self.forward(model_inputs, **forward_params)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/base.py", line 1163, in forward
    model_outputs = self._forward(model_inputs, **forward_params)
                    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/pipelines/text_generation.py", line 403, in _forward
    output = self.model.generate(input_ids=input_ids, attention_mask=attention_mask, **generate_kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/utils/_contextlib.py", line 120, in decorate_context
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 2543, in generate
    result = decoding_method(
             ^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 2736, in _sample
    outputs = self._prefill(
              ^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/generation/utils.py", line 3768, in _prefill
    return self(**model_inputs, return_dict=True)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/generic.py", line 876, in wrapper
    output = func(self, *args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 649, in forward
    outputs: MoeModelOutputWithPast = self.model(
                                      ^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/generic.py", line 952, in wrapper
    output = func(self, *args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/utils/output_capturing.py", line 248, in wrapper
    outputs = func(self, *args, **kwargs)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 490, in forward
    hidden_states = decoder_layer(
                    ^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/modeling_layers.py", line 93, in __call__
    return super().__call__(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/models/gpt_oss/modeling_gpt_oss.py", line 384, in forward
    hidden_states, _ = self.mlp(hidden_states)  # diff with llama: router scores
                       ^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/integrations/mxfp4.py", line 508, in mlp_forward
    routed_out = self.experts(hidden_states, routing_data, gather_idx, scatter_idx=scatter_idx)
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/transformers/integrations/mxfp4.py", line 400, in forward
    intermediate_cache1 = matmul_ogs(
                          ^^^^^^^^^^^
  File "/root/.cache/huggingface/hub/models--kernels-community--gpt-oss-triton-kernels/snapshots/76c23fb9a6607cd5c62c1e6b8e7f436ec5385517/build/torch-rocm/matmul_ogs.py", line 531, in matmul_ogs
    (kernels._p_matmul_ogs if opt_flags.is_persistent else kernels._matmul_ogs)[(grid,)](
  File "/opt/venv/lib/python3.12/site-packages/triton/runtime/jit.py", line 390, in <lambda>
    return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
                                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/triton/runtime/jit.py", line 594, in run
    kernel = self.compile(src, target=target, options=options.__dict__)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/triton/compiler/compiler.py", line 359, in compile
    next_module = compile_ir(module, metadata)
                  ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.py", line 446, in <lambda>
    stages["llir"] = lambda src, metadata: self.make_llir(src, metadata, options)
                                           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/opt/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.py", line 330, in make_llir
    pm.run(mod)
RuntimeError: PassManager::run failed