[SYCL][libdevice] Add IMF support in devicelib #5999

jinge90 · 2022-04-11T11:50:13Z

Signed-off-by: jinge90 [email protected]

To provide support for math APIs and functions, not included in standard or <CL/sycl.hpp>.
These math APIs are implemented in "imf" (Intel math functions) SYCL device library. All function names are decorated with __imf_ prefix. These functions may be further categorized into:

Different accuracy implementation for some math functions such as log, exp, cos, sin…
Type cast util functions for float, double, half, bfloat16
Integer type util functions
SIMD util functions
Half type util functions
Bfloat16 type util functions

This PR is 1st of a series of patches to add “imf” SYCL device library, it includes type cast util functions for float, double and integer type util functions.

Signed-off-by: jinge90 <[email protected]>

akolesov-nv

Thank you, JIn Ge, I will look at the IMF part. Try to add me to reviewer using this username.

akolesov-nv · 2022-05-05T22:58:43Z

libdevice/device_imf.hpp

+#endif
+}
+
+class _iml_half {


Consider to split this source by several ones because

There are logically different parts: one for half emulation class and routines, another ones for device independent basic plugins (__fma, ...)

Basic plugins list are going to be expanded in future (__fma, __floor, ....)

Hi, @akolesov-intel
I split the original device_imf.hpp into 2 separate files. All _iml_half related parts are in imf_half.hpp, everything else is still in device_imf.hpp. And device_imf.hpp will include imf_half.hpp, so we only need to include device_imf.hpp in source code.

Thank you, @jinge90 . Now it is more clear and easy to find - having all fp16 emulation stuff in the separate file.

libdevice/device_imf.hpp

akolesov-nv · 2022-05-05T23:08:59Z

libdevice/imf_utils/saturatef.cpp

+#ifdef __LIBDEVICE_IMF_ENABLED__
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_imf_saturatef(float x) { return __fclamp(x, .0f, 1.f); }


Numerics team owns a number of trivial math functions like __devicelib_imf_copysign, __devicelib_imf_fabs, __devicelib_imf_rsqrt which can be either mapped to just one device independent builtin or implemented by a couple of simple operations - these functions are better to be inlined from header file than called from binary. We would like to gather them all under the ./imf folder which our team owns, promotes and supports. All trivial functions may combined within just one file. Could you please add such header inside ./imf, say, ./imf/imf_inline.hpp, with one function example (floor?)? Thanks!

Hi, @akolesov-intel
If we want to inline those trivial functions in header, then the header should be exposed to users, users will include it in their source code. And in the header file, we need to include:

declarations for all imf functions provided by us.

Some trivial functions inline definition.

I am fine if these functions are going to be inlined as IR, not as sources in header. I am just suggesting to gather them all into one file, not to spread by many separate tiny sources. I did that in my package through imf_inline.cpp source (can be named differently).

Hi, @akolesov-intel
I updated the patch to integrate all "trivial inline" functions, those functions are inlined as IR.
If we inline in LLVM IR level, we need also split imf_inline.cpp into imf_inline_fp32.cpp and imf_inline_fp64.cpp, I did this in the latest patch and basic tests for fp32, fp64, fp16 passed. Could you help review?
Thanks very much.

Thank you, @jinge90 , the split looks good to me.

Would be better to move saturatef implementation to imf_inline_fp32.cpp to gather all trivial functions mapped to math builtins in one place. And this file may be deleted.

libdevice/imf_utils/saturatef.cpp

libdevice/imf_wrapper.cpp

libdevice/cmake/modules/SYCLLibdevice.cmake

akolesov-nv · 2022-05-05T23:33:17Z

libdevice/cmake/modules/SYCLLibdevice.cmake

+set(fallback-imf-src imf_utils/saturatef.cpp
+                     imf_utils/float_convert.cpp
+                     imf_utils/half_convert.cpp
+                     imf/exp_s_la.cpp


Just curious is it possible to take the section with all Numerics function out of this file in a separate small cname include file and move it under the ../imf folder? This case Numerics team is going to be responsible to expand this list after all new function additions and no needs for Compiler team assistance required this case.

Hi, @akolesov-intel
Moving the section with all Numerics function out of this file in a separate small cmake is OK. However, moving it to ../imf folder doesn't align with libdevice's building system.

I am Ok not to move the section under ./imf, but we would need to agree on promotion process in details and document it somewhere.

xtian-github

LGTM

libdevice/device_imf.hpp

libdevice/imf/_imf_include.hpp

jinge90 · 2022-05-10T09:18:36Z

Hi, all.
I am focusing on implement type cast functions for float, double in this PR. When those functions are done, I will focus on addressing your comments. Type cast for fp16 and bfloat16 will be in future PR.
Thanks very much for the review.

libdevice/device_imf.hpp

libdevice/imf/_imf_include.hpp

xtian-github · 2022-05-12T16:37:04Z

With all the latest changes, the patch LGTM. But, we need to define the promotion/merge process between CE and MKL teams. Thanks.

xtian-github · 2022-05-13T04:44:50Z

@jinge90 *fp32.hpp and *fp64.hpp splitting is good. as fp64 does not work for ATS-M.

zettai-reido

Thanks for addressing the comments.
LGTM as it of now.

clang/lib/Driver/Driver.cpp

Signed-off-by: jinge90 <[email protected]>

libdevice/device_imf.hpp

Signed-off-by: jinge90 <[email protected]>

libdevice/device_imf.hpp

…ice_imf.hpp Signed-off-by: jinge90 <[email protected]>

steffenlarsen

LGTM!

jinge90 · 2022-06-09T13:22:46Z

Hi, all.
We made a little change in comments and added one compile-time macro check, could you help review again? The patch is ready to merge now.

Thanks very much.

jinge90 · 2022-06-10T01:31:14Z

Hi, @pvchupin
We have finished review for this patch, could you help merge it?
Thanks very much.

pvchupin · 2022-06-10T18:43:22Z

@jinge90, can you please add PR description
@asudarsa, feel free to reapprove.

asudarsa

LGTM. Thanks

xtian-github · 2022-06-10T22:26:47Z

@jinge90 is @zettai-reido going to do merge? or you?

pvchupin · 2022-06-10T22:39:44Z

@xtian-github, I'm going to merge as soon as somebody update PR description above (to be used as commit message).

xtian-github · 2022-06-10T22:46:32Z

@pvchupin Thanks Pavel!

Signed-off-by: jinge90 <[email protected]>

jinge90 · 2022-06-12T01:05:09Z

Hi, @pvchupin
I just updated the PR description and also made a little change to the patch, since the latest change only includes some little bugfix, it is OK to merge.

Thanks very much.

zettai-reido

LGTM

We added some imf(intel math) APIs in libdevice: #5999, #6327, those APIs are C interface which means we don't expect users to call them directly in SYCL kernel since SYCL programming uses cpp. This PR will provide cpp wrapper for those imf APIs which users can call directly in source, the approach is similar to STL , std math functions will call corresponding C math functions or compiler math builtins. The user scenario should be following: cpp #include <sycl/ext/intel/math.hpp> ... float x, y, z; z = sycl::ext::intel::math::copysign(x, y); ... ``

jinge90 requested review from a team as code owners April 11, 2022 11:50

jinge90 requested a review from v-klochkov April 11, 2022 11:50

jinge90 marked this pull request as draft April 11, 2022 11:50

jinge90 added 4 commits April 11, 2022 19:59

[SYCL][libdevice] Add sycl imf devicelib

9a69b78

Signed-off-by: jinge90 <[email protected]>

fix clang format

1065621

Signed-off-by: jinge90 <[email protected]>

Merge remote-tracking branch 'origin/sycl' into libdevice_imf

ca5f4ae

Enable imf sycl devicelib on Windows.

63a6aac

Signed-off-by: jinge90 <[email protected]>

v-klochkov removed their request for review April 19, 2022 18:28

jinge90 requested review from xtian-github and zettai-reido May 5, 2022 12:40

akolesov-nv reviewed May 5, 2022

View reviewed changes

xtian-github previously approved these changes May 6, 2022

View reviewed changes

jinge90 dismissed xtian-github’s stale review via 03d39bf May 7, 2022 09:11

akolesov-nv suggested changes May 9, 2022

View reviewed changes

libdevice/device_imf.hpp Outdated Show resolved Hide resolved

libdevice/imf/_imf_include.hpp Outdated Show resolved Hide resolved

zettai-reido reviewed May 10, 2022

View reviewed changes

jinge90 marked this pull request as ready for review May 12, 2022 10:39

zettai-reido previously approved these changes May 13, 2022

View reviewed changes

jinge90 dismissed zettai-reido’s stale review via 9fae9d6 May 13, 2022 12:59

mdtoguchi reviewed May 16, 2022

View reviewed changes

clang/lib/Driver/Driver.cpp Show resolved Hide resolved

jinge90 closed this May 17, 2022

jinge90 reopened this May 17, 2022

jinge90 force-pushed the libdevice_imf branch from 9fae9d6 to 63a6aac Compare May 17, 2022 01:40

restore imf utils

e483162

Signed-off-by: jinge90 <[email protected]>

akolesov-nv previously approved these changes Jun 8, 2022

View reviewed changes

steffenlarsen reviewed Jun 8, 2022

View reviewed changes

libdevice/device_imf.hpp Show resolved Hide resolved

libdevice/device_imf.hpp Show resolved Hide resolved

jinge90 dismissed stale reviews from akolesov-nv and asudarsa via 61b50ed June 9, 2022 04:03

Add TODO for bitcast

61b50ed

Signed-off-by: jinge90 <[email protected]>

jinge90 commented Jun 9, 2022

View reviewed changes

libdevice/device_imf.hpp Show resolved Hide resolved

Double confirm __SPIR__ or __LIBDEVICE_HOST_IMPL__ is defined for dev…

c47c071

…ice_imf.hpp Signed-off-by: jinge90 <[email protected]>

steffenlarsen previously approved these changes Jun 9, 2022

View reviewed changes

mdtoguchi previously approved these changes Jun 9, 2022

View reviewed changes

akolesov-nv previously approved these changes Jun 9, 2022

View reviewed changes

asudarsa previously approved these changes Jun 10, 2022

View reviewed changes

jinge90 dismissed stale reviews from asudarsa, akolesov-nv, mdtoguchi, and steffenlarsen via 6dfa250 June 11, 2022 13:45

jinge90 added 2 commits June 11, 2022 21:58

Fix bugs in brev, byte_perm

6dfa250

Signed-off-by: jinge90 <[email protected]>

Fix brevll prototype

d85f86f

Signed-off-by: jinge90 <[email protected]>

zettai-reido approved these changes Jun 12, 2022

View reviewed changes

akolesov-nv approved these changes Jun 13, 2022

View reviewed changes

pvchupin merged commit a310952 into intel:sycl Jun 13, 2022

jinge90 mentioned this pull request Aug 20, 2022

[SYCL] Add sycl cpp header for imf libdevice API #6487

Merged

jinge90 deleted the libdevice_imf branch August 26, 2022 03:09

[SYCL][libdevice] Add IMF support in devicelib #5999

[SYCL][libdevice] Add IMF support in devicelib #5999

Uh oh!

Conversation

jinge90 commented Apr 11, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

akolesov-nv left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

xtian-github left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

jinge90 commented May 10, 2022

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

xtian-github commented May 12, 2022

Uh oh!

xtian-github commented May 13, 2022

Uh oh!

zettai-reido left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

jinge90 commented Jun 9, 2022

Uh oh!

jinge90 commented Jun 10, 2022

Uh oh!

pvchupin commented Jun 10, 2022

Uh oh!

asudarsa left a comment

Choose a reason for hiding this comment

Uh oh!

xtian-github commented Jun 10, 2022

jinge90 commented Apr 11, 2022 •

edited

Loading