add clz builtin #965

tdavidcl · 2023-03-07T15:10:41Z

Hi,

I've tried to implement the clz builtin in OpenSYCL by mimicking the way it was done for mul24.

The part where I'm unsure whether the definition is correct are :

HIPSYCL_DEFINE_BUILTIN(clz, HIPSYCL_BUILTIN_OVERLOAD_SET_GENINTEGER,
                       HIPSYCL_BUILTIN_GENERATOR_UNARY_T)

and

template<class T>
HIPSYCL_BUILTIN T __hipsycl_clz(T x) noexcept {
  HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
}

Also do you have a mechanism in place to check the correctness of builtin functions ?

illuhad · 2023-03-07T20:01:49Z

Nice, more awesome PRs! :D

HIPSYCL_DEFINE_BUILTIN(clz, HIPSYCL_BUILTIN_OVERLOAD_SET_GENINTEGER,
HIPSYCL_BUILTIN_GENERATOR_UNARY_T)

This looks plausible to me.

template
HIPSYCL_BUILTIN T __hipsycl_clz(T x) noexcept {
HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
}

That too.

Also do you have a mechanism in place to check the correctness of builtin functions ?

You can add a test case. We do test builtin correctness for math functions here: https://github.com/OpenSYCL/OpenSYCL/blob/develop/tests/sycl/math.cpp
But we don't yet have a similar file for integer builtins. You can add it :)

There seems to be an issue with the current enable_if logic in the host path. Maybe it is easier just using overloads?

src/libkernel/sscp/ptx/integer.cpp

include/hipSYCL/sycl/libkernel/sscp/builtins.hpp

tdavidcl · 2023-03-09T14:54:30Z

There seems to be an issue with the current enable_if logic in the host path. Maybe it is easier just using overloads?

I've changed the declaration now the checks passes on my computer at least

Currently on the host side I use a declaration like this

template <class T,
          std::enable_if_t<
              (std::is_same_v<T, unsigned int> || std::is_same_v<T, int> ||
               std::is_same_v<T, unsigned short> || std::is_same_v<T, short> ||
               std::is_same_v<T, unsigned char> ||
               std::is_same_v<T, signed char> || std::is_same_v<T, char>),
              int> = 0>
HIPSYCL_BUILTIN T __hipsycl_clz(T x) noexcept {
  return __builtin_clz(x);
}

template <class T, std::enable_if_t<(std::is_same_v<T, unsigned long> ||
                                     std::is_same_v<T, long>),
                                    int> = 0>
HIPSYCL_BUILTIN T __hipsycl_clz(T x) noexcept {
  return __builtin_clzl(x);
}

template <class T, std::enable_if_t<(std::is_same_v<T, unsigned long long> ||
                                     std::is_same_v<T, long long>),
                                    int> = 0>
HIPSYCL_BUILTIN T __hipsycl_clz(T x) noexcept {
  return __builtin_clzll(x);
}

which does work but is rather long. Would this be ok ?

illuhad · 2023-03-09T15:18:00Z

I feel like this could be simplified. Would maybe something like this work?

template<class T>
T __hipsycl_clz(T x) {
  if constexpr(std::is_same_v<T, long long> || std::is_same_v<T, unsigned long long>)
    return __builtin_clzll(x);
  else if constexpr(std::is_same_v<T, long> || std::is_same_v<T unsigned long>)
    return __builtin_clzl(x);
  else
   return __builtin_clz(x); // do we need static_cast<T>(...) for the return statements?
}

illuhad · 2023-03-09T15:21:14Z

CI says nvc++ is not happy:

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (char)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=char]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (signed char)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=signed char]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (unsigned char)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=unsigned char]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (short)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=short]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (unsigned short)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=unsigned short]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (long)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=long]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

"/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtin_interface.hpp", line 534: error: no instance of overloaded function "hipsycl::sycl::detail::hiplike_builtins::__hipsycl_clz" matches the argument list
            argument types are: (unsigned long)
    HIPSYCL_RETURN_DISPATCH_BUILTIN(__hipsycl_clz, x);
    ^
          detected during instantiation of "T hipsycl::sycl::detail::__hipsycl_clz(T) noexcept [with T=unsigned long]" at line 788 of "/home/runner/work/OpenSYCL/OpenSYCL/build/install/bin/../include/CL/../hipSYCL/sycl/libkernel/builtins.hpp"

include/hipSYCL/sycl/libkernel/host/builtins.hpp

tdavidcl · 2023-03-09T16:35:54Z

CI says nvc++ is not happy:

