Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Improve program checks #394

Open
wants to merge 6 commits into
base: main
Choose a base branch
from

Conversation

ksimpson-work
Copy link
Contributor

@ksimpson-work ksimpson-work commented Jan 14, 2025

This review addresses the suggestions in 313 for improving some of the runtime checks in program.
close #313

Copy link
Contributor

copy-pr-bot bot commented Jan 14, 2025

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@ksimpson-work ksimpson-work self-assigned this Jan 14, 2025
@ksimpson-work ksimpson-work added P1 Medium priority - Should do cuda.core Everything related to the cuda.core module labels Jan 14, 2025
@ksimpson-work ksimpson-work added this to the cuda.core beta 3 milestone Jan 14, 2025
@ksimpson-work ksimpson-work force-pushed the improve-program-checks branch from 17f13b7 to b2495d3 Compare January 21, 2025 00:03
@ksimpson-work
Copy link
Contributor Author

/ok to test

Copy link

@ksimpson-work ksimpson-work marked this pull request as ready for review January 22, 2025 17:01
@ksimpson-work ksimpson-work requested a review from leofang January 22, 2025 17:01
@ksimpson-work
Copy link
Contributor Author

/ok to test

options.arch is not None and int(options.arch.split("_")[-1]) not in supported_archs
)
default_arch_not_supported = (
options.arch is None
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is no longer possible now that we default to the current arch when creating ProgramOptions without arch being set, so this logic here and below can be removed/simplified.

Comment on lines +461 to +466
version = handle_return(nvrtc.nvrtcVersion())
if handle_return(driver.cuDriverGetVersion()) > version[0] * 1000 + version[1] * 10:
raise RuntimeError(
"The CUDA driver version is newer than the NVRTC version. "
"Please update your NVRTC library to match the CUDA driver version."
)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Unfortunately this makes our run-time requirements unnecessarily strict. Driver is always backward compatible so new driver + old CTK (on a device that the old CTK supports) should work just fine.

In the test file there's a removed comment. If this part is meant to follow it, it should check whether target_type is PTX and then compare the driver and NVRTC versions. PTX generated by new NVRTC/NVVM/nvJitLink cannot be loaded (JIT-compiled) by old driver.

I think we should just move around that can_load_generated_ptx check to the codebase, and reuse it here. Then, in the test suite if the local/CI environment does not meet the requirement, we handle the raised exception by the check, instead of skipping the test. Does it make sense?

@leofang leofang added the enhancement Any code-related improvements label Feb 5, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
cuda.core Everything related to the cuda.core module enhancement Any code-related improvements P1 Medium priority - Should do
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Investigate improving the error message for PTX_COMPILE error
2 participants