Disable cl_khr_fp16 for the Environments unsupporting cl_khr_fp16 #266

vorj · 2020-02-05T16:44:12Z

Tasks of #264 , and #265 .

ClPy in this PR checks that the device supports cl_khr_fp16 or not, then enable or disable cl_khr_fp16 .
When cl_khr_fp16 isn't supported on the device, the ClPy uses a fallback implementation of half , which doesn't require cl_khr_fp16 .

This PR also contains below changes:

Ultima supports member functions.
- "Member functions" includes user-defined conversion operator.
Giving up to support type casting to half
- Currently we doesn't support half type, so it makes little impact.

jenkins-maekawa

Test (commit 56efdba) failed in vega.

============================= test session starts ==============================
platform linux -- Python 3.6.5, pytest-4.3.0, py-1.8.1, pluggy-0.13.1
rootdir: /home/clpy-jenkins-slave/workspace/clpy_testing_vega, inifile: setup.cfg
collected 148 items

test_api.py .                                                            [  0%]
test_atomicAdd.py ..                                                     [  2%]
test_carray.py .                                                         [  2%]
test_clblast.py ......................                                   [ 17%]
test_concatenate.py ........                                             [ 22%]
test_elementwise.py .                                                    [ 23%]
test_exception.py ..                                                     [ 25%]
test_linalg.py ...                                                       [ 27%]
test_memory.py .............                                             [ 35%]
test_ndarray.py ....................                                     [ 49%]
test_reduction.py ..............                                         [ 58%]
test_rollaxis.py ...                                                     [ 60%]
test_sample_rand.py ...                                                  [ 62%]
headercvt_tests/test_headercvt_funcdecl.py ..                            [ 64%]
headercvt_tests/test_headercvt_preproc_defines.py ..                     [ 65%]
headercvt_tests/test_headercvt_types.py ..............                   [ 75%]
ultima_tests/test_cast.py ....                                           [ 77%]
ultima_tests/test_cindexer.py ..                                         [ 79%]
ultima_tests/test_constructor.py ...........                             [ 86%]
ultima_tests/test_half.py .FF.                                           [ 89%]
ultima_tests/test_overload.py ....                                       [ 91%]
ultima_tests/test_template.py ........                                   [ 97%]
ultima_tests/test_ultima_carray.py ....                                  [100%]

=================================== FAILURES ===================================
______________________ TestUltimaHalfTrick.test_clpy_half ______________________

self = <test_half.TestUltimaHalfTrick testMethod=test_clpy_half>

        def test_clpy_half(self):
            x = '''
    void f()
    {
        __clpy__half half_ = 42.F;
        int __clpy__half = half_;
    }
    '''
            y = clpy.backend.ultima.exec_ultima(
                '''
                void f(){
                  __clpy__half half_ = 42.f;
                  int __clpy__half = half_;
                }
>               ''')

... and more 72 lines

jenkins-maekawa

Test (commit 56efdba) failed in titanv.

============================= test session starts ==============================
platform linux -- Python 3.6.5, pytest-4.3.0, py-1.8.1, pluggy-0.13.1
rootdir: /home/clpy-jenkins-slave/workspace/clpy_testing_titan, inifile: setup.cfg
collected 148 items

test_api.py .                                                            [  0%]
test_atomicAdd.py ..                                                     [  2%]
test_carray.py .                                                         [  2%]
test_clblast.py ......................                                   [ 17%]
test_concatenate.py ........                                             [ 22%]
test_elementwise.py .                                                    [ 23%]
test_exception.py ..                                                     [ 25%]
test_linalg.py ...                                                       [ 27%]
test_memory.py .............                                             [ 35%]
test_ndarray.py ....................                                     [ 49%]
test_reduction.py ..............                                         [ 58%]
test_rollaxis.py ...                                                     [ 60%]
test_sample_rand.py ...                                                  [ 62%]
headercvt_tests/test_headercvt_funcdecl.py ..                            [ 64%]
headercvt_tests/test_headercvt_preproc_defines.py ..                     [ 65%]
headercvt_tests/test_headercvt_types.py ..............                   [ 75%]
ultima_tests/test_cast.py ....                                           [ 77%]
ultima_tests/test_cindexer.py ..                                         [ 79%]
ultima_tests/test_constructor.py ...........                             [ 86%]
ultima_tests/test_half.py .FF.                                           [ 89%]
ultima_tests/test_overload.py ....                                       [ 91%]
ultima_tests/test_template.py ........                                   [ 97%]
ultima_tests/test_ultima_carray.py ....                                  [100%]

