Skip to content

[libclc] Move clcmacro.h to CLC library. NFC #114845

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

Merged
merged 1 commit into from
Nov 4, 2024
Merged

[libclc] Move clcmacro.h to CLC library. NFC #114845

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
2 changes: 1 addition & 1 deletion libclc/amdgcn/lib/integer/popcount.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include <utils.h>
#include <clc/utils.h>
#include <integer/popcount.h>

#define __CLC_BODY "popcount.inc"
Expand Down
3 changes: 1 addition & 2 deletions libclc/amdgcn/lib/math/fmax.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,5 @@
#include <clc/clc.h>

#include "../../../generic/lib/clcmacro.h"
#include <clc/clcmacro.h>

_CLC_DEF _CLC_OVERLOAD float fmax(float x, float y)
{
Expand Down
3 changes: 1 addition & 2 deletions libclc/amdgcn/lib/math/fmin.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,5 @@
#include <clc/clc.h>

#include "../../../generic/lib/clcmacro.h"
#include <clc/clcmacro.h>

_CLC_DEF _CLC_OVERLOAD float fmin(float x, float y)
{
Expand Down
3 changes: 1 addition & 2 deletions libclc/amdgcn/lib/math/ldexp.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,7 @@
*/

#include <clc/clc.h>

#include "../../../generic/lib/clcmacro.h"
#include <clc/clcmacro.h>

#ifdef __HAS_LDEXPF__
#define BUILTINF __builtin_amdgcn_ldexpf
Expand Down
2 changes: 1 addition & 1 deletion libclc/amdgpu/lib/math/half_native_unary.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
#include <utils.h>
#include <clc/utils.h>

#define __CLC_HALF_FUNC(x) __CLC_CONCAT(half_, x)
#define __CLC_NATIVE_FUNC(x) __CLC_CONCAT(native_, x)
Expand Down
2 changes: 1 addition & 1 deletion libclc/amdgpu/lib/math/nextafter.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "../lib/clcmacro.h"
#include <clc/clcmacro.h>
#include <math/clc_nextafter.h>

_CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)
Expand Down
4 changes: 2 additions & 2 deletions libclc/amdgpu/lib/math/sqrt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,9 +20,9 @@
* THE SOFTWARE.
*/

#include <clc/clc.h>
#include "../../../generic/lib/clcmacro.h"
#include "math/clc_sqrt.h"
#include <clc/clc.h>
#include <clc/clcmacro.h>

_CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)

Expand Down
214 changes: 116 additions & 98 deletions libclc/generic/lib/clcmacro.h → libclc/clc/include/clc/clcmacro.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,92 +1,105 @@
#include <clc/clc.h>
#include <utils.h>

#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
#ifndef __CLC_CLCMACRO_H__
#define __CLC_CLCMACRO_H__

#include <clc/internal/clc.h>
#include <clc/utils.h>

#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
}

#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \
return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
FUNCTION(x.z, y.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
ARG2_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \
return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
FUNCTION(x.z, y.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
}

#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
FUNCTION(x, y.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\

#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE, ARG3_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, ARG3_TYPE##2 z) { \
return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, ARG3_TYPE##3 z) { \
return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
FUNCTION(x.z, y.z, z.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, ARG3_TYPE##4 z) { \
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, ARG3_TYPE##8 z) { \
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, ARG3_TYPE##16 z) { \
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
ARG2_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
FUNCTION(x, y.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
}

#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
ARG2_TYPE, ARG3_TYPE) \
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, \
ARG3_TYPE##2 z) { \
return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
} \
\
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, \
ARG3_TYPE##3 z) { \
return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
FUNCTION(x.z, y.z, z.z)); \
} \
\
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, \
ARG3_TYPE##4 z) { \
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), \
FUNCTION(x.hi, y.hi, z.hi)); \
} \
\
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, \
ARG3_TYPE##8 z) { \
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), \
FUNCTION(x.hi, y.hi, z.hi)); \
} \
\
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, \
ARG3_TYPE##16 z) { \
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), \
FUNCTION(x.hi, y.hi, z.hi)); \
}

#define _CLC_V_S_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
Expand Down Expand Up @@ -161,21 +174,24 @@
ARG2_TYPE, 8) *)((ADDR_SPACE ARG2_TYPE *)y + 8))); \
}