actually after checking the hip/cuda doc specify clz for 32 bit ints and clzll for 64 bit ints.

So I've changed it to :

template<class T>
HIPSYCL_HIPLIKE_BUILTIN T __hipsycl_clz(T x) noexcept {

  // use __clzll or __clz by checking the bit lenght because
  // the nvidia/hip documentation mention clz as 32 bits and clzll as 64
  
  if constexpr (sizeof(T)*CHAR_BIT == 64){
    return __clzll(static_cast<__hipsycl_int64>(x));
  }

  return __clz(static_cast<__hipsycl_int32>(x));
}

It should also solve the instantiation issues.

include/hipSYCL/sycl/libkernel/generic/hiplike/builtins.hpp

include/hipSYCL/sycl/libkernel/host/builtins.hpp

include/hipSYCL/sycl/libkernel/generic/hiplike/builtins.hpp

Co-authored-by: Ronan Keryell <ronan@keryell.fr>

illuhad · 2023-04-05T06:40:53Z

Also do you have a mechanism in place to check the correctness of builtin functions ?

Can you confirm whether you intend to also add a test case, or if this PR should be considered for final review/testing as is?

tdavidcl · 2023-04-05T23:09:19Z

I think I won't be able to work on this one for a few weeks, so I won't be able to add testing to this as of now...

Can you confirm whether you intend to also add a test case, or if this PR should be considered for final review/testing as is?

So do you prefer to wait until i add some test case or do you want to review it now ?

include/hipSYCL/sycl/libkernel/sscp/builtins.hpp

src/libkernel/sscp/amdgpu/integer.cpp

src/libkernel/sscp/spirv/integer.cpp

illuhad · 2023-04-26T19:56:35Z

So do you prefer to wait until i add some test case or do you want to review it now ?

I'll start reviewing now and merge once tests are there too :-)

tdavidcl · 2023-05-03T12:25:44Z

I've added the test, and cleaned the clz implementation.
The tests pass on omp mode, but i don't have hardware/config to test on amd & nvidia right now.

illuhad · 2023-05-03T12:34:06Z

Thank you, I have just merged our new self-hosted CI with NVIDIA and AMD GPUs - so if you rebase on current develop, the tests will be run on GPUs :)

illuhad · 2023-05-04T12:31:04Z

I fixed the typo for nvcc, but i'm suprised that nvcc doesn't have __builtin_clz, would __clz, __clzll work since it is compiling some CUDA ?

We don't support nvcc, only nvc++ :-) (those are two very different compilers).
I'm not sure if the CUDA intrinsic would work in the host path - most likely not, so we'd have to use the fallback clz.

tdavidcl · 2023-05-04T13:00:54Z

I fixed the typo for nvcc, but i'm suprised that nvcc doesn't have __builtin_clz, would __clz, __clzll work since it is compiling some CUDA ?

We don't support nvcc, only nvc++ :-) (those are two very different compilers). I'm not sure if the CUDA intrinsic would work in the host path - most likely not, so we'd have to use the fallback clz.

yeah my bad i meant nvc++. The test in math.cpp with nvc++ passes, with the fix, tho i'm not sure if the __clz or fallback_clz version was used

edit : after checking in the doc it says __device__ int __clz ( int x ), so yes not on the host

tdavidcl · 2023-05-04T13:10:30Z

So i reverted to the previous commit. With the test passing on all config (+ self hosted), is everything ok for final review ?

include/hipSYCL/sycl/libkernel/host/builtins.hpp

illuhad

Thanks!

illuhad · 2023-05-04T14:53:00Z

Hm.. it seems that there might still be an issue on CPU: https://github.com/OpenSYCL/OpenSYCL/actions/runs/4883593430/jobs/8716435004?pr=965

tdavidcl · 2023-05-04T16:28:58Z

Yes, but i can not reproduce the issue ...
The previous version with the issue in the null case was working, so i've rolled back to it with an additional check on x==0

https://godbolt.org/z/ToKsdoTKM

illuhad · 2023-05-04T16:42:03Z

The only potential issue I see is that it would probably evalute !bset[sizeof(T)*CHAR_BIT - idx -1] before it would evaluate the idx<sizeof(T)*CHAR_BIT. So in the step when the second condition evaluates to false, it could happen that it accessed the invalid memory one time (thus triggering UB) for the first condition before finding out that it should not have done that.

Let's see if it works now :)

tdavidcl · 2023-05-04T17:25:51Z

