[cfe-commits] [libclc] r157590 - in /libclc/trunk: generic/include/clc/clc.h generic/include/clc/convert.h generic/lib/SOURCES generic/lib/convert.cl test/convert.cl
Peter Collingbourne
peter at pcc.me.uk
Mon May 28 13:42:55 PDT 2012
Author: pcc
Date: Mon May 28 15:42:54 2012
New Revision: 157590
URL: http://llvm.org/viewvc/llvm-project?rev=157590&view=rev
Log:
Explicit conversions.
Added:
libclc/trunk/generic/include/clc/convert.h
libclc/trunk/generic/lib/convert.cl
libclc/trunk/test/convert.cl
Modified:
libclc/trunk/generic/include/clc/clc.h
libclc/trunk/generic/lib/SOURCES
Modified: libclc/trunk/generic/include/clc/clc.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/clc.h?rev=157590&r1=157589&r2=157590&view=diff
==============================================================================
--- libclc/trunk/generic/include/clc/clc.h (original)
+++ libclc/trunk/generic/include/clc/clc.h Mon May 28 15:42:54 2012
@@ -17,6 +17,9 @@
/* 6.1 Supported Data Types */
#include <clc/clctypes.h>
+/* 6.2.3 Explicit Conversions */
+#include <clc/convert.h>
+
/* 6.2.4.2 Reinterpreting Types Using as_type() and as_typen() */
#include <clc/as_type.h>
Added: libclc/trunk/generic/include/clc/convert.h
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/convert.h?rev=157590&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/convert.h (added)
+++ libclc/trunk/generic/include/clc/convert.h Mon May 28 15:42:54 2012
@@ -0,0 +1,82 @@
+#define _CLC_CONVERT_DECL(FROM_TYPE, TO_TYPE, SUFFIX) \
+ _CLC_OVERLOAD _CLC_DECL TO_TYPE convert_##TO_TYPE##SUFFIX(FROM_TYPE x);
+
+_CLC_CONVERT_DECL(long, char, )
+_CLC_CONVERT_DECL(ulong, uchar, )
+_CLC_CONVERT_DECL(long, short, )
+_CLC_CONVERT_DECL(ulong, ushort, )
+_CLC_CONVERT_DECL(long, int, )
+_CLC_CONVERT_DECL(ulong, uint, )
+_CLC_CONVERT_DECL(long, long, )
+_CLC_CONVERT_DECL(ulong, ulong, )
+#ifdef cl_khr_fp64
+_CLC_CONVERT_DECL(double, float, )
+_CLC_CONVERT_DECL(double, double, )
+#else
+_CLC_CONVERT_DECL(float, float, )
+#endif
+
+_CLC_CONVERT_DECL(long, char, _sat)
+_CLC_CONVERT_DECL(ulong, uchar, _sat)
+_CLC_CONVERT_DECL(long, short, _sat)
+_CLC_CONVERT_DECL(ulong, ushort, _sat)
+_CLC_CONVERT_DECL(long, int, _sat)
+_CLC_CONVERT_DECL(ulong, uint, _sat)
+_CLC_CONVERT_DECL(long, long, _sat)
+_CLC_CONVERT_DECL(ulong, ulong, _sat)
+#ifdef cl_khr_fp64
+_CLC_CONVERT_DECL(double, float, _sat)
+_CLC_CONVERT_DECL(double, double, _sat)
+#else
+_CLC_CONVERT_DECL(float, float, _sat)
+#endif
+
+#define _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, TO_TYPE, SUFFIX) \
+ _CLC_CONVERT_DECL(FROM_TYPE##2, TO_TYPE##2, SUFFIX) \
+ _CLC_CONVERT_DECL(FROM_TYPE##3, TO_TYPE##3, SUFFIX) \
+ _CLC_CONVERT_DECL(FROM_TYPE##4, TO_TYPE##4, SUFFIX) \
+ _CLC_CONVERT_DECL(FROM_TYPE##8, TO_TYPE##8, SUFFIX) \
+ _CLC_CONVERT_DECL(FROM_TYPE##16, TO_TYPE##16, SUFFIX)
+
+#define _CLC_VECTOR_CONVERT_FROM1(FROM_TYPE, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, char, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, uchar, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, int, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, uint, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, short, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, ushort, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, long, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, ulong, SUFFIX) \
+ _CLC_VECTOR_CONVERT_DECL(FROM_TYPE, float, SUFFIX)
+
+#ifdef cl_khr_fp64
+#define _CLC_VECTOR_CONVERT_FROM(FROM_TYPE, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM1(FROM_TYPE, SUFFIX) \
+ _CLC_VECTOR_CONVERT(FROM_TYPE, double, SUFFIX)
+#else
+#define _CLC_VECTOR_CONVERT_FROM(FROM_TYPE, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM1(FROM_TYPE, SUFFIX)
+#endif
+
+#define _CLC_VECTOR_CONVERT_TO1(SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(char, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(uchar, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(int, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(uint, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(short, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(ushort, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(long, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(ulong, SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(float, SUFFIX)
+
+#ifdef cl_khr_fp64
+#define _CLC_VECTOR_CONVERT_TO(SUFFIX) \
+ _CLC_VECTOR_CONVERT_TO1(SUFFIX) \
+ _CLC_VECTOR_CONVERT_FROM(double, SUFFIX)
+#else
+#define _CLC_VECTOR_CONVERT_TO(SUFFIX) \
+ _CLC_VECTOR_CONVERT_TO1(SUFFIX)
+#endif
+
+_CLC_VECTOR_CONVERT_TO()
+_CLC_VECTOR_CONVERT_TO(_sat)
Modified: libclc/trunk/generic/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/SOURCES?rev=157590&r1=157589&r2=157590&view=diff
==============================================================================
--- libclc/trunk/generic/lib/SOURCES (original)
+++ libclc/trunk/generic/lib/SOURCES Mon May 28 15:42:54 2012
@@ -1,3 +1,4 @@
+convert.cl
geometric/cross.cl
geometric/dot.cl
geometric/length.cl
Added: libclc/trunk/generic/lib/convert.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/convert.cl?rev=157590&view=auto
==============================================================================
--- libclc/trunk/generic/lib/convert.cl (added)
+++ libclc/trunk/generic/lib/convert.cl Mon May 28 15:42:54 2012
@@ -0,0 +1,122 @@
+#include <clc/clc.h>
+
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+
+#define CONVERT_ID(FROM_TYPE, TO_TYPE, SUFFIX) \
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE convert_##TO_TYPE##SUFFIX(FROM_TYPE x) { \
+ return x; \
+ }
+
+#define CONVERT_VECTORIZE(FROM_TYPE, TO_TYPE, SUFFIX) \
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE##2 convert_##TO_TYPE##2##SUFFIX(FROM_TYPE##2 x) { \
+ return (TO_TYPE##2)(convert_##TO_TYPE##SUFFIX(x.x), convert_##TO_TYPE##SUFFIX(x.y)); \
+ } \
+\
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE##3 convert_##TO_TYPE##3##SUFFIX(FROM_TYPE##3 x) { \
+ return (TO_TYPE##3)(convert_##TO_TYPE##SUFFIX(x.x), convert_##TO_TYPE##SUFFIX(x.y), convert_##TO_TYPE##SUFFIX(x.z)); \
+ } \
+\
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE##4 convert_##TO_TYPE##4##SUFFIX(FROM_TYPE##4 x) { \
+ return (TO_TYPE##4)(convert_##TO_TYPE##2##SUFFIX(x.lo), convert_##TO_TYPE##2##SUFFIX(x.hi)); \
+ } \
+\
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE##8 convert_##TO_TYPE##8##SUFFIX(FROM_TYPE##8 x) { \
+ return (TO_TYPE##8)(convert_##TO_TYPE##4##SUFFIX(x.lo), convert_##TO_TYPE##4##SUFFIX(x.hi)); \
+ } \
+\
+ _CLC_OVERLOAD _CLC_DEF TO_TYPE##16 convert_##TO_TYPE##16##SUFFIX(FROM_TYPE##16 x) { \
+ return (TO_TYPE##16)(convert_##TO_TYPE##8##SUFFIX(x.lo), convert_##TO_TYPE##8##SUFFIX(x.hi)); \
+ }
+
+CONVERT_ID(long, char, )
+CONVERT_ID(ulong, uchar, )
+CONVERT_ID(long, short, )
+CONVERT_ID(ulong, ushort, )
+CONVERT_ID(long, int, )
+CONVERT_ID(ulong, uint, )
+CONVERT_ID(long, long, )
+CONVERT_ID(ulong, ulong, )
+#ifdef cl_khr_fp64
+CONVERT_ID(double, float, )
+CONVERT_ID(double, double, )
+#else
+CONVERT_ID(float, float, )
+#endif
+
+_CLC_OVERLOAD _CLC_DEF char convert_char_sat(long l) {
+ return l > 127 ? 127 : l < -128 ? -128 : l;
+}
+
+_CLC_OVERLOAD _CLC_DEF uchar convert_uchar_sat(ulong l) {
+ return l > 255U ? 255U : l;
+}
+
+_CLC_OVERLOAD _CLC_DEF short convert_short_sat(long l) {
+ return l > 32767 ? 32767 : l < -32768 ? -32768 : l;
+}
+
+_CLC_OVERLOAD _CLC_DEF ushort convert_ushort_sat(ulong l) {
+ return l > 65535U ? 65535U : l;
+}
+
+_CLC_OVERLOAD _CLC_DEF int convert_int_sat(long l) {
+ return l > ((1L<<31)-1) ? ((1L<<31L)-1) : l < -(1L<<31) ? -(1L<<31) : l;
+}
+
+_CLC_OVERLOAD _CLC_DEF uint convert_uint_sat(ulong l) {
+ return l > ((1UL<<32)-1) ? ((1UL<<32)-1) : l;
+}
+
+CONVERT_ID(long, long, _sat)
+CONVERT_ID(ulong, ulong, _sat)
+#ifdef cl_khr_fp64
+CONVERT_ID(double, float, _sat)
+CONVERT_ID(double, double, _sat)
+#else
+CONVERT_ID(float, float, _sat)
+#endif
+
+#define CONVERT_VECTORIZE_FROM1(FROM_TYPE, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, char, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, uchar, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, int, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, uint, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, short, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, ushort, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, long, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, ulong, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, float, SUFFIX)
+
+#ifdef cl_khr_fp64
+#define CONVERT_VECTORIZE_FROM(FROM_TYPE, SUFFIX) \
+ CONVERT_VECTORIZE_FROM1(FROM_TYPE, SUFFIX) \
+ CONVERT_VECTORIZE(FROM_TYPE, double, SUFFIX)
+#else
+#define CONVERT_VECTORIZE_FROM(FROM_TYPE, SUFFIX) \
+ CONVERT_VECTORIZE_FROM1(FROM_TYPE, SUFFIX)
+#endif
+
+#define CONVERT_VECTORIZE_TO1(SUFFIX) \
+ CONVERT_VECTORIZE_FROM(char, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(uchar, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(int, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(uint, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(short, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(ushort, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(long, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(ulong, SUFFIX) \
+ CONVERT_VECTORIZE_FROM(float, SUFFIX)
+
+#ifdef cl_khr_fp64
+#define CONVERT_VECTORIZE_TO(SUFFIX) \
+ CONVERT_VECTORIZE_TO1(SUFFIX) \
+ CONVERT_VECTORIZE_FROM(double, SUFFIX)
+#else
+#define CONVERT_VECTORIZE_TO(SUFFIX) \
+ CONVERT_VECTORIZE_TO1(SUFFIX)
+#endif
+
+CONVERT_VECTORIZE_TO()
+CONVERT_VECTORIZE_TO(_sat)
Added: libclc/trunk/test/convert.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/test/convert.cl?rev=157590&view=auto
==============================================================================
--- libclc/trunk/test/convert.cl (added)
+++ libclc/trunk/test/convert.cl Mon May 28 15:42:54 2012
@@ -0,0 +1,3 @@
+__kernel void foo(int4 *x, float4 *y) {
+ *x = convert_int4(*y);
+}
More information about the cfe-commits
mailing list