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.

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

[SYCL] [NATIVECPU] Implement missing math builtins for scalar data types #11321

Merged
merged 12 commits into from
Oct 24, 2023
Merged

[SYCL] [NATIVECPU] Implement missing math builtins for scalar data types #11321

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
806f49a
Implement missing math builtins for scalar data types
PietroGhg Sep 27, 2023
ec01ed2
formatting
PietroGhg Sep 27, 2023
908425c
use clang builtins
PietroGhg Oct 2, 2023
c9d1bda
formatting
PietroGhg Oct 2, 2023
126f265
Merge branch 'sycl' into pietro/math_builtins
PietroGhg Oct 9, 2023
a08c279
formatting
PietroGhg Oct 9, 2023
b8f667a
Remove commented code
PietroGhg Oct 9, 2023
a9d6d09
Remove math builtins test
PietroGhg Oct 9, 2023
7556444
Use NativeCPU AS map when building libclc
PietroGhg Oct 9, 2023
edfe058
Add FE test
PietroGhg Oct 11, 2023
5009f36
Skip test on windows
PietroGhg Oct 11, 2023
45b75c5
undef IS_NATIVE
PietroGhg Oct 24, 2023
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
1 change: 1 addition & 0 deletions libclc/x86_64-unknown-linux/libspirv/SOURCES
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,3 +16,4 @@ math/native_sqrt.cl
math/rint.cl
math/round.cl
math/trunc.cl
shared/helpers.ll
9 changes: 9 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/integer/helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,9 @@
#include "func.h"

#define GEN_UNARY_BUILTIN_T(NAME, TYPE) \
_CLC_OVERLOAD TYPE __##NAME##_helper(TYPE); \
_CLC_OVERLOAD TYPE __spirv_ocl_##NAME(TYPE n) { return __##NAME##_helper(n); }

#define GEN_UNARY_BUILTIN(NAME) \
GEN_UNARY_BUILTIN_T(NAME, int) \
GEN_UNARY_BUILTIN_T(NAME, signed char)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/integer/popcount.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(popcount)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/ceil.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(ceil)
4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/fabs.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_FABS
#include "helpers.h"

GEN_UNARY_BUILTIN(fabs)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/floor.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(floor)
4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/fma.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#include "helpers.h"

GEN_TERNARY_BUILTIN(fma);

82 changes: 82 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,82 @@
#include "func.h"
#include "types.h"

#ifdef NO_CLANG_BUILTINS

#define GEN_UNARY_BUILTIN_T(NAME, TYPE) \
_CLC_OVERLOAD TYPE __##NAME##_helper(TYPE); \
_CLC_OVERLOAD TYPE __spirv_ocl_##NAME(TYPE n) { return __##NAME##_helper(n); }

#define GEN_TERNARY_BUILTIN_T(NAME, TYPE) \
_CLC_OVERLOAD TYPE __##NAME##_helper(TYPE, TYPE, TYPE); \
_CLC_OVERLOAD TYPE __spirv_ocl_##NAME(TYPE a, TYPE b, TYPE c) { \
return __##NAME##_helper(a, b, c); \
}
#define GEN_UNARY_BUILTIN(NAME) \
GEN_UNARY_BUILTIN_T(NAME, float) \
GEN_UNARY_BUILTIN_T(NAME, double)

#define GEN_TERNARY_BUILTIN(NAME) \
GEN_TERNARY_BUILTIN_T(NAME, float) \
GEN_TERNARY_BUILTIN_T(NAME, double)

#else

#ifndef IS_NATIVE
#define GETNAME(ID) __spirv_ocl_##ID
#else
#define GETNAME(ID) __spirv_ocl_native_##ID
#endif

// Todo: fabs is the only builtin whose vector version is not named
// __builtin_elementwise_##NAME
#ifndef IS_FABS
#define GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, NUM) \
_CLC_OVERLOAD TYPE##NUM GETNAME(NAME)(TYPE##NUM n) { \
return __builtin_elementwise_##NAME(n); \
}
#else
#define GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, NUM) \
_CLC_OVERLOAD TYPE##NUM GETNAME(NAME)(TYPE##NUM n) { \
return __builtin_elementwise_abs(n); \
}
#endif

#define GEN_UNARY_VECTOR_BUILTIN_T(NAME, TYPE) \
GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, 2) \
GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, 3) \
GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, 4) \
GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, 8) \
GEN_UNARY_VECTOR_BUILTIN(NAME, TYPE, 16)

#define GEN_UNARY_BUILTIN(NAME) \
_CLC_OVERLOAD float GETNAME(NAME)(float n) { \
return __builtin_##NAME##f(n); \
} \
_CLC_OVERLOAD double GETNAME(NAME)(double n) { return __builtin_##NAME(n); } \
GEN_UNARY_VECTOR_BUILTIN_T(NAME, float) \
GEN_UNARY_VECTOR_BUILTIN_T(NAME, double)

#define GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, NUM) \
_CLC_OVERLOAD TYPE##NUM GETNAME(NAME)(TYPE##NUM n1, TYPE##NUM n2, \
TYPE##NUM n3) { \
return __builtin_elementwise_##NAME(n1, n2, n3); \
Copy link
Contributor

Choose a reason for hiding this comment

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

Are these __builtin* specific to nativecpu?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

They are part of clang see here

Copy link
Contributor

Choose a reason for hiding this comment

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

That's what I was thinking. It seems a bit strange to me to be calling generic clang builtins in a target specific libclc backend. Is there a good reason to do it like this?

Maybe these clang builtins better than the ones in the libclc/generic backend? Such generic impls that would be called if you don't add a target specific impl. Or could the clang builtins used here just be added to the libclc/generic impls?

Thanks

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The clang builtins should lead to the same LLVM intrinsics used in libclc/generic, but the way the SPIRV builtins are implemented in generic doesn't play well with the x86 ABI on Linux (see #10970). I think you are right and it's worth considering changing the implementation in generic, but I'm not 100% sure about the implications of doing so.

Copy link
Contributor

Choose a reason for hiding this comment

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

OK I see, makes sense, thanks for the clarification.

}

#define GEN_TERNARY_VECTOR_BUILTIN_T(NAME, TYPE) \
GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, 2) \
GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, 3) \
GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, 4) \
GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, 8) \
GEN_TERNARY_VECTOR_BUILTIN(NAME, TYPE, 16)

#define GEN_TERNARY_BUILTIN(NAME) \
_CLC_OVERLOAD float GETNAME(NAME)(float n1, float n2, float n3) { \
return __builtin_##NAME##f(n1, n2, n3); \
} \
_CLC_OVERLOAD double GETNAME(NAME)(double n1, double n2, double n3) { \
return __builtin_##NAME(n1, n2, n3); \
} \
GEN_TERNARY_VECTOR_BUILTIN_T(NAME, float) \
GEN_TERNARY_VECTOR_BUILTIN_T(NAME, double)
#endif
4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_cos.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(cos)
4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_exp.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(exp)
Copy link
Contributor

Choose a reason for hiding this comment

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

I think IS_NATIVE could be staying defined when you don't want it to?

Suggested change
GEN_UNARY_BUILTIN(exp)
GEN_UNARY_BUILTIN(exp)
#undef IS_NATIVE

etc

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Good point, that's done, thank you

4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_exp2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(exp2)
4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_log.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(log)
5 changes: 5 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_log10.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,5 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(log10)

5 changes: 5 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_log2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,5 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(log2)

5 changes: 5 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_sin.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,5 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(sin)

4 changes: 4 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/native_sqrt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
#define IS_NATIVE
#include "helpers.h"

GEN_UNARY_BUILTIN(sqrt)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/rint.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(rint)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/round.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(round)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/sqrt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(sqrt)
3 changes: 3 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/math/trunc.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
#include "helpers.h"

GEN_UNARY_BUILTIN(trunc)
17 changes: 17 additions & 0 deletions libclc/x86_64-unknown-linux/libspirv/shared/helpers.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
declare i32 @llvm.ctpop.i32(i32 %n)
declare i8 @llvm.ctpop.i8(i8 %n)
npmiller marked this conversation as resolved.
Show resolved Hide resolved


define dso_local i32 @_Z17__popcount_helperi(i32 %x) {
entry:
%call = call i32 @llvm.ctpop.i32(i32 %x)
ret i32 %call
}


define dso_local i8 @_Z17__popcount_helpera(i8 %x) {
entry:
%call = call i8 @llvm.ctpop.i8(i8 %x)
ret i8 %call
}

13 changes: 13 additions & 0 deletions llvm/lib/SYCLLowerIR/RenameKernelSYCLNativeCPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,5 +45,18 @@ RenameKernelSYCLNativeCPUPass::run(Module &M, ModuleAnalysisManager &MAM) {
}
}

//<<<<<<< HEAD
// for (auto &F : CalledSet) {
// auto NewName = sycl::utils::addSYCLNativeCPUSuffix(F->getName());
// F->setName(NewName);
// auto Comdat = F->getComdat();
// if (Comdat) {
// auto NewComdat = M.getOrInsertComdat(NewName.str());
// F->setComdat(NewComdat);
// }
// ModuleChanged |= true;
// }
//=======
//>>>>>>> sycl
Copy link
Contributor

Choose a reason for hiding this comment

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

A merge gone wrong?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes :) That's fixed now, thanks for spotting it

return ModuleChanged ? PreservedAnalyses::none() : PreservedAnalyses::all();
}
Loading