it's odd that the macos test doesn't work, it was ok here : https://github.com/OpenSYCL/OpenSYCL/actions/runs/4883180762/jobs/8714601525

and the only difference is this line :

if(x==0){return sizeof(T)*CHAR_BIT;}

and the test pass on my MAC (M1) with appleClang :'(

illuhad · 2023-05-04T17:34:51Z

Mac also failed previously, before the fix. So I wonder why changing to if(x==0) still has not resolved the issue entirely... For reference, this is how Mac CI works: https://github.com/OpenSYCL/OpenSYCL/blob/develop/.github/workflows/macos.yml
If you have a Mac, can you try following these commands to see if you can reproduce? Or maybe try address sanitizer to see if it's still doing some invalid access.

tdavidcl · 2023-05-04T18:08:21Z

it is still correct on my Mac, I've added bunch of print statements, can you try running it only the Mac OS workflow?

illuhad · 2023-05-04T18:59:55Z

I've invited you into the OpenSYCL organization; the workflows will then run automatically for you.

tdavidcl · 2023-05-04T20:00:50Z

I've invited you into the OpenSYCL organization; the workflows will then run automatically for you.

Oh thanks a lot, it will speed up the process for sure ^^

tdavidcl · 2023-05-04T21:45:27Z

ok I found the issue : in the failing tests you have this ...

castvar = 0
__builtin_clz(castvar) = 73911 (unsigned int)
__builtin_clz(0) = 73896 (unsigned int)

the clz builtin looks broken in the Mac OS workflow, also it does not look consistent and hard to check ...

also I found this : https://stackoverflow.com/questions/19527897/how-undefined-are-builtin-ctz0-or-builtin-clz0
love it ...

illuhad · 2023-05-04T21:57:39Z

Good find!

So basically we should always

if(x == 0)
  return sizeof(T)*CHAR_BIT
else
  return __builtin_clz(x)

whenever using the builtin?

I do wonder though when this case is so undefined, whether the SYCL specification even mandates a specific result here? I can kind of see that the concept of leading or trailing zeros falls apart somewhat in the case of a value of 0.

Good thing we noticed this edge case :)

tdavidcl · 2023-05-04T22:05:09Z

The SYCL standard actually says :

geninteger clz(geninteger x)

Returns the number of leading 0-bits in x, starting at the most significant bit position. 
If x is 0, returns the size in bits of the type of x or component type of x, if x is a vector type.

So I've added the check for __builtin_clz

In cuda :

__device__  int __clz ( int  x )
Return the number of consecutive high-order zero bits in a 32-bit integer.

so this one is fine

I'm not 100% sure for the others

illuhad · 2023-05-04T22:08:44Z

My guess is that SPIR-V will probably be aligned with OpenCL, which is probably aligned with SYCL. No idea for AMD. So maybe also add the check for those two backends to be sure?

illuhad · 2023-05-04T22:10:20Z

... Or maybe just add the check to the high-level builtin interface, so that the backends don't need to be concerned about this case anymore?

tdavidcl · 2023-05-04T22:16:26Z

My guess is that SPIR-V will probably be aligned with OpenCL, which is probably aligned with SYCL. No idea for AMD. So maybe also add the check for those two backends to be sure?

As you guessed
From https://registry.khronos.org/SPIR-V/specs/unified1/OpenCL.ExtendedInstructionSet.100.html

clz

Returns the number of leading 0 bits in x, starting at the most significant bit position. 
If x is 0, returns the size in bits of the type of x or component type of x, if x is a vector.

So spir-v should be fine

... Or maybe just add the check to the high-level builtin interface, so that the backends don't need to be concerned about this case anymore?

Actually we can just add the check in tests, also I don't know also why in the first place the value was 0 in the inputs, it shouldn't when i'm looking at the test.

tdavidcl · 2023-05-04T22:21:40Z

From RadeonCompute gits:

https://github.com/RadeonOpenCompute/clang/blob/a09d37e345861d68f9768939e485d265f4fcb0ce/include/clang/Basic/TargetInfo.h#L776

Looks like they are aware and check for it :), so looks like we only have to be careful with the host.

add clz builtin

4b5a43d

illuhad reviewed Mar 7, 2023

View reviewed changes

src/libkernel/sscp/ptx/integer.cpp Outdated Show resolved Hide resolved

include/hipSYCL/sycl/libkernel/sscp/builtins.hpp Outdated Show resolved Hide resolved

tdavidcl added 3 commits March 8, 2023 15:06

fix host signature

a8fb714

use lidevice header for nv builtins

d910069

change sscp clz code path

c78e774

