[NFC] FileCheck tests check all overloads by ashermancinelli · Pull Request #354 · NVIDIA/numba-cuda

ashermancinelli · 2025-07-26T00:16:39Z

Clean up type hints
On failure, dump the failing IR to a file and provide a reproducer command
- When updating tests, this is extremely helpful in my opinion. Right now, it is enabled with an environment variable, and I am not sure if this is the right option. It is not very discoverable.
When multiple overloaded IR is available for a kernel and no signature has been passed, check all of the overloads

TODO:

Need to decide: is an environment variable how we want to toggle the dumping? I don't think it should happen by default in case multiple filecheck tests run and fail in parallel.
- We will use a pytest command line option registered in conftest.py.

copy-pr-bot · 2025-07-26T00:16:42Z

Auto-sync is disabled for ready for review pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

ashermancinelli · 2025-07-26T00:16:56Z

/ok to test 66f2f8c

numba_cuda/numba/cuda/testing.py

gmarkall · 2025-07-28T07:54:09Z

numba_cuda/numba/cuda/testing.py

-            check_prefixes=check_prefixes,
-            **extra_filecheck_options,
+        ir_contents = self._getIRContent(ir_producer.inspect_asm(), signature)
+        assert ir_contents


Hitting a bare assert is usually not the nicest debug experience. What is this guarding against?

That the IR given to us by the user was not empty - Are assertions in the assert methods of a test case still used to indicate a violated precondition, and self.fail/skip is used to indicate the result of the assertion represented by the method? Assuming so, I just added messages to these assertions. Please let me know if you had something else in mind.

numba_cuda/numba/cuda/testing.py

gmarkall · 2025-07-28T08:17:57Z

Need to decide: is an environment variable how we want to toggle the dumping? I don't think it should happen by default in case multiple filecheck tests run and fail in parallel.

Now we've switched over to pytest, it might be nice to add an option with https://docs.pytest.org/en/stable/reference/reference.html#pytest.hookspec.pytest_addoption so that user can run e.g.:

pytest --dump-failed-filechecks --pyargs numba.cuda.tests

I don't think the name of the file should be hardcoded - perhaps file names could be derived from the name of the test case (and potentially suffixed with a numeric index if there is a name clash due to things like subtests?

ashermancinelli · 2025-07-28T14:20:37Z

Need to decide: is an environment variable how we want to toggle the dumping? I don't think it should happen by default in case multiple filecheck tests run and fail in parallel.

Now we've switched over to pytest, it might be nice to add an option with https://docs.pytest.org/en/stable/reference/reference.html#pytest.hookspec.pytest_addoption so that user can run e.g.:
pytest --dump-failed-filechecks --pyargs numba.cuda.tests
I don't think the name of the file should be hardcoded - perhaps file names could be derived from the name of the test case (and potentially suffixed with a numeric index if there is a name clash due to things like subtests?

I love this suggestion. I added it in the latest revision, so now the failing IR and checks are dumped to files like this: numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.ll. Is that what you had in mind? Do you think we should try to rename it with a numeric suffix if it already exists, or is the name so specific that we don't mind overwriting a dump from a previous failing run?

ashermancinelli · 2025-07-28T15:45:49Z

/ok to test 67642ed

ashermancinelli · 2025-07-28T16:09:01Z

/ok to test 67642ed

ashermancinelli · 2025-07-28T20:55:01Z

/ok to test 49a007a

isVoid · 2025-07-28T21:53:17Z

numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.ll. Is that what you had in mind? Do you think we should try to rename it with a numeric suffix if it already exists, or is the name so specific that we don't mind overwriting a dump from a previous failing run?

One idea of separating them from one test run from another is placing the dump files in a folder named numba-cuda-tests_YYYY:MM:DD:HH:MM:SS and place the dump there. But second level internal may not be sufficient to isolate repeated runs.

I think isolation is necessary when pytest --count is used to test bugs which is not deterministic.

ashermancinelli · 2025-07-28T22:22:15Z

numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.ll. Is that what you had in mind? Do you think we should try to rename it with a numeric suffix if it already exists, or is the name so specific that we don't mind overwriting a dump from a previous failing run?

One idea of separating them from one test run from another is placing the dump files in a folder named numba-cuda-tests_YYYY:MM:DD:HH:MM:SS and place the dump there. But second level internal may not be sufficient to isolate repeated runs.

I think isolation is necessary when pytest --count is used to test bugs which is not deterministic.

I like that idea! Feel free to let me know what's most intuitive and I'll do it. If I had to guess, I think the workflow will look like this:

Make a commit that changes the IR we produce somehow
Run the tests and dump the failing IR
Make sure the differences are what we expected
Regenerate and update the checks in one test
Rerun, and if any tests are still failing, goto step 2

I don't envision myself needing to keep multiple directories with IR dumps from failing tests. I think a directory with a timestamp in the name would be a great way to accomplish that though.

isVoid · 2025-07-28T22:37:55Z

@ashermancinelli The workflow sounds right. I think my original language is misleading as "second level interval" is not referring to multiple directories. I meant using second as the smallest time interval may not be sufficient as pytest --count can rerun tests in sub-second frequency. What do you think of adding millisecond interval in the timestamp?

ashermancinelli · 2025-07-28T22:53:20Z

Okay, the IR is now dumped in this format on error:


> filecheck --check-prefixes=LLVM numba-ir-2025_07_28_15_49_27/numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.checks --input-file=numba-ir-2025_07_28_15_49_27/numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.ll
numba-ir-2025_07_28_15_49_27/numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.checks:9: error: Couldn't match "bar".
Current position at numba-ir-2025_07_28_15_49_27/numba_cuda_tests_cudapy_test_inspect_TestInspect_test_monotyped.ll:9:1
}
^

It is very verbose, but we will be unlikely to see collisions. Thanks for the review @isVoid!

ashermancinelli · 2025-07-28T22:53:43Z

/ok to test 10e93cc

On failure, IR is dumped to a directory based on the current date and time and the file is based on the test name. This way, the IR is preserved between failures and multiple failures in parallel in different tests do not cause collisions.

ashermancinelli · 2025-07-29T15:35:44Z

/ok to test d41fb5e

ashermancinelli · 2025-07-30T02:22:59Z

/ok to test b0bed24

gmarkall · 2025-07-30T14:06:52Z

numba_cuda/numba/cuda/testing.py

+                    _ = ir_file.write(ir_content + "\n")
+                    _ = checks_file.write(check_patterns)


I'm a little confused here: if you don't want the return value, why assign it?

gmarkall · 2025-07-30T14:10:09Z

