Author: Fraser Cormack Date: 2025-07-08T10:37:06+01:00 New Revision: 9d11bd0db8bb96a106b2ea9018be52b20582e2f0
URL: https://github.com/llvm/llvm-project/commit/9d11bd0db8bb96a106b2ea9018be52b20582e2f0 DIFF: https://github.com/llvm/llvm-project/commit/9d11bd0db8bb96a106b2ea9018be52b20582e2f0.diff LOG: [libclc] Remove catch-all opencl/clc.h (#147490) This commit finishes the work started in #146840 and #147276. It makes each OpenCL header self-contained and each implementation file include only the headers it needs. It removes the need for a catch-all include file of all OpenCL builtin declarations. Added: Modified: libclc/opencl/include/clc/opencl/as_type.h libclc/opencl/include/clc/opencl/async/async_work_group_copy.h libclc/opencl/include/clc/opencl/async/async_work_group_strided_copy.h libclc/opencl/include/clc/opencl/async/prefetch.h libclc/opencl/include/clc/opencl/async/wait_group_events.h libclc/opencl/include/clc/opencl/atomic/atom_add.h libclc/opencl/include/clc/opencl/atomic/atom_and.h libclc/opencl/include/clc/opencl/atomic/atom_cmpxchg.h libclc/opencl/include/clc/opencl/atomic/atom_dec.h libclc/opencl/include/clc/opencl/atomic/atom_inc.h libclc/opencl/include/clc/opencl/atomic/atom_max.h libclc/opencl/include/clc/opencl/atomic/atom_min.h libclc/opencl/include/clc/opencl/atomic/atom_or.h libclc/opencl/include/clc/opencl/atomic/atom_sub.h libclc/opencl/include/clc/opencl/atomic/atom_xchg.h libclc/opencl/include/clc/opencl/atomic/atom_xor.h libclc/opencl/include/clc/opencl/atomic/atomic_add.h libclc/opencl/include/clc/opencl/atomic/atomic_and.h libclc/opencl/include/clc/opencl/atomic/atomic_cmpxchg.h libclc/opencl/include/clc/opencl/atomic/atomic_dec.h libclc/opencl/include/clc/opencl/atomic/atomic_inc.h libclc/opencl/include/clc/opencl/atomic/atomic_max.h libclc/opencl/include/clc/opencl/atomic/atomic_min.h libclc/opencl/include/clc/opencl/atomic/atomic_or.h libclc/opencl/include/clc/opencl/atomic/atomic_sub.h libclc/opencl/include/clc/opencl/atomic/atomic_xchg.h libclc/opencl/include/clc/opencl/atomic/atomic_xor.h libclc/opencl/include/clc/opencl/convert.h libclc/opencl/include/clc/opencl/explicit_fence/explicit_memory_fence.h libclc/opencl/include/clc/opencl/image/image.h libclc/opencl/include/clc/opencl/image/image_defines.h libclc/opencl/include/clc/opencl/misc/shuffle.h libclc/opencl/include/clc/opencl/misc/shuffle2.h libclc/opencl/include/clc/opencl/shared/clamp.h libclc/opencl/include/clc/opencl/shared/max.h libclc/opencl/include/clc/opencl/shared/min.h libclc/opencl/include/clc/opencl/shared/vload.h libclc/opencl/include/clc/opencl/shared/vstore.h libclc/opencl/include/clc/opencl/synchronization/barrier.h libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h libclc/opencl/include/clc/opencl/workitem/get_global_id.h libclc/opencl/include/clc/opencl/workitem/get_global_offset.h libclc/opencl/include/clc/opencl/workitem/get_global_size.h libclc/opencl/include/clc/opencl/workitem/get_group_id.h libclc/opencl/include/clc/opencl/workitem/get_local_id.h libclc/opencl/include/clc/opencl/workitem/get_local_size.h libclc/opencl/include/clc/opencl/workitem/get_num_groups.h libclc/opencl/include/clc/opencl/workitem/get_work_dim.h libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl libclc/opencl/lib/amdgcn-amdhsa/workitem/get_num_groups.cl libclc/opencl/lib/amdgcn/mem_fence/fence.cl libclc/opencl/lib/amdgcn/synchronization/barrier.cl libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl libclc/opencl/lib/amdgcn/workitem/get_global_size.cl libclc/opencl/lib/amdgcn/workitem/get_group_id.cl libclc/opencl/lib/amdgcn/workitem/get_local_id.cl libclc/opencl/lib/amdgcn/workitem/get_local_size.cl libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl libclc/opencl/lib/clspv/shared/vstore_half.cl libclc/opencl/lib/clspv/subnormal_config.cl libclc/opencl/lib/generic/async/async_work_group_copy.cl libclc/opencl/lib/generic/async/async_work_group_strided_copy.cl libclc/opencl/lib/generic/async/prefetch.cl libclc/opencl/lib/generic/async/wait_group_events.cl libclc/opencl/lib/generic/atomic/atom_add.cl libclc/opencl/lib/generic/atomic/atom_and.cl libclc/opencl/lib/generic/atomic/atom_int32_binary.inc libclc/opencl/lib/generic/atomic/atom_max.cl libclc/opencl/lib/generic/atomic/atom_min.cl libclc/opencl/lib/generic/atomic/atom_or.cl libclc/opencl/lib/generic/atomic/atom_sub.cl libclc/opencl/lib/generic/atomic/atom_xchg.cl libclc/opencl/lib/generic/atomic/atom_xor.cl libclc/opencl/lib/generic/atomic/atomic_add.cl libclc/opencl/lib/generic/atomic/atomic_and.cl libclc/opencl/lib/generic/atomic/atomic_cmpxchg.cl libclc/opencl/lib/generic/atomic/atomic_dec.cl libclc/opencl/lib/generic/atomic/atomic_inc.cl libclc/opencl/lib/generic/atomic/atomic_max.cl libclc/opencl/lib/generic/atomic/atomic_min.cl libclc/opencl/lib/generic/atomic/atomic_or.cl libclc/opencl/lib/generic/atomic/atomic_sub.cl libclc/opencl/lib/generic/atomic/atomic_xchg.cl libclc/opencl/lib/generic/atomic/atomic_xor.cl libclc/opencl/lib/generic/misc/shuffle.cl libclc/opencl/lib/generic/misc/shuffle2.cl libclc/opencl/lib/generic/shared/clamp.cl libclc/opencl/lib/generic/shared/max.cl libclc/opencl/lib/generic/shared/min.cl libclc/opencl/lib/generic/shared/vload.cl libclc/opencl/lib/generic/shared/vstore.cl libclc/opencl/lib/generic/subnormal_config.cl libclc/opencl/lib/generic/workitem/get_global_id.cl libclc/opencl/lib/generic/workitem/get_global_size.cl libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl libclc/opencl/lib/r600/image/get_image_channel_data_type.cl libclc/opencl/lib/r600/image/get_image_channel_order.cl libclc/opencl/lib/r600/image/get_image_depth.cl libclc/opencl/lib/r600/image/get_image_dim.cl libclc/opencl/lib/r600/image/get_image_height.cl libclc/opencl/lib/r600/image/get_image_width.cl libclc/opencl/lib/r600/image/read_imagef.cl libclc/opencl/lib/r600/image/read_imagei.cl libclc/opencl/lib/r600/image/read_imageui.cl libclc/opencl/lib/r600/image/write_imagef.cl libclc/opencl/lib/r600/image/write_imagei.cl libclc/opencl/lib/r600/image/write_imageui.cl libclc/opencl/lib/r600/synchronization/barrier.cl libclc/opencl/lib/r600/workitem/get_global_offset.cl libclc/opencl/lib/r600/workitem/get_global_size.cl libclc/opencl/lib/r600/workitem/get_group_id.cl libclc/opencl/lib/r600/workitem/get_local_id.cl libclc/opencl/lib/r600/workitem/get_local_size.cl libclc/opencl/lib/r600/workitem/get_num_groups.cl libclc/opencl/lib/r600/workitem/get_work_dim.cl libclc/opencl/lib/spirv/subnormal_config.cl libclc/utils/gen_convert.py Removed: libclc/opencl/include/clc/opencl/clc.h ################################################################################ diff --git a/libclc/opencl/include/clc/opencl/as_type.h b/libclc/opencl/include/clc/opencl/as_type.h index a68e8a96f252b..abe643df81046 100644 --- a/libclc/opencl/include/clc/opencl/as_type.h +++ b/libclc/opencl/include/clc/opencl/as_type.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_AS_TYPE_H__ +#define __CLC_OPENCL_AS_TYPE_H__ + +#include <clc/opencl/opencl-base.h> + #define as_char(x) __builtin_astype(x, char) #define as_uchar(x) __builtin_astype(x, uchar) #define as_short(x) __builtin_astype(x, short) @@ -83,3 +88,5 @@ #define as_half8(x) __builtin_astype(x, half8) #define as_half16(x) __builtin_astype(x, half16) #endif + +#endif // __CLC_OPENCL_AS_TYPE_H__ diff --git a/libclc/opencl/include/clc/opencl/async/async_work_group_copy.h b/libclc/opencl/include/clc/opencl/async/async_work_group_copy.h index 14420b3c9eeac..1bd16a411554a 100644 --- a/libclc/opencl/include/clc/opencl/async/async_work_group_copy.h +++ b/libclc/opencl/include/clc/opencl/async/async_work_group_copy.h @@ -6,6 +6,9 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_COPY_H__ +#define __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_COPY_H__ + #define __CLC_DST_ADDR_SPACE local #define __CLC_SRC_ADDR_SPACE global #define __CLC_BODY <clc/opencl/async/async_work_group_copy.inc> @@ -23,3 +26,5 @@ #include <clc/math/gentype.inc> #undef __CLC_DST_ADDR_SPACE #undef __CLC_SRC_ADDR_SPACE + +#endif // __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_COPY_H__ diff --git a/libclc/opencl/include/clc/opencl/async/async_work_group_strided_copy.h b/libclc/opencl/include/clc/opencl/async/async_work_group_strided_copy.h index fd37e35f18c5b..5929ce73a9317 100644 --- a/libclc/opencl/include/clc/opencl/async/async_work_group_strided_copy.h +++ b/libclc/opencl/include/clc/opencl/async/async_work_group_strided_copy.h @@ -6,6 +6,9 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_STRIDED_COPY_H__ +#define __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_STRIDED_COPY_H__ + #define __CLC_DST_ADDR_SPACE local #define __CLC_SRC_ADDR_SPACE global #define __CLC_BODY <clc/opencl/async/async_work_group_strided_copy.inc> @@ -23,3 +26,5 @@ #include <clc/math/gentype.inc> #undef __CLC_DST_ADDR_SPACE #undef __CLC_SRC_ADDR_SPACE + +#endif // __CLC_OPENCL_ASYNC_ASYNC_WORK_GROUP_STRIDED_COPY_H__ diff --git a/libclc/opencl/include/clc/opencl/async/prefetch.h b/libclc/opencl/include/clc/opencl/async/prefetch.h index 5075dcaee1409..1ba4414d9eb10 100644 --- a/libclc/opencl/include/clc/opencl/async/prefetch.h +++ b/libclc/opencl/include/clc/opencl/async/prefetch.h @@ -6,8 +6,13 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ASYNC_PREFETCH_H__ +#define __CLC_OPENCL_ASYNC_PREFETCH_H__ + #define __CLC_BODY <clc/opencl/async/prefetch.inc> #include <clc/integer/gentype.inc> #define __CLC_BODY <clc/opencl/async/prefetch.inc> #include <clc/math/gentype.inc> + +#endif // __CLC_OPENCL_ASYNC_PREFETCH_H__ diff --git a/libclc/opencl/include/clc/opencl/async/wait_group_events.h b/libclc/opencl/include/clc/opencl/async/wait_group_events.h index 6aab22e5406ea..d3ace9c34a25d 100644 --- a/libclc/opencl/include/clc/opencl/async/wait_group_events.h +++ b/libclc/opencl/include/clc/opencl/async/wait_group_events.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ASYNC_WAIT_GROUP_EVENTS_H__ +#define __CLC_OPENCL_ASYNC_WAIT_GROUP_EVENTS_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD void wait_group_events(int num_events, event_t *event_list); + +#endif // __CLC_OPENCL_ASYNC_WAIT_GROUP_EVENTS_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_add.h b/libclc/opencl/include/clc/opencl/atomic/atom_add.h index 8b3b8372ca9ff..7f68fca0fc3a5 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_add.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_add.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_ADD_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_ADD_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_base_atomics #define FUNCTION atom_add #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_add #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_ADD_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_and.h b/libclc/opencl/include/clc/opencl/atomic/atom_and.h index 8eda27bd0732f..5c40aa4db70ad 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_and.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_and.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_AND_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_AND_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_extended_atomics #define FUNCTION atom_and #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_and #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_extended_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_AND_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_cmpxchg.h b/libclc/opencl/include/clc/opencl/atomic/atom_cmpxchg.h index 92d922b0d2766..83ee99f414241 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_cmpxchg.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_cmpxchg.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_CMPXCHG_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_CMPXCHG_H__ + +#include <clc/opencl/opencl-base.h> + #include <clc/clcfunc.h> #include <clc/clctypes.h> @@ -37,3 +42,5 @@ _CLC_OVERLOAD _CLC_DECL unsigned long atom_cmpxchg(volatile local unsigned long *p, unsigned long cmp, unsigned long val); #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_CMPXCHG_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_dec.h b/libclc/opencl/include/clc/opencl/atomic/atom_dec.h index afb829bb83159..f57d9c6afca5d 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_dec.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_dec.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_DEC_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_DEC_H__ + +#include <clc/opencl/opencl-base.h> + #include <clc/clcfunc.h> #include <clc/clctypes.h> @@ -26,3 +31,5 @@ atom_dec(volatile global unsigned long *p); _CLC_OVERLOAD _CLC_DECL long atom_dec(volatile local long *p); _CLC_OVERLOAD _CLC_DECL unsigned long atom_dec(volatile local unsigned long *p); #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_DEC_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_inc.h b/libclc/opencl/include/clc/opencl/atomic/atom_inc.h index 227538b153ab5..96d3faa593fc9 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_inc.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_inc.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_INC_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_INC_H__ + +#include <clc/opencl/opencl-base.h> + #include <clc/clcfunc.h> #include <clc/clctypes.h> @@ -26,3 +31,5 @@ atom_inc(volatile global unsigned long *p); _CLC_OVERLOAD _CLC_DECL long atom_inc(volatile local long *p); _CLC_OVERLOAD _CLC_DECL unsigned long atom_inc(volatile local unsigned long *p); #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_INC_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_max.h b/libclc/opencl/include/clc/opencl/atomic/atom_max.h index 778085138f795..87c70f20058d9 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_max.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_max.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_MAX_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_MAX_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_extended_atomics #define FUNCTION atom_max #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_max #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_extended_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_MAX_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_min.h b/libclc/opencl/include/clc/opencl/atomic/atom_min.h index d3169397b89fe..487364c617ab5 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_min.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_min.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_MIN_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_MIN_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_extended_atomics #define FUNCTION atom_min #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_min #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_extended_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_MIN_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_or.h b/libclc/opencl/include/clc/opencl/atomic/atom_or.h index 9a058c813b999..1c3ae6984eead 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_or.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_or.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_OR_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_OR_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_extended_atomics #define FUNCTION atom_or #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_or #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_extended_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_OR_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_sub.h b/libclc/opencl/include/clc/opencl/atomic/atom_sub.h index 7cbcce5c8bf51..eb17aa8ce73fe 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_sub.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_sub.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_SUB_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_SUB_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_base_atomics #define FUNCTION atom_sub #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_sub #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_SUB_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_xchg.h b/libclc/opencl/include/clc/opencl/atomic/atom_xchg.h index 75d9a391d1c7b..df17d31d6e67a 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_xchg.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_xchg.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_XCHG_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_XCHG_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_base_atomics #define FUNCTION atom_xchg #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_xchg #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_base_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_XCHG_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atom_xor.h b/libclc/opencl/include/clc/opencl/atomic/atom_xor.h index b01b1fba23787..8a97818821582 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atom_xor.h +++ b/libclc/opencl/include/clc/opencl/atomic/atom_xor.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOM_XOR_H__ +#define __CLC_OPENCL_ATOMIC_ATOM_XOR_H__ + +#include <clc/opencl/opencl-base.h> + #ifdef cl_khr_global_int32_extended_atomics #define FUNCTION atom_xor #define __CLC_ADDRESS_SPACE global @@ -22,3 +27,5 @@ #define FUNCTION atom_xor #include <clc/opencl/atomic/atom_decl_int64.inc> #endif // cl_khr_int64_extended_atomics + +#endif // __CLC_OPENCL_ATOMIC_ATOM_XOR_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_add.h b/libclc/opencl/include/clc/opencl/atomic/atomic_add.h index 74b3d8ffd8c35..821ae7aab05bf 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_add.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_add.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_ADD_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_ADD_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_add #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_ADD_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_and.h b/libclc/opencl/include/clc/opencl/atomic/atomic_and.h index aec521e07bc20..d10cfed9b581a 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_and.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_and.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_AND_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_AND_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_and #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_AND_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_cmpxchg.h b/libclc/opencl/include/clc/opencl/atomic/atomic_cmpxchg.h index fc7975910cb24..f784984299344 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_cmpxchg.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_cmpxchg.h @@ -6,7 +6,14 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_CMPXCHG_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_CMPXCHG_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_OVERLOAD _CLC_DECL int atomic_cmpxchg(volatile local int *, int, int); _CLC_OVERLOAD _CLC_DECL int atomic_cmpxchg(volatile global int *, int, int); _CLC_OVERLOAD _CLC_DECL uint atomic_cmpxchg(volatile local uint *, uint, uint); _CLC_OVERLOAD _CLC_DECL uint atomic_cmpxchg(volatile global uint *, uint, uint); + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_CMPXCHG_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_dec.h b/libclc/opencl/include/clc/opencl/atomic/atomic_dec.h index 354a721003f56..253a64f491fb4 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_dec.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_dec.h @@ -6,7 +6,14 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_DEC_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_DEC_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_OVERLOAD _CLC_DECL int atomic_dec(volatile local int *); _CLC_OVERLOAD _CLC_DECL int atomic_dec(volatile global int *); _CLC_OVERLOAD _CLC_DECL uint atomic_dec(volatile local uint *); _CLC_OVERLOAD _CLC_DECL uint atomic_dec(volatile global uint *); + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_DEC_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_inc.h b/libclc/opencl/include/clc/opencl/atomic/atomic_inc.h index 2b81c23bf0e5e..75b3fe076cf21 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_inc.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_inc.h @@ -6,7 +6,14 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_INC_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_INC_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_OVERLOAD _CLC_DECL int atomic_inc(volatile local int *); _CLC_OVERLOAD _CLC_DECL int atomic_inc(volatile global int *); _CLC_OVERLOAD _CLC_DECL uint atomic_inc(volatile local uint *); _CLC_OVERLOAD _CLC_DECL uint atomic_inc(volatile global uint *); + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_INC_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_max.h b/libclc/opencl/include/clc/opencl/atomic/atomic_max.h index 4f23f3858531c..667fa36f16f9d 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_max.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_max.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_MAX_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_MAX_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_max #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_MAX_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_min.h b/libclc/opencl/include/clc/opencl/atomic/atomic_min.h index 1d99600f6d13f..91bb636eec875 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_min.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_min.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_MIN_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_MIN_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_min #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_MIN_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_or.h b/libclc/opencl/include/clc/opencl/atomic/atomic_or.h index 7cb920df99760..5c03fd157a2bc 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_or.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_or.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_OR_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_OR_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_or #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_OR_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_sub.h b/libclc/opencl/include/clc/opencl/atomic/atomic_sub.h index f12e77523770b..25ffe9ff4a9b7 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_sub.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_sub.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_SUB_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_SUB_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_sub #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_SUB_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_xchg.h b/libclc/opencl/include/clc/opencl/atomic/atomic_xchg.h index ef5ad542a2d45..6b4206dedb820 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_xchg.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_xchg.h @@ -6,8 +6,15 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_XCHG_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_XCHG_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_xchg _CLC_OVERLOAD _CLC_DECL float FUNCTION(volatile local float *, float); _CLC_OVERLOAD _CLC_DECL float FUNCTION(volatile global float *, float); #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_XCHG_H__ diff --git a/libclc/opencl/include/clc/opencl/atomic/atomic_xor.h b/libclc/opencl/include/clc/opencl/atomic/atomic_xor.h index fb6498d9d0b32..e94560cb6b9ed 100644 --- a/libclc/opencl/include/clc/opencl/atomic/atomic_xor.h +++ b/libclc/opencl/include/clc/opencl/atomic/atomic_xor.h @@ -6,5 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_ATOMIC_ATOMIC_XOR_H__ +#define __CLC_OPENCL_ATOMIC_ATOMIC_XOR_H__ + +#include <clc/opencl/opencl-base.h> + #define FUNCTION atomic_xor #include <clc/opencl/atomic/atomic_decl.inc> + +#endif // __CLC_OPENCL_ATOMIC_ATOMIC_XOR_H__ diff --git a/libclc/opencl/include/clc/opencl/clc.h b/libclc/opencl/include/clc/opencl/clc.h deleted file mode 100644 index 7d34f00113517..0000000000000 --- a/libclc/opencl/include/clc/opencl/clc.h +++ /dev/null @@ -1,256 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 -// -//===----------------------------------------------------------------------===// - -#ifndef __CLC_OPENCL_CLC_H__ -#define __CLC_OPENCL_CLC_H__ - -#ifndef cl_clang_storage_class_specifiers -#error Implementation requires cl_clang_storage_class_specifiers extension! -#endif - -#pragma OPENCL EXTENSION cl_clang_storage_class_specifiers : enable - -#include <clc/opencl/opencl-base.h> - -/* 6.2.3 Explicit Conversions */ -#include <clc/opencl/convert.h> - -/* 6.2.4.2 Reinterpreting Types Using as_type() and as_typen() */ -#include <clc/opencl/as_type.h> - -/* 6.11.1 Work-Item Functions */ -#include <clc/opencl/workitem/get_global_id.h> -#include <clc/opencl/workitem/get_global_offset.h> -#include <clc/opencl/workitem/get_global_size.h> -#include <clc/opencl/workitem/get_group_id.h> -#include <clc/opencl/workitem/get_local_id.h> -#include <clc/opencl/workitem/get_local_size.h> -#include <clc/opencl/workitem/get_num_groups.h> -#include <clc/opencl/workitem/get_work_dim.h> - -/* 6.11.2 Math Functions */ -#include <clc/opencl/math/acos.h> -#include <clc/opencl/math/acosh.h> -#include <clc/opencl/math/acospi.h> -#include <clc/opencl/math/asin.h> -#include <clc/opencl/math/asinh.h> -#include <clc/opencl/math/asinpi.h> -#include <clc/opencl/math/atan.h> -#include <clc/opencl/math/atan2.h> -#include <clc/opencl/math/atan2pi.h> -#include <clc/opencl/math/atanh.h> -#include <clc/opencl/math/atanpi.h> -#include <clc/opencl/math/cbrt.h> -#include <clc/opencl/math/ceil.h> -#include <clc/opencl/math/copysign.h> -#include <clc/opencl/math/cos.h> -#include <clc/opencl/math/cosh.h> -#include <clc/opencl/math/cospi.h> -#include <clc/opencl/math/erf.h> -#include <clc/opencl/math/erfc.h> -#include <clc/opencl/math/exp.h> -#include <clc/opencl/math/exp10.h> -#include <clc/opencl/math/exp2.h> -#include <clc/opencl/math/expm1.h> -#include <clc/opencl/math/fabs.h> -#include <clc/opencl/math/fdim.h> -#include <clc/opencl/math/floor.h> -#include <clc/opencl/math/fma.h> -#include <clc/opencl/math/fmax.h> -#include <clc/opencl/math/fmin.h> -#include <clc/opencl/math/fmod.h> -#include <clc/opencl/math/fract.h> -#include <clc/opencl/math/frexp.h> -#include <clc/opencl/math/half_cos.h> -#include <clc/opencl/math/half_divide.h> -#include <clc/opencl/math/half_exp.h> -#include <clc/opencl/math/half_exp10.h> -#include <clc/opencl/math/half_exp2.h> -#include <clc/opencl/math/half_log.h> -#include <clc/opencl/math/half_log10.h> -#include <clc/opencl/math/half_log2.h> -#include <clc/opencl/math/half_powr.h> -#include <clc/opencl/math/half_recip.h> -#include <clc/opencl/math/half_rsqrt.h> -#include <clc/opencl/math/half_sin.h> -#include <clc/opencl/math/half_sqrt.h> -#include <clc/opencl/math/half_tan.h> -#include <clc/opencl/math/hypot.h> -#include <clc/opencl/math/ilogb.h> -#include <clc/opencl/math/ldexp.h> -#include <clc/opencl/math/lgamma.h> -#include <clc/opencl/math/lgamma_r.h> -#include <clc/opencl/math/log.h> -#include <clc/opencl/math/log10.h> -#include <clc/opencl/math/log1p.h> -#include <clc/opencl/math/log2.h> -#include <clc/opencl/math/logb.h> -#include <clc/opencl/math/mad.h> -#include <clc/opencl/math/maxmag.h> -#include <clc/opencl/math/minmag.h> -#include <clc/opencl/math/modf.h> -#include <clc/opencl/math/nan.h> -#include <clc/opencl/math/native_cos.h> -#include <clc/opencl/math/native_divide.h> -#include <clc/opencl/math/native_exp.h> -#include <clc/opencl/math/native_exp10.h> -#include <clc/opencl/math/native_exp2.h> -#include <clc/opencl/math/native_log.h> -#include <clc/opencl/math/native_log10.h> -#include <clc/opencl/math/native_log2.h> -#include <clc/opencl/math/native_powr.h> -#include <clc/opencl/math/native_recip.h> -#include <clc/opencl/math/native_rsqrt.h> -#include <clc/opencl/math/native_sin.h> -#include <clc/opencl/math/native_sqrt.h> -#include <clc/opencl/math/native_tan.h> -#include <clc/opencl/math/nextafter.h> -#include <clc/opencl/math/pow.h> -#include <clc/opencl/math/pown.h> -#include <clc/opencl/math/powr.h> -#include <clc/opencl/math/remainder.h> -#include <clc/opencl/math/remquo.h> -#include <clc/opencl/math/rint.h> -#include <clc/opencl/math/rootn.h> -#include <clc/opencl/math/round.h> -#include <clc/opencl/math/rsqrt.h> -#include <clc/opencl/math/sin.h> -#include <clc/opencl/math/sincos.h> -#include <clc/opencl/math/sinh.h> -#include <clc/opencl/math/sinpi.h> -#include <clc/opencl/math/sqrt.h> -#include <clc/opencl/math/tan.h> -#include <clc/opencl/math/tanh.h> -#include <clc/opencl/math/tanpi.h> -#include <clc/opencl/math/tgamma.h> -#include <clc/opencl/math/trunc.h> - -/* 6.11.2.1 Floating-point macros */ -#include <clc/float/definitions.h> - -/* 6.11.3 Integer Functions */ -#include <clc/opencl/integer/abs.h> -#include <clc/opencl/integer/abs_ diff .h> -#include <clc/opencl/integer/add_sat.h> -#include <clc/opencl/integer/clz.h> -#include <clc/opencl/integer/ctz.h> -#include <clc/opencl/integer/hadd.h> -#include <clc/opencl/integer/mad24.h> -#include <clc/opencl/integer/mad_hi.h> -#include <clc/opencl/integer/mad_sat.h> -#include <clc/opencl/integer/mul24.h> -#include <clc/opencl/integer/mul_hi.h> -#include <clc/opencl/integer/popcount.h> -#include <clc/opencl/integer/rhadd.h> -#include <clc/opencl/integer/rotate.h> -#include <clc/opencl/integer/sub_sat.h> -#include <clc/opencl/integer/upsample.h> - -/* 6.11.3 Integer Definitions */ -#include <clc/integer/definitions.h> - -/* 6.11.2 and 6.11.3 Shared Integer/Math Functions */ -#include <clc/opencl/shared/clamp.h> -#include <clc/opencl/shared/max.h> -#include <clc/opencl/shared/min.h> -#include <clc/opencl/shared/vload.h> -#include <clc/opencl/shared/vstore.h> - -/* 6.11.4 Common Functions */ -#include <clc/opencl/common/degrees.h> -#include <clc/opencl/common/mix.h> -#include <clc/opencl/common/radians.h> -#include <clc/opencl/common/sign.h> -#include <clc/opencl/common/smoothstep.h> -#include <clc/opencl/common/step.h> - -/* 6.11.5 Geometric Functions */ -#include <clc/opencl/geometric/cross.h> -#include <clc/opencl/geometric/distance.h> -#include <clc/opencl/geometric/dot.h> -#include <clc/opencl/geometric/fast_distance.h> -#include <clc/opencl/geometric/fast_length.h> -#include <clc/opencl/geometric/fast_normalize.h> -#include <clc/opencl/geometric/length.h> -#include <clc/opencl/geometric/normalize.h> - -/* 6.11.6 Relational Functions */ -#include <clc/opencl/relational/all.h> -#include <clc/opencl/relational/any.h> -#include <clc/opencl/relational/bitselect.h> -#include <clc/opencl/relational/isequal.h> -#include <clc/opencl/relational/isfinite.h> -#include <clc/opencl/relational/isgreater.h> -#include <clc/opencl/relational/isgreaterequal.h> -#include <clc/opencl/relational/isinf.h> -#include <clc/opencl/relational/isless.h> -#include <clc/opencl/relational/islessequal.h> -#include <clc/opencl/relational/islessgreater.h> -#include <clc/opencl/relational/isnan.h> -#include <clc/opencl/relational/isnormal.h> -#include <clc/opencl/relational/isnotequal.h> -#include <clc/opencl/relational/isordered.h> -#include <clc/opencl/relational/isunordered.h> -#include <clc/opencl/relational/select.h> -#include <clc/opencl/relational/signbit.h> - -#include <clc/opencl/synchronization/cl_mem_fence_flags.h> - -/* 6.11.8 Synchronization Functions */ -#include <clc/opencl/synchronization/barrier.h> - -/* 6.11.9 Explicit Memory Fence Functions */ -#include <clc/opencl/explicit_fence/explicit_memory_fence.h> - -/* 6.11.10 Async Copy and Prefetch Functions */ -#include <clc/opencl/async/async_work_group_copy.h> -#include <clc/opencl/async/async_work_group_strided_copy.h> -#include <clc/opencl/async/prefetch.h> -#include <clc/opencl/async/wait_group_events.h> - -/* 6.11.11 Atomic Functions */ -#include <clc/opencl/atomic/atomic_add.h> -#include <clc/opencl/atomic/atomic_and.h> -#include <clc/opencl/atomic/atomic_cmpxchg.h> -#include <clc/opencl/atomic/atomic_dec.h> -#include <clc/opencl/atomic/atomic_inc.h> -#include <clc/opencl/atomic/atomic_max.h> -#include <clc/opencl/atomic/atomic_min.h> -#include <clc/opencl/atomic/atomic_or.h> -#include <clc/opencl/atomic/atomic_sub.h> -#include <clc/opencl/atomic/atomic_xchg.h> -#include <clc/opencl/atomic/atomic_xor.h> - -/* cl_khr_global_int32_base_atomics, cl_khr_local_int32_base_atomics and - * cl_khr_int64_base_atomics Extension Functions */ -#include <clc/opencl/atomic/atom_add.h> -#include <clc/opencl/atomic/atom_cmpxchg.h> -#include <clc/opencl/atomic/atom_dec.h> -#include <clc/opencl/atomic/atom_inc.h> -#include <clc/opencl/atomic/atom_sub.h> -#include <clc/opencl/atomic/atom_xchg.h> - -/* cl_khr_global_int32_extended_atomics, cl_khr_local_int32_extended_atomics and - * cl_khr_int64_extended_atomics Extension Functions */ -#include <clc/opencl/atomic/atom_and.h> -#include <clc/opencl/atomic/atom_max.h> -#include <clc/opencl/atomic/atom_min.h> -#include <clc/opencl/atomic/atom_or.h> -#include <clc/opencl/atomic/atom_xor.h> - -/* 6.12.12 Miscellaneous Vector Functions */ -#include <clc/opencl/misc/shuffle.h> -#include <clc/opencl/misc/shuffle2.h> - -/* 6.11.13 Image Read and Write Functions */ -#include <clc/opencl/image/image.h> -#include <clc/opencl/image/image_defines.h> - -#pragma OPENCL EXTENSION all : disable - -#endif // __CLC_OPENCL_CLC_H__ diff --git a/libclc/opencl/include/clc/opencl/convert.h b/libclc/opencl/include/clc/opencl/convert.h index cdd7cc7b1503c..1f71f673619b9 100644 --- a/libclc/opencl/include/clc/opencl/convert.h +++ b/libclc/opencl/include/clc/opencl/convert.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_CONVERT_H__ +#define __CLC_OPENCL_CONVERT_H__ + +#include <clc/opencl/opencl-base.h> + #define _CLC_CONVERT_DECL(FROM_TYPE, TO_TYPE, SUFFIX) \ _CLC_OVERLOAD _CLC_DECL TO_TYPE convert_##TO_TYPE##SUFFIX(FROM_TYPE x); @@ -99,3 +104,5 @@ _CLC_VECTOR_CONVERT_TO_SUFFIX() #undef _CLC_VECTOR_CONVERT_FROM1 #undef _CLC_VECTOR_CONVERT_DECL #undef _CLC_CONVERT_DECL + +#endif // __CLC_OPENCL_CONVERT_H__ diff --git a/libclc/opencl/include/clc/opencl/explicit_fence/explicit_memory_fence.h b/libclc/opencl/include/clc/opencl/explicit_fence/explicit_memory_fence.h index 3ad3757d43e7d..3ade970a33c61 100644 --- a/libclc/opencl/include/clc/opencl/explicit_fence/explicit_memory_fence.h +++ b/libclc/opencl/include/clc/opencl/explicit_fence/explicit_memory_fence.h @@ -6,6 +6,14 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_EXPLICIT_FENCE_EXPLICIT_MEMORY_FENCE_H__ +#define __CLC_OPENCL_EXPLICIT_FENCE_EXPLICIT_MEMORY_FENCE_H__ + +#include <clc/opencl/opencl-base.h> +#include <clc/opencl/synchronization/cl_mem_fence_flags.h> + _CLC_DECL _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags); _CLC_DECL _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags); _CLC_DECL _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags); + +#endif // __CLC_OPENCL_EXPLICIT_FENCE_EXPLICIT_MEMORY_FENCE_H__ diff --git a/libclc/opencl/include/clc/opencl/image/image.h b/libclc/opencl/include/clc/opencl/image/image.h index 6435f3430e2fe..7a248304acff5 100644 --- a/libclc/opencl/include/clc/opencl/image/image.h +++ b/libclc/opencl/include/clc/opencl/image/image.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_IMAGE_H__ +#define __CLC_OPENCL_IMAGE_H__ + +#include <clc/opencl/opencl-base.h> + #if defined(__opencl_c_images) _CLC_OVERLOAD _CLC_DECL int get_image_width(image2d_t image); @@ -46,3 +51,5 @@ _CLC_OVERLOAD _CLC_DECL uint4 read_imageui(image2d_t image, sampler_t sampler, float2 coord); #endif + +#endif // __CLC_OPENCL_IMAGE_H__ diff --git a/libclc/opencl/include/clc/opencl/image/image_defines.h b/libclc/opencl/include/clc/opencl/image/image_defines.h index f3589fdc4a89a..75e08dec3f044 100644 --- a/libclc/opencl/include/clc/opencl/image/image_defines.h +++ b/libclc/opencl/include/clc/opencl/image/image_defines.h @@ -6,6 +6,9 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_IMAGE_IMAGE_DEFINES_H__ +#define __CLC_OPENCL_IMAGE_IMAGE_DEFINES_H__ + /* get_image_channel_data_type flags */ #define CLK_SNORM_INT8 0x10D0 #define CLK_SNORM_INT16 0x10D1 @@ -55,3 +58,5 @@ #define CLK_FILTER_NEAREST 0x0000 #define CLK_FILTER_LINEAR 0x0010 #define __CLC_FILTER_MASK 0x0010 + +#endif // __CLC_OPENCL_IMAGE_IMAGE_DEFINES_H__ diff --git a/libclc/opencl/include/clc/opencl/misc/shuffle.h b/libclc/opencl/include/clc/opencl/misc/shuffle.h index 3c3be84a5baf3..03c2718b0c4ff 100644 --- a/libclc/opencl/include/clc/opencl/misc/shuffle.h +++ b/libclc/opencl/include/clc/opencl/misc/shuffle.h @@ -6,6 +6,9 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_MISC_SHUFFLE_H__ +#define __CLC_OPENCL_MISC_SHUFFLE_H__ + #define FUNCTION shuffle // Integer-type decls @@ -17,3 +20,5 @@ #include <clc/math/gentype.inc> #undef FUNCTION + +#endif // __CLC_OPENCL_MISC_SHUFFLE_H__ diff --git a/libclc/opencl/include/clc/opencl/misc/shuffle2.h b/libclc/opencl/include/clc/opencl/misc/shuffle2.h index 865de53029be8..5b6ba99686516 100644 --- a/libclc/opencl/include/clc/opencl/misc/shuffle2.h +++ b/libclc/opencl/include/clc/opencl/misc/shuffle2.h @@ -6,6 +6,9 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_MISC_SHUFFLE2_H__ +#define __CLC_OPENCL_MISC_SHUFFLE2_H__ + #define FUNCTION shuffle2 // Integer-type decls @@ -17,3 +20,5 @@ #include <clc/math/gentype.inc> #undef FUNCTION + +#endif // __CLC_OPENCL_MISC_SHUFFLE2_H__ diff --git a/libclc/opencl/include/clc/opencl/shared/clamp.h b/libclc/opencl/include/clc/opencl/shared/clamp.h index 4fe7897f56016..4bd365e9dd006 100644 --- a/libclc/opencl/include/clc/opencl/shared/clamp.h +++ b/libclc/opencl/include/clc/opencl/shared/clamp.h @@ -6,8 +6,13 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SHARED_CLAMP_H__ +#define __CLC_OPENCL_SHARED_CLAMP_H__ + #define __CLC_BODY <clc/opencl/shared/clamp.inc> #include <clc/integer/gentype.inc> #define __CLC_BODY <clc/opencl/shared/clamp.inc> #include <clc/math/gentype.inc> + +#endif // __CLC_OPENCL_SHARED_CLAMP_H__ diff --git a/libclc/opencl/include/clc/opencl/shared/max.h b/libclc/opencl/include/clc/opencl/shared/max.h index 5ff01ba8082a2..512e696885817 100644 --- a/libclc/opencl/include/clc/opencl/shared/max.h +++ b/libclc/opencl/include/clc/opencl/shared/max.h @@ -6,8 +6,13 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SHARED_MAX_H__ +#define __CLC_OPENCL_SHARED_MAX_H__ + #define __CLC_BODY <clc/opencl/shared/max.inc> #include <clc/integer/gentype.inc> #define __CLC_BODY <clc/opencl/shared/max.inc> #include <clc/math/gentype.inc> + +#endif // __CLC_OPENCL_SHARED_MAX_H__ diff --git a/libclc/opencl/include/clc/opencl/shared/min.h b/libclc/opencl/include/clc/opencl/shared/min.h index 573f837685be1..e20545210dcd2 100644 --- a/libclc/opencl/include/clc/opencl/shared/min.h +++ b/libclc/opencl/include/clc/opencl/shared/min.h @@ -6,8 +6,13 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SHARED_MIN_H__ +#define __CLC_OPENCL_SHARED_MIN_H__ + #define __CLC_BODY <clc/opencl/shared/min.inc> #include <clc/integer/gentype.inc> #define __CLC_BODY <clc/opencl/shared/min.inc> #include <clc/math/gentype.inc> + +#endif // __CLC_OPENCL_SHARED_MIN_H__ diff --git a/libclc/opencl/include/clc/opencl/shared/vload.h b/libclc/opencl/include/clc/opencl/shared/vload.h index 7b4092c0a1c32..2ec087d4f3c18 100644 --- a/libclc/opencl/include/clc/opencl/shared/vload.h +++ b/libclc/opencl/include/clc/opencl/shared/vload.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SHARED_VLOAD_H__ +#define __CLC_OPENCL_SHARED_VLOAD_H__ + +#include <clc/opencl/opencl-base.h> + #define _CLC_VLOAD_DECL(SUFFIX, MEM_TYPE, VEC_TYPE, WIDTH, ADDR_SPACE) \ _CLC_OVERLOAD _CLC_DECL VEC_TYPE vload##SUFFIX##WIDTH( \ size_t offset, const ADDR_SPACE MEM_TYPE *x); @@ -79,3 +84,5 @@ _CLC_VLOAD_DECL(a_half, half, float, , __generic) #undef _CLC_VECTOR_VLOAD_PRIM3 #undef _CLC_VECTOR_VLOAD_PRIM1 #undef _CLC_VECTOR_VLOAD_GENERIC_DECL + +#endif // __CLC_OPENCL_SHARED_VLOAD_H__ diff --git a/libclc/opencl/include/clc/opencl/shared/vstore.h b/libclc/opencl/include/clc/opencl/shared/vstore.h index c1b78b734138f..00eaff41f77fb 100644 --- a/libclc/opencl/include/clc/opencl/shared/vstore.h +++ b/libclc/opencl/include/clc/opencl/shared/vstore.h @@ -6,6 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SHARED_VSTORE_H__ +#define __CLC_OPENCL_SHARED_VSTORE_H__ + +#include <clc/opencl/opencl-base.h> + #define _CLC_VSTORE_DECL(SUFFIX, PRIM_TYPE, VEC_TYPE, WIDTH, ADDR_SPACE, RND) \ _CLC_OVERLOAD _CLC_DECL void vstore##SUFFIX##WIDTH##RND( \ VEC_TYPE vec, size_t offset, ADDR_SPACE PRIM_TYPE *out); @@ -82,3 +87,5 @@ _CLC_VECTOR_VSTORE_PRIM1(half) #undef _CLC_VECTOR_VSTORE_PRIM3 #undef _CLC_VECTOR_VSTORE_PRIM1 #undef _CLC_VECTOR_VSTORE_GENERIC_DECL + +#endif // __CLC_OPENCL_SHARED_VSTORE_H__ diff --git a/libclc/opencl/include/clc/opencl/synchronization/barrier.h b/libclc/opencl/include/clc/opencl/synchronization/barrier.h index 5efe25c3194f7..728ba4ae8a532 100644 --- a/libclc/opencl/include/clc/opencl/synchronization/barrier.h +++ b/libclc/opencl/include/clc/opencl/synchronization/barrier.h @@ -6,4 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SYNCHRONIZATION_BARRIER_H__ +#define __CLC_OPENCL_SYNCHRONIZATION_BARRIER_H__ + +#include <clc/opencl/opencl-base.h> +#include <clc/opencl/synchronization/cl_mem_fence_flags.h> + _CLC_DECL _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags); + +#endif // __CLC_OPENCL_SYNCHRONIZATION_BARRIER_H__ diff --git a/libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h b/libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h index f98b1e2a2bc51..6636515fca47d 100644 --- a/libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h +++ b/libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h @@ -6,7 +6,12 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_SYNCHRONIZATION_CL_MEM_FENCE_FLAGS_H__ +#define __CLC_OPENCL_SYNCHRONIZATION_CL_MEM_FENCE_FLAGS_H__ + typedef uint cl_mem_fence_flags; #define CLK_LOCAL_MEM_FENCE 1 #define CLK_GLOBAL_MEM_FENCE 2 + +#endif // __CLC_OPENCL_SYNCHRONIZATION_CL_MEM_FENCE_FLAGS_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h index c60e9bc8b9b78..98217c0ba6b4d 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__ +#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h index 7f06476048c52..8c820331e6ab2 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__ +#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h index e235d990c79fb..5c7426047d7d0 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__ +#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h index 78b78e8e56922..3d0f51aeb8dc2 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__ +#define __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h index 82b569380d471..31fe9d2943f48 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__ +#define __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h index 9458ba3923f7b..c93a7a27bd7cf 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__ +#define __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h index 3f0d3cb2c4fbd..0c47a5ef462b3 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__ +#define __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim); + +#endif // __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__ diff --git a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h index dc6ae4e9f93bf..c4c3c1b224a9a 100644 --- a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h +++ b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h @@ -6,4 +6,11 @@ // //===----------------------------------------------------------------------===// +#ifndef __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__ +#define __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__ + +#include <clc/opencl/opencl-base.h> + _CLC_DECL _CLC_OVERLOAD uint get_work_dim(void); + +#endif // __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__ diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl index 2d7fea91233aa..8f9b47978c6af 100644 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl +++ b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_size.h> #if __clang_major__ >= 8 #define CONST_AS __constant diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl index 84552e2e08fd7..fd30e6ab7a47c 100644 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl +++ b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_size.h> #if __clang_major__ >= 8 #define CONST_AS __constant diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_num_groups.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_num_groups.cl index 1ed16736376f6..a88dae97c7958 100644 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_num_groups.cl +++ b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_num_groups.cl @@ -6,7 +6,9 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_size.h> +#include <clc/opencl/workitem/get_local_size.h> +#include <clc/opencl/workitem/get_num_groups.h> _CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { size_t global_size = get_global_size(dim); diff --git a/libclc/opencl/lib/amdgcn/mem_fence/fence.cl b/libclc/opencl/lib/amdgcn/mem_fence/fence.cl index 5c5029a63cfc7..88b953005aae6 100644 --- a/libclc/opencl/lib/amdgcn/mem_fence/fence.cl +++ b/libclc/opencl/lib/amdgcn/mem_fence/fence.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/explicit_fence/explicit_memory_fence.h> void __clc_amdgcn_s_waitcnt(unsigned flags); diff --git a/libclc/opencl/lib/amdgcn/synchronization/barrier.cl b/libclc/opencl/lib/amdgcn/synchronization/barrier.cl index f7c41aafa83c9..5203db72f484c 100644 --- a/libclc/opencl/lib/amdgcn/synchronization/barrier.cl +++ b/libclc/opencl/lib/amdgcn/synchronization/barrier.cl @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/explicit_fence/explicit_memory_fence.h> +#include <clc/opencl/synchronization/barrier.h> _CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) { mem_fence(flags); diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl b/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl index a1b3ce4192793..55c09a0a2d4e9 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_offset.h> #if __clang_major__ >= 8 #define CONST_AS __constant diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl index 8f1507765f934..f49d26eb684ac 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_size.h> _CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl b/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl index 446cc63ab759d..d0fae52f6f535 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_group_id.h> _CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl b/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl index cd07d8645cd53..5b8b8f40d095d 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_id.h> _CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl index cec2d358df8ac..aeaa383ce22f4 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_size.h> uint __clc_amdgcn_get_local_size_x(void) __asm("llvm.r600.read.local.size.x"); uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y"); diff --git a/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl b/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl index 0d8f0b8736585..3d602fb821dcd 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_num_groups.h> uint __clc_amdgcn_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x"); uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y"); diff --git a/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl b/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl index 8ca8b0b61ce54..f7af4ee57432f 100644 --- a/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl +++ b/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_work_dim.h> #if __clang_major__ >= 8 #define CONST_AS __constant diff --git a/libclc/opencl/lib/clspv/shared/vstore_half.cl b/libclc/opencl/lib/clspv/shared/vstore_half.cl index 1694ebeebe689..df5f30a711df5 100644 --- a/libclc/opencl/lib/clspv/shared/vstore_half.cl +++ b/libclc/opencl/lib/clspv/shared/vstore_half.cl @@ -6,7 +6,15 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/float/definitions.h> +#include <clc/opencl/as_type.h> +#include <clc/opencl/math/copysign.h> +#include <clc/opencl/math/fabs.h> +#include <clc/opencl/math/nextafter.h> +#include <clc/opencl/relational/isinf.h> +#include <clc/opencl/relational/isnan.h> +#include <clc/opencl/shared/min.h> +#include <clc/opencl/shared/vstore.h> #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable diff --git a/libclc/opencl/lib/clspv/subnormal_config.cl b/libclc/opencl/lib/clspv/subnormal_config.cl index 77d60ab9cbbde..114aabb2e9435 100644 --- a/libclc/opencl/lib/clspv/subnormal_config.cl +++ b/libclc/opencl/lib/clspv/subnormal_config.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/math/clc_subnormal_config.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/opencl-base.h> _CLC_DEF bool __clc_fp16_subnormals_supported() { return false; } diff --git a/libclc/opencl/lib/generic/async/async_work_group_copy.cl b/libclc/opencl/lib/generic/async/async_work_group_copy.cl index a0b9d7d1f52ec..67eadf5e2569a 100644 --- a/libclc/opencl/lib/generic/async/async_work_group_copy.cl +++ b/libclc/opencl/lib/generic/async/async_work_group_copy.cl @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/async/async_work_group_copy.h> +#include <clc/opencl/async/async_work_group_strided_copy.h> #define __CLC_BODY <async_work_group_copy.inc> #include <clc/integer/gentype.inc> diff --git a/libclc/opencl/lib/generic/async/async_work_group_strided_copy.cl b/libclc/opencl/lib/generic/async/async_work_group_strided_copy.cl index ac4d1da3bc215..2c17e2a685e8a 100644 --- a/libclc/opencl/lib/generic/async/async_work_group_strided_copy.cl +++ b/libclc/opencl/lib/generic/async/async_work_group_strided_copy.cl @@ -6,7 +6,9 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/async/async_work_group_strided_copy.h> +#include <clc/opencl/workitem/get_local_id.h> +#include <clc/opencl/workitem/get_local_size.h> #define __CLC_BODY <async_work_group_strided_copy.inc> #include <clc/integer/gentype.inc> diff --git a/libclc/opencl/lib/generic/async/prefetch.cl b/libclc/opencl/lib/generic/async/prefetch.cl index 50abbc5083428..8ab2101205868 100644 --- a/libclc/opencl/lib/generic/async/prefetch.cl +++ b/libclc/opencl/lib/generic/async/prefetch.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/async/clc_prefetch.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/async/prefetch.h> #define __CLC_BODY <prefetch.inc> #include <clc/integer/gentype.inc> diff --git a/libclc/opencl/lib/generic/async/wait_group_events.cl b/libclc/opencl/lib/generic/async/wait_group_events.cl index e353e48cfe09f..00b8ddf3342fb 100644 --- a/libclc/opencl/lib/generic/async/wait_group_events.cl +++ b/libclc/opencl/lib/generic/async/wait_group_events.cl @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/async/wait_group_events.h> +#include <clc/opencl/synchronization/barrier.h> _CLC_DEF _CLC_OVERLOAD void wait_group_events(int num_events, event_t *event_list) { diff --git a/libclc/opencl/lib/generic/atomic/atom_add.cl b/libclc/opencl/lib/generic/atomic/atom_add.cl index 9f60881c2717d..65d7ad595ace6 100644 --- a/libclc/opencl/lib/generic/atomic/atom_add.cl +++ b/libclc/opencl/lib/generic/atomic/atom_add.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_add.h> +#include <clc/opencl/atomic/atomic_add.h> #ifdef cl_khr_global_int32_base_atomics #define __CLC_ATOMIC_OP add diff --git a/libclc/opencl/lib/generic/atomic/atom_and.cl b/libclc/opencl/lib/generic/atomic/atom_and.cl index 2653296d29e4f..371e5f0aa4d0d 100644 --- a/libclc/opencl/lib/generic/atomic/atom_and.cl +++ b/libclc/opencl/lib/generic/atomic/atom_and.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_and.h> +#include <clc/opencl/atomic/atomic_and.h> #ifdef cl_khr_global_int32_extended_atomics #define __CLC_ATOMIC_OP and diff --git a/libclc/opencl/lib/generic/atomic/atom_int32_binary.inc b/libclc/opencl/lib/generic/atomic/atom_int32_binary.inc index bee20d9146420..0afc2190263fa 100644 --- a/libclc/opencl/lib/generic/atomic/atom_int32_binary.inc +++ b/libclc/opencl/lib/generic/atomic/atom_int32_binary.inc @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/opencl-base.h> #include <clc/utils.h> #define __CLC_ATOM_IMPL(AS, TYPE) \ diff --git a/libclc/opencl/lib/generic/atomic/atom_max.cl b/libclc/opencl/lib/generic/atomic/atom_max.cl index 744f39f6cb494..2542191d04f53 100644 --- a/libclc/opencl/lib/generic/atomic/atom_max.cl +++ b/libclc/opencl/lib/generic/atomic/atom_max.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_max.h> +#include <clc/opencl/atomic/atomic_max.h> #ifdef cl_khr_global_int32_extended_atomics #define __CLC_ATOMIC_OP max diff --git a/libclc/opencl/lib/generic/atomic/atom_min.cl b/libclc/opencl/lib/generic/atomic/atom_min.cl index 75e5fb02e1289..4e62804824d82 100644 --- a/libclc/opencl/lib/generic/atomic/atom_min.cl +++ b/libclc/opencl/lib/generic/atomic/atom_min.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_min.h> +#include <clc/opencl/atomic/atomic_min.h> #ifdef cl_khr_global_int32_extended_atomics #define __CLC_ATOMIC_OP min diff --git a/libclc/opencl/lib/generic/atomic/atom_or.cl b/libclc/opencl/lib/generic/atomic/atom_or.cl index 4b1ecac95d878..30ad8e52ca232 100644 --- a/libclc/opencl/lib/generic/atomic/atom_or.cl +++ b/libclc/opencl/lib/generic/atomic/atom_or.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_or.h> +#include <clc/opencl/atomic/atomic_or.h> #ifdef cl_khr_global_int32_extended_atomics #define __CLC_ATOMIC_OP or diff --git a/libclc/opencl/lib/generic/atomic/atom_sub.cl b/libclc/opencl/lib/generic/atomic/atom_sub.cl index 24c7e85ca73a2..f6ee94db7469f 100644 --- a/libclc/opencl/lib/generic/atomic/atom_sub.cl +++ b/libclc/opencl/lib/generic/atomic/atom_sub.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_sub.h> +#include <clc/opencl/atomic/atomic_sub.h> #ifdef cl_khr_global_int32_base_atomics #define __CLC_ATOMIC_OP sub diff --git a/libclc/opencl/lib/generic/atomic/atom_xchg.cl b/libclc/opencl/lib/generic/atomic/atom_xchg.cl index f14b18070deb0..3fa17c950a82e 100644 --- a/libclc/opencl/lib/generic/atomic/atom_xchg.cl +++ b/libclc/opencl/lib/generic/atomic/atom_xchg.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_xchg.h> +#include <clc/opencl/atomic/atomic_xchg.h> #ifdef cl_khr_global_int32_base_atomics #define __CLC_ATOMIC_OP xchg diff --git a/libclc/opencl/lib/generic/atomic/atom_xor.cl b/libclc/opencl/lib/generic/atomic/atom_xor.cl index 9373fae2487de..ac08ef63b1a4b 100644 --- a/libclc/opencl/lib/generic/atomic/atom_xor.cl +++ b/libclc/opencl/lib/generic/atomic/atom_xor.cl @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/opencl/atomic/atom_xor.h> +#include <clc/opencl/atomic/atomic_xor.h> #ifdef cl_khr_global_int32_extended_atomics #define __CLC_ATOMIC_OP xor diff --git a/libclc/opencl/lib/generic/atomic/atomic_add.cl b/libclc/opencl/lib/generic/atomic/atomic_add.cl index cb3935cc4f8a1..a0effced7dc62 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_add.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_add.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_add.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_add(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_and.cl b/libclc/opencl/lib/generic/atomic/atomic_and.cl index f5b6e97246c6b..629e6638d3bc2 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_and.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_and.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_and.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_and(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_cmpxchg.cl b/libclc/opencl/lib/generic/atomic/atomic_cmpxchg.cl index c2227ea5207df..db0495b004c1a 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_cmpxchg.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_cmpxchg.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_cmpxchg.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_cmpxchg(volatile AS TYPE *p, TYPE cmp, \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_dec.cl b/libclc/opencl/lib/generic/atomic/atomic_dec.cl index d67f6382fd07d..6f18cdf13428a 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_dec.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_dec.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_dec.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_dec(volatile AS TYPE *p) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_inc.cl b/libclc/opencl/lib/generic/atomic/atomic_inc.cl index 989c45c0e5d8a..13349e5432e5c 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_inc.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_inc.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_inc.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_inc(volatile AS TYPE *p) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_max.cl b/libclc/opencl/lib/generic/atomic/atomic_max.cl index dbd5aa049c5e7..ae929f0523dc2 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_max.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_max.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_max.h> #define IMPL(TYPE, AS, OP) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_max(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_min.cl b/libclc/opencl/lib/generic/atomic/atomic_min.cl index d43f74c82ffc5..c7ebe71ee8295 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_min.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_min.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_min.h> #define IMPL(TYPE, AS, OP) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_min(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_or.cl b/libclc/opencl/lib/generic/atomic/atomic_or.cl index dac7eb66033b4..45fb865689655 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_or.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_or.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_or.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_or(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_sub.cl b/libclc/opencl/lib/generic/atomic/atomic_sub.cl index 5ebe1f0b13694..74977f15155fb 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_sub.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_sub.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_sub.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_sub(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/atomic/atomic_xchg.cl b/libclc/opencl/lib/generic/atomic/atomic_xchg.cl index b3f631c21aec3..883132caf9fbb 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_xchg.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_xchg.cl @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/as_type.h> +#include <clc/opencl/atomic/atomic_xchg.h> _CLC_OVERLOAD _CLC_DEF float atomic_xchg(volatile global float *p, float val) { return as_float(atomic_xchg((volatile global uint *)p, as_uint(val))); diff --git a/libclc/opencl/lib/generic/atomic/atomic_xor.cl b/libclc/opencl/lib/generic/atomic/atomic_xor.cl index 864e22b774538..246752c258374 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_xor.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_xor.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/atomic/atomic_xor.h> #define IMPL(TYPE, AS) \ _CLC_OVERLOAD _CLC_DEF TYPE atomic_xor(volatile AS TYPE *p, TYPE val) { \ diff --git a/libclc/opencl/lib/generic/misc/shuffle.cl b/libclc/opencl/lib/generic/misc/shuffle.cl index f8c5d16d21946..15295bf7d907a 100644 --- a/libclc/opencl/lib/generic/misc/shuffle.cl +++ b/libclc/opencl/lib/generic/misc/shuffle.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/misc/clc_shuffle.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/misc/shuffle.h> #define FUNCTION shuffle diff --git a/libclc/opencl/lib/generic/misc/shuffle2.cl b/libclc/opencl/lib/generic/misc/shuffle2.cl index 74487d07db3d9..7e4c5f0922682 100644 --- a/libclc/opencl/lib/generic/misc/shuffle2.cl +++ b/libclc/opencl/lib/generic/misc/shuffle2.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/misc/clc_shuffle2.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/misc/shuffle2.h> #define FUNCTION shuffle2 diff --git a/libclc/opencl/lib/generic/shared/clamp.cl b/libclc/opencl/lib/generic/shared/clamp.cl index 4d00f6f8c3a68..6e2bedc0704ff 100644 --- a/libclc/opencl/lib/generic/shared/clamp.cl +++ b/libclc/opencl/lib/generic/shared/clamp.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/shared/clamp.h> #include <clc/shared/clc_clamp.h> #define __CLC_BODY <clamp.inc> diff --git a/libclc/opencl/lib/generic/shared/max.cl b/libclc/opencl/lib/generic/shared/max.cl index 59f0b7e3b4c8e..5cb5ff124bd01 100644 --- a/libclc/opencl/lib/generic/shared/max.cl +++ b/libclc/opencl/lib/generic/shared/max.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/shared/max.h> #include <clc/shared/clc_max.h> #define __CLC_BODY <max.inc> diff --git a/libclc/opencl/lib/generic/shared/min.cl b/libclc/opencl/lib/generic/shared/min.cl index 6c491c5eb1cbf..2a1ded3aa0288 100644 --- a/libclc/opencl/lib/generic/shared/min.cl +++ b/libclc/opencl/lib/generic/shared/min.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/shared/min.h> #include <clc/shared/clc_min.h> #define __CLC_BODY <min.inc> diff --git a/libclc/opencl/lib/generic/shared/vload.cl b/libclc/opencl/lib/generic/shared/vload.cl index ad22839580132..6f1ca06ccb8be 100644 --- a/libclc/opencl/lib/generic/shared/vload.cl +++ b/libclc/opencl/lib/generic/shared/vload.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/shared/vload.h> #include <clc/shared/clc_vload.h> #define __CLC_BODY "vload.inc" diff --git a/libclc/opencl/lib/generic/shared/vstore.cl b/libclc/opencl/lib/generic/shared/vstore.cl index 145658f873dc5..1d62de810ff75 100644 --- a/libclc/opencl/lib/generic/shared/vstore.cl +++ b/libclc/opencl/lib/generic/shared/vstore.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/shared/vstore.h> #include <clc/shared/clc_vstore.h> #define __CLC_BODY "vstore.inc" diff --git a/libclc/opencl/lib/generic/subnormal_config.cl b/libclc/opencl/lib/generic/subnormal_config.cl index a32af9b0cf2e7..aa2e30935e5f0 100644 --- a/libclc/opencl/lib/generic/subnormal_config.cl +++ b/libclc/opencl/lib/generic/subnormal_config.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/math/clc_subnormal_config.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/opencl-base.h> _CLC_DEF bool __clc_fp16_subnormals_supported() { return false; } diff --git a/libclc/opencl/lib/generic/workitem/get_global_id.cl b/libclc/opencl/lib/generic/workitem/get_global_id.cl index 26c3bf528cd4d..5e417bef0d2bd 100644 --- a/libclc/opencl/lib/generic/workitem/get_global_id.cl +++ b/libclc/opencl/lib/generic/workitem/get_global_id.cl @@ -6,7 +6,11 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_id.h> +#include <clc/opencl/workitem/get_global_offset.h> +#include <clc/opencl/workitem/get_group_id.h> +#include <clc/opencl/workitem/get_local_id.h> +#include <clc/opencl/workitem/get_local_size.h> _CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) { return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + diff --git a/libclc/opencl/lib/generic/workitem/get_global_size.cl b/libclc/opencl/lib/generic/workitem/get_global_size.cl index 747115d524885..d62e9e91888ac 100644 --- a/libclc/opencl/lib/generic/workitem/get_global_size.cl +++ b/libclc/opencl/lib/generic/workitem/get_global_size.cl @@ -6,7 +6,9 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_size.h> +#include <clc/opencl/workitem/get_local_size.h> +#include <clc/opencl/workitem/get_num_groups.h> _CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { return get_num_groups(dim) * get_local_size(dim); diff --git a/libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl b/libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl index 284a0517d3308..d24569ecda1bc 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/explicit_fence/explicit_memory_fence.h> _CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) { if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE)) diff --git a/libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl b/libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl index 173d771ccaa39..7c57478795dda 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/synchronization/barrier.h> _CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) { __syncthreads(); diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl index 2f1f7cd4250fe..bcfddbe70cb63 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_id.h> +#include <clc/opencl/workitem/get_group_id.h> +#include <clc/opencl/workitem/get_local_id.h> +#include <clc/opencl/workitem/get_local_size.h> _CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) { return get_group_id(dim) * get_local_size(dim) + get_local_id(dim); diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl index 0dad4c2061fe6..7b3e463fd423a 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_group_id.h> _CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl index 199b4610bdb7b..b5133848e9a2e 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_id.h> _CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl index a93fa0a3c9649..814ac462a2bc2 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_size.h> _CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl index 4c934968df865..e9e4e4d06930a 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl +++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_num_groups.h> _CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl b/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl index 813572ae5d2d7..f28b7d458ca4f 100644 --- a/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl +++ b/libclc/opencl/lib/r600/image/get_image_channel_data_type.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL int __clc_get_image_channel_data_type_2d(image2d_t); _CLC_DECL int __clc_get_image_channel_data_type_3d(image3d_t); diff --git a/libclc/opencl/lib/r600/image/get_image_channel_order.cl b/libclc/opencl/lib/r600/image/get_image_channel_order.cl index 0e5882b155079..da6d783e7f01a 100644 --- a/libclc/opencl/lib/r600/image/get_image_channel_order.cl +++ b/libclc/opencl/lib/r600/image/get_image_channel_order.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL int __clc_get_image_channel_order_2d(image2d_t); _CLC_DECL int __clc_get_image_channel_order_3d(image3d_t); diff --git a/libclc/opencl/lib/r600/image/get_image_depth.cl b/libclc/opencl/lib/r600/image/get_image_depth.cl index 431889b6ca258..1731d85dc4c3d 100644 --- a/libclc/opencl/lib/r600/image/get_image_depth.cl +++ b/libclc/opencl/lib/r600/image/get_image_depth.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL int __clc_get_image_depth_3d(image3d_t); diff --git a/libclc/opencl/lib/r600/image/get_image_dim.cl b/libclc/opencl/lib/r600/image/get_image_dim.cl index 5a123425a2910..6929d12ad1a5a 100644 --- a/libclc/opencl/lib/r600/image/get_image_dim.cl +++ b/libclc/opencl/lib/r600/image/get_image_dim.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_OVERLOAD _CLC_DEF int2 get_image_dim(image2d_t image) { return (int2)(get_image_width(image), get_image_height(image)); diff --git a/libclc/opencl/lib/r600/image/get_image_height.cl b/libclc/opencl/lib/r600/image/get_image_height.cl index 0c498bed8e608..14a21823a826d 100644 --- a/libclc/opencl/lib/r600/image/get_image_height.cl +++ b/libclc/opencl/lib/r600/image/get_image_height.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL int __clc_get_image_height_2d(image2d_t); _CLC_DECL int __clc_get_image_height_3d(image3d_t); diff --git a/libclc/opencl/lib/r600/image/get_image_width.cl b/libclc/opencl/lib/r600/image/get_image_width.cl index c54f66176e764..fa54ec3628f2d 100644 --- a/libclc/opencl/lib/r600/image/get_image_width.cl +++ b/libclc/opencl/lib/r600/image/get_image_width.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL int __clc_get_image_width_2d(image2d_t); _CLC_DECL int __clc_get_image_width_3d(image3d_t); diff --git a/libclc/opencl/lib/r600/image/read_imagef.cl b/libclc/opencl/lib/r600/image/read_imagef.cl index 54983363e33c5..fc597b762e52e 100644 --- a/libclc/opencl/lib/r600/image/read_imagef.cl +++ b/libclc/opencl/lib/r600/image/read_imagef.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); diff --git a/libclc/opencl/lib/r600/image/read_imagei.cl b/libclc/opencl/lib/r600/image/read_imagei.cl index 918a444f90a11..bc54a631eabb5 100644 --- a/libclc/opencl/lib/r600/image/read_imagei.cl +++ b/libclc/opencl/lib/r600/image/read_imagei.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); diff --git a/libclc/opencl/lib/r600/image/read_imageui.cl b/libclc/opencl/lib/r600/image/read_imageui.cl index 8236b542ea0e8..8c266b9f2f854 100644 --- a/libclc/opencl/lib/r600/image/read_imageui.cl +++ b/libclc/opencl/lib/r600/image/read_imageui.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2); diff --git a/libclc/opencl/lib/r600/image/write_imagef.cl b/libclc/opencl/lib/r600/image/write_imagef.cl index d9784f5322c94..0e828acad0823 100644 --- a/libclc/opencl/lib/r600/image/write_imagef.cl +++ b/libclc/opencl/lib/r600/image/write_imagef.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL void __clc_write_imagef_2d(image2d_t image, int2 coord, float4 color); diff --git a/libclc/opencl/lib/r600/image/write_imagei.cl b/libclc/opencl/lib/r600/image/write_imagei.cl index ced65184071b3..88ea35634abcd 100644 --- a/libclc/opencl/lib/r600/image/write_imagei.cl +++ b/libclc/opencl/lib/r600/image/write_imagei.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL void __clc_write_imagei_2d(image2d_t image, int2 coord, int4 color); diff --git a/libclc/opencl/lib/r600/image/write_imageui.cl b/libclc/opencl/lib/r600/image/write_imageui.cl index ef2ee5a835ec7..40f3f33345836 100644 --- a/libclc/opencl/lib/r600/image/write_imageui.cl +++ b/libclc/opencl/lib/r600/image/write_imageui.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/image/image.h> _CLC_DECL void __clc_write_imageui_2d(image2d_t image, int2 coord, uint4 color); diff --git a/libclc/opencl/lib/r600/synchronization/barrier.cl b/libclc/opencl/lib/r600/synchronization/barrier.cl index 83438efa26498..69e87d9497fde 100644 --- a/libclc/opencl/lib/r600/synchronization/barrier.cl +++ b/libclc/opencl/lib/r600/synchronization/barrier.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/explicit_fence/explicit_memory_fence.h> _CLC_DEF void __clc_r600_barrier(void) __asm("llvm.r600.group.barrier"); diff --git a/libclc/opencl/lib/r600/workitem/get_global_offset.cl b/libclc/opencl/lib/r600/workitem/get_global_offset.cl index cc42fe4beef0c..988e482de2391 100644 --- a/libclc/opencl/lib/r600/workitem/get_global_offset.cl +++ b/libclc/opencl/lib/r600/workitem/get_global_offset.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_offset.h> _CLC_DEF _CLC_OVERLOAD uint get_global_offset(uint dim) { __attribute__((address_space(7))) uint *ptr = diff --git a/libclc/opencl/lib/r600/workitem/get_global_size.cl b/libclc/opencl/lib/r600/workitem/get_global_size.cl index 1c93568f80c74..298487f87dc1e 100644 --- a/libclc/opencl/lib/r600/workitem/get_global_size.cl +++ b/libclc/opencl/lib/r600/workitem/get_global_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_global_size.h> uint __clc_r600_get_global_size_x(void) __asm("llvm.r600.read.global.size.x"); uint __clc_r600_get_global_size_y(void) __asm("llvm.r600.read.global.size.y"); diff --git a/libclc/opencl/lib/r600/workitem/get_group_id.cl b/libclc/opencl/lib/r600/workitem/get_group_id.cl index c771e835cda12..196f0f21241e8 100644 --- a/libclc/opencl/lib/r600/workitem/get_group_id.cl +++ b/libclc/opencl/lib/r600/workitem/get_group_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_group_id.h> _CLC_DEF _CLC_OVERLOAD uint get_group_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/r600/workitem/get_local_id.cl b/libclc/opencl/lib/r600/workitem/get_local_id.cl index 9d8e044812f24..1321e1e8362e2 100644 --- a/libclc/opencl/lib/r600/workitem/get_local_id.cl +++ b/libclc/opencl/lib/r600/workitem/get_local_id.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_id.h> _CLC_DEF _CLC_OVERLOAD uint get_local_id(uint dim) { switch (dim) { diff --git a/libclc/opencl/lib/r600/workitem/get_local_size.cl b/libclc/opencl/lib/r600/workitem/get_local_size.cl index 3755cf5e8b0e6..c0e4e15542769 100644 --- a/libclc/opencl/lib/r600/workitem/get_local_size.cl +++ b/libclc/opencl/lib/r600/workitem/get_local_size.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_local_size.h> uint __clc_r600_get_local_size_x(void) __asm("llvm.r600.read.local.size.x"); uint __clc_r600_get_local_size_y(void) __asm("llvm.r600.read.local.size.y"); diff --git a/libclc/opencl/lib/r600/workitem/get_num_groups.cl b/libclc/opencl/lib/r600/workitem/get_num_groups.cl index 088a055aabfce..098a9cb86a6dc 100644 --- a/libclc/opencl/lib/r600/workitem/get_num_groups.cl +++ b/libclc/opencl/lib/r600/workitem/get_num_groups.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_num_groups.h> uint __clc_r600_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x"); uint __clc_r600_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y"); diff --git a/libclc/opencl/lib/r600/workitem/get_work_dim.cl b/libclc/opencl/lib/r600/workitem/get_work_dim.cl index b59dd3af6c334..a7917b864e407 100644 --- a/libclc/opencl/lib/r600/workitem/get_work_dim.cl +++ b/libclc/opencl/lib/r600/workitem/get_work_dim.cl @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/clc.h> +#include <clc/opencl/workitem/get_work_dim.h> _CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) { __attribute__((address_space(7))) uint *ptr = diff --git a/libclc/opencl/lib/spirv/subnormal_config.cl b/libclc/opencl/lib/spirv/subnormal_config.cl index 77d60ab9cbbde..114aabb2e9435 100644 --- a/libclc/opencl/lib/spirv/subnormal_config.cl +++ b/libclc/opencl/lib/spirv/subnormal_config.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include <clc/math/clc_subnormal_config.h> -#include <clc/opencl/clc.h> +#include <clc/opencl/opencl-base.h> _CLC_DEF bool __clc_fp16_subnormals_supported() { return false; } diff --git a/libclc/utils/gen_convert.py b/libclc/utils/gen_convert.py index 02893dbad6a37..737ab9d1a5cd6 100644 --- a/libclc/utils/gen_convert.py +++ b/libclc/utils/gen_convert.py @@ -167,7 +167,7 @@ def conditional_guard(src, dst): nl = "\n" includes = [] if not clc: - includes = ["<clc/opencl/clc.h>"] + includes = ["<clc/opencl/convert.h>"] else: includes = sorted( [ _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
