Fix opus gemm aiter check

csrc/gemm_a16w16/README.md

@@ -60,17 +60,24 @@ If you have built kernels before tuning, add `AITER_REBUILD=1` to rebuild with n
 - **Type**: Comma-separated string
 - **Default**: `all`
 - **Choices**: `all`, `asm`, `hipblaslt`, `triton`, `flydsl`, `torch`, `skinny`, `opus`
-- **Description**: Choose which backends to tune. `all` includes every backend **except** hipblaslt. To include hipblaslt, you must also pass `--with-hipblaslt`.
+- **Description**: Choose which backends to tune. hipblaslt requires **both** `--libtype all` (or `--libtype hipblaslt`) **and** `--with-hipblaslt` to run.
 
 **Example**:
 ```bash
+# tune asm and opus only
 --libtype asm,opus
+
+# tune all backends including hipblaslt
+--with-hipblaslt
+
+# tune hipblaslt only
+--libtype hipblaslt --with-hipblaslt
 ```
 
 ### `--with-hipblaslt`
 - **Type**: Flag
 - **Default**: disabled
-- **Description**: Include hipblaslt solutions in tuning (imports from gradlib). Without this flag, hipblaslt is **never** run — even with `--libtype all` or `--libtype hipblaslt`. When using this flag, it is recommended to run via `gemm_tuner.py` (the subprocess wrapper) so that GPU-level crashes from hipblaslt are retried automatically.
+- **Description**: Enable hipblaslt backend (imports from gradlib). This is a **gate switch** — hipblaslt is never run without it, regardless of `--libtype`. With the default `--libtype all`, adding `--with-hipblaslt` is sufficient to include hipblaslt alongside all other backends. When using this flag, it is recommended to run via `gemm_tuner.py` (the subprocess wrapper) so that GPU-level crashes from hipblaslt are retried automatically.
 
 ### `--indtype` / `--outdtype`
 - **Choices**: `f32`, `f16`, `bf16`, `fp8`

csrc/gemm_a16w16/gemm_a16w16_tune.py

@@ -85,6 +85,17 @@
     HipblasltGemm = None
     HIPBLASLT_TUNE_ERROR = str(_hipb_exc)
 
+# ---------------------------------------------------------------------------
+# Tolerance helpers
+# ---------------------------------------------------------------------------
+
+
+def _default_tol(outdtype):
+    """Return (rtol, atol) for the given output dtype."""
+    tol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+    return tol, tol
+
+
 # ---------------------------------------------------------------------------
 # Data generation & reference
 # ---------------------------------------------------------------------------
@@ -505,8 +516,7 @@ def run_config(self, args):
                     eval(indtype),
                     eval(outdtype),
                 )
-                _atol = 5e-2 if eval(outdtype) == torch.bfloat16 else 1e-2
-                _rtol = _atol
+                _rtol, _atol = _default_tol(eval(outdtype))
                 err_ratio = checkAllclose(
                     out, ref, atol=_atol, rtol=_rtol, msg=f"run_config {shape_str}"
                 )
