Fix CUDA kernel index data type in vllm/csrc/quantization/gptq_marlin/awq_marlin_repack.cu +10

[houseroad]

Summary:
CUDA kernel variables matching the type (thread|block|grid).(Idx|Dim).(x|y|z) have the data type uint.

Many programmers mistakenly use implicit casts to turn these data types into int. In fact, the CUDA Programming Guide it self is inconsistent and incorrect in its use of data types in programming examples.

The result of these implicit casts is that our kernels may give unexpected results when exposed to large datasets, i.e., those exceeding >~2B items.

While we now have linters in place to prevent simple mistakes (D71236150), our codebase has many problematic instances. This diff fixes some of them.

Differential Revision: D71355454

[github-actions]

👋 Hi! Thank you for contributing to the vLLM project.

💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in #pr-reviews, coordinate on features in #feat- channels, or join special interest groups in #sig- channels.

Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors. You can run other CI tests on top of those by going to your fastcheck build on Buildkite UI (linked in the PR checks section) and unblock them. If you do not have permission to unblock, ping simon-mo or khluu to add you in our Buildkite org.

Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging.

To run CI, PR reviewers can either: Add ready label to the PR or enable auto-merge.

🚀

[facebook-github-bot]

This pull request was exported from Phabricator. Differential Revision: D71355454

[comaniac]

LGTM. Also cc @tlrmchlsmth

[DarkLight1337]

Can you update this after #15282 has been merged?

[houseroad]

rebased. hope the CI can be fixed. Keep running into timeout... sigh.

[DarkLight1337]

Kernels tests aren't failing on main, so maybe there is indeed something wrong with this PR?

[houseroad]

kernel test-2:

attempt 1, failed on

[2025-03-22T08:06:54Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors1-128-quant_type1-256-128] PASSED

  | [2025-03-22T09:27:08Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors2--1-quant_type1-256-128] # Received cancellation signal, interrupting

attempt 2, failed on

[2025-03-23T00:44:47Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors1-128-quant_type1-256-128] PASSED

  | [2025-03-23T02:03:09Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors2--1-quant_type1-256-128] # Received cancellation signal, interrupting
  | [2025-03-23T02:03:11Z] 🚨 Error: The command exited with status -1

attempt3 failed on

[2025-03-23T06:29:36Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors1-64-quant_type0-256-128] PASSED

  | [2025-03-23T06:29:36Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors1-128-quant_type1-256-128] PASSED
  | [2025-03-23T07:50:10Z] kernels/test_marlin_gemm.py::test_gptq_marlin_gemm[False-False-False-False-mnk_factors2--1-quant_type1-256-128] # Received cancellation signal, interrupting
  | [2025-03-23T07:50:12Z] 🚨 Error: The command exited with status -1

Seems problematic, let me take a close look

[houseroad]

I think I found the problematic place:

int delta_first = iters * blockIdx.x - col_first; // this shouldn't be changed to auto. Since the type deduce to uint, where iters and col_first are int, and blockIdx.x is uint.

[houseroad]

The problems should be fixed, cc: @DarkLight1337

[DarkLight1337]

Pre-commit is failing

[houseroad]

It shows all 60 checks passed

[DarkLight1337]

Weird, the mobile app showed that it's failing. I'm back on my PC now and everything looks fine, sorry for the confusion!