how to add self-math-function in sycl？

[wangzy0327]

How to extend a custom mathematical function in sycl, what code should be added to which file? Can you tell me how to extend a custom math function? For example, extend cuda's built-in functions into sycl

math-function

[wangzy0327]

for example, If I want to add sincospi() func in sycl on cuda backend, how could I do ? @sommerlukas

[JackAKirk]

A few things. Firstly you need to either append or add a new DPC++ extension. In this case I don't think such an addition maps to an existing extension, since I think this function would be placed in sycl math functions (rather than sycl native math functions). However a very similar extension that you could use as a template is:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_native_math.asciidoc

Be aware also of the extension template that has some advice and expectations:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/template.asciidoc

Math functions in the cuda backend are either called from clang builtins or asm ptx or libdevice (via libclc). In the case of sincospi it is defined in libdevice:

https://docs.nvidia.com/cuda/libdevice-users-guide/__nv_sincospi.html#__nv_sincospi

Therefore you should create a new file very similar to this one for sincos, but for sincospi

https://github.com/intel/llvm/blob/ec7fb7c4aebae3ea642a269e8cc4d4ab57f721ef/libclc/ptx-nvidiacl/libspirv/math/sincos.cl

The full implementation needs to wire this up to DPC++ runtime. An example PR doing this for a different new libclc function is:
250c498

Hope this is helpful. sincospi seems to be commonly used as an optimized high precision builtin for box-muller etc, so it seems suitable to be added to SYCL.

Note that the above describes only the cuda implementation. You would either have to state the extension is only currently supported on cuda, or also add a simple generic implementation (see other math functions for examples), via sin(cospi()) or otherwise.

[wangzy0327]

A few things. Firstly you need to either append or add a new DPC++ extension. In this case I don't think such an addition maps to an existing extension, since I think this function would be placed in sycl math functions (rather than sycl native math functions). However a very similar extension that you could use as a template is: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_native_math.asciidoc

Be aware also of the extension template that has some advice and expectations: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/template.asciidoc

Math functions in the cuda backend are either called from clang builtins or asm ptx or libdevice (via libclc). In the case of sincospi it is defined in libdevice:

https://docs.nvidia.com/cuda/libdevice-users-guide/__nv_sincospi.html#__nv_sincospi

Therefore you should create a new file very similar to this one for sincos, but for sincospi

https://github.com/intel/llvm/blob/ec7fb7c4aebae3ea642a269e8cc4d4ab57f721ef/libclc/ptx-nvidiacl/libspirv/math/sincos.cl

The full implementation needs to wire this up to DPC++ runtime. An example PR doing this for a different new libclc function is: 250c498

Hope this is helpful. sincospi seems to be commonly used as an optimized high precision builtin for box-muller etc, so it seems suitable to be added to SYCL.

Note that the above describes only the cuda implementation. You would either have to state the extension is only currently supported on cuda, or also add a simple generic implementation (see other math functions for examples), via sin(cospi()) or otherwise.

Thank you for reply. I would like to follow up and ask, currently libclc's libdevice related implementations are scalar math functions, such as float abs_f32(float x). How can we add vector math functions through libdevice, such as void vector_abs_f32(int n, float *y, const float *x)?

[jinz2014]

Thank you for the question and answer. Could Jack's answer turn into a blog or technical guide for users/developers to add more functions ? There are a lot of changes in [SYCL][libclc][CUDA] Add native math extension, I hope the PR developer could write a document explaining the extension.

[JackAKirk]

Thank you for the question and answer. Could Jack's answer turn into a blog or technical guide for users/developers to add more functions ? There are a lot of changes in [SYCL][libclc][CUDA] Add native math extension, I hope the PR developer could write a document explaining the extension.

This is an interesting idea. After some discussion we came to the conclusion that the best option is to check that the project docs are correctly signposting contributors to the appropriate documentation if they wish to contribute extensions. Generally I think existing docs are quite good and comprehensive, however I made a small PR to try to improve things slightly:
#12332

Application developers often have useful insight into missing features so I think it is a good idea to try to make any such contributors lives easier. The concern with writing specific technical blogs, or even blogs on processes, is that they could quickly become out of date. The established solution is to have good up to date project documentation.

[JackAKirk]