=================================== FAILURES ===================================
______________________ TestUltimaHalfTrick.test_clpy_half ______________________

self = <test_half.TestUltimaHalfTrick testMethod=test_clpy_half>

        def test_clpy_half(self):
            x = '''
    void f()
    {
        __clpy__half half_ = 42.F;
        int __clpy__half = half_;
    }
    '''
            y = clpy.backend.ultima.exec_ultima(
                '''
                void f(){
                  __clpy__half half_ = 42.f;
                  int __clpy__half = half_;
                }
>               ''')

... and more 72 lines

This reverts commit a1d49aa.

jenkins-maekawa

Test (commit b3482fe) failed in vega.

============================= test session starts ==============================
platform linux -- Python 3.6.5, pytest-4.3.0, py-1.8.1, pluggy-0.13.1
rootdir: /home/clpy-jenkins-slave/workspace/clpy_testing_vega, inifile: setup.cfg
collected 148 items

test_api.py .                                                            [  0%]
test_atomicAdd.py ..                                                     [  2%]
test_carray.py .                                                         [  2%]
test_clblast.py ......................                                   [ 17%]
test_concatenate.py ........                                             [ 22%]
test_elementwise.py .                                                    [ 23%]
test_exception.py ..                                                     [ 25%]
test_linalg.py ...                                                       [ 27%]
test_memory.py .............                                             [ 35%]
test_ndarray.py ....................                                     [ 49%]
test_reduction.py ..............                                         [ 58%]
test_rollaxis.py ...                                                     [ 60%]
test_sample_rand.py ...                                                  [ 62%]
headercvt_tests/test_headercvt_funcdecl.py ..                            [ 64%]
headercvt_tests/test_headercvt_preproc_defines.py ..                     [ 65%]
headercvt_tests/test_headercvt_types.py ..............                   [ 75%]
ultima_tests/test_cast.py ....                                           [ 77%]
ultima_tests/test_cindexer.py ..                                         [ 79%]
ultima_tests/test_constructor.py ...........                             [ 86%]
ultima_tests/test_half.py .FF.                                           [ 89%]
ultima_tests/test_overload.py ....                                       [ 91%]
ultima_tests/test_template.py ........                                   [ 97%]
ultima_tests/test_ultima_carray.py ....                                  [100%]

=================================== FAILURES ===================================
______________________ TestUltimaHalfTrick.test_clpy_half ______________________

