[TIRX] Bind parallel loops to GPU threads before VerifyMemory

[zhils]

VerifyMemory on GPU targets treats direct accesses outside thread environments as illegal. In the ScatterValue CUDA lowering path, topi.scatter_elements emits ForKind::kParallel loops without explicit thread bindings, which triggers false host-memory access failures (e.g. "Did you forget to bind?") during TIR verification.

This change adds a new tirx pass (BindParallelLoopsToThreads) and inserts it before VerifyMemory in the s_tir pipelines (including adreno). The pass rewrites parallel loops into blockIdx.x/threadIdx.x thread-extent regions, substitutes loop vars with global thread indices, and adds bounds checks for non-divisible extents. This preserves correctness while ensuring GPU kernels pass memory verification for this path.

[gemini-code-assist]

Code Review

This pull request introduces the BindParallelLoopsToThreads pass, which converts ForKind::kParallel loops into GPU block and thread bindings, and integrates this pass into the S-TIR pipelines. Additionally, it provides a configuration option to allow unsupported host compilers for NVCC on Windows and adds a functional test for scatter operations on CUDA. Review feedback identifies a critical issue regarding the handling of nested parallel loops which could lead to invalid GPU register bindings, an inconsistency in GPU device type definitions between files, and a minor code redundancy in the loop variable substitution logic.

[tlopex]

A few things to address here:

1. Please remove .tmp_scatter_cuda_check.py. This looks like a local debugging script and should not be committed. It should be replaced with a proper pytest test instead.

2. This PR currently mixes several unrelated changes: the NVCC -allow-unsupported-compiler workaround, the loop_partition.cc Var{} fix, and the kDLWebGPU change in verify_memory.cc are all independent and should be split into separate PRs.

3. Please clarify the semantics for nested parallel loops. When in_parallel_loop_ is already true, inner parallel loops are left as kParallel, but once they are inside a GPU kernel they effectively become serial. If that is intentional, it would be good to document it explicitly; otherwise it may be better to reject this case.

4. The else branch in IfThenElse looks unnecessary. Evaluate(IntImm(DataType::Int(32), 0)) is a no-op, and TIR already supports IfThenElse without an else, so this can just be IfThenElse(global_idx < extent, mapped_body).

5. For target-less PrimFuncs, please avoid falling back to Target::Current(). If the function has no target attribute, it is better to leave it unchanged rather than guessing from ambient context. That would also be consistent with how other recent passes handle target-less functions.

6. This PR also needs proper pytest coverage for the new pass.

[zhils]

A few things to address here:

1. Please remove .tmp_scatter_cuda_check.py. This looks like a local debugging script and should not be committed. It should be replaced with a proper pytest test instead.
2. This PR currently mixes several unrelated changes: the NVCC -allow-unsupported-compiler workaround, the loop_partition.cc Var{} fix, and the kDLWebGPU change in verify_memory.cc are all independent and should be split into separate PRs.
3. Please clarify the semantics for nested parallel loops. When in_parallel_loop_ is already true, inner parallel loops are left as kParallel, but once they are inside a GPU kernel they effectively become serial. If that is intentional, it would be good to document it explicitly; otherwise it may be better to reject this case.
4. The else branch in IfThenElse looks unnecessary. Evaluate(IntImm(DataType::Int(32), 0)) is a no-op, and TIR already supports IfThenElse without an else, so this can just be IfThenElse(global_idx < extent, mapped_body).
5. For target-less PrimFuncs, please avoid falling back to Target::Current(). If the function has no target attribute, it is better to leave it unchanged rather than guessing from ambient context. That would also be consistent with how other recent passes handle target-less functions.
6. This PR also needs proper pytest coverage for the new pass.

Thanks for the review. I removed .tmp_scatter_cuda_check.py and added tests/python/tirx-transform/test_tir_transform_bind_parallel_loops_to_threads.py for the pass. The NVCC workaround, loop_partition Var{} fix, and verify_memory WebGPU change are reverted here and will be submitted as separate PRs. Nested T.parallel is now rejected with a clear error; the pass doc comment explains semantics (including leaving kParallel unchanged under an existing thread env). IfThenElse no longer has a no-op else. PrimFuncs without kTarget are left unchanged (no Target::Current fallback).