Skip to content

[libclc] Add new function implementations in math. #4849

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

Closed
wants to merge 1 commit into from
Closed

[libclc] Add new function implementations in math. #4849

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
44 changes: 42 additions & 2 deletions libclc/amdgcn-amdhsa/libspirv/SOURCES
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,17 +1,57 @@
atomic/loadstore_helpers.ll
cl_khr_int64_extended_atomics/minmax_helpers.ll
synchronization/barrier.cl

math/cos.cl
math/sin.cl
math/sqrt.cl
math/atan.cl
math/tan.cl
math/atanh.cl
math/cbrt.cl
math/acos.cl
math/asin.cl
math/asinh.cl
math/acosh.cl
math/cosh.cl
math/cospi.cl
math/sinpi.cl
math/atan2.cl
math/exp.cl
math/exp2.cl
math/erfc.cl
math/erf.cl
math/fdim.cl
math/fmin.cl
math/ceil.cl
math/pow.cl
math/fmod.cl
math/exp10.cl
math/expm1.cl
math/floor.cl
math/nextafter.cl
math/round.cl
math/sincos.cl
math/fabs.cl
math/fmax.cl
math/fmod.cl
math/frexp.cl
math/hypot.cl
math/lgamma.cl
math/ldexp.cl
math/log.cl
math/modf.cl
math/trunc.cl
math/copysign.cl
math/log10.cl
math/logb.cl
math/log2.cl
math/log1p.cl
math/rsqrt.cl
workitem/get_global_size.cl
workitem/get_local_size.cl
workitem/get_num_groups.cl
workitem/get_max_sub_group_size.cl
workitem/get_num_sub_groups.cl
workitem/get_sub_group_id.cl
Copy link
Contributor

Choose a reason for hiding this comment

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

Why workitem/get_sub_group_local_id.cl and misc/sub_group_shuffle.cl are removed?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Sorry, the wrong file was selected when submitting, new submission has modified this.

workitem/get_sub_group_local_id.cl
workitem/get_sub_group_size.cl
misc/sub_group_shuffle.cl
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/acos.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
Copy link
Contributor

Choose a reason for hiding this comment

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

Please, add this comment to all new files:

//===----------------------------------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Sorry，I delete my repository by mistake, and I forked again. But I don't know how to update this request now, should I close this one ,and create a new pull request?

#include <clcmacro.h>

double __ocml_acos_f64(double);
float __ocml_acos_f32(float);

#define __CLC_FUNCTION __spirv_ocl_acos
#define __CLC_BUILTIN __ocml_acos
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/acosh.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_acosh_f64(double);
float __ocml_acosh_f32(float);

#define __CLC_FUNCTION __spirv_ocl_acosh
#define __CLC_BUILTIN __ocml_acosh
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
13 changes: 13 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/asin.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_asin_f64(double);
float __ocml_asin_f32(float);

Copy link
Contributor

Choose a reason for hiding this comment

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

Could you please unify the white spaces so the files have the same layout?


#define __CLC_FUNCTION __spirv_ocl_asin
#define __CLC_BUILTIN __ocml_asin

#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/asinh.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_asinh_f64(double);
float __ocml_asinh_f32(float);

#define __CLC_FUNCTION __spirv_ocl_asinh
#define __CLC_BUILTIN __ocml_asinh
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/atan2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_atan2_f64(double,double);
float __ocml_atan2_f32(float,float);

#define __CLC_FUNCTION __spirv_ocl_atan2
#define __CLC_BUILTIN __ocml_atan2
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/atanh.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_atanh_f64(double);
float __ocml_atanh_f32(float);

#define __CLC_FUNCTION __spirv_ocl_atanh
#define __CLC_BUILTIN __ocml_atanh
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/ceil.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_ceil_f64(double);
float __ocml_ceil_f32(float);

#define __CLC_FUNCTION __spirv_ocl_ceil
#define __CLC_BUILTIN __ocml_ceil
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
13 changes: 13 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/clc_sw_binary.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
//===----------------------------------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include <utils.h>

_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __CLC_FUNC(__CLC_GENTYPE x,
Copy link
Contributor

Choose a reason for hiding this comment

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

Not sure if that's really needed, as only used once?

__CLC_GENTYPE y) {
return __CLC_SW_FUNC(x, y);
}
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/copysign.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_copysign_f64(double, double);
float __ocml_copysign_f32(float, float);

#define __CLC_FUNCTION __spirv_ocl_copysign
#define __CLC_BUILTIN __ocml_copysign
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
20 changes: 15 additions & 5 deletions libclc/amdgcn-amdhsa/libspirv/math/cos.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,9 +6,19 @@
//
//===----------------------------------------------------------------------===//

#include <spirv/spirv.h>
//#include <spirv/spirv.h>

_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
__spirv_ocl_cos(__clc_fp32_t In) {
return __builtin_amdgcn_cosf(In);
}
//_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
//__spirv_ocl_cos(__clc_fp32_t In) {
// return __builtin_amdgcn_cosf(In);
//}

#include <spirv/spirv.h>
#include <clcmacro.h>
double __ocml_cos_f64(double);
float __ocml_cos_f32(float);
#define __CLC_FUNCTION __spirv_ocl_cos
#define __CLC_BUILTIN __ocml_cos
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/cosh.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_cosh_f64(double);
float __ocml_cosh_f32(float);

#define __CLC_FUNCTION __spirv_ocl_cosh
#define __CLC_BUILTIN __ocml_cosh
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/cospi.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_cospi_f64(double);
float __ocml_cospi_f32(float);

#define __CLC_FUNCTION __spirv_ocl_cospi
#define __CLC_BUILTIN __ocml_cospi
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/erf.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_erf_f64(double);
float __ocml_erf_f32(float);

#define __CLC_FUNCTION __spirv_ocl_erf
#define __CLC_BUILTIN __ocml_erf
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/erfc.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_erfc_f64(double);
float __ocml_erfc_f32(float);

#define __CLC_FUNCTION __spirv_ocl_erfc
#define __CLC_BUILTIN __ocml_erfc
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/exp.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_exp_f64(double);
float __ocml_exp_f32(float);

#define __CLC_FUNCTION __spirv_ocl_exp
#define __CLC_BUILTIN __ocml_exp
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/exp10.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_exp10_f64(double);
float __ocml_exp10_f32(float);

#define __CLC_FUNCTION __spirv_ocl_exp10
#define __CLC_BUILTIN __ocml_exp10
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/exp2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_exp2_f64(double);
float __ocml_exp2_f32(float);

#define __CLC_FUNCTION __spirv_ocl_exp2
#define __CLC_BUILTIN __ocml_exp2
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/expm1.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_expm1_f64(double);
float __ocml_expm1_f32(float);

#define __CLC_FUNCTION __spirv_ocl_expm1
#define __CLC_BUILTIN __ocml_expm1
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/fabs.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_fabs_f64(double);
float __ocml_fabs_f32(float);

#define __CLC_FUNCTION __spirv_ocl_fabs
#define __CLC_BUILTIN __ocml_fabs
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/fdim.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_fdim_f64(double,double);
float __ocml_fdim_f32(float,float);

#define __CLC_FUNCTION __spirv_ocl_fdim
#define __CLC_BUILTIN __ocml_fdim
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/floor.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_floor_f64(double);
float __ocml_floor_f32(float);

#define __CLC_FUNCTION __spirv_ocl_floor
#define __CLC_BUILTIN __ocml_floor
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/unary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/fmax.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_fmax_f64(double,double);
float __ocml_fmax_f32(float,float);

#define __CLC_FUNCTION __spirv_ocl_fmax
#define __CLC_BUILTIN __ocml_fmax
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/fmin.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_fmin_f64(double,double);
float __ocml_fmin_f32(float,float);

#define __CLC_FUNCTION __spirv_ocl_fmin
#define __CLC_BUILTIN __ocml_fmin
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
11 changes: 11 additions & 0 deletions libclc/amdgcn-amdhsa/libspirv/math/fmod.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <spirv/spirv.h>
#include <clcmacro.h>

double __ocml_fmod_f64(double,double);
float __ocml_fmod_f32(float,float);

#define __CLC_FUNCTION __spirv_ocl_fmod
#define __CLC_BUILTIN __ocml_fmod
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
#include <math/binary_builtin.inc>
Loading