Skip to content

Conversation

@jiel-nv
Copy link
Contributor

@jiel-nv jiel-nv commented Feb 21, 2025

The customized CUDA front end type 'GridGroup' maps to the back end type 'ir.IntType(64)' which is defined in numba.cuda.models.

The numba DIBuilder emit variable types for simple numeric types by looking into back end type which falls through the 'DW_ATE_float' because the 'datamodel.fe_type' is not 'types.Integer'.

So emitting this type through the CUDADIBuilder is needed. Since the type represents the handle of a group object, giving it the encoding of 'DW_ATE_unsigned' should be reasonable.

A test is also added in this change.

This change addresses #107

The customized CUDA front end type 'GridGroup' maps to the back end
type 'ir.IntType(64)' which is defined in numba.cuda.models.

The numba DIBuilder emit variable types for simple numeric types by
looking into back end type which falls through the 'DW_ATE_float'
because the 'datamodel.fe_type' is not 'types.Integer'.

So emitting this type through the CUDADIBuilder is needed. Since the
type represents the handle of a group object, giving it the encoding
of 'DW_ATE_unsigned' should be reasonable.

A test is also added in this change.

This change addresses NVIDIA#107
@copy-pr-bot
Copy link

copy-pr-bot bot commented Feb 21, 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.

@gmarkall gmarkall added the 3 - Ready for Review Ready for review by team label Feb 21, 2025
@gmarkall
Copy link
Contributor

/ok to test

Copy link
Contributor

@gmarkall gmarkall left a comment

Choose a reason for hiding this comment

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

Many thanks! Just waiting for CI.

@gmarkall gmarkall linked an issue Feb 21, 2025 that may be closed by this pull request
@gmarkall gmarkall merged commit 3e9e705 into NVIDIA:main Feb 21, 2025
31 checks passed
gmarkall added a commit to gmarkall/numba-cuda that referenced this pull request Mar 6, 2025
- Fix linking of external code from callees (NVIDIA#137)
- Try using a newer branch workflow (NVIDIA#148)
- Move publish step out of `wheels-build.yaml` (NVIDIA#147)
- Upload wheels to PyPI from GitHub-hosted runner (NVIDIA#142)
- Add paddle to interoperability chapter (NVIDIA#144)
- Fix the debug info of GridGroup type (NVIDIA#131)
- Remove dead `prepare_cuda_kernel()` (NVIDIA#130)
- Add a CUDA DI Builder (NVIDIA#104)
- dont launch extra kernels when stats counting is disabled (NVIDIA#127)
- Fixup debug metadata in kernel fixup (NVIDIA#97)
- Implement debuginfo bool name fix (numba/numba#9888) in numba-cuda (NVIDIA#106)
@gmarkall gmarkall mentioned this pull request Mar 6, 2025
gmarkall added a commit that referenced this pull request Mar 6, 2025
- Fix linking of external code from callees (#137)
- Try using a newer branch workflow (#148)
- Move publish step out of `wheels-build.yaml` (#147)
- Upload wheels to PyPI from GitHub-hosted runner (#142)
- Add paddle to interoperability chapter (#144)
- Fix the debug info of GridGroup type (#131)
- Remove dead `prepare_cuda_kernel()` (#130)
- Add a CUDA DI Builder (#104)
- dont launch extra kernels when stats counting is disabled (#127)
- Fixup debug metadata in kernel fixup (#97)
- Implement debuginfo bool name fix (numba/numba#9888) in numba-cuda (#106)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

3 - Ready for Review Ready for review by team

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[BUG] Debug: grid groups appear to be floats, in cuda-gdb

2 participants