A few things. Firstly you need to either append or add a new DPC++ extension. In this case I don't think such an addition maps to an existing extension, since I think this function would be placed in sycl math functions (rather than sycl native math functions). However a very similar extension that you could use as a template is: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_native_math.asciidoc
Be aware also of the extension template that has some advice and expectations: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/template.asciidoc
Math functions in the cuda backend are either called from clang builtins or asm ptx or libdevice (via libclc). In the case of sincospi it is defined in libdevice:
https://docs.nvidia.com/cuda/libdevice-users-guide/__nv_sincospi.html#__nv_sincospi
Therefore you should create a new file very similar to this one for sincos, but for sincospi
https://github.com/intel/llvm/blob/ec7fb7c4aebae3ea642a269e8cc4d4ab57f721ef/libclc/ptx-nvidiacl/libspirv/math/sincos.cl
The full implementation needs to wire this up to DPC++ runtime. An example PR doing this for a different new libclc function is: 250c498
Hope this is helpful. sincospi seems to be commonly used as an optimized high precision builtin for box-muller etc, so it seems suitable to be added to SYCL.
Note that the above describes only the cuda implementation. You would either have to state the extension is only currently supported on cuda, or also add a simple generic implementation (see other math functions for examples), via sin(cospi()) or otherwise.

Thank you for reply. I would like to follow up and ask, currently libclc's libdevice related implementations are scalar math functions, such as float abs_f32(float x). How can we add vector math functions through libdevice, such as void vector_abs_f32(int n, float *y, const float *x)?

If you look at the various math functions in libclc, you will see that there are many examples of vector math functions that call libdevice. There are also (a small number of) cuda instructions that are vector instructions, that are correctly called through libclc (see for example the bfloat16 function files in nvptx libclc).

[keryell]

This kind of API has already been discussed in the SYCL committee meeting.
The presented sincos has a C-like API with pointer-through side effects, which makes unfortunately sense for CUDA which is almost 20 year old.
But SYCL is based on modern C++, so we could expect a more mathematical streamlined API like

auto [sin, cos] = sycl::sincos(angle);

[wangzy0327]

Could you be more specific and give an example of a vector math func which have already accomplished by pvptx libdevice based on which sycl version?

…

------------------ Original ------------------ From: JackAKirk ***@***.***> Date: Tue,Jan 9,2024 7:44 PM To: intel/llvm ***@***.***> Cc: wangzy0327 ***@***.***>, Author ***@***.***> Subject: Re: [intel/llvm] how to add self-math-function in sycl？ (Issue #12251) A few things. Firstly you need to either append or add a new DPC++ extension. In this case I don't think such an addition maps to an existing extension, since I think this function would be placed in sycl math functions (rather than sycl native math functions). However a very similar extension that you could use as a template is: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_native_math.asciidoc Be aware also of the extension template that has some advice and expectations: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/template.asciidoc Math functions in the cuda backend are either called from clang builtins or asm ptx or libdevice (via libclc). In the case of sincospi it is defined in libdevice: https://docs.nvidia.com/cuda/libdevice-users-guide/__nv_sincospi.html#__nv_sincospi Therefore you should create a new file very similar to this one for sincos, but for sincospi https://github.com/intel/llvm/blob/ec7fb7c4aebae3ea642a269e8cc4d4ab57f721ef/libclc/ptx-nvidiacl/libspirv/math/sincos.cl The full implementation needs to wire this up to DPC++ runtime. An example PR doing this for a different new libclc function is: 250c498 Hope this is helpful. sincospi seems to be commonly used as an optimized high precision builtin for box-muller etc, so it seems suitable to be added to SYCL. Note that the above describes only the cuda implementation. You would either have to state the extension is only currently supported on cuda, or also add a simple generic implementation (see other math functions for examples), via sin(cospi()) or otherwise. Thank you for reply. I would like to follow up and ask, currently libclc's libdevice related implementations are scalar math functions, such as float abs_f32(float x). How can we add vector math functions through libdevice, such as void vector_abs_f32(int n, float *y, const float *x)? If you look at the various math functions in libclc, you will see that there are many examples of vector math functions that call libdevice. There are also (a small number of) cuda instructions that are vector instructions, that are correctly called through libclc (see for example the bfloat16 function files in nvptx libclc). — Reply to this email directly, view it on GitHub, or unsubscribe. You are receiving this because you authored the thread.Message ID: ***@***.***>

[wangzy0327]

This kind of API has already been discussed in the SYCL committee meeting. The presented sincos has a C-like API with pointer-through side effects, which makes unfortunately sense for CUDA which is almost 20 year old. But SYCL is based on modern C++, so we could expect a more mathematical streamlined API like

auto [sin, cos] = sycl::sincos(angle);

Ok Thanks. If I want to add a SPIR-V vector math function interface to the device vector math function mapping, what should I do? Do I need to customize interfaces like spirv_ocl_vec_xxx?