[KMCompiler]feat(conj_physical): add hygon runtime and ascend tune config

benchmark/test_conj_physical.py

@@ -1,16 +1,4 @@
-# Copyright 2026 FlagOS Contributors
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-#     http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
+import time
 
 import pytest
 import torch
@@ -34,7 +22,6 @@ def _input_fn(shape, dtype, device):
 
 
 class Conj_physicalBenchmark(base.GenericBenchmarkExcluse3D):
-    # TODO(Qiming): Check if this is necessary
     def __init__(self, *args, **kwargs):
         super().__init__(*args, **kwargs)
 
@@ -52,10 +39,27 @@ def set_shapes(self, shape_file_path=None):
     def set_more_shapes(self):
         return None
 
+    def get_latency(self, op, *args, **kwargs):
+        """Override for Ascend NPU: wall-clock timing to avoid event timer bug."""
+        if "npu" in str(self.device) or "ascend" in str(self.device).lower():
+            fn = lambda: op(*args, **kwargs)
+            torch.npu.synchronize()
+            start = time.time()
+            for _ in range(10):
+                fn()
+            torch.npu.synchronize()
+            end = time.time()
+            return (end - start) / 10 * 1000
+        return super().get_latency(op, *args, **kwargs)
+
 
 @pytest.mark.conj_physical
 def test_conj_physical():
-    dtypes = consts.FLOAT_DTYPES + consts.INT_DTYPES + consts.COMPLEX_DTYPES
+    # Ascend NPU Triton backend does not support complex64
+    if "npu" in flag_gems.device or "ascend" in flag_gems.device.lower():
+        dtypes = consts.FLOAT_DTYPES + consts.INT_DTYPES
+    else:
+        dtypes = consts.FLOAT_DTYPES + consts.INT_DTYPES + consts.COMPLEX_DTYPES
 
     bench = Conj_physicalBenchmark(
         input_fn=_input_fn,

src/flag_gems/runtime/backend/_ascend/tune_configs.yaml

@@ -651,6 +651,35 @@ var_mean:
   block_n:
   - 128
 
+conj_physical:
+  - META:
+      BLOCK_SIZE: 64
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 128
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 256
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 512
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 64
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 128
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 256
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 512
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 1024
+    num_warps: 8
+
 conv2d_forward:
 - META:
     BLOCK_NI_HO_WO: 32

src/flag_gems/runtime/backend/_hygon/ops/__init__.py

@@ -21,6 +21,7 @@
     scaled_dot_product_attention_backward,
     scaled_dot_product_attention_forward,
 )
+from .conj_physical import conj_physical
 from .div import (
     div_mode,
     div_mode_,
@@ -72,6 +73,7 @@
 
 __all__ = [
     "_unique2",
+    "conj_physical",
     "ScaleDotProductAttention",
     "SUPPORTED_FP8_DTYPE",
     "any",

src/flag_gems/runtime/backend/_hygon/ops/conj_physical.py

@@ -0,0 +1,49 @@
+import logging
+
+import torch
+import triton
+import triton.language as tl
+
+from flag_gems.utils import libentry
+
+logger = logging.getLogger(__name__)
+
+
+@libentry()
+@triton.jit
+def conj_physical_kernel(in_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
+    pid = tl.program_id(0)
+    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
+    mask = offsets < n_elements
+
+    base = offsets * 2
+    real = tl.load(in_ptr + base, mask=mask)
+    imag = tl.load(in_ptr + base + 1, mask=mask)
+
+    tl.store(out_ptr + base, real, mask=mask)
+    tl.store(out_ptr + base + 1, -imag, mask=mask)
+
+
+def conj_physical(input: torch.Tensor) -> torch.Tensor:
+    logger.debug("GEMS_HYGON CONJ_PHYSICAL")
+    if not input.is_complex():
+        return input
+
+    n_elements = input.numel()
+    src = input if input.is_contiguous() else input.contiguous()
+    output = torch.empty_like(src)
+    in_real_ptr = torch.view_as_real(src)
+    out_real_ptr = torch.view_as_real(output)
+
+    BLOCK_SIZE = 256
+    grid = (triton.cdiv(n_elements, BLOCK_SIZE),)
+
+    conj_physical_kernel[grid](
+        in_real_ptr,
+        out_real_ptr,
+        n_elements,
+        BLOCK_SIZE=BLOCK_SIZE,
+        num_warps=4,
+    )
+
+    return output

src/flag_gems/runtime/backend/_hygon/tune_configs.yaml

@@ -1167,6 +1167,35 @@ var_mean:
   - 8
   - 16
 
+conj_physical:
+  - META:
+      BLOCK_SIZE: 64
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 128
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 256
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 512
+    num_warps: 4
+  - META:
+      BLOCK_SIZE: 64
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 128
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 256
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 512
+    num_warps: 8
+  - META:
+      BLOCK_SIZE: 1024
+    num_warps: 8
+
 conv2d_forward:
 - META:
     BLOCK_NI_HO_WO: 32