From aa9ede0dd2466ac49c85fae4e4f4272ef0e8168d Mon Sep 17 00:00:00 2001 From: rankaiyx Date: Fri, 4 Jul 2025 09:02:07 +0800 Subject: [PATCH 1/3] Update scan.py to fix pascal error. --- python/tvm/topi/gpu/scan.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/python/tvm/topi/gpu/scan.py b/python/tvm/topi/gpu/scan.py index 5be4033e4575..ae32c707f5a1 100644 --- a/python/tvm/topi/gpu/scan.py +++ b/python/tvm/topi/gpu/scan.py @@ -40,6 +40,8 @@ def _can_use_scan_thrust(binop): target = tvm.target.Target.current() if target is None: return False + if target.arch == "sm_60": + return False # pylint: disable=comparison-with-callable return binop == tvm.tir.generic.add and any( [ From 7d672a1ab8e6f29c73ee5b97bd8729d31e92ca1f Mon Sep 17 00:00:00 2001 From: rankaiyx Date: Tue, 26 Aug 2025 08:37:35 +0800 Subject: [PATCH 2/3] Update scan.py to fix Pascal P40 error. --- python/tvm/topi/gpu/scan.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tvm/topi/gpu/scan.py b/python/tvm/topi/gpu/scan.py index ae32c707f5a1..2a0814d69d8d 100644 --- a/python/tvm/topi/gpu/scan.py +++ b/python/tvm/topi/gpu/scan.py @@ -40,7 +40,7 @@ def _can_use_scan_thrust(binop): target = tvm.target.Target.current() if target is None: return False - if target.arch == "sm_60": + if target.arch in ("sm_60", "sm_61"): return False # pylint: disable=comparison-with-callable return binop == tvm.tir.generic.add and any( From e26bf686fa9df2d4e6683b69c8a0e32ed937bdfc Mon Sep 17 00:00:00 2001 From: rankaiyx Date: Tue, 26 Aug 2025 10:18:31 +0800 Subject: [PATCH 3/3] Update scan.py --- python/tvm/topi/gpu/scan.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/tvm/topi/gpu/scan.py b/python/tvm/topi/gpu/scan.py index 2a0814d69d8d..a4019403bb6e 100644 --- a/python/tvm/topi/gpu/scan.py +++ b/python/tvm/topi/gpu/scan.py @@ -40,6 +40,7 @@ def _can_use_scan_thrust(binop): target = tvm.target.Target.current() if target is None: return False + # Excluding Pascal architecture, use fallback path code if target.arch in ("sm_60", "sm_61"): return False # pylint: disable=comparison-with-callable