#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \
return BUILTIN(x, y); \
} \
_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \
ARG2_TYPE) \
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \
return BUILTIN(x, y); \
} \
_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, \
ARG2_TYPE)

#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG( \
RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \
ARG2_TYPE) \
_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, \
FUNCTION, ARG1_TYPE, ARG2_TYPE)

#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { \
return BUILTIN(x); \
} \
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { return BUILTIN(x); } \
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)

#ifdef cl_khr_fp16

Expand All @@ -199,3 +215,5 @@ _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
#define _CLC_DEFINE_BINARY_BUILTIN_FP16(FUNCTION)

#endif

#endif // __CLC_CLCMACRO_H__
4 changes: 2 additions & 2 deletions libclc/clspv/lib/math/fma.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,9 +24,9 @@
// (__clc_sw_fma), but avoids the use of ulong in favor of uint2. The logic has
// been updated as appropriate.

#include <clc/clc.h>
#include "../../../generic/lib/clcmacro.h"
#include "../../../generic/lib/math/math.h"
#include <clc/clc.h>
#include <clc/clcmacro.h>

struct fp {
uint2 mantissa;
Expand Down
3 changes: 3 additions & 0 deletions libclc/generic/include/math/clc_sqrt.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
#include <clc/clcfunc.h>
#include <clc/clctypes.h>

#define __CLC_FUNCTION __clc_sqrt
#define __CLC_BODY <clc/math/unary_decl.inc>
#include <clc/math/gentype.inc>
Expand Down
10 changes: 0 additions & 10 deletions libclc/generic/include/utils.h
View file
Open in desktop

This file was deleted.

2 changes: 1 addition & 1 deletion libclc/generic/lib/atom_int32_binary.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "utils.h"
#include <clc/utils.h>

#define __CLC_ATOM_IMPL(AS, TYPE) \
_CLC_OVERLOAD _CLC_DEF TYPE __CLC_XCONCAT(atom_, __CLC_ATOMIC_OP) (volatile AS TYPE *p, TYPE val) { \
Expand Down
3 changes: 1 addition & 2 deletions libclc/generic/lib/common/degrees.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,7 @@
*/

#include <clc/clc.h>

#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF float degrees(float radians) {
// 180/pi = ~57.29577951308232087685 or 0x1.ca5dc1a63c1f8p+5 or 0x1.ca5dc2p+5F
Expand Down
3 changes: 1 addition & 2 deletions libclc/generic/lib/common/radians.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,7 @@
*/

#include <clc/clc.h>

#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF float radians(float degrees) {
// pi/180 = ~0.01745329251994329577 or 0x1.1df46a2529d39p-6 or 0x1.1df46ap-6F
Expand Down
2 changes: 1 addition & 1 deletion libclc/generic/lib/common/sign.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "../clcmacro.h"
#include <clc/clcmacro.h>

#define SIGN(TYPE, F) \
_CLC_DEF _CLC_OVERLOAD TYPE sign(TYPE x) { \
Expand Down
3 changes: 1 addition & 2 deletions libclc/generic/lib/common/smoothstep.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,7 @@
*/

#include <clc/clc.h>

#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF float smoothstep(float edge0, float edge1, float x) {
float t = clamp((x - edge0) / (edge1 - edge0), 0.0f, 1.0f);
Expand Down
3 changes: 1 addition & 2 deletions libclc/generic/lib/common/step.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,7 @@
*/

#include <clc/clc.h>

#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF float step(float edge, float x) {
return x < edge ? 0.0f : 1.0f;
Expand Down
2 changes: 1 addition & 1 deletion libclc/generic/lib/integer/add_sat.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "../clcmacro.h"
#include <clc/clcmacro.h>

// From add_sat.ll
_CLC_DECL char __clc_add_sat_s8(char, char);
Expand Down
2 changes: 1 addition & 1 deletion libclc/generic/lib/integer/clz.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF char clz(char x) {
return clz((ushort)(uchar)x) - 8;
Expand Down
2 changes: 1 addition & 1 deletion libclc/generic/lib/integer/mad_sat.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#include <clc/clc.h>
#include "../clcmacro.h"
#include <clc/clcmacro.h>

_CLC_OVERLOAD _CLC_DEF char mad_sat(char x, char y, char z) {
return clamp((short)mad24((short)x, (short)y, (short)z), (short)CHAR_MIN, (short) CHAR_MAX);
Expand Down
Loading
Loading