You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "Aminsed (via GitHub)" <gi...@apache.org> on 2025/09/27 05:22:29 UTC
[PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Aminsed opened a new pull request, #18347:
URL: https://github.com/apache/tvm/pull/18347
This PR adds a deterministic, schedule-free CUDA test to validate shared memory limits using tir.analysis.verify_gpu_code.\n\n- Adds a direct TIR kernel test that allocates a 64 KB shared buffer (T.launch_thread + T.allocate) and asserts verification fails when max_shared_memory_per_block=48 KB.\n- Keeps changes scoped to tests; no default behavior altered.\n- Documents an opt-in pipeline toggle (off by default) and a per-device cap hint in tvm/python/tvm/tir/pipeline.py.\n Usage:\n \n\nReferences: #18094
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on PR #18347:
URL: https://github.com/apache/tvm/pull/18347#issuecomment-3342151102
given this is something that can be composed, we will likely not need this in the default pipeline
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on code in PR #18347:
URL: https://github.com/apache/tvm/pull/18347#discussion_r2383895791
##########
python/tvm/tir/pipeline.py:
##########
@@ -120,6 +120,21 @@ def _pipeline(mod: tvm.ir.IRModule, _ctx: tvm.transform.PassContext) -> tvm.ir.I
tir.transform.LowerDeviceKernelLaunch(),
]
)
+ # Optional GPU verification based on PassContext configuration.
+ # Usage example:
+ #
+ # with tvm.transform.PassContext(config={
+ # "tir.verify_gpu_code": True,
+ # # Optional per-device cap override; for Ampere (e.g., RTX A6000, SM 86),
+ # # you may choose 96 KB (98304 bytes). Default is conservative 48 KB.
+ # # "tir.cuda.max_shared_memory_per_block": 96 * 1024,
+ # }):
+ # lib = tvm.tir.build(mod, target="cuda")
+ #
+ # This check is opt-in and does not change defaults.
+ if bool(config.get("tir.verify_gpu_code", False)):
+ cap = int(config.get("tir.cuda.max_shared_memory_per_block", 48 * 1024))
Review Comment:
The `bool()` and `int()` type casts are redundant. `config.get()` with a default value already returns the correct type, and the value from the first `get` can be used directly in the `if` condition.
```suggestion
if config.get("tir.verify_gpu_code", False):
cap = config.get("tir.cuda.max_shared_memory_per_block", 48 * 1024)
```
##########
tests/python/relay/test_cuda_shared_mem_overflow.py:
##########
@@ -0,0 +1,45 @@
+"""
+Tests for verifying shared memory limit using a deterministic TIR kernel.
+
+To enable optional pipeline verification during normal builds (opt-in, off by default):
+
+ with tvm.transform.PassContext(config={
+ "tir.verify_gpu_code": True,
+ # For Ampere (e.g., RTX A6000, SM 86), you may set a higher cap, e.g., 96 KB:
+ # "tir.cuda.max_shared_memory_per_block": 96 * 1024,
+ # By default, leave unset to use a conservative 48 KB.
+ }):
+ lib = tvm.tir.build(mod, target="cuda")
+
+This test avoids schedule/lowering variability by using a direct kernel that allocates a
+64 KB shared buffer and asserts the verifier fails when the cap is 48 KB.
+"""
+
+import pytest
+
+import tvm
+from tvm import tir
+from tvm.script import tir as T
+
+
+@T.prim_func
+def _pf_direct_kernel_shared_large(A: T.handle) -> None:
+ T.func_attr({"global_symbol": "_pf_direct_kernel_shared_large", "tir.noalias": True})
+ A_buf = T.match_buffer(A, (1,), dtype="float32")
+ blockIdx_x = T.launch_thread("blockIdx.x", 1)
+ threadIdx_x = T.launch_thread("threadIdx.x", 1)
+ # 16384 float32 elements = 64 KB shared allocation
+ sh = T.allocate([16384], "float32", "shared")
+ s = T.float32(0)
+ for t in T.serial(0, 16384):
+ s = s + T.float32(1)
Review Comment:
The allocated shared memory buffer `sh` is not used in the loop. While this works for the verifier, using the buffer would make the test's intent clearer and more robust against potential dead-code elimination if this function were used in a context with more optimizations. For example:
```python
for t in T.serial(0, 16384):
sh[t] = T.float32(1)
s = s + sh[t]
```
##########
tests/python/relay/test_cuda_shared_mem_overflow.py:
##########
@@ -0,0 +1,45 @@
+"""
Review Comment:
This test appears to be for TIR analysis rather than Relay. To improve project structure and discoverability, consider moving this file to a more appropriate location, such as under `tests/python/tir/`.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "Aminsed (via GitHub)" <gi...@apache.org>.
Aminsed commented on PR #18347:
URL: https://github.com/apache/tvm/pull/18347#issuecomment-3344251224
Thanks for the guidance. I have removed the optional VerifyGPUCode hook from the default TIR pipeline to keep the pipeline minimal. The PR now recommends composing verification explicitly after lowering.\n\nIn particular, users can:\n- Call the analysis helper directly on a PrimFunc: , or\n- Compose a short post-pipeline sequence including .\n\nThe deterministic test remains and demonstrates the analysis on a tiny kernel without changing defaults. Please let me know if you prefer any further adjustments.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen commented on PR #18347:
URL: https://github.com/apache/tvm/pull/18347#issuecomment-3346561926
Thanks, given the remaining changes can be done in downstream, we will close this for now
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "gemini-code-assist[bot] (via GitHub)" <gi...@apache.org>.
gemini-code-assist[bot] commented on PR #18347:
URL: https://github.com/apache/tvm/pull/18347#issuecomment-3341215296
## Summary of Changes
Hello @Aminsed, I'm Gemini Code Assist[^1]! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!
This pull request introduces a robust, deterministic CUDA test to validate the `tir.analysis.verify_gpu_code` utility's ability to detect shared memory overflows. By creating a direct TIR kernel that attempts to allocate 64KB of shared memory and asserting its failure against a 48KB limit, the PR ensures the verification mechanism functions as expected. Additionally, it integrates an opt-in pipeline toggle and documentation for configuring per-device shared memory caps, enhancing the flexibility and diagnostic capabilities for GPU code without altering any default system behaviors.
### Highlights
* **New Deterministic CUDA Test**: A new, deterministic CUDA test has been added to validate shared memory limits using `tir.analysis.verify_gpu_code`.
* **Shared Memory Overflow Verification**: The test specifically verifies that a 64KB shared buffer allocation fails verification when the maximum allowed shared memory per block is set to 48KB.
* **Opt-in GPU Code Verification**: An optional, opt-in GPU code verification step has been integrated into the TIR pipeline, allowing users to enable shared memory limit checks via `PassContext` configuration.
* **Configurable Shared Memory Cap**: The `tir.cuda.max_shared_memory_per_block` configuration option allows overriding the per-device shared memory cap for verification, with a default conservative limit of 48KB.
* **No Default Behavior Altered**: All changes are confined to testing and optional configurations, ensuring no existing default behaviors are modified.
<details>
<summary><b>Using Gemini Code Assist</b></summary>
<br>
The full guide for Gemini Code Assist can be found on our [documentation page](https://developers.google.com/gemini-code-assist/docs/review-github-code), here are some quick tips.
<b>Invoking Gemini</b>
You can request assistance from Gemini at any point by creating a comment using either `/gemini <command>` or `@gemini-code-assist <command>`. Below is a summary of the supported commands on the current page.
Feature | Command | Description
--- | --- | ---
Code Review | `/gemini review` | Performs a code review for the current pull request in its current state.
Pull Request Summary | `/gemini summary` | Provides a summary of the current pull request in its current state.
Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in pull request comments and review comments.
Help | `/gemini help` | Displays a list of available commands.
<b>Customization</b>
To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a `.gemini/` folder in the base of the repository. Detailed instructions can be found [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github).
<b>Limitations & Feedback</b>
Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with :thumbsup: and :thumbsdown: on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up [here](https://google.qualtrics.com/jfe/form/SV_2cyuGuTWsEw84yG).
<b>You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the [Gemini Code Assist IDE Extension](https://cloud.google.com/products/gemini/code-assist).</b>
</details>
[^1]: Review the [Privacy Notices](https://policies.google.com/privacy), [Generative AI Prohibited Use Policy](https://policies.google.com/terms/generative-ai/use-policy), [Terms of Service](https://policies.google.com/terms), and learn how to configure Gemini Code Assist in GitHub [here](https://developers.google.com/gemini-code-assist/docs/customize-gemini-behavior-github). Gemini can make mistakes, so double check it and [use code with caution](https://support.google.com/legal/answer/13505487).
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB) [tvm]
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen closed pull request #18347: tests(cuda): deterministic VerifyGPUCode test for shared-memory overflow (64KB vs 48KB)
URL: https://github.com/apache/tvm/pull/18347
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org