illuhad reviewed Mar 9, 2023

View reviewed changes

include/hipSYCL/sycl/libkernel/host/builtins.hpp Outdated Show resolved Hide resolved

tdavidcl added 2 commits March 9, 2023 17:25

fix hiplike clz by checking bit lenght

f1ccfbf

add host clz fallback

cbcfcab

illuhad reviewed Mar 10, 2023

View reviewed changes

tdavidcl added 2 commits March 13, 2023 18:33

cleaning clz implementation

81d7e6e

fall back clz

00e5acb

illuhad reviewed Mar 16, 2023

View reviewed changes

include/hipSYCL/sycl/libkernel/host/builtins.hpp Outdated Show resolved Hide resolved

keryell reviewed Mar 17, 2023

View reviewed changes

include/hipSYCL/sycl/libkernel/generic/hiplike/builtins.hpp Outdated Show resolved Hide resolved

include/hipSYCL/sycl/libkernel/generic/hiplike/builtins.hpp Outdated Show resolved Hide resolved

tdavidcl and others added 2 commits March 18, 2023 16:06

fix comment typo

2489bd9

Co-authored-by: Ronan Keryell <ronan@keryell.fr>

remove __ for the fallback clz

01bbb25

illuhad reviewed Apr 26, 2023

View reviewed changes

include/hipSYCL/sycl/libkernel/sscp/builtins.hpp Show resolved Hide resolved

src/libkernel/sscp/amdgpu/integer.cpp Show resolved Hide resolved

src/libkernel/sscp/spirv/integer.cpp Outdated Show resolved Hide resolved

tdavidcl added 4 commits May 3, 2023 11:09

Merge branch 'OpenSYCL:develop' into feature/clz-math-builtin

ebd33bf

implement u16 & u8 variant of clz

404f43e

ptx clz u8 & u16

4e826b0

clz testing

3c9d419

remove the __clz host for nvc++ and fallback to fallback_clz

5688ee3

illuhad reviewed May 4, 2023

View reviewed changes

include/hipSYCL/sycl/libkernel/host/builtins.hpp Show resolved Hide resolved

fix issue in fallback clz

0babda5

illuhad approved these changes May 4, 2023

View reviewed changes

changing fallback implementation

107ddbb

print test

842fb9f

additional prints

aab81b5

tdavidcl added 3 commits May 4, 2023 22:26

additional cast layer

0033961

printing clz(0)

d6ed518

fix build

eceb269

remove prints & add check to avoid clz(0)

0e616c2

illuhad merged commit a54d87b into AdaptiveCpp:develop May 5, 2023
17 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

add clz builtin #965

add clz builtin #965

tdavidcl commented Mar 7, 2023 •

edited

illuhad commented Mar 7, 2023

tdavidcl commented Mar 9, 2023

illuhad commented Mar 9, 2023

illuhad commented Mar 9, 2023

tdavidcl commented Mar 9, 2023 •

edited

illuhad commented Apr 5, 2023

tdavidcl commented Apr 5, 2023

illuhad commented Apr 26, 2023

tdavidcl commented May 3, 2023

illuhad commented May 3, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 •

edited

tdavidcl commented May 4, 2023

illuhad left a comment

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 •

edited

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 •

edited

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

tdavidcl commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 •

edited

illuhad commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

tdavidcl commented May 4, 2023

add clz builtin #965

add clz builtin #965

Conversation

tdavidcl commented Mar 7, 2023 • edited

illuhad commented Mar 7, 2023

tdavidcl commented Mar 9, 2023

illuhad commented Mar 9, 2023

illuhad commented Mar 9, 2023

tdavidcl commented Mar 9, 2023 • edited

illuhad commented Apr 5, 2023

tdavidcl commented Apr 5, 2023

illuhad commented Apr 26, 2023

tdavidcl commented May 3, 2023

illuhad commented May 3, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 • edited

tdavidcl commented May 4, 2023

illuhad left a comment

Choose a reason for hiding this comment

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 • edited

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 • edited

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

tdavidcl commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023 • edited

illuhad commented May 4, 2023

illuhad commented May 4, 2023

tdavidcl commented May 4, 2023

tdavidcl commented May 4, 2023

tdavidcl commented Mar 7, 2023 •

edited

tdavidcl commented Mar 9, 2023 •

edited

tdavidcl commented May 4, 2023 •

edited

tdavidcl commented May 4, 2023 •

edited

tdavidcl commented May 4, 2023 •

edited

tdavidcl commented May 4, 2023 •

edited