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

[GPU/OpenCL] Check fp16(half) support #2608

Merged
merged 1 commit into from
Jun 1, 2024

Conversation

s-debadri
Copy link
Contributor

Added check for OpenCL fp16 half support on device.
Will show proper message if not found.

To use fp16 enable cl_khr_fp16 as following with the kernels.

#pragma OPENCL EXTENSION cl_khr_fp16 : enable
    
 __kernel void kernel_name(...) { } 

Signed-off-by: Debadri Samaddar [email protected]

Added check for fp16 support on OpenCL.
Will show proper message if not found.

Signed-off-by: Debadri Samaddar <[email protected]>
@taos-ci
Copy link

taos-ci commented May 28, 2024

📝 TAOS-CI Version: 1.5.20200925. Thank you for submitting PR #2608. Please a submit 1commit/1PR (one commit per one PR) policy to get comments quickly from reviewers. Your PR must pass all verificiation processes of cibot before starting a review process from reviewers. If you are new member to join this project, please read manuals in documentation folder and wiki page. In order to monitor a progress status of your PR in more detail, visit http://ci.nnstreamer.ai/.

@taos-ci
Copy link

taos-ci commented May 28, 2024

:octocat: cibot: @s-debadri, nntrainer/opencl/opencl_context_manager.cpp includes bug(s). Please fix incorrect coding constructs in your commit before entering a review process.

Copy link

@taos-ci taos-ci left a comment

Choose a reason for hiding this comment

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

@s-debadri, 💯 All CI checkers are successfully verified. Thanks.

@@ -155,6 +155,31 @@ bool ContextManager::CreateDefaultGPUDevice() {
device_id_ = devices[0];
platform_id_ = platform_id_;

#ifdef ENABLE_FP16
// check for fp16 (half) support available on device
Copy link
Contributor

Choose a reason for hiding this comment

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

which half-precision does it support? _Float16, __fp16, or half

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Ideally all of _Float16, __fp16 and half can be used interchangeably inside kernels when cl_khr_fp16 extension is enabled. I have tested with __fp16 and half myself, they worked exactly the same. However I did not test with __Float16 but it is supported by OpenCL 2.1 or later.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

half is specifically designed for use in OpenCL kernels. Some are architecture specific (e.g: __fp16 for ARM). So half will be more portable and benefit from further optimizations..

Copy link
Contributor

Choose a reason for hiding this comment

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

thanks for the clarification! if using the half type has advantages, should it be added as a tensor data type?

#ifdef ENABLE_FP16
#ifdef USE__FP16
#define _FP16 __fp16
#else
#define _FP16 _Float16
#endif
#endif

Copy link
Contributor Author

Choose a reason for hiding this comment

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

half will only work inside OpenCL kernels. So I think adding it to tensor data type won't be useful in this scenario.

@myungjoo myungjoo merged commit 2a66923 into nnstreamer:main Jun 1, 2024
33 of 35 checks passed
@s-debadri s-debadri deleted the opencl_fp16_support branch June 5, 2024 05:15
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants