Skip to content

Regenerate float16 bindings #436

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.

Sign up for GitHub

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

Merged
isVoid merged 5 commits into from
Aug 29, 2025
Merged

Regenerate float16 bindings #436

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
9d97bd4
Generate fp16 bindings using numbast
lakshayg Aug 28, 2025
a8c5c6e
Alias __half to numba.core.types.float16
lakshayg Aug 28, 2025
ac499a2
Replace Numbast generated operators
lakshayg Aug 28, 2025
26ef143
Replace Numbast generated lower_casts
lakshayg Aug 28, 2025
ab7d9fb
Add performance warning
lakshayg Aug 28, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions configs/cuda_fp16.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,9 @@ Version: 0.0.1
Entry Point: ./numba_cuda/numba/cuda/include/12/cuda_fp16.h
File List:
- ./numba_cuda/numba/cuda/include/12/cuda_fp16.h
Exclude:
Struct:
- __half2
Types:
__half: Number
__half2: Type
Expand All @@ -13,3 +16,6 @@ Data Models:
__half2: StructModel
Shim Include Override: "\"cuda_fp16.h\""
Require Pynvjitlink: False
Use Separate Registry: True
GPU Arch:
- sm_86
Copy link
Contributor

@isVoid isVoid Aug 28, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

GPU Arch defines the macro CUDA_ARCH when clang is used to parse fp16 headers. We don't have a good guidance of this value yet. For float16, after sm70 it is feature complete. So any value after 70 should give us all of the API that's in the header.

Copy link
Contributor

@isVoid isVoid Aug 28, 2025

Choose a reason for hiding this comment

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

Ideally - Numbast generated bindings should generate different bindings for different arches based on supported arches the library is generating for. Then use difftools to conditionally enable/disable these features.

Loading