[Testing] Restore tma gather scatter fallback tests (#7619)

Fixes #7596

Author: peterbell10

python/test/unit/language/test_tensor_descriptor.py

@@ -4,7 +4,7 @@
 
 import triton
 import triton.language as tl
-from triton._internal_testing import is_blackwell, is_hopper, is_interpreter, numpy_random, to_triton, unwrap_tensor, tma_dtypes, to_numpy
+from triton._internal_testing import is_hopper, is_interpreter, numpy_random, to_triton, unwrap_tensor, tma_dtypes, to_numpy
 from triton.tools.mxfp import MXFP4Tensor, MXScaleTensor
 from typing import Optional
 from triton._internal_testing import is_cuda, is_hip, is_hip_cdna3
@@ -1315,7 +1315,7 @@ def torch_gather_rows(input, idx, y, block_y):
 @pytest.mark.parametrize("BLOCK_X, BLOCK_Y", [(32, 32), (64, 128), (16, 128), (512, 16)])
 @pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.int8])
 @pytest.mark.parametrize("y", [0, 32, 48])
-@pytest.mark.skipif(not is_blackwell(), reason="TMA Gather requires blackwell")
+@pytest.mark.skipif(is_hopper(), reason="TMA Scatter is not supported on hopper")
 def test_tma_gather(X, Y, BLOCK_X, BLOCK_Y, dtype, y, device):
     if BLOCK_X > X or y + BLOCK_Y > Y:
         pytest.skip()
@@ -1367,7 +1367,7 @@ def tma_gather_dot_pipeline(  #
 @pytest.mark.interpreter
 @pytest.mark.parametrize("BLOCK_M, BLOCK_N, BLOCK_K", [(16, 16, 16)])
 @pytest.mark.parametrize("K", [128])
-@pytest.mark.skipif(not is_blackwell(), reason="TMA Gather requires blackwell")
+@pytest.mark.skipif(is_hopper(), reason="TMA Scatter is not supported on hopper")
 def test_tma_gather_dot_pipeline(BLOCK_M, BLOCK_N, BLOCK_K, K, device):
 
     def alloc_fn(size: int, align: int, steam):
@@ -1414,7 +1414,7 @@ def tma_scatter_rows_kernel(out_ptr, in_ptr, idx_ptr, y, X: tl.constexpr, Y: tl.
 @pytest.mark.parametrize("BLOCK_X, BLOCK_Y", [(32, 32), (64, 128), (16, 128), (512, 16)])
 @pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.int8])
 @pytest.mark.parametrize("y", [0, 32, 48])
-@pytest.mark.skipif(not is_blackwell(), reason="TMA Scatter requires blackwell")
+@pytest.mark.skipif(is_hopper(), reason="TMA Scatter is not supported on hopper")
 def test_tma_scatter(X, Y, BLOCK_X, BLOCK_Y, dtype, y):
     if BLOCK_X > X or y + BLOCK_Y > Y:
         pytest.skip()