self = <test_half.TestUltimaHalfTrick testMethod=test_clpy_half>

        def test_clpy_half(self):
            supports_cl_khr_fp16 = clpy.backend.opencl.env.supports_cl_khr_fp16()
            options = ('-D__CLPY_ENABLE_CL_KHR_FP16'
                       if supports_cl_khr_fp16
                       else '', )
            x = clpy.backend.ultima.exec_ultima('', '#include <cupy/carray.hpp>') + ('''
    void f()
    {
        __clpy__half half_ = 42.F;
        int __clpy__half = half_;
    }
    ''' if supports_cl_khr_fp16 else '''
    void f()
    {
        __clpy__half half_;constructor___clpy__half___left_paren____clpy__half_float__right_paren__(&half_, 42.F);
        int __clpy__half = operatorfloat___clpy__half_(&half_);

... and more 53 lines

jenkins-maekawa

Test (commit b3482fe) passed in titanv.

jenkins-maekawa

Test (commit 33338e8) passed in titanv.

jenkins-maekawa

Test (commit 33338e8) passed in vega.

vorj · 2020-02-05T18:53:51Z

@LWisteria

The CI has passed.
This PR contains some changes to test cases, but affected test cases are only in ultima_tests .
I didn't touch the tests to originate from CuPy.

I hear that @ybsh will return to our development today (welcome back!).
He works on our secondary machine, so we need to fix current situation before he will restart to work.
This PR fix the issue, so I want you to review this as soon as possible.
However I know that you are extremely busy, so if you have no time, you can take below steps instead of reviewing this PR:

Give the Mainter role for ClPy to @nsakabe-fixstars
Change reviewer of this PR to @nsakabe-fixstars

LWisteria

@vorj I will make reviews for PRs to solve #264 because I opened it.

I also welcome other's (including @nsakabe-fixstars 's) review and/or comments.

LWisteria · 2020-02-05T22:27:57Z

clpy/core/include/clpy/carray.clh

@@ -1,5 +1,7 @@
 #pragma once
+#ifdef __CLPY_ENABLE_CL_KHR_FP16


Why didn't you just remove cl_khr_fp16 for all environments but disable it only if it doesn't support it?

Do you have any definite disadvantages or problems when you treat half as int16 on a cl_khr_fp16 machine? If yes, please leave comments on the code, otherwise (yet) please make the code simple.

I think that it's better in performance to use half implemented as hardware. If we will support all half operations, the performance given from native half in supported environments is more important.

If we will support all half operations

We have not proven the ability neither decided to do that, must continue to discuss on #265 .

Then, until to decide it, the half performance is not important even if it has native half support.
Adding cl_khr_fp16 is now too superfluous to just disable it #264 .

must continue to discuss on #265 .

Sure.

Adding cl_khr_fp16 is now too superfluous to just disable it #264 .

So, you want to just remove cl_khr_fp16 whether the extension is available or not, right?

clpy/backend/opencl/env.pyx

LWisteria · 2020-02-05T22:47:04Z

clpy/core/include/clpy/carray.clh

+  ushort v;
+  __clpy__half_() = default;
+  __clpy__half_(float f):v{convert_float_to_half_ushort(f)}{}
+  operator float()const{return convert_half_ushort_to_float(v);}


I understand Ultima can solve the problem.

But I don't understand why you need Ultima. As I told you in the some old issue, you tend to solve everything by using Ultima but we must avoid it unless there's no solution without Ultima. Minimize Ultima to minimize maintenance difficulties.

Actually, I need the changes in this PR to ultima only for passing ultima_tests .
If we decide to change the behavior of half identifier, we doesn't need it.
However, we will need it to support all half operations, so I proposed this changes.

If we decide to change the behavior of half

As of #266 (comment), no half variable will be constructed by users. So you don't need Ultima.

As of #266 (comment), no half variable will be constructed by users.

Until to decide to support all half operations . OK.

Then, when you

decide to change the behavior of half identifier

? Again, we need to decide to change the behavior of half identifier from current ClPy to disable cl_khr_fp16 without Ultima. It means that we need to fix or disable some tests. If you have already decided it in your own, you should indicate your opinion here.

jenkins-maekawa

Test (commit 7ee88ab) passed in titanv.

jenkins-maekawa

Test (commit 7ee88ab) passed in vega.

vorj · 2020-02-06T06:06:25Z

@LWisteria I commented your questions and fixed the code you reviewed.

LWisteria · 2020-02-18T08:48:20Z

@vorj After our offline discussion, we concluded that we must focus on just removing cl_khr_fp16 until we finish to investigate and decide to support all half operations.

However, your commits on this PR are valuable to store our commit history. So I want you not to make new branch/PR but to continue on this branch/PR and remove inessentials.

vorj · 2020-02-21T13:59:13Z

@LWisteria

However, your commits on this PR are valuable to store our commit history. So I want you not to make new branch/PR but to continue on this branch/PR and remove inessentials.

It is quite different to just remove cl_khr_fp16 than this PR, so if I will continue on this PR, almost all commits will be reverted.
I propose that we will just close this PR and preserve the branch, then we will fix #264 by new PR.

LWisteria · 2020-02-22T00:33:19Z

After #269, this PR is obsolete. I'm closing but not deleting this branch for the further development.

jenkins-maekawa suggested changes Feb 5, 2020

View reviewed changes

vorj added 6 commits February 6, 2020 03:08

Revert "disable member function"

396cc35

This reverts commit a1d49aa.

support user-defined conversion

d024318

implement clpy.backend.opencl.env.supports_cl_khr_fp16

324126b

ultima supports options

93b0e71

define __CLPY_ENABLE_CL_KHR_FP16 when the device supports cl_khr_fp16

aa89256

disable cl_khr_fp16 when __CLPY_ENABLE_CL_KHR_FP16 is not defined

909ec19

vorj force-pushed the 264-remove-cl_khr_fp16 branch from 9cb6676 to ae09728 Compare February 5, 2020 18:17

jenkins-maekawa suggested changes Feb 5, 2020

View reviewed changes

jenkins-maekawa approved these changes Feb 5, 2020

View reviewed changes

fix test cases

d5e4865

vorj force-pushed the 264-remove-cl_khr_fp16 branch from ae09728 to d5e4865 Compare February 5, 2020 18:39

jenkins-maekawa approved these changes Feb 5, 2020

View reviewed changes

vorj marked this pull request as ready for review February 5, 2020 18:51

vorj assigned LWisteria Feb 5, 2020

vorj added the bug Something isn't working label Feb 5, 2020

LWisteria requested changes Feb 5, 2020

View reviewed changes

LWisteria assigned vorj and unassigned LWisteria Feb 5, 2020

fix partial matching on check_device_extension

aaf590c

jenkins-maekawa approved these changes Feb 6, 2020

View reviewed changes

vorj assigned LWisteria and unassigned vorj Feb 6, 2020

LWisteria assigned vorj and unassigned LWisteria Feb 9, 2020

vorj assigned LWisteria and unassigned vorj Feb 12, 2020

LWisteria closed this Feb 22, 2020

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Disable cl_khr_fp16 for the Environments unsupporting cl_khr_fp16 #266

Disable cl_khr_fp16 for the Environments unsupporting cl_khr_fp16 #266

vorj commented Feb 5, 2020 •

edited

Loading

jenkins-maekawa left a comment

jenkins-maekawa left a comment

jenkins-maekawa left a comment

jenkins-maekawa left a comment

jenkins-maekawa left a comment

jenkins-maekawa left a comment

vorj commented Feb 5, 2020

LWisteria left a comment

LWisteria Feb 5, 2020

vorj Feb 6, 2020

LWisteria Feb 9, 2020 •

edited

Loading

vorj Feb 12, 2020

LWisteria Feb 5, 2020

vorj Feb 6, 2020

LWisteria Feb 9, 2020 •

edited

Loading

vorj Feb 12, 2020

jenkins-maekawa left a comment

jenkins-maekawa left a comment

vorj commented Feb 6, 2020

LWisteria commented Feb 18, 2020

vorj commented Feb 21, 2020

LWisteria commented Feb 22, 2020

Disable cl_khr_fp16 for the Environments unsupporting cl_khr_fp16 #266

Disable cl_khr_fp16 for the Environments unsupporting cl_khr_fp16 #266

Conversation

vorj commented Feb 5, 2020 • edited Loading

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

vorj commented Feb 5, 2020

LWisteria left a comment

Choose a reason for hiding this comment

LWisteria Feb 5, 2020

Choose a reason for hiding this comment

vorj Feb 6, 2020

Choose a reason for hiding this comment

LWisteria Feb 9, 2020 • edited Loading

Choose a reason for hiding this comment

vorj Feb 12, 2020

Choose a reason for hiding this comment

LWisteria Feb 5, 2020

Choose a reason for hiding this comment

vorj Feb 6, 2020

Choose a reason for hiding this comment

LWisteria Feb 9, 2020 • edited Loading

Choose a reason for hiding this comment

vorj Feb 12, 2020

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

jenkins-maekawa left a comment

Choose a reason for hiding this comment

vorj commented Feb 6, 2020

LWisteria commented Feb 18, 2020

vorj commented Feb 21, 2020

LWisteria commented Feb 22, 2020

vorj commented Feb 5, 2020 •

edited

Loading

LWisteria Feb 9, 2020 •

edited

Loading

LWisteria Feb 9, 2020 •

edited

Loading