Skip to content

Update _launcher.py to fix vector_add.py #592

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

Open
wants to merge 1 commit into
base: main
Choose a base branch
from

Conversation

markwhen
Copy link

Description

closes

Checklist

  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

Copy link
Contributor

copy-pr-bot bot commented Apr 30, 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.

Copy link
Member

@leofang leofang left a comment

Choose a reason for hiding this comment

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

Thanks for catching this bug @markwhen. The fix is not quite right, though, see below. Also, I am curious if you are using driver <11.8 by any chance? This can be confirmed via nvidia-smi.

@rwgk this seems like an overlook from our earlier refactoring. I suspect our CI did not catch it because we don't have runners with old 11.x driver that could trigger this path.

@@ -146,6 +146,6 @@ def launch(stream, config, kernel, *kernel_args):
# TODO: check if config has any unsupported attrs
handle_return(
driver.cuLaunchKernel(
int(kernel._handle), *config.grid, *config.block, config.shmem_size, stream._handle, args_ptr, 0
int(kernel._handle), *config.grid, *config.block, config.shmem_size, stream._ctx_handle, args_ptr, 0
Copy link
Member

Choose a reason for hiding this comment

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

We need a stream here, not context:

Suggested change
int(kernel._handle), *config.grid, *config.block, config.shmem_size, stream._ctx_handle, args_ptr, 0
int(kernel._handle), *config.grid, *config.block, config.shmem_size, stream.handle, args_ptr, 0

@leofang leofang added bug Something isn't working P0 High priority - Must do! cuda.core Everything related to the cuda.core module labels Apr 30, 2025
@leofang leofang added this to the cuda.core beta 4 milestone Apr 30, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working cuda.core Everything related to the cuda.core module P0 High priority - Must do!
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants