[FRONTEND] Better error for num_cta > 1 on sm < 90 (#7812)

Fixes triton-lang/triton#7811

Author: lezcano

python/test/unit/language/test_core.py

@@ -7190,6 +7190,26 @@ def simple(data, out):
         assert amdgcn_gfx.group(1) == arch
 
 
+def test_num_ctas_pre_sm90(device):
+    if not is_cuda() and not is_hip():
+        pytest.skip("Only supported on CUDA and HIP")
+
+    @triton.jit
+    def _kernel(src):
+        pass
+
+    src = torch.empty(1, device=device)
+    if is_cuda():
+        arch = "sm80"
+        msg = r"num_ctas > 1 requires NVIDIA SM90\+ \(Hopper\)"
+    else:
+        arch = "gfx942"
+        msg = r"num_ctas > 1 not supported for AMD GPUs"
+
+    with pytest.raises(ValueError, match=msg):
+        _kernel.warmup(src, grid=(1, ), num_ctas=2, arch=arch)
+
+
 # -----------------------
 # test propagate_nan
 # -----------------------

third_party/amd/backend/compiler.py

@@ -110,6 +110,9 @@ def get_target_name(self, options) -> str:
     def parse_options(self, opts) -> Any:
         args = {'arch': knobs.runtime.override_arch or self.target.arch}
 
+        if opts.get("num_ctas", 1) > 1:
+            raise ValueError("num_ctas > 1 not supported for AMD GPUs")
+
         # Enable XF32 (TF32) for CDNA3 GPUs
         if self.target.arch == 'gfx942':
             allowed_dot_input_precisions = set(HIPOptions.allowed_dot_input_precisions)

third_party/nvidia/backend/compiler.py

@@ -174,6 +174,11 @@ def parse_options(self, opts) -> Any:
         args.update({k: opts[k] for k in CUDAOptions.__dataclass_fields__.keys() if k in opts if opts[k] is not None})
         capability = int(self._parse_arch(args["arch"]))
 
+        if args.get("num_ctas", 1) > 1 and capability < 90:
+            raise ValueError((f"num_ctas > 1 requires NVIDIA SM90+ (Hopper). "
+                              f"Current target is sm_{capability}. This configuration will fail. "
+                              f"Please set num_ctas=1 or target an SM90+ GPU."))
+
         if "supported_fp8_dtypes" not in args:
             supported_fp8_dtypes = set(CUDAOptions.supported_fp8_dtypes)
             if capability >= 89: