gemini-code-assist[bot] commented on code in PR #19985:
URL: https://github.com/apache/tvm/pull/19985#discussion_r3564815288
##########
tests/python/tirx/codegen/test_codegen_cuda.py:
##########
@@ -24,10 +24,11 @@
from tvm.testing import env
-def _get_source(func: tvm.tirx.PrimFunc) -> str:
- target = tvm.target.Target("cuda")
+def _get_source(func: tvm.tirx.PrimFunc, target="cuda") -> str:
Review Comment:
The return type annotation of `_get_source` is declared as `-> str`, but the
function actually returns a tuple of `(src, mod)` (i.e., `tuple[str,
tvm.IRModule]`). This mismatch can cause confusion and type-checking issues.
```suggestion
def _get_source(func: tvm.tirx.PrimFunc, target="cuda") -> tuple[str,
tvm.IRModule]:
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy/test_reg.py:
##########
@@ -338,6 +338,7 @@ def run_and_check():
tvm.testing.run_with_gpu_lock(run_and_check)
[email protected]
def test_reg_copy_wg_local_to_swizzled_shared_uses_swizzle_fastpath():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_90a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_reg_copy_wg_local_to_swizzled_shared_uses_swizzle_fastpath():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy_async/test_tma.py:
##########
@@ -1576,6 +1577,7 @@ def copy_async_dynamic_mask(A_ptr: T.handle) -> None:
assert "multicast" in src, "Expected multicast TMA instruction in
generated code"
[email protected]
def test_copy_tma_uint32_shape_extent():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_90a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_copy_tma_uint32_shape_extent():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/gemm/test_gemm_mma_m16n8k_.py:
##########
@@ -599,6 +601,7 @@ def run_and_check():
(2, 2, 3, 8), # k8, every dim tiled
],
)
[email protected]
def test_cuda_gemm_mma_lowers_tiled(Mt, Nt, Kt, kinst):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and modify the `_lower`
helper function to use a target with a concrete architecture (e.g., `"cuda
-arch=sm_80"`) to avoid querying the host GPU.
```suggestion
def test_cuda_gemm_mma_lowers_tiled(Mt, Nt, Kt, kinst):
```
##########
tests/python/tirx/operator/tile_primitive/cuda/permute_layout/test_permute_layout.py:
##########
@@ -421,6 +421,7 @@ def f(A: T.handle, B: T.handle):
@pytest.mark.parametrize("dtype", ["uint32", "float32"])
[email protected]
def test_shared_to_shared_uses_direct_ldst(dtype):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_90a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_shared_to_shared_uses_direct_ldst(dtype):
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy_async/test_smem_tmem.py:
##########
@@ -447,6 +447,7 @@ def test_dispatch_rejects_bad_inputs(bad):
tvm.compile(tvm.IRModule({"main": kernel}), target=target,
tir_pipeline="tirx")
[email protected]
def test_multi_cp_encodes_descriptor_once_and_patches_addr():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_100a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_multi_cp_encodes_descriptor_once_and_patches_addr():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy/test_ld_stmatrix.py:
##########
@@ -469,6 +469,7 @@ def run_and_check():
tvm.testing.run_with_gpu_lock(run_and_check)
[email protected]
def test_ldstmatrix_tcgen05_warpgroup_atom_emits_ldmatrix():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
To allow it to run in regular CPU-only CI (which aligns with the PR's goal of
running compile-only tests in regular CI), we should remove the
`@pytest.mark.gpu` decorator. To avoid querying the host GPU architecture
during target creation, we can modify the `_compile_src` helper function to use
a target with a specific architecture (e.g., `"cuda -arch=sm_100a"`).
```suggestion
def test_ldstmatrix_tcgen05_warpgroup_atom_emits_ldmatrix():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy_async/test_tma.py:
##########
@@ -1003,6 +1003,7 @@ def _tma_case(
# fmt: on
[email protected]
@pytest.mark.parametrize("case", TMA_CASES)
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
Furthermore, the underlying helper `_make_tma_call` already specifies a
concrete target architecture (`"sm_90a"`), meaning it does not query the host
GPU at all. We should remove the `@pytest.mark.gpu` decorator so that this test
can run in regular CPU-only CI.
```suggestion
@pytest.mark.parametrize("case", TMA_CASES)
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy_async/test_tmem_16xnb.py:
##########
@@ -826,6 +826,7 @@ def run_and_check():
("16x256b", 64, 64), # .16x256b.x8 fp32
],
)
[email protected]
def test_alloc_tcgen05_frag_wrapper_compiles(shape, frag_rows, K_cols):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_100a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_alloc_tcgen05_frag_wrapper_compiles(shape, frag_rows, K_cols):
```
##########
tests/python/tirx/operator/tile_primitive/cuda/copy_async/test_tma.py:
##########
@@ -1609,6 +1611,7 @@ def tma_load(n: T.uint32, a_ptr: T.handle, o_ptr:
T.handle) -> None:
tvm.compile(tvm.IRModule({"main": tma_load}), target=target,
tir_pipeline="tirx")
[email protected]
def test_copy_tma_uint32_slice_base():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_90a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_copy_tma_uint32_slice_base():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/gemm/test_gemm_mma_m16n8k_.py:
##########
@@ -389,6 +390,7 @@ def test_cuda_gemm_mma_lowers_to_mma_sync(dtype):
assert f"b_local[{r}]" in script
[email protected]
def test_cuda_gemm_mma_accumulates_c_when_beta_one():
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and modify the `_lower`
helper function to use a target with a concrete architecture (e.g., `"cuda
-arch=sm_80"`) to avoid querying the host GPU.
```suggestion
def test_cuda_gemm_mma_accumulates_c_when_beta_one():
```
##########
tests/python/tirx/operator/tile_primitive/cuda/gemm/test_gemm_mma_m16n8k_.py:
##########
@@ -369,6 +369,7 @@ def test_cuda_gemm_mma_variant_is_registered():
@pytest.mark.parametrize("dtype", ["bfloat16", "float16"])
[email protected]
def test_cuda_gemm_mma_lowers_to_mma_sync(dtype):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and modify the `_lower`
helper function to use a target with a concrete architecture (e.g., `"cuda
-arch=sm_80"`) to avoid querying the host GPU.
```suggestion
@pytest.mark.parametrize("dtype", ["bfloat16", "float16"])
def test_cuda_gemm_mma_lowers_to_mma_sync(dtype):
```
##########
tests/python/tirx/operator/tile_primitive/cuda/gemm/test_gemm_mma_m16n8k_.py:
##########
@@ -644,6 +647,7 @@ def test_cuda_gemm_mma_codegen_issue_count(Mt, Nt, Kt,
kinst):
"transpose_A, transpose_B",
[(False, False), (True, False), (False, True), (True, True)],
)
[email protected]
def test_cuda_gemm_mma_lowers_transpose(transpose_A, transpose_B):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and modify the `_lower`
helper function to use a target with a concrete architecture (e.g., `"cuda
-arch=sm_80"`) to avoid querying the host GPU.
```suggestion
def test_cuda_gemm_mma_lowers_transpose(transpose_A, transpose_B):
```
##########
tests/python/tirx/operator/tile_primitive/cuda/gemm_async/test_gemm_async.py:
##########
@@ -2285,6 +2285,7 @@ def gemm_async(A_ptr: T.handle, B_ptr: T.handle, C_ptr:
T.handle) -> None:
@pytest.mark.parametrize("smem_desc", ["hoist", "recompute"])
[email protected]
def test_gemm_smem_desc_hoist_vs_recompute(smem_desc):
Review Comment:
This test is compile-only and does not execute any code on a physical GPU.
We should remove the `@pytest.mark.gpu` decorator and specify a target with a
concrete architecture (e.g., `"cuda -arch=sm_100a"`) inside the test to avoid
querying the host GPU.
```suggestion
def test_gemm_smem_desc_hoist_vs_recompute(smem_desc):
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