I forgot to mention - I think the dumps for subtests will still overwrite each other if they don't form part of the test's id() (I'm not sure if they do or not) but I don't think this is a massive problem - one case is sufficient to identify an issue, and after fixing it if a different subtest still has a problem that will show up on a re-run.

- [NFC] FileCheck tests check all overloads (NVIDIA#354) - [REVIEW][NFC] Vendor in serialize to allow for future CUDA-specific refactoring and changes (NVIDIA#349) - Vendor in usecases used in testing (NVIDIA#359) - Add thirdparty tests of numba extensions (NVIDIA#348) - Support running tests in parallel (NVIDIA#350) - Add more debuginfo tests (NVIDIA#358) - [REVIEW][NFC] Vendor in the Cache, CacheImpl used by CUDACache and CUDACacheImpl to allow for future CUDA-specific refactoring and changes (NVIDIA#334) - [NFC] Vendor in Dispatcher as CUDADispatcher to allow for future CUDA-specific customization (NVIDIA#338) - Vendor in BaseNativeLowering and BaseLower for CUDA-specific customizations (NVIDIA#329) - [REVIEW] Vendor in the CompilerBase used by CUDACompiler to allow for future CUDA-specific refactoring and changes (NVIDIA#322) - Vendor in Codegen and CodeLibrary for CUDA-specific customization (NVIDIA#327) - Disable tests that deadlock due to NVIDIA#317 (NVIDIA#356) - FIX: Add type check for shape elements in DeviceNDArrayBase constructor (NVIDIA#352) - Merge pull request NVIDIA#265 from lakshayg/fp16-support - Add performance warning - Fix tests - Create and register low++ bindings for float16 - Create typing/target registries for float16 - Replace Numbast generated lower_casts - Replace Numbast generated operators - Alias __half to numba.core.types.float16 - Generate fp16 bindings using numbast - Remove existing fp16 logic - [REVIEW][NFC] Vendor in the utils and cgutils to allow for future CUDA-specific refactoring and changes (NVIDIA#340) - [RFC,TESTING] Add filecheck test infrastructure (NVIDIA#342) - Migrate test infra to pytest (NVIDIA#347) - Add .vscode to gitignore (NVIDIA#344) - [NFC] Add dev dependencies to project config (NVIDIA#341) - Allow Inspection of Link-Time Optimized PTX (NVIDIA#326) - [NFC] Vendor in DIBuilder used by CUDADIBuilder (NVIDIA#332) - Add guidance on setting up pre-commit (NVIDIA#339) - [Refactor][NFC] Vendor in MinimalCallConv (NVIDIA#333) - [Refactor][NFC] Vendor in BaseCallConv (NVIDIA#324) - [REVIEW] Vendor in CompileResult as CUDACompileResult to allow for future CUDA-specific customizations (NVIDIA#325)

- [NFC] FileCheck tests check all overloads (#354) - [REVIEW][NFC] Vendor in serialize to allow for future CUDA-specific refactoring and changes (#349) - Vendor in usecases used in testing (#359) - Add thirdparty tests of numba extensions (#348) - Support running tests in parallel (#350) - Add more debuginfo tests (#358) - [REVIEW][NFC] Vendor in the Cache, CacheImpl used by CUDACache and CUDACacheImpl to allow for future CUDA-specific refactoring and changes (#334) - [NFC] Vendor in Dispatcher as CUDADispatcher to allow for future CUDA-specific customization (#338) - Vendor in BaseNativeLowering and BaseLower for CUDA-specific customizations (#329) - [REVIEW] Vendor in the CompilerBase used by CUDACompiler to allow for future CUDA-specific refactoring and changes (#322) - Vendor in Codegen and CodeLibrary for CUDA-specific customization (#327) - Disable tests that deadlock due to #317 (#356) - FIX: Add type check for shape elements in DeviceNDArrayBase constructor (#352) - Merge pull request #265 from lakshayg/fp16-support - Add performance warning - Fix tests - Create and register low++ bindings for float16 - Create typing/target registries for float16 - Replace Numbast generated lower_casts - Replace Numbast generated operators - Alias __half to numba.core.types.float16 - Generate fp16 bindings using numbast - Remove existing fp16 logic - [REVIEW][NFC] Vendor in the utils and cgutils to allow for future CUDA-specific refactoring and changes (#340) - [RFC,TESTING] Add filecheck test infrastructure (#342) - Migrate test infra to pytest (#347) - Add .vscode to gitignore (#344) - [NFC] Add dev dependencies to project config (#341) - Allow Inspection of Link-Time Optimized PTX (#326) - [NFC] Vendor in DIBuilder used by CUDADIBuilder (#332) - Add guidance on setting up pre-commit (#339) - [Refactor][NFC] Vendor in MinimalCallConv (#333) - [Refactor][NFC] Vendor in BaseCallConv (#324) - [REVIEW] Vendor in CompileResult as CUDACompileResult to allow for future CUDA-specific customizations (#325)

ashermancinelli requested a review from gmarkall July 26, 2025 00:16

ashermancinelli self-assigned this Jul 26, 2025

ashermancinelli added the 3 - Ready for Review Ready for review by team label Jul 26, 2025

gmarkall reviewed Jul 28, 2025

View reviewed changes

numba_cuda/numba/cuda/testing.py Outdated Show resolved Hide resolved

gmarkall reviewed Jul 28, 2025

View reviewed changes

numba_cuda/numba/cuda/testing.py Outdated Show resolved Hide resolved

ashermancinelli force-pushed the ajm/filecheck-fixes branch from 23ddaab to 67642ed Compare July 28, 2025 14:32

ashermancinelli requested a review from gmarkall July 28, 2025 15:49

ashermancinelli requested a review from isVoid July 28, 2025 23:28

ashermancinelli added 5 commits July 29, 2025 08:27

[NFC] FileCheck tests check all overloads

795bc76

Use command line option for dumping failing IR

9af5319

Remove unnecessary comment

fd4a7bf

Dump IR to timestamp directory

c6fd1af

On failure, IR is dumped to a directory based on the current date and time and the file is based on the test name. This way, the IR is preserved between failures and multiple failures in parallel in different tests do not cause collisions.

Integrate main branch

d41fb5e

ashermancinelli force-pushed the ajm/filecheck-fixes branch from 10e93cc to d41fb5e Compare July 29, 2025 15:35

Merge branch 'main' into ajm/filecheck-fixes

b0bed24

gmarkall reviewed Jul 30, 2025

View reviewed changes

gmarkall approved these changes Jul 30, 2025

View reviewed changes

gmarkall merged commit f762fce into NVIDIA:main Jul 30, 2025
41 checks passed

gmarkall mentioned this pull request Jul 31, 2025

Bump version to 0.18.0 #365

Merged

		_ = ir_file.write(ir_content + "\n")
		_ = checks_file.write(check_patterns)

Conversation

ashermancinelli commented Jul 26, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot bot commented Jul 26, 2025

Uh oh!

ashermancinelli commented Jul 26, 2025

Uh oh!

Uh oh!

gmarkall Jul 28, 2025

Choose a reason for hiding this comment

Uh oh!

ashermancinelli Jul 28, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

gmarkall commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

isVoid commented Jul 28, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

isVoid commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 28, 2025

Uh oh!

ashermancinelli commented Jul 29, 2025

Uh oh!

ashermancinelli commented Jul 30, 2025

Uh oh!

gmarkall Jul 30, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

gmarkall commented Jul 30, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

ashermancinelli commented Jul 26, 2025 •

edited

Loading

isVoid commented Jul 28, 2025 •

edited

Loading