@@ -582,8 +592,7 @@ def _get_asm_tasks(
             return []
         asm_kernel_list_csv = f"{get_asm_dir()}/bf16gemm/bf16gemm_fp32bf16.csv"
         asm_kernels = get_asm_kernels(asm_kernel_list_csv, is_shuffle)
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
-        atol = rtol
+        rtol, atol = _default_tol(outdtype)
         tasks = []
         solidx = 0
         for key in asm_kernels.keys():
@@ -652,8 +661,7 @@ def _get_opus_tasks(
             return []
         M, N, K = info_keys[2], info_keys[3], info_keys[4]
         cu_num = get_cu_num()
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
-        atol = rtol
+        rtol, atol = _default_tol(outdtype)
         cand_kids = _opus_candidate_kids_for_shape(M, N, K, has_bias, cu_num)
         tasks = []
         for kid in sorted(cand_kids):
@@ -705,8 +713,7 @@ def _get_flydsl_tasks(
         if scaleAB or indtype != dtypes.bf16:
             return []
         M, N, K = info_keys[2], info_keys[3], info_keys[4]
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
-        atol = rtol
+        rtol, atol = _default_tol(outdtype)
         flydsl_catalog = get_flydsl_bf16_catalog(M, N, K)
         weight_key = "shuffleweights" if is_shuffle else "weights"
         min_tile_m = min((c["tile_m"] for _, _, c in flydsl_catalog), default=16)
@@ -769,7 +776,7 @@ def _get_skinny_tasks(
         _, _, native_is_skinny = get_native_gemm_funcs()
         if not native_is_skinny(M, N, K, indtype):
             return []
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+        rtol, _ = _default_tol(outdtype)
         info = (info_keys, 2, 0, "sol2", "skinny", is_shuffle)
         return [
             (
@@ -796,7 +803,7 @@ def _get_torch_tasks(
         if indtype not in [dtypes.fp16, dtypes.bf16, dtypes.fp8]:
             return []
         M, N, K = info_keys[2], info_keys[3], info_keys[4]
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+        rtol, _ = _default_tol(outdtype)
         info = (info_keys, 0, 0, "native", "torch", is_shuffle)
         return [
             (
@@ -821,7 +828,7 @@ def _get_triton_tasks(
         if scaleAB or is_shuffle or outdtype == dtypes.fp32 or indtype != dtypes.bf16:
             return []
         M, N, K = info_keys[2], info_keys[3], info_keys[4]
-        rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+        rtol, _ = _default_tol(outdtype)
         info = (info_keys, 0, 0, "auto", "triton", is_shuffle)
         return [
             (
@@ -858,6 +865,7 @@ def _run_hipblaslt(self, ds, args):
             return []
         indtype = eval(ds["dtype"])
         outdtype = eval(ds["outdtype"])
+        rtol, atol = _default_tol(outdtype)
         gfx = self.get_gfx()
         cu_num = self.get_cu_num()
         gemmobj = HipblasltGemm(
@@ -875,6 +883,8 @@ def _run_hipblaslt(self, ds, args):
             num_warmup=10,
             timeout=args.timeout,
             verbose=args.verbose,
+            rtol=rtol,
+            atol=atol,
         )
         rets = gemmobj.run_solutions()
         gemmobj.cleanup()

csrc/opus_gemm/opus_gemm.cu

@@ -108,6 +108,7 @@ void opus_gemm(
   std::optional<aiter_tensor_t> w_scale,
   std::optional<aiter_tensor_t> bias)
 {
+  aiter_detail::g_aiter_can_throw = true;
   AITER_CHECK(XQ.dim() == 3, "XQ must be 3D [batch, M, K]");
   AITER_CHECK(WQ.dim() == 3, "WQ must be 3D [batch, N, K]");
   AITER_CHECK(Y.dim() == 3, "Y must be 3D [batch, M, N]");
@@ -248,6 +249,7 @@ void opus_gemm_a16w16_tune(
     int kernelId,
     int splitK)
 {
+  aiter_detail::g_aiter_can_throw = true;
   AITER_CHECK(XQ.dim() == 3, "XQ must be 3D [batch, M, K]");
   AITER_CHECK(WQ.dim() == 3, "WQ must be 3D [batch, N, K]");
   AITER_CHECK(Y.dim() == 3, "Y must be 3D [batch, M, N]");

gradlib/README.md

@@ -22,28 +22,51 @@ AITER_TUNE_GEMM=1 python {workload_tests}
 Captured shapes are written to `aiter/configs/bf16_untuned_gemm.csv`.
 
 ### 2) Tune GEMMs
-Run the tuner:
+
+The multi-backend tuner lives at `csrc/gemm_a16w16/gemm_a16w16_tune.py` and supports
+asm, opus, flydsl, triton, skinny, torch, and hipblaslt (opt-in) backends:
 
 ```bash
-python3 gradlib/gradlib/gemm_tuner.py \
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
   --tuned_file aiter/configs/bf16_tuned_gemm.csv \
   --input_file aiter/configs/bf16_untuned_gemm.csv
 ```
 
+To tune with a specific backend:
+
+```bash
+# opus only
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
+  --input_file aiter/configs/bf16_untuned_gemm.csv \
+  --libtype opus
+
+# all backends including hipblaslt (opt-in, requires gradlib)
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
+  --input_file aiter/configs/bf16_untuned_gemm.csv \
+  --with-hipblaslt
+
+# hipblaslt only
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
+  --input_file aiter/configs/bf16_untuned_gemm.csv \
+  --libtype hipblaslt --with-hipblaslt
+```
+
+The legacy hipblaslt-only tuner is still available at `gradlib/gradlib/gemm_tuner.py`.
+
 Tuned results are saved to `aiter/configs/bf16_tuned_gemm.csv`.
 
 Example columns:
 
-|**cu_num**|**M**|**N**|**K**|**bias**|**dtype**|**outdtype**|**scaleAB**|**bpreshuffle**|**libtype**|**solidx**|**splitK**|**soltimes**|**kernelName**|**tflops**|**bw**|
-|----------|-----|-----|-----|--------|---------|-----------|-----------|---------------|-----------|----------|----------|------------|--------------|----------|------|
-|80|128|1536|7168|False|torch.bfloat16|torch.float32|False|False|hipblaslt|667788|0|10.6|xxxxxxx|xx|xx|
+|**cu_num**|**M**|**N**|**K**|**bias**|**dtype**|**outdtype**|**scaleAB**|**bpreshuffle**|**libtype**|**solidx**|**splitK**|**us**|**kernelName**|**tflops**|**bw**|
+|----------|-----|-----|-----|--------|---------|-----------|-----------|---------------|-----------|----------|----------|------|--------------|----------|------|
+|80|128|1536|7168|False|torch.bfloat16|torch.float32|False|False|asm|5|2|10.6|bf16gemm_fp32bf16_tn_64x64_splitk_clean|xx|xx|
 
 Notes:
 - `cu_num`: compute units for current GPU.
 - `bpreshuffle`: whether weight is shuffled.
 - `dtype`: input dtype (`hipblaslt` supports fp8/bf16/fp16; asm/triton supports bf16/fp16).
-- `libtype`: kernel backend (`hipblaslt` / `rocblas` / `asm` / `triton`).
-- `splitK`: valid when `libtype == asm`.
+- `libtype`: kernel backend (`asm` / `opus` / `flydsl` / `triton` / `skinny` / `torch` / `hipblaslt`).
+- `splitK`: split-K factor (backend-dependent).
 - `tflops`: throughput in TFLOPS.
 - `bw`: bandwidth in GB/s.
 
@@ -75,6 +98,29 @@ After tuning, run your model/tests as usual.
 
 ### Tuning Configuration
 
+#### `--libtype`
+- **Type**: String (comma-separated list)
+- **Default**: `all`
+- **Description**: Choose which backends to tune: `all`, `asm`, `opus`, `flydsl`, `triton`, `skinny`, `torch`, `hipblaslt`. `hipblaslt` requires `--with-hipblaslt`.
+
+**Example**:
+```bash
+--libtype asm
+--libtype opus,flydsl
+--libtype hipblaslt --with-hipblaslt
+```
+
+#### `--with-hipblaslt`
+- **Type**: Flag (boolean)
+- **Default**: `False`
+- **Description**: Include hipblaslt in tuning (disabled by default). hipblaslt tuning is also available standalone via `gradlib/gradlib/gemm_tuner.py`.
+
+#### `--outdtype`
+- **Type**: String
+- **Default**: None (uses CSV column value or bf16)
+- **Choices**: `f32`, `f16`, `bf16`, `fp8`
+- **Description**: Output dtype override.
+
 #### `--errRatio`
 - **Type**: Float
 - **Default**: `0.05` (5%)
@@ -118,12 +164,12 @@ After tuning, run your model/tests as usual.
 **Examples**:
 ```bash
 # benchmark tuned kernels from specified tuned config
-python3 gradlib/gradlib/gemm_tuner.py \
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
   --input_file aiter/configs/bf16_untuned_gemm.csv \
   --run_config aiter/configs/bf16_tuned_gemm.csv
 
 # benchmark default kernels using shapes from --input_file
-python3 gradlib/gradlib/gemm_tuner.py \
+python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
   --input_file aiter/configs/bf16_untuned_gemm.csv \
   --run_config
 ```
@@ -179,4 +225,3 @@ export HIP_ONLINE_TUNING=1
 ```
 
 The one-time overhead can take several minutes. Results are saved to `hip_online_tuning_res.csv`.
-

gradlib/gradlib/GemmTuner.py

@@ -153,6 +153,8 @@ def __init__(
         num_warmup=10,
         timeout=None,
         verbose=False,
+        rtol=None,
+        atol=None,
     ):
         torch.cuda.empty_cache()
         self.m = m
@@ -173,8 +175,9 @@ def __init__(
         self.blob = torch.ones(128 * 1024 * 1024, dtype=dtypes.fp32, device="cuda")
         self.topn = 20
         self.hipb_sols = []
-        self.rtol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
-        self.atol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+        _tol = 5e-2 if outdtype == dtypes.bf16 else 1e-2
+        self.rtol = rtol if rtol is not None else _tol
+        self.atol = atol if atol is not None else _tol
         self.check_err_ratio = err_ratio
         self.profile_file = profile_file
         self.hipb_prefer_ratio = 0.995
@@ -525,8 +528,9 @@ def run_config(self, args):
                     eval(indtype),
                     eval(outdtype),
                 )
-                _atol = 5e-2 if eval(outdtype) == torch.bfloat16 else 1e-2
-                _rtol = 5e-2 if eval(outdtype) == torch.bfloat16 else 1e-2
+                _tol = 5e-2 if eval(outdtype) == torch.bfloat16 else 1e-2
+                _atol = _tol
+                _rtol = _tol
                 err_ratio = checkAllclose(
                     out, ref, atol=_atol, rtol=_rtol, msg=f"run_config {shape_str}"
                 )