elars/copnix
1
0 Fork 0
mirror of https://github.com/elwrcl/copnix.git synced 2026-10-31 14:28:02 +03:00
Code Issues Packages Projects Releases Wiki Activity

local drivers

Browse Source
This commit is contained in:
elwrcl committed 2026-10-09 19:06:03 +03:00
1 parent 7e965fe195
commit a71d522e5a
76 files changed
+37344 -18
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

drivers/pocl-gen7/share/pocl/include/_builtin_renames.h
+1168
View File
File diff suppressed because it is too large. Load diff
drivers/pocl-gen7/share/pocl/include/_clang_opencl.h
+770
View File
@@ -0,0 +1,770 @@
/* This file adds declarations missing from opencl-c.h.
Copyright (c) 2011-2017 Pekka Jääskeläinen / TUT
2024 Pekka Jääskeläinen / Intel Finland Oy
Copyright (c) 2017 Michal Babej / Tampere University of Technology
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to
deal in the Software without restriction, including without limitation the
rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
sell copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
IN THE SOFTWARE.
*/
#ifndef _CLANG_OPENCL_H
#define _CLANG_OPENCL_H
#ifndef SKIP_HEADER_INCLUDE
/* Use the declarations shipped with Clang. */
#ifndef _OPENCL_H_
/* Check for _OPENCL_H already here because the kernel compiler loads the
header beforehand, but cannot find the file due to include paths not
set up. */
#include "opencl-c.h"
#endif
#endif
/* We can include this header directly (to patch up opencl-c.h without
including _kernel.h, thus need to define the attribute macros here. */
#ifndef _CL_OVERLOADABLE
#if __has_attribute(__overloadable__)
# define _CL_OVERLOADABLE __attribute__((__overloadable__))
#else
# define _CL_OVERLOADABLE
#endif
#endif
#ifndef _CL_READNONE
#if __has_attribute(__const__)
# define _CL_READNONE __attribute__((__const__))
#else
# define _CL_READNONE
#endif
#endif
#ifndef _CL_CONV
#if __has_attribute(__convergent__)
# define _CL_CONV __attribute__((__convergent__))
#else
# define _CL_CONV
#endif
#endif
/* Some of the geometric builtins are defined only up to 4 vectors, but we
implement them all: */
#ifdef cl_khr_fp16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
half _CL_OVERLOADABLE _CL_READNONE length (half8 p);
half _CL_OVERLOADABLE _CL_READNONE length (half16 p);
half _CL_OVERLOADABLE _CL_READNONE fast_length (half8 p);
half _CL_OVERLOADABLE _CL_READNONE fast_length (half16 p);
half8 _CL_OVERLOADABLE _CL_READNONE normalize (half8 p);
half16 _CL_OVERLOADABLE _CL_READNONE normalize (half16 p);
half8 _CL_OVERLOADABLE _CL_READNONE fast_normalize (half8 p);
half16 _CL_OVERLOADABLE _CL_READNONE fast_normalize (half16 p);
half _CL_OVERLOADABLE _CL_READNONE dot (half8 p0, half8 p1);
half _CL_OVERLOADABLE _CL_READNONE dot (half16 p0, half16 p1);
#endif
float _CL_OVERLOADABLE _CL_READNONE length (float8 p);
float _CL_OVERLOADABLE _CL_READNONE length (float16 p);
float _CL_OVERLOADABLE _CL_READNONE fast_length (float8 p);
float _CL_OVERLOADABLE _CL_READNONE fast_length (float16 p);
float8 _CL_OVERLOADABLE _CL_READNONE normalize (float8 p);
float16 _CL_OVERLOADABLE _CL_READNONE normalize (float16 p);
float8 _CL_OVERLOADABLE _CL_READNONE fast_normalize (float8 p);
float16 _CL_OVERLOADABLE _CL_READNONE fast_normalize (float16 p);
float _CL_OVERLOADABLE _CL_READNONE dot (float8 p0, float8 p1);
float _CL_OVERLOADABLE _CL_READNONE dot (float16 p0, float16 p1);
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
double _CL_OVERLOADABLE _CL_READNONE length (double8 p);
double _CL_OVERLOADABLE _CL_READNONE length (double16 p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double2 p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double3 p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double4 p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double8 p);
double _CL_OVERLOADABLE _CL_READNONE fast_length (double16 p);
double8 _CL_OVERLOADABLE _CL_READNONE normalize (double8 p);
double16 _CL_OVERLOADABLE _CL_READNONE normalize (double16 p);
double8 _CL_OVERLOADABLE _CL_READNONE fast_normalize (double8 p);
double16 _CL_OVERLOADABLE _CL_READNONE fast_normalize (double16 p);
double _CL_OVERLOADABLE _CL_READNONE dot (double8 p0, double8 p1);
double _CL_OVERLOADABLE _CL_READNONE dot (double16 p0, double16 p1);
#endif
#if defined(cl_khr_kernel_clock)
#if defined(__opencl_c_kernel_clock_scope_device)
ulong _CL_OVERLOADABLE _CL_CONV clock_read_device(void);
uint2 _CL_OVERLOADABLE _CL_CONV clock_read_hilo_device(void);
#endif // __opencl_c_kernel_clock_scope_device
#if defined(__opencl_c_kernel_clock_scope_work_group)
ulong _CL_OVERLOADABLE _CL_CONV clock_read_work_group(void);
uint2 _CL_OVERLOADABLE _CL_CONV clock_read_hilo_work_group(void);
#endif // __opencl_c_kernel_clock_scope_work_group
#if defined(__opencl_c_kernel_clock_scope_sub_group)
ulong _CL_OVERLOADABLE _CL_CONV clock_read_sub_group(void);
uint2 _CL_OVERLOADABLE _CL_CONV clock_read_hilo_sub_group(void);
#endif // __opencl_c_kernel_clock_scope_sub_group
#endif // cl_khr_kernel_clock
/**************************************************************************/
#if defined(cl_intel_subgroups_char)
/* The char version is missing from opencl-c.h */
char _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (char, uint sub_group_local_id);
char2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (char2, uint sub_group_local_id);
char3 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (char3, uint sub_group_local_id);
char4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (char4, uint sub_group_local_id);
char8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (char8, uint sub_group_local_id);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (uchar, uint sub_group_local_id);
uchar2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (uchar2, uint sub_group_local_id);
uchar3 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (uchar3, uint sub_group_local_id);
uchar4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (uchar4, uint sub_group_local_id);
uchar8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (uchar8, uint sub_group_local_id);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char, uint);
char2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char2, uint);
char3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char3, uint);
char4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char4, uint);
char8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char8, uint);
char16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (char16, uint);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar, uint);
uchar2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar2, uint);
uchar3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar3, uint);
uchar4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar4, uint);
uchar8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar8, uint);
uchar16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (uchar16, uint);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char cur,
char next,
uint);
char2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char2 cur,
char2 next,
uint);
char3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char3 cur,
char3 next,
uint);
char4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char4 cur,
char4 next,
uint);
char8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char8 cur,
char8 next,
uint);
char16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (char16 cur,
char16 next,
uint);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (uchar cur,
uchar next,
uint);
uchar2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (uchar2 cur,
uchar2 next,
uint);
uchar3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (uchar3 cur,
uchar3 next,
uint);
uchar4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (uchar4 cur,
uchar4 next,
uint);
uchar8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (uchar8 cur,
uchar8 next,
uint);
uchar16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_shuffle_down (uchar16 cur, uchar16 next, uint);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char cur,
char next,
uint);
char2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char2 cur,
char2 next,
uint);
char3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char3 cur,
char3 next,
uint);
char4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char4 cur,
char4 next,
uint);
char8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char8 cur,
char8 next,
uint);
char16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (char16 cur,
char16 next,
uint);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar cur,
uchar next,
uint);
uchar2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar2 cur,
uchar2 next,
uint);
uchar3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar3 cur,
uchar3 next,
uint);
uchar4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar4 cur,
uchar4 next,
uint);
uchar8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar8 cur,
uchar8 next,
uint);
uchar16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (uchar16 cur,
uchar16 next,
uint);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char, uint);
char2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char2, uint);
char3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char3, uint);
char4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char4, uint);
char8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char8, uint);
char16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (char16,
uint);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar, uint);
uchar2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar2,
uint);
uchar3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar3,
uint);
uchar4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar4,
uint);
uchar8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar8,
uint);
uchar16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (uchar16,
uint);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_add (char x);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_add (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_min (char x);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_min (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_max (char x);
uchar _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_max (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_add (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_add (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_min (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_min (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_max (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_max (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_add (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_add (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_min (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_min (uchar x);
char _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_max (char x);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_max (uchar x);
#if defined(__opencl_c_images)
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc (read_only image2d_t, int2);
uchar2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc2 (read_only image2d_t, int2);
uchar4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc4 (read_only image2d_t, int2);
uchar8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc8 (read_only image2d_t, int2);
uchar16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc16 (read_only image2d_t, int2);
#if defined(__opencl_c_read_write_images)
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc (read_write image2d_t, int2);
uchar2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc2 (read_write image2d_t, int2);
uchar4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc4 (read_write image2d_t, int2);
uchar8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc8 (read_write image2d_t, int2);
uchar16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc16 (read_write image2d_t, int2);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc (const global uchar *p);
uchar2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc2 (const global uchar *p);
uchar4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc4 (const global uchar *p);
uchar8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc8 (const global uchar *p);
uchar16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc16 (const global uchar *p);
uchar _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc (const local uchar *p);
uchar2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc2 (const local uchar *p);
uchar4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc4 (const local uchar *p);
uchar8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc8 (const local uchar *p);
uchar16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_uc16 (const local uchar *p);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc (global uchar *p,
uchar data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc2 (global uchar *p,
uchar2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc4 (global uchar *p,
uchar4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc8 (global uchar *p,
uchar8 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc16 (global uchar *p,
uchar16 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc (local uchar *p,
uchar data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc2 (local uchar *p,
uchar2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc4 (local uchar *p,
uchar4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc8 (local uchar *p,
uchar8 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_uc16 (local uchar *p,
uchar16 data);
#if defined(__opencl_c_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc (write_only image2d_t, int2, uchar);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc2 (write_only image2d_t, int2, uchar2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc4 (write_only image2d_t, int2, uchar4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc8 (write_only image2d_t, int2, uchar8);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc16 (write_only image2d_t, int2, uchar16);
#if defined(__opencl_c_read_write_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc (read_write image2d_t, int2, uchar);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc2 (read_write image2d_t, int2, uchar2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc4 (read_write image2d_t, int2, uchar4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc8 (read_write image2d_t, int2, uchar8);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_uc16 (read_write image2d_t, int2, uchar16);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
#endif // cl_intel_subgroups_char
/**************************************************************************/
#if defined(cl_intel_subgroups_short)
/* The short version is missing from opencl-c.h */
short _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (short, uint sub_group_local_id);
short2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (short2, uint sub_group_local_id);
short3 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (short3, uint sub_group_local_id);
short4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (short4, uint sub_group_local_id);
short8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (short8, uint sub_group_local_id);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (ushort, uint sub_group_local_id);
ushort2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (ushort2, uint sub_group_local_id);
ushort3 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (ushort3, uint sub_group_local_id);
ushort4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (ushort4, uint sub_group_local_id);
ushort8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_broadcast (ushort8, uint sub_group_local_id);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short, uint);
short2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short2, uint);
short3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short3, uint);
short4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short4, uint);
short8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short8, uint);
short16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (short16, uint);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort, uint);
ushort2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort2, uint);
ushort3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort3, uint);
ushort4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort4, uint);
ushort8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort8, uint);
ushort16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle (ushort16, uint);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short cur,
short next,
uint);
short2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short2 cur,
short2 next,
uint);
short3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short3 cur,
short3 next,
uint);
short4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short4 cur,
short4 next,
uint);
short8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short8 cur,
short8 next,
uint);
short16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (short16 cur,
short16 next,
uint);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (ushort cur,
ushort next,
uint);
ushort2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (ushort2 cur,
ushort2 next,
uint);
ushort3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (ushort3 cur,
ushort3 next,
uint);
ushort4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (ushort4 cur,
ushort4 next,
uint);
ushort8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_down (ushort8 cur,
ushort8 next,
uint);
ushort16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_shuffle_down (ushort16 cur, ushort16 next, uint);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short cur,
short next,
uint);
short2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short2 cur,
short2 next,
uint);
short3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short3 cur,
short3 next,
uint);
short4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short4 cur,
short4 next,
uint);
short8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short8 cur,
short8 next,
uint);
short16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (short16 cur,
short16 next,
uint);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort cur,
ushort next,
uint);
ushort2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort2 cur,
ushort2 next,
uint);
ushort3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort3 cur,
ushort3 next,
uint);
ushort4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort4 cur,
ushort4 next,
uint);
ushort8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort8 cur,
ushort8 next,
uint);
ushort16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_up (ushort16 cur,
ushort16 next,
uint);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short, uint);
short2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short2, uint);
short3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short3, uint);
short4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short4, uint);
short8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short8, uint);
short16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (short16,
uint);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort, uint);
ushort2 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort2,
uint);
ushort3 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort3,
uint);
ushort4 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort4,
uint);
ushort8 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort8,
uint);
ushort16 _CL_OVERLOADABLE _CL_CONV intel_sub_group_shuffle_xor (ushort16,
uint);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_add (short x);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_add (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_min (short x);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_min (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_max (short x);
ushort _CL_OVERLOADABLE _CL_CONV intel_sub_group_reduce_max (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_add (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_add (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_min (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_min (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_exclusive_max (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_exclusive_max (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_add (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_add (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_min (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_min (ushort x);
short _CL_OVERLOADABLE _CL_CONV intel_sub_group_scan_inclusive_max (short x);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_scan_inclusive_max (ushort x);
#if defined(__opencl_c_images)
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us (read_only image2d_t, int2);
ushort2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us2 (read_only image2d_t, int2);
ushort4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us4 (read_only image2d_t, int2);
ushort8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us8 (read_only image2d_t, int2);
ushort16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us16 (read_only image2d_t, int2);
#if defined(__opencl_c_read_write_images)
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us (read_write image2d_t, int2);
ushort2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us2 (read_write image2d_t, int2);
ushort4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us4 (read_write image2d_t, int2);
ushort8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us8 (read_write image2d_t, int2);
ushort16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us16 (read_write image2d_t, int2);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us (const global ushort *p);
ushort2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us2 (const global ushort *p);
ushort4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us4 (const global ushort *p);
ushort8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us8 (const global ushort *p);
ushort16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us16 (const global ushort *p);
ushort _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us (const local ushort *p);
ushort2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us2 (const local ushort *p);
ushort4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us4 (const local ushort *p);
ushort8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us8 (const local ushort *p);
ushort16 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_us16 (const local ushort *p);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us (global ushort *p,
ushort data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us2 (global ushort *p,
ushort2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us4 (global ushort *p,
ushort4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us8 (global ushort *p,
ushort8 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us16 (global ushort *p,
ushort16 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us (local ushort *p,
ushort data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us2 (local ushort *p,
ushort2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us4 (local ushort *p,
ushort4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us8 (local ushort *p,
ushort8 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_us16 (local ushort *p,
ushort16 data);
#if defined(__opencl_c_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us (write_only image2d_t, int2, ushort);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us2 (write_only image2d_t, int2, ushort2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us4 (write_only image2d_t, int2, ushort4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us8 (write_only image2d_t, int2, ushort8);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us16 (write_only image2d_t, int2, ushort16);
#if defined(__opencl_c_read_write_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us (read_write image2d_t, int2, ushort);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us2 (read_write image2d_t, int2, ushort2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us4 (read_write image2d_t, int2, ushort4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us8 (read_write image2d_t, int2, ushort8);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_us16 (read_write image2d_t, int2, ushort16);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
#endif // cl_intel_subgroups_short
/**************************************************************************/
#if defined(cl_intel_subgroups_long)
#if defined(__opencl_c_images)
ulong _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul (read_only image2d_t, int2);
ulong2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul2 (read_only image2d_t, int2);
ulong4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul4 (read_only image2d_t, int2);
ulong8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul8 (read_only image2d_t, int2);
#if defined(__opencl_c_read_write_images)
ulong _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul (read_write image2d_t, int2);
ulong2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul2 (read_write image2d_t, int2);
ulong4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul4 (read_write image2d_t, int2);
ulong8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul8 (read_write image2d_t, int2);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
ulong _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul (const global ulong *p);
ulong2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul2 (const global ulong *p);
ulong4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul4 (const global ulong *p);
ulong8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul8 (const global ulong *p);
ulong _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul (const local ulong *p);
ulong2 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul2 (const local ulong *p);
ulong4 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul4 (const local ulong *p);
ulong8 _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_read_ul8 (const local ulong *p);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul (global ulong *p,
ulong data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul2 (global ulong *p,
ulong2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul4 (global ulong *p,
ulong4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul8 (global ulong *p,
ulong8 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul (local ulong *p,
ulong data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul2 (local ulong *p,
ulong2 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul4 (local ulong *p,
ulong4 data);
void _CL_OVERLOADABLE _CL_CONV intel_sub_group_block_write_ul8 (local ulong *p,
ulong8 data);
#if defined(__opencl_c_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul (write_only image2d_t, int2, ulong);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul2 (write_only image2d_t, int2, ulong2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul4 (write_only image2d_t, int2, ulong4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul8 (write_only image2d_t, int2, ulong8);
#if defined(__opencl_c_read_write_images)
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul (read_write image2d_t, int2, ulong);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul2 (read_write image2d_t, int2, ulong2);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul4 (read_write image2d_t, int2, ulong4);
void _CL_OVERLOADABLE _CL_CONV
intel_sub_group_block_write_ul8 (read_write image2d_t, int2, ulong8);
#endif // defined(__opencl_c_read_write_images)
#endif // defined(__opencl_c_images)
#endif // cl_intel_subgroups_long
/**************************************************************************/
#ifdef cl_intel_split_work_group_barrier
void _CL_OVERLOADABLE
intel_work_group_barrier_arrive (cl_mem_fence_flags flags);
void _CL_OVERLOADABLE intel_work_group_barrier_wait (cl_mem_fence_flags flags);
// For OpenCL C 2.0 or newer:
#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
void _CL_OVERLOADABLE
intel_work_group_barrier_arrive (cl_mem_fence_flags flags, memory_scope scope);
void _CL_OVERLOADABLE intel_work_group_barrier_wait (cl_mem_fence_flags flags,
memory_scope scope);
#endif
#endif // cl_intel_split_work_group_barrier
#endif
drivers/pocl-gen7/share/pocl/include/_enable_all_exts.h
+53
View File
@@ -0,0 +1,53 @@
/* Enable all extensions known to pocl, which a device supports.
* This is required at the start of include/_kernel.h for prototypes,
* then at kernel lib compilation phase (because _kernel.h disables
* everything at the end).
*/
/* OpenCL 1.0-only extensions */
#if (__OPENCL_C_VERSION__ < 110)
#ifdef cl_khr_global_int32_base_atomics
# pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#endif
#ifdef cl_khr_global_int32_extended_atomics
# pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
#endif
#ifdef cl_khr_local_int32_base_atomics
# pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#endif
#ifdef cl_khr_local_int32_extended_atomics
# pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#endif
#ifdef cl_khr_byte_addressable_store
# pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#endif
#endif
/* all versions */
#ifdef cl_khr_fp16
# pragma OPENCL EXTENSION cl_khr_fp16: enable
#endif
#ifdef cl_khr_fp64
# pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif
#ifdef cl_khr_int64_base_atomics
# pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#endif
#ifdef cl_khr_int64_extended_atomics
# pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
#endif
#ifdef cl_khr_3d_image_writes
# pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#endif
drivers/pocl-gen7/share/pocl/include/_kernel.h
+179
View File
@@ -0,0 +1,179 @@
/* pocl/_kernel.h - OpenCL types and runtime library
functions declarations. This should be included only from OpenCL C files.
Copyright (c) 2011 Universidad Rey Juan Carlos
Copyright (c) 2011-2017 Pekka Jääskeläinen / TUT
Copyright (c) 2011-2013 Erik Schnetter <[email protected]>
Perimeter Institute for Theoretical Physics
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* If the -cl-std build option is not specified, the highest OpenCL C 1.x
* language version supported by each device is used as the version of
* OpenCL C when compiling the program for each device.
*/
#ifndef __OPENCL_C_VERSION__
#define __OPENCL_C_VERSION__ 120
#endif
#if (__OPENCL_C_VERSION__ > 99)
#define CL_VERSION_1_0 100
#endif
#if (__OPENCL_C_VERSION__ > 109)
#define CL_VERSION_1_1 110
#endif
#if (__OPENCL_C_VERSION__ > 119)
#define CL_VERSION_1_2 120
#endif
#if (__OPENCL_C_VERSION__ > 199)
#define CL_VERSION_2_0 200
#endif
#include "_enable_all_exts.h"
#include "_builtin_renames.h"
/* Define some feature test macros to help write generic code. These are used
* mostly in _pocl_opencl.h header + some .cl files in kernel library */
#ifdef cl_khr_int64
# define __IF_INT64(x) x
#else
# define __IF_INT64(x)
#endif
#ifdef cl_khr_fp16
# define __IF_FP16(x) x
#else
# define __IF_FP16(x)
#endif
#ifdef cl_khr_fp64
# define __IF_FP64(x) x
#else
# define __IF_FP64(x)
#endif
#ifdef cl_khr_int64_base_atomics
#define __IF_BA64(x) x
#else
#define __IF_BA64(x)
#endif
#ifdef cl_khr_int64_extended_atomics
#define __IF_EA64(x) x
#else
#define __IF_EA64(x)
#endif
/****************************************************************************/
/* Function/type attributes supported by Clang/SPIR */
#if __has_attribute(__always_inline__)
# define _CL_ALWAYSINLINE __attribute__((__always_inline__))
#else
# define _CL_ALWAYSINLINE
#endif
#if __has_attribute(__noinline__)
# define _CL_NOINLINE __attribute__((__noinline__))
#else
# define _CL_NOINLINE
#endif
#if __has_attribute(__overloadable__)
# define _CL_OVERLOADABLE __attribute__((__overloadable__))
#else
# define _CL_OVERLOADABLE
#endif
#if __has_attribute(__pure__)
# define _CL_READONLY __attribute__((__pure__))
#else
# define _CL_READONLY
#endif
#if __has_attribute(__const__)
# define _CL_READNONE __attribute__((__const__))
#else
# define _CL_READNONE
#endif
#if __has_attribute(convergent)
# define _CL_CONVERGENT __attribute__((convergent))
#else
# define _CL_CONVERGENT
#endif
/****************************************************************************/
/* A static assert statement to catch inconsistencies at build time */
#if __has_extension(__c_static_assert__)
# define _CL_STATIC_ASSERT(_t, _x) _Static_assert(_x, #_t)
#else
# define _CL_STATIC_ASSERT(_t, _x) typedef int __cl_ai##_t[(x) ? 1 : -1];
#endif
/****************************************************************************/
#define IMG_RO_AQ __read_only
#define IMG_WO_AQ __write_only
#ifdef __opencl_c_read_write_images
#define CLANG_HAS_RW_IMAGES
#define IMG_RW_AQ __read_write
#else
#undef CLANG_HAS_RW_IMAGES
#define IMG_RW_AQ __RW_IMAGES_UNSUPPORTED_BEFORE_CL_20
#endif
/****************************************************************************/
/* use Clang opencl header for definitions. */
#ifdef POCL_DEVICE_ADDRESS_BITS
/* If we wish to override the Clang set __SIZE_TYPE__ for this target,
let's do it here so the opencl-c.h sets size_t to the wanted type. */
#ifdef __SIZE_TYPE__
#undef __SIZE_TYPE__
#endif
#if POCL_DEVICE_ADDRESS_BITS == 32
#define __SIZE_TYPE__ uint
#elif POCL_DEVICE_ADDRESS_BITS == 64
#define __SIZE_TYPE__ ulong
#else
#error Unsupported POCL_DEVICE_ADDRESS_BITS value.
#endif
#endif
#include "_clang_opencl.h"
/****************************************************************************/
/* GNU's libm seems to use INT_MIN here while the Clang's header uses
INT_MAX. Both are allowed by the OpenCL specs, but we want them to
be unified to avoid failing tests. */
#undef FP_ILOGBNAN
#undef FP_ILOGB0
#define FP_ILOGB0 INT_MIN
#define FP_ILOGBNAN INT_MAX
/****************************************************************************/
#include "pocl_image_types.h"
#pragma OPENCL EXTENSION all : disable
drivers/pocl-gen7/share/pocl/include/_kernel_c.h
+199
View File
@@ -0,0 +1,199 @@
/* pocl/_kernel_c.h - C compatible OpenCL types and runtime library
functions declarations for kernel builtin implementations using C.
Copyright (c) 2011 Universidad Rey Juan Carlos
Copyright (c) 2011-2017 Pekka Jääskeläinen / TUT
Copyright (c) 2011-2013 Erik Schnetter <[email protected]>
Perimeter Institute for Theoretical Physics
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* Header that can be implemented in C compiled implementations of
* built-in functions to introduce the OpenCL C compatible types etc.
*/
#ifndef _KERNEL_C_H
#define _KERNEL_C_H
#include "pocl_types.h"
/* cannot include "_builtin_renames.h" because it overrides some C/C++ macros,
but we need the prefixed functions in the builtin library */
#define POCL_BUILTIN_PREFIX(FUNC) _cl_##FUNC
#include "_kernel_constants.h"
/* Function/type attributes supported by Clang/SPIR */
#if __has_attribute(__always_inline__)
# define _CL_ALWAYSINLINE __attribute__((__always_inline__))
#else
# define _CL_ALWAYSINLINE
#endif
#if __has_attribute(__noinline__)
# define _CL_NOINLINE __attribute__((__noinline__))
#else
# define _CL_NOINLINE
#endif
#if __has_attribute(__overloadable__)
# define _CL_OVERLOADABLE __attribute__((__overloadable__))
#else
# define _CL_OVERLOADABLE
#endif
#if __has_attribute(__pure__)
# define _CL_READONLY __attribute__((__pure__))
#else
# define _CL_READONLY
#endif
#if __has_attribute(__const__)
# define _CL_READNONE __attribute__((__const__))
#else
# define _CL_READNONE
#endif
#if __has_attribute(convergent)
# define _CL_CONVERGENT __attribute__((convergent))
#else
# define _CL_CONVERGENT
#endif
#if __has_attribute(optnone)
#define _CL_OPTNONE __attribute__ ((optnone))
#else
#define _CL_OPTNONE
#endif
typedef char char2 __attribute__((__ext_vector_type__(2)));
typedef char char3 __attribute__((__ext_vector_type__(3)));
typedef char char4 __attribute__((__ext_vector_type__(4)));
typedef char char8 __attribute__((__ext_vector_type__(8)));
typedef char char16 __attribute__((__ext_vector_type__(16)));
typedef uchar uchar2 __attribute__((__ext_vector_type__(2)));
typedef uchar uchar3 __attribute__((__ext_vector_type__(3)));
typedef uchar uchar4 __attribute__((__ext_vector_type__(4)));
typedef uchar uchar8 __attribute__((__ext_vector_type__(8)));
typedef uchar uchar16 __attribute__((__ext_vector_type__(16)));
typedef short short2 __attribute__((__ext_vector_type__(2)));
typedef short short3 __attribute__((__ext_vector_type__(3)));
typedef short short4 __attribute__((__ext_vector_type__(4)));
typedef short short8 __attribute__((__ext_vector_type__(8)));
typedef short short16 __attribute__((__ext_vector_type__(16)));
typedef ushort ushort2 __attribute__((__ext_vector_type__(2)));
typedef ushort ushort3 __attribute__((__ext_vector_type__(3)));
typedef ushort ushort4 __attribute__((__ext_vector_type__(4)));
typedef ushort ushort8 __attribute__((__ext_vector_type__(8)));
typedef ushort ushort16 __attribute__((__ext_vector_type__(16)));
typedef int int2 __attribute__((__ext_vector_type__(2)));
typedef int int3 __attribute__((__ext_vector_type__(3)));
typedef int int4 __attribute__((__ext_vector_type__(4)));
typedef int int8 __attribute__((__ext_vector_type__(8)));
typedef int int16 __attribute__((__ext_vector_type__(16)));
typedef uint uint2 __attribute__((__ext_vector_type__(2)));
typedef uint uint3 __attribute__((__ext_vector_type__(3)));
typedef uint uint4 __attribute__((__ext_vector_type__(4)));
typedef uint uint8 __attribute__((__ext_vector_type__(8)));
typedef uint uint16 __attribute__((__ext_vector_type__(16)));
#if defined(__CBUILD__) && defined(cl_khr_fp16)
#ifdef _HAS_FLOAT16_TYPE
typedef _Float16 half;
#else
typedef __fp16 half;
#endif
typedef half half2 __attribute__((__ext_vector_type__(2)));
typedef half half3 __attribute__((__ext_vector_type__(3)));
typedef half half4 __attribute__((__ext_vector_type__(4)));
typedef half half8 __attribute__((__ext_vector_type__(8)));
typedef half half16 __attribute__((__ext_vector_type__(16)));
#endif
typedef float float2 __attribute__((__ext_vector_type__(2)));
typedef float float3 __attribute__((__ext_vector_type__(3)));
typedef float float4 __attribute__((__ext_vector_type__(4)));
typedef float float8 __attribute__((__ext_vector_type__(8)));
typedef float float16 __attribute__((__ext_vector_type__(16)));
#ifdef cl_khr_fp64
# ifndef __CBUILD__
# pragma OPENCL EXTENSION cl_khr_fp64 : enable
# endif
typedef double double2 __attribute__((__ext_vector_type__(2)));
typedef double double3 __attribute__((__ext_vector_type__(3)));
typedef double double4 __attribute__((__ext_vector_type__(4)));
typedef double double8 __attribute__((__ext_vector_type__(8)));
typedef double double16 __attribute__((__ext_vector_type__(16)));
#endif
#ifdef cl_khr_int64
typedef long long2 __attribute__((__ext_vector_type__(2)));
typedef long long3 __attribute__((__ext_vector_type__(3)));
typedef long long4 __attribute__((__ext_vector_type__(4)));
typedef long long8 __attribute__((__ext_vector_type__(8)));
typedef long long16 __attribute__((__ext_vector_type__(16)));
typedef ulong ulong2 __attribute__((__ext_vector_type__(2)));
typedef ulong ulong3 __attribute__((__ext_vector_type__(3)));
typedef ulong ulong4 __attribute__((__ext_vector_type__(4)));
typedef ulong ulong8 __attribute__((__ext_vector_type__(8)));
typedef ulong ulong16 __attribute__((__ext_vector_type__(16)));
#endif
#if defined(__TCE__)
#define POCL_ADDRESS_SPACE_PRIVATE 0
#define POCL_ADDRESS_SPACE_GLOBAL 1
#define POCL_ADDRESS_SPACE_LOCAL 3
#define POCL_ADDRESS_SPACE_CONSTANT 2
#define POCL_ADDRESS_SPACE_GENERIC 6
#endif
typedef uint cl_mem_fence_flags;
/* Integer Constants */
#if defined(__CBUILD__)
#define CHAR_BIT 8
#define CHAR_MAX SCHAR_MAX
#define CHAR_MIN SCHAR_MIN
#define INT_MAX 2147483647
#define INT_MIN (-2147483647 - 1)
#ifdef cl_khr_int64
#define LONG_MAX 0x7fffffffffffffffL
#define LONG_MIN (-0x7fffffffffffffffL - 1)
#endif
#define SCHAR_MAX 127
#define SCHAR_MIN (-127 - 1)
#define SHRT_MAX 32767
#define SHRT_MIN (-32767 - 1)
#define UCHAR_MAX 255
#define USHRT_MAX 65535
#define UINT_MAX 0xffffffff
#ifdef cl_khr_int64
#define ULONG_MAX 0xffffffffffffffffUL
#endif
#endif /* __CBUILD__ */
#endif
drivers/pocl-gen7/share/pocl/include/_kernel_constants.h
+93
View File
@@ -0,0 +1,93 @@
/* pocl/_kernel_constants.h - C compatible OpenCL types and runtime library
constants declarations.
Copyright (c) 2011 Universidad Rey Juan Carlos
Copyright (c) 2011-2013 Pekka Jääskeläinen / TUT
Copyright (c) 2011-2013 Erik Schnetter <[email protected]>
Perimeter Institute for Theoretical Physics
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* Header that can be implemented in C compiled implementations of
* built-in functions to introduce the OpenCL C compatible constants.
*/
#ifndef _KERNEL_CONSTANTS_H
#define _KERNEL_CONSTANTS_H
/* clang's header defines these */
#ifndef _OPENCL_H_
/* cl_channel_order */
#define CLK_R 0x10B0
#define CLK_A 0x10B1
#define CLK_RG 0x10B2
#define CLK_RA 0x10B3
#define CLK_RGB 0x10B4
#define CLK_RGBA 0x10B5
#define CLK_BGRA 0x10B6
#define CLK_ARGB 0x10B7
#define CLK_INTENSITY 0x10B8
#define CLK_LUMINANCE 0x10B9
#define CLK_Rx 0x10BA
#define CLK_RGx 0x10BB
#define CLK_RGBx 0x10BC
#define CLK_DEPTH 0x10BD
#define CLK_DEPTH_STENCIL 0x10BE
/* cl_channel_type */
#define CLK_SNORM_INT8 0x10D0
#define CLK_SNORM_INT16 0x10D1
#define CLK_UNORM_INT8 0x10D2
#define CLK_UNORM_INT16 0x10D3
#define CLK_UNORM_SHORT_565 0x10D4
#define CLK_UNORM_SHORT_555 0x10D5
#define CLK_UNORM_INT_101010 0x10D6
#define CLK_SIGNED_INT8 0x10D7
#define CLK_SIGNED_INT16 0x10D8
#define CLK_SIGNED_INT32 0x10D9
#define CLK_UNSIGNED_INT8 0x10DA
#define CLK_UNSIGNED_INT16 0x10DB
#define CLK_UNSIGNED_INT32 0x10DC
#define CLK_HALF_FLOAT 0x10DD
#define CLK_FLOAT 0x10DE
#define CLK_UNORM_INT24 0x10DF
/* cl_addressing _mode */
#define CLK_ADDRESS_NONE 0x00
#define CLK_ADDRESS_CLAMP_TO_EDGE 0x02
#define CLK_ADDRESS_CLAMP 0x04
#define CLK_ADDRESS_REPEAT 0x06
#define CLK_ADDRESS_MIRRORED_REPEAT 0x08
/* cl_sampler_info */
#define CLK_NORMALIZED_COORDS_FALSE 0x00
#define CLK_NORMALIZED_COORDS_TRUE 0x01
/* filter_mode */
#define CLK_FILTER_NEAREST 0x10
#define CLK_FILTER_LINEAR 0x20
/* barrier() flags */
#define CLK_LOCAL_MEM_FENCE 0x01
#define CLK_GLOBAL_MEM_FENCE 0x02
#endif
#endif
drivers/pocl-gen7/share/pocl/include/opencl-c-base.h
+830
View File
@@ -0,0 +1,830 @@
//===----- opencl-c-base.h - OpenCL C language base definitions -----------===//
//
// 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 _OPENCL_BASE_H_
#define _OPENCL_BASE_H_
// Define extension macros
#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
// For SPIR and SPIR-V all extensions are supported.
#if defined(__SPIR__) || defined(__SPIRV__)
#define cl_khr_subgroup_extended_types 1
#define cl_khr_subgroup_non_uniform_vote 1
#define cl_khr_subgroup_ballot 1
#define cl_khr_subgroup_non_uniform_arithmetic 1
#define cl_khr_subgroup_shuffle 1
#define cl_khr_subgroup_shuffle_relative 1
#define cl_khr_subgroup_clustered_reduce 1
#define cl_khr_subgroup_rotate 1
#define cl_khr_extended_bit_ops 1
#define cl_khr_integer_dot_product 1
#define __opencl_c_integer_dot_product_input_4x8bit 1
#define __opencl_c_integer_dot_product_input_4x8bit_packed 1
#define cl_ext_float_atomics 1
#ifdef cl_khr_fp16
#define __opencl_c_ext_fp16_global_atomic_load_store 1
#define __opencl_c_ext_fp16_local_atomic_load_store 1
#define __opencl_c_ext_fp16_global_atomic_add 1
#define __opencl_c_ext_fp16_local_atomic_add 1
#define __opencl_c_ext_fp16_global_atomic_min_max 1
#define __opencl_c_ext_fp16_local_atomic_min_max 1
#endif
#ifdef cl_khr_fp64
#define __opencl_c_ext_fp64_global_atomic_add 1
#define __opencl_c_ext_fp64_local_atomic_add 1
#define __opencl_c_ext_fp64_global_atomic_min_max 1
#define __opencl_c_ext_fp64_local_atomic_min_max 1
#endif
#define __opencl_c_ext_fp32_global_atomic_add 1
#define __opencl_c_ext_fp32_local_atomic_add 1
#define __opencl_c_ext_fp32_global_atomic_min_max 1
#define __opencl_c_ext_fp32_local_atomic_min_max 1
#endif // defined(__SPIR__) || defined(__SPIRV__)
#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
// Define feature macros for OpenCL C 2.0
#if (__OPENCL_CPP_VERSION__ == 100 || __OPENCL_C_VERSION__ == 200)
#define __opencl_c_pipes 1
#define __opencl_c_generic_address_space 1
#define __opencl_c_work_group_collective_functions 1
#define __opencl_c_atomic_order_acq_rel 1
#define __opencl_c_atomic_order_seq_cst 1
#define __opencl_c_atomic_scope_device 1
#define __opencl_c_atomic_scope_all_devices 1
#define __opencl_c_device_enqueue 1
#define __opencl_c_read_write_images 1
#define __opencl_c_program_scope_global_variables 1
#define __opencl_c_images 1
#endif
// Define header-only feature macros for OpenCL C 3.0.
#if (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300)
// For the SPIR and SPIR-V target all features are supported.
#if defined(__SPIR__) || defined(__SPIRV__)
#define __opencl_c_work_group_collective_functions 1
#define __opencl_c_atomic_order_seq_cst 1
#define __opencl_c_atomic_scope_device 1
#define __opencl_c_atomic_scope_all_devices 1
#define __opencl_c_read_write_images 1
#endif // defined(__SPIR__)
// Undefine any feature macros that have been explicitly disabled using
// an __undef_<feature> macro.
#ifdef __undef___opencl_c_work_group_collective_functions
#undef __opencl_c_work_group_collective_functions
#endif
#ifdef __undef___opencl_c_atomic_order_seq_cst
#undef __opencl_c_atomic_order_seq_cst
#endif
#ifdef __undef___opencl_c_atomic_scope_device
#undef __opencl_c_atomic_scope_device
#endif
#ifdef __undef___opencl_c_atomic_scope_all_devices
#undef __opencl_c_atomic_scope_all_devices
#endif
#ifdef __undef___opencl_c_read_write_images
#undef __opencl_c_read_write_images
#endif
#endif // (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300)
#if !defined(__opencl_c_generic_address_space)
// Internal feature macro to provide named (global, local, private) address
// space overloads for builtin functions that take a pointer argument.
#define __opencl_c_named_address_space_builtins 1
#endif // !defined(__opencl_c_generic_address_space)
#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups) || defined(__opencl_c_subgroups)
// Internal feature macro to provide subgroup builtins.
#define __opencl_subgroup_builtins 1
#endif
// built-in scalar data types:
/**
* An unsigned 8-bit integer.
*/
typedef unsigned char uchar;
/**
* An unsigned 16-bit integer.
*/
typedef unsigned short ushort;
/**
* An unsigned 32-bit integer.
*/
typedef unsigned int uint;
/**
* An unsigned 64-bit integer.
*/
typedef unsigned long ulong;
/**
* The unsigned integer type of the result of the sizeof operator. This
* is a 32-bit unsigned integer if CL_DEVICE_ADDRESS_BITS
* defined in table 4.3 is 32-bits and is a 64-bit unsigned integer if
* CL_DEVICE_ADDRESS_BITS is 64-bits.
*/
typedef __SIZE_TYPE__ size_t;
/**
* A signed integer type that is the result of subtracting two pointers.
* This is a 32-bit signed integer if CL_DEVICE_ADDRESS_BITS
* defined in table 4.3 is 32-bits and is a 64-bit signed integer if
* CL_DEVICE_ADDRESS_BITS is 64-bits.
*/
typedef __PTRDIFF_TYPE__ ptrdiff_t;
/**
* A signed integer type with the property that any valid pointer to
* void can be converted to this type, then converted back to pointer
* to void, and the result will compare equal to the original pointer.
*/
typedef __INTPTR_TYPE__ intptr_t;
/**
* An unsigned integer type with the property that any valid pointer to
* void can be converted to this type, then converted back to pointer
* to void, and the result will compare equal to the original pointer.
*/
typedef __UINTPTR_TYPE__ uintptr_t;
// built-in vector data types:
typedef char char2 __attribute__((ext_vector_type(2))) __attribute__((aligned(2)));
typedef char char3 __attribute__((ext_vector_type(3))) __attribute__((aligned(4)));
typedef char char4 __attribute__((ext_vector_type(4))) __attribute__((aligned(4)));
typedef char char8 __attribute__((ext_vector_type(8))) __attribute__((aligned(8)));
typedef char char16 __attribute__((ext_vector_type(16))) __attribute__((aligned(16)));
typedef uchar uchar2 __attribute__((ext_vector_type(2))) __attribute__((aligned(2)));
typedef uchar uchar3 __attribute__((ext_vector_type(3))) __attribute__((aligned(4)));
typedef uchar uchar4 __attribute__((ext_vector_type(4))) __attribute__((aligned(4)));
typedef uchar uchar8 __attribute__((ext_vector_type(8))) __attribute__((aligned(8)));
typedef uchar uchar16 __attribute__((ext_vector_type(16))) __attribute__((aligned(16)));
typedef short short2 __attribute__((ext_vector_type(2))) __attribute__((aligned(4)));
typedef short short3 __attribute__((ext_vector_type(3))) __attribute__((aligned(8)));
typedef short short4 __attribute__((ext_vector_type(4))) __attribute__((aligned(8)));
typedef short short8 __attribute__((ext_vector_type(8))) __attribute__((aligned(16)));
typedef short short16 __attribute__((ext_vector_type(16))) __attribute__((aligned(32)));
typedef ushort ushort2 __attribute__((ext_vector_type(2))) __attribute__((aligned(4)));
typedef ushort ushort3 __attribute__((ext_vector_type(3))) __attribute__((aligned(8)));
typedef ushort ushort4 __attribute__((ext_vector_type(4))) __attribute__((aligned(8)));
typedef ushort ushort8 __attribute__((ext_vector_type(8))) __attribute__((aligned(16)));
typedef ushort ushort16 __attribute__((ext_vector_type(16))) __attribute__((aligned(32)));
typedef int int2 __attribute__((ext_vector_type(2))) __attribute__((aligned(8)));
typedef int int3 __attribute__((ext_vector_type(3))) __attribute__((aligned(16)));
typedef int int4 __attribute__((ext_vector_type(4))) __attribute__((aligned(16)));
typedef int int8 __attribute__((ext_vector_type(8))) __attribute__((aligned(32)));
typedef int int16 __attribute__((ext_vector_type(16))) __attribute__((aligned(64)));
typedef uint uint2 __attribute__((ext_vector_type(2))) __attribute__((aligned(8)));
typedef uint uint3 __attribute__((ext_vector_type(3))) __attribute__((aligned(16)));
typedef uint uint4 __attribute__((ext_vector_type(4))) __attribute__((aligned(16)));
typedef uint uint8 __attribute__((ext_vector_type(8))) __attribute__((aligned(32)));
typedef uint uint16 __attribute__((ext_vector_type(16))) __attribute__((aligned(64)));
typedef long long2 __attribute__((ext_vector_type(2))) __attribute__((aligned(16)));
typedef long long3 __attribute__((ext_vector_type(3))) __attribute__((aligned(32)));
typedef long long4 __attribute__((ext_vector_type(4))) __attribute__((aligned(32)));
typedef long long8 __attribute__((ext_vector_type(8))) __attribute__((aligned(64)));
typedef long long16 __attribute__((ext_vector_type(16))) __attribute__((aligned(128)));
typedef ulong ulong2 __attribute__((ext_vector_type(2))) __attribute__((aligned(16)));
typedef ulong ulong3 __attribute__((ext_vector_type(3))) __attribute__((aligned(32)));
typedef ulong ulong4 __attribute__((ext_vector_type(4))) __attribute__((aligned(32)));
typedef ulong ulong8 __attribute__((ext_vector_type(8))) __attribute__((aligned(64)));
typedef ulong ulong16 __attribute__((ext_vector_type(16))) __attribute__((aligned(128)));
typedef float float2 __attribute__((ext_vector_type(2))) __attribute__((aligned(8)));
typedef float float3 __attribute__((ext_vector_type(3))) __attribute__((aligned(16)));
typedef float float4 __attribute__((ext_vector_type(4))) __attribute__((aligned(16)));
typedef float float8 __attribute__((ext_vector_type(8))) __attribute__((aligned(32)));
typedef float float16 __attribute__((ext_vector_type(16))) __attribute__((aligned(64)));
#ifdef cl_khr_fp16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
typedef half half2 __attribute__((ext_vector_type(2))) __attribute__((aligned(4)));
typedef half half3 __attribute__((ext_vector_type(3))) __attribute__((aligned(8)));
typedef half half4 __attribute__((ext_vector_type(4))) __attribute__((aligned(8)));
typedef half half8 __attribute__((ext_vector_type(8))) __attribute__((aligned(16)));
typedef half half16 __attribute__((ext_vector_type(16))) __attribute__((aligned(32)));
#endif
#ifdef cl_khr_fp64
#if __OPENCL_C_VERSION__ < CL_VERSION_1_2
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#endif
typedef double double2 __attribute__((ext_vector_type(2))) __attribute__((aligned(16)));
typedef double double3 __attribute__((ext_vector_type(3))) __attribute__((aligned(32)));
typedef double double4 __attribute__((ext_vector_type(4))) __attribute__((aligned(32)));
typedef double double8 __attribute__((ext_vector_type(8))) __attribute__((aligned(64)));
typedef double double16 __attribute__((ext_vector_type(16))) __attribute__((aligned(128)));
#endif
// An internal alias for half, for use by OpenCLBuiltins.td.
#define __half half
#if defined(__OPENCL_CPP_VERSION__)
#define NULL nullptr
#elif defined(__OPENCL_C_VERSION__)
#define NULL ((void*)0)
#endif
/**
* Value of maximum non-infinite single-precision floating-point
* number.
*/
#define MAXFLOAT 0x1.fffffep127f
/**
* A positive float constant expression. HUGE_VALF evaluates
* to +infinity. Used as an error value returned by the built-in
* math functions.
*/
#define HUGE_VALF (__builtin_huge_valf())
/**
* A positive double constant expression. HUGE_VAL evaluates
* to +infinity. Used as an error value returned by the built-in
* math functions.
*/
#define HUGE_VAL (__builtin_huge_val())
/**
* A constant expression of type float representing positive or
* unsigned infinity.
*/
#define INFINITY (__builtin_inff())
/**
* A constant expression of type float representing a quiet NaN.
*/
#define NAN as_float(INT_MAX)
#define FP_ILOGB0 INT_MIN
#define FP_ILOGBNAN INT_MAX
#define FLT_DIG 6
#define FLT_MANT_DIG 24
#define FLT_MAX_10_EXP +38
#define FLT_MAX_EXP +128
#define FLT_MIN_10_EXP -37
#define FLT_MIN_EXP -125
#define FLT_RADIX 2
#define FLT_MAX 0x1.fffffep127f
#define FLT_MIN 0x1.0p-126f
#define FLT_EPSILON 0x1.0p-23f
#define M_E_F 2.71828182845904523536028747135266250f
#define M_LOG2E_F 1.44269504088896340735992468100189214f
#define M_LOG10E_F 0.434294481903251827651128918916605082f
#define M_LN2_F 0.693147180559945309417232121458176568f
#define M_LN10_F 2.30258509299404568401799145468436421f
#define M_PI_F 3.14159265358979323846264338327950288f
#define M_PI_2_F 1.57079632679489661923132169163975144f
#define M_PI_4_F 0.785398163397448309615660845819875721f
#define M_1_PI_F 0.318309886183790671537767526745028724f
#define M_2_PI_F 0.636619772367581343075535053490057448f
#define M_2_SQRTPI_F 1.12837916709551257389615890312154517f
#define M_SQRT2_F 1.41421356237309504880168872420969808f
#define M_SQRT1_2_F 0.707106781186547524400844362104849039f
#define DBL_DIG 15
#define DBL_MANT_DIG 53
#define DBL_MAX_10_EXP +308
#define DBL_MAX_EXP +1024
#define DBL_MIN_10_EXP -307
#define DBL_MIN_EXP -1021
#define DBL_RADIX 2
#define DBL_MAX 0x1.fffffffffffffp1023
#define DBL_MIN 0x1.0p-1022
#define DBL_EPSILON 0x1.0p-52
#define M_E 0x1.5bf0a8b145769p+1
#define M_LOG2E 0x1.71547652b82fep+0
#define M_LOG10E 0x1.bcb7b1526e50ep-2
#define M_LN2 0x1.62e42fefa39efp-1
#define M_LN10 0x1.26bb1bbb55516p+1
#define M_PI 0x1.921fb54442d18p+1
#define M_PI_2 0x1.921fb54442d18p+0
#define M_PI_4 0x1.921fb54442d18p-1
#define M_1_PI 0x1.45f306dc9c883p-2
#define M_2_PI 0x1.45f306dc9c883p-1
#define M_2_SQRTPI 0x1.20dd750429b6dp+0
#define M_SQRT2 0x1.6a09e667f3bcdp+0
#define M_SQRT1_2 0x1.6a09e667f3bcdp-1
#ifdef cl_khr_fp16
#define HALF_DIG 3
#define HALF_MANT_DIG 11
#define HALF_MAX_10_EXP +4
#define HALF_MAX_EXP +16
#define HALF_MIN_10_EXP -4
#define HALF_MIN_EXP -13
#define HALF_RADIX 2
#define HALF_MAX ((0x1.ffcp15h))
#define HALF_MIN ((0x1.0p-14h))
#define HALF_EPSILON ((0x1.0p-10h))
#define M_E_H 2.71828182845904523536028747135266250h
#define M_LOG2E_H 1.44269504088896340735992468100189214h
#define M_LOG10E_H 0.434294481903251827651128918916605082h
#define M_LN2_H 0.693147180559945309417232121458176568h
#define M_LN10_H 2.30258509299404568401799145468436421h
#define M_PI_H 3.14159265358979323846264338327950288h
#define M_PI_2_H 1.57079632679489661923132169163975144h
#define M_PI_4_H 0.785398163397448309615660845819875721h
#define M_1_PI_H 0.318309886183790671537767526745028724h
#define M_2_PI_H 0.636619772367581343075535053490057448h
#define M_2_SQRTPI_H 1.12837916709551257389615890312154517h
#define M_SQRT2_H 1.41421356237309504880168872420969808h
#define M_SQRT1_2_H 0.707106781186547524400844362104849039h
#endif //cl_khr_fp16
#define CHAR_BIT 8
#define SCHAR_MAX 127
#define SCHAR_MIN (-128)
#define UCHAR_MAX 255
#define CHAR_MAX SCHAR_MAX
#define CHAR_MIN SCHAR_MIN
#define USHRT_MAX 65535
#define SHRT_MAX 32767
#define SHRT_MIN (-32768)
#define UINT_MAX 0xffffffff
#define INT_MAX 2147483647
#define INT_MIN (-2147483647-1)
#define ULONG_MAX 0xffffffffffffffffUL
#define LONG_MAX 0x7fffffffffffffffL
#define LONG_MIN (-0x7fffffffffffffffL-1)
// OpenCL v1.1 s6.11.8, v1.2 s6.12.8, v2.0 s6.13.8 - Synchronization Functions
// Flag type and values for barrier, mem_fence, read_mem_fence, write_mem_fence
typedef uint cl_mem_fence_flags;
/**
* Queue a memory fence to ensure correct
* ordering of memory operations to local memory
*/
#define CLK_LOCAL_MEM_FENCE 0x01
/**
* Queue a memory fence to ensure correct
* ordering of memory operations to global memory
*/
#define CLK_GLOBAL_MEM_FENCE 0x02
#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
typedef enum memory_scope {
memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
#if defined(__opencl_c_atomic_scope_all_devices)
memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
#if (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 || __OPENCL_CPP_VERSION__ >= 202100)
memory_scope_all_devices = memory_scope_all_svm_devices,
#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 || __OPENCL_CPP_VERSION__ >= 202100)
#endif // defined(__opencl_c_atomic_scope_all_devices)
/**
* Subgroups have different requirements on forward progress, so just test
* all the relevant macros.
* CL 3.0 sub-groups "they are not guaranteed to make independent forward progress"
* KHR subgroups "Subgroups within a workgroup are independent, make forward progress with respect to each other"
*/
#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups) || defined(__opencl_c_subgroups)
memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
#endif
} memory_scope;
/**
* Queue a memory fence to ensure correct ordering of memory
* operations between work-items of a work-group to
* image memory.
*/
#define CLK_IMAGE_MEM_FENCE 0x04
#ifndef ATOMIC_VAR_INIT
#define ATOMIC_VAR_INIT(x) (x)
#endif //ATOMIC_VAR_INIT
#define ATOMIC_FLAG_INIT 0
// enum values aligned with what clang uses in EmitAtomicExpr()
typedef enum memory_order
{
memory_order_relaxed = __ATOMIC_RELAXED,
memory_order_acquire = __ATOMIC_ACQUIRE,
memory_order_release = __ATOMIC_RELEASE,
memory_order_acq_rel = __ATOMIC_ACQ_REL,
#if defined(__opencl_c_atomic_order_seq_cst)
memory_order_seq_cst = __ATOMIC_SEQ_CST
#endif
} memory_order;
#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
// OpenCL v1.1 s6.11.3, v1.2 s6.12.14, v2.0 s6.13.14 - Image Read and Write Functions
// These values need to match the runtime equivalent
//
// Addressing Mode.
//
#define CLK_ADDRESS_NONE 0
#define CLK_ADDRESS_CLAMP_TO_EDGE 2
#define CLK_ADDRESS_CLAMP 4
#define CLK_ADDRESS_REPEAT 6
#define CLK_ADDRESS_MIRRORED_REPEAT 8
//
// Coordination Normalization
//
#define CLK_NORMALIZED_COORDS_FALSE 0
#define CLK_NORMALIZED_COORDS_TRUE 1
//
// Filtering Mode.
//
#define CLK_FILTER_NEAREST 0x10
#define CLK_FILTER_LINEAR 0x20
#ifdef cl_khr_gl_msaa_sharing
#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable
#endif //cl_khr_gl_msaa_sharing
//
// Channel Datatype.
//
#define CLK_SNORM_INT8 0x10D0
#define CLK_SNORM_INT16 0x10D1
#define CLK_UNORM_INT8 0x10D2
#define CLK_UNORM_INT16 0x10D3
#define CLK_UNORM_SHORT_565 0x10D4
#define CLK_UNORM_SHORT_555 0x10D5
#define CLK_UNORM_INT_101010 0x10D6
#define CLK_SIGNED_INT8 0x10D7
#define CLK_SIGNED_INT16 0x10D8
#define CLK_SIGNED_INT32 0x10D9
#define CLK_UNSIGNED_INT8 0x10DA
#define CLK_UNSIGNED_INT16 0x10DB
#define CLK_UNSIGNED_INT32 0x10DC
#define CLK_HALF_FLOAT 0x10DD
#define CLK_FLOAT 0x10DE
#define CLK_UNORM_INT24 0x10DF
#if __OPENCL_C_VERSION__ >= CL_VERSION_3_0
#define CLK_UNORM_INT_101010_2 0x10E0
#endif // __OPENCL_C_VERSION__ >= CL_VERSION_3_0
// Channel order, numbering must be aligned with cl_channel_order in cl.h
//
#define CLK_R 0x10B0
#define CLK_A 0x10B1
#define CLK_RG 0x10B2
#define CLK_RA 0x10B3
#define CLK_RGB 0x10B4
#define CLK_RGBA 0x10B5
#define CLK_BGRA 0x10B6
#define CLK_ARGB 0x10B7
#define CLK_INTENSITY 0x10B8
#define CLK_LUMINANCE 0x10B9
#define CLK_Rx 0x10BA
#define CLK_RGx 0x10BB
#define CLK_RGBx 0x10BC
#define CLK_DEPTH 0x10BD
#define CLK_DEPTH_STENCIL 0x10BE
#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
#define CLK_sRGB 0x10BF
#define CLK_sRGBx 0x10C0
#define CLK_sRGBA 0x10C1
#define CLK_sBGRA 0x10C2
#define CLK_ABGR 0x10C3
#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
// OpenCL v2.0 s6.13.16 - Pipe Functions
#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
#define CLK_NULL_RESERVE_ID (__builtin_astype(((void*)(__SIZE_MAX__)), reserve_id_t))
// OpenCL v2.0 s6.13.17 - Enqueue Kernels
#define CL_COMPLETE 0x0
#define CL_RUNNING 0x1
#define CL_SUBMITTED 0x2
#define CL_QUEUED 0x3
#define CLK_SUCCESS 0
#define CLK_ENQUEUE_FAILURE -101
#define CLK_INVALID_QUEUE -102
#define CLK_INVALID_NDRANGE -160
#define CLK_INVALID_EVENT_WAIT_LIST -57
#define CLK_DEVICE_QUEUE_FULL -161
#define CLK_INVALID_ARG_SIZE -51
#define CLK_EVENT_ALLOCATION_FAILURE -100
#define CLK_OUT_OF_RESOURCES -5
#define CLK_NULL_QUEUE 0
#define CLK_NULL_EVENT (__builtin_astype(((__SIZE_MAX__)), clk_event_t))
// execution model related definitions
#define CLK_ENQUEUE_FLAGS_NO_WAIT 0x0
#define CLK_ENQUEUE_FLAGS_WAIT_KERNEL 0x1
#define CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP 0x2
typedef int kernel_enqueue_flags_t;
typedef int clk_profiling_info;
// Profiling info name (see capture_event_profiling_info)
#define CLK_PROFILING_COMMAND_EXEC_TIME 0x1
#define MAX_WORK_DIM 3
#ifdef __opencl_c_device_enqueue
typedef struct {
unsigned int workDimension;
size_t globalWorkOffset[MAX_WORK_DIM];
size_t globalWorkSize[MAX_WORK_DIM];
size_t localWorkSize[MAX_WORK_DIM];
} ndrange_t;
#endif // __opencl_c_device_enqueue
#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
/**
* OpenCL v1.1/1.2/2.0 s6.2.4.2 - as_type operators
* Reinterprets a data type as another data type of the same size
*/
#define as_char(x) __builtin_astype((x), char)
#define as_char2(x) __builtin_astype((x), char2)
#define as_char3(x) __builtin_astype((x), char3)
#define as_char4(x) __builtin_astype((x), char4)
#define as_char8(x) __builtin_astype((x), char8)
#define as_char16(x) __builtin_astype((x), char16)
#define as_uchar(x) __builtin_astype((x), uchar)
#define as_uchar2(x) __builtin_astype((x), uchar2)
#define as_uchar3(x) __builtin_astype((x), uchar3)
#define as_uchar4(x) __builtin_astype((x), uchar4)
#define as_uchar8(x) __builtin_astype((x), uchar8)
#define as_uchar16(x) __builtin_astype((x), uchar16)
#define as_short(x) __builtin_astype((x), short)
#define as_short2(x) __builtin_astype((x), short2)
#define as_short3(x) __builtin_astype((x), short3)
#define as_short4(x) __builtin_astype((x), short4)
#define as_short8(x) __builtin_astype((x), short8)
#define as_short16(x) __builtin_astype((x), short16)
#define as_ushort(x) __builtin_astype((x), ushort)
#define as_ushort2(x) __builtin_astype((x), ushort2)
#define as_ushort3(x) __builtin_astype((x), ushort3)
#define as_ushort4(x) __builtin_astype((x), ushort4)
#define as_ushort8(x) __builtin_astype((x), ushort8)
#define as_ushort16(x) __builtin_astype((x), ushort16)
#define as_int(x) __builtin_astype((x), int)
#define as_int2(x) __builtin_astype((x), int2)
#define as_int3(x) __builtin_astype((x), int3)
#define as_int4(x) __builtin_astype((x), int4)
#define as_int8(x) __builtin_astype((x), int8)
#define as_int16(x) __builtin_astype((x), int16)
#define as_uint(x) __builtin_astype((x), uint)
#define as_uint2(x) __builtin_astype((x), uint2)
#define as_uint3(x) __builtin_astype((x), uint3)
#define as_uint4(x) __builtin_astype((x), uint4)
#define as_uint8(x) __builtin_astype((x), uint8)
#define as_uint16(x) __builtin_astype((x), uint16)
#define as_long(x) __builtin_astype((x), long)
#define as_long2(x) __builtin_astype((x), long2)
#define as_long3(x) __builtin_astype((x), long3)
#define as_long4(x) __builtin_astype((x), long4)
#define as_long8(x) __builtin_astype((x), long8)
#define as_long16(x) __builtin_astype((x), long16)
#define as_ulong(x) __builtin_astype((x), ulong)
#define as_ulong2(x) __builtin_astype((x), ulong2)
#define as_ulong3(x) __builtin_astype((x), ulong3)
#define as_ulong4(x) __builtin_astype((x), ulong4)
#define as_ulong8(x) __builtin_astype((x), ulong8)
#define as_ulong16(x) __builtin_astype((x), ulong16)
#define as_float(x) __builtin_astype((x), float)
#define as_float2(x) __builtin_astype((x), float2)
#define as_float3(x) __builtin_astype((x), float3)
#define as_float4(x) __builtin_astype((x), float4)
#define as_float8(x) __builtin_astype((x), float8)
#define as_float16(x) __builtin_astype((x), float16)
#ifdef cl_khr_fp64
#define as_double(x) __builtin_astype((x), double)
#define as_double2(x) __builtin_astype((x), double2)
#define as_double3(x) __builtin_astype((x), double3)
#define as_double4(x) __builtin_astype((x), double4)
#define as_double8(x) __builtin_astype((x), double8)
#define as_double16(x) __builtin_astype((x), double16)
#endif // cl_khr_fp64
#ifdef cl_khr_fp16
#define as_half(x) __builtin_astype((x), half)
#define as_half2(x) __builtin_astype((x), half2)
#define as_half3(x) __builtin_astype((x), half3)
#define as_half4(x) __builtin_astype((x), half4)
#define as_half8(x) __builtin_astype((x), half8)
#define as_half16(x) __builtin_astype((x), half16)
#endif // cl_khr_fp16
#define as_size_t(x) __builtin_astype((x), size_t)
#define as_ptrdiff_t(x) __builtin_astype((x), ptrdiff_t)
#define as_intptr_t(x) __builtin_astype((x), intptr_t)
#define as_uintptr_t(x) __builtin_astype((x), uintptr_t)
// C++ for OpenCL - __remove_address_space
#if defined(__OPENCL_CPP_VERSION__)
template <typename _Tp> struct __remove_address_space { using type = _Tp; };
#if defined(__opencl_c_generic_address_space)
template <typename _Tp> struct __remove_address_space<__generic _Tp> {
using type = _Tp;
};
#endif
template <typename _Tp> struct __remove_address_space<__global _Tp> {
using type = _Tp;
};
template <typename _Tp> struct __remove_address_space<__private _Tp> {
using type = _Tp;
};
template <typename _Tp> struct __remove_address_space<__local _Tp> {
using type = _Tp;
};
template <typename _Tp> struct __remove_address_space<__constant _Tp> {
using type = _Tp;
};
#endif
// OpenCL v1.1 s6.9, v1.2/2.0 s6.10 - Function qualifiers
#define __kernel_exec(X, typen) __kernel \
__attribute__((work_group_size_hint(X, 1, 1))) \
__attribute__((vec_type_hint(typen)))
#define kernel_exec(X, typen) __kernel \
__attribute__((work_group_size_hint(X, 1, 1))) \
__attribute__((vec_type_hint(typen)))
#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
// OpenCL v1.2 s6.12.13, v2.0 s6.13.13 - printf
int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
#endif
#ifdef cl_intel_device_side_avc_motion_estimation
#define CLK_AVC_ME_MAJOR_16x16_INTEL 0x0
#define CLK_AVC_ME_MAJOR_16x8_INTEL 0x1
#define CLK_AVC_ME_MAJOR_8x16_INTEL 0x2
#define CLK_AVC_ME_MAJOR_8x8_INTEL 0x3
#define CLK_AVC_ME_MINOR_8x8_INTEL 0x0
#define CLK_AVC_ME_MINOR_8x4_INTEL 0x1
#define CLK_AVC_ME_MINOR_4x8_INTEL 0x2
#define CLK_AVC_ME_MINOR_4x4_INTEL 0x3
#define CLK_AVC_ME_MAJOR_FORWARD_INTEL 0x0
#define CLK_AVC_ME_MAJOR_BACKWARD_INTEL 0x1
#define CLK_AVC_ME_MAJOR_BIDIRECTIONAL_INTEL 0x2
#define CLK_AVC_ME_PARTITION_MASK_ALL_INTEL 0x0
#define CLK_AVC_ME_PARTITION_MASK_16x16_INTEL 0x7E
#define CLK_AVC_ME_PARTITION_MASK_16x8_INTEL 0x7D
#define CLK_AVC_ME_PARTITION_MASK_8x16_INTEL 0x7B
#define CLK_AVC_ME_PARTITION_MASK_8x8_INTEL 0x77
#define CLK_AVC_ME_PARTITION_MASK_8x4_INTEL 0x6F
#define CLK_AVC_ME_PARTITION_MASK_4x8_INTEL 0x5F
#define CLK_AVC_ME_PARTITION_MASK_4x4_INTEL 0x3F
#define CLK_AVC_ME_SLICE_TYPE_PRED_INTEL 0x0
#define CLK_AVC_ME_SLICE_TYPE_BPRED_INTEL 0x1
#define CLK_AVC_ME_SLICE_TYPE_INTRA_INTEL 0x2
#define CLK_AVC_ME_SEARCH_WINDOW_EXHAUSTIVE_INTEL 0x0
#define CLK_AVC_ME_SEARCH_WINDOW_SMALL_INTEL 0x1
#define CLK_AVC_ME_SEARCH_WINDOW_TINY_INTEL 0x2
#define CLK_AVC_ME_SEARCH_WINDOW_EXTRA_TINY_INTEL 0x3
#define CLK_AVC_ME_SEARCH_WINDOW_DIAMOND_INTEL 0x4
#define CLK_AVC_ME_SEARCH_WINDOW_LARGE_DIAMOND_INTEL 0x5
#define CLK_AVC_ME_SEARCH_WINDOW_RESERVED0_INTEL 0x6
#define CLK_AVC_ME_SEARCH_WINDOW_RESERVED1_INTEL 0x7
#define CLK_AVC_ME_SEARCH_WINDOW_CUSTOM_INTEL 0x8
#define CLK_AVC_ME_SAD_ADJUST_MODE_NONE_INTEL 0x0
#define CLK_AVC_ME_SAD_ADJUST_MODE_HAAR_INTEL 0x2
#define CLK_AVC_ME_SUBPIXEL_MODE_INTEGER_INTEL 0x0
#define CLK_AVC_ME_SUBPIXEL_MODE_HPEL_INTEL 0x1
#define CLK_AVC_ME_SUBPIXEL_MODE_QPEL_INTEL 0x3
#define CLK_AVC_ME_COST_PRECISION_QPEL_INTEL 0x0
#define CLK_AVC_ME_COST_PRECISION_HPEL_INTEL 0x1
#define CLK_AVC_ME_COST_PRECISION_PEL_INTEL 0x2
#define CLK_AVC_ME_COST_PRECISION_DPEL_INTEL 0x3
#define CLK_AVC_ME_BIDIR_WEIGHT_QUARTER_INTEL 0x10
#define CLK_AVC_ME_BIDIR_WEIGHT_THIRD_INTEL 0x15
#define CLK_AVC_ME_BIDIR_WEIGHT_HALF_INTEL 0x20
#define CLK_AVC_ME_BIDIR_WEIGHT_TWO_THIRD_INTEL 0x2B
#define CLK_AVC_ME_BIDIR_WEIGHT_THREE_QUARTER_INTEL 0x30
#define CLK_AVC_ME_BORDER_REACHED_LEFT_INTEL 0x0
#define CLK_AVC_ME_BORDER_REACHED_RIGHT_INTEL 0x2
#define CLK_AVC_ME_BORDER_REACHED_TOP_INTEL 0x4
#define CLK_AVC_ME_BORDER_REACHED_BOTTOM_INTEL 0x8
#define CLK_AVC_ME_INTRA_16x16_INTEL 0x0
#define CLK_AVC_ME_INTRA_8x8_INTEL 0x1
#define CLK_AVC_ME_INTRA_4x4_INTEL 0x2
#define CLK_AVC_ME_SKIP_BLOCK_PARTITION_16x16_INTEL 0x0
#define CLK_AVC_ME_SKIP_BLOCK_PARTITION_8x8_INTEL 0x4000
#define CLK_AVC_ME_SKIP_BLOCK_16x16_FORWARD_ENABLE_INTEL (0x1 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_16x16_BACKWARD_ENABLE_INTEL (0x2 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_16x16_DUAL_ENABLE_INTEL (0x3 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_FORWARD_ENABLE_INTEL (0x55 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_BACKWARD_ENABLE_INTEL (0xAA << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_DUAL_ENABLE_INTEL (0xFF << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_0_FORWARD_ENABLE_INTEL (0x1 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_0_BACKWARD_ENABLE_INTEL (0x2 << 24)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_1_FORWARD_ENABLE_INTEL (0x1 << 26)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_1_BACKWARD_ENABLE_INTEL (0x2 << 26)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_2_FORWARD_ENABLE_INTEL (0x1 << 28)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_2_BACKWARD_ENABLE_INTEL (0x2 << 28)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_3_FORWARD_ENABLE_INTEL (0x1 << 30)
#define CLK_AVC_ME_SKIP_BLOCK_8x8_3_BACKWARD_ENABLE_INTEL (0x2 << 30)
#define CLK_AVC_ME_BLOCK_BASED_SKIP_4x4_INTEL 0x00
#define CLK_AVC_ME_BLOCK_BASED_SKIP_8x8_INTEL 0x80
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_ALL_INTEL 0x0
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_16x16_INTEL 0x6
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_8x8_INTEL 0x5
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_4x4_INTEL 0x3
#define CLK_AVC_ME_INTRA_NEIGHBOR_LEFT_MASK_ENABLE_INTEL 0x60
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_MASK_ENABLE_INTEL 0x10
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_RIGHT_MASK_ENABLE_INTEL 0x8
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_LEFT_MASK_ENABLE_INTEL 0x4
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_INTEL 0x0
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_INTEL 0x1
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DC_INTEL 0x2
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DIAGONAL_DOWN_LEFT_INTEL 0x3
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DIAGONAL_DOWN_RIGHT_INTEL 0x4
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_PLANE_INTEL 0x4
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_RIGHT_INTEL 0x5
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_DOWN_INTEL 0x6
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_LEFT_INTEL 0x7
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_UP_INTEL 0x8
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_DC_INTEL 0x0
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_HORIZONTAL_INTEL 0x1
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_VERTICAL_INTEL 0x2
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_PLANE_INTEL 0x3
#define CLK_AVC_ME_FRAME_FORWARD_INTEL 0x1
#define CLK_AVC_ME_FRAME_BACKWARD_INTEL 0x2
#define CLK_AVC_ME_FRAME_DUAL_INTEL 0x3
#define CLK_AVC_ME_INTERLACED_SCAN_TOP_FIELD_INTEL 0x0
#define CLK_AVC_ME_INTERLACED_SCAN_BOTTOM_FIELD_INTEL 0x1
#define CLK_AVC_ME_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_PAYLOAD_INITIALIZE_INTEL 0x0
#define CLK_AVC_REF_PAYLOAD_INITIALIZE_INTEL 0x0
#define CLK_AVC_SIC_PAYLOAD_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_RESULT_INITIALIZE_INTEL 0x0
#define CLK_AVC_REF_RESULT_INITIALIZE_INTEL 0x0
#define CLK_AVC_SIC_RESULT_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMOUT_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMIN_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMOUT_INITIALIZE_INTEL 0x0
#define CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMIN_INITIALIZE_INTEL 0x0
#endif // cl_intel_device_side_avc_motion_estimation
// Disable any extensions we may have enabled previously.
#pragma OPENCL EXTENSION all : disable
#endif //_OPENCL_BASE_H_
drivers/pocl-gen7/share/pocl/include/opencl-c.h
+18375
View File
File diff suppressed because it is too large. Load diff
drivers/pocl-gen7/share/pocl/include/pocl.h
+617
View File
@@ -0,0 +1,617 @@
/* pocl.h - global pocl declarations for the host side runtime.
Copyright (c) 2011 Universidad Rey Juan Carlos
2011-2019 Pekka Jääskeläinen
2024 Pekka Jääskeläinen / Intel Finland Oy
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to
deal in the Software without restriction, including without limitation the
rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
sell copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
IN THE SOFTWARE.
*/
/**
* @file pocl.h
*
* The declarations in this file are such that are used both in the
* libpocl implementation CL and the kernel compiler. Others should be
* moved to pocl_cl.h of lib/CL or under the kernel compiler dir.
* @todo Check if there are extra declarations here that could be moved.
*/
#ifndef POCL_H
#define POCL_H
#include <stdint.h>
/* The running number used for identifying PoCL runtime allocated objects.
0 marks an invalid/undefined object. */
typedef uint64_t pocl_obj_id_t;
#if defined(EXPORT_POCL_LIB) && defined(_MSC_VER)
/* To successfully export OpenCL API, the dllexport must be placed on
* both the declaration and the definition as MSVC requires that both
* of them have the same attributes.
*
* NOTE: Trouble is expected if CL/opencl.h is not included for the
* first time here.
*/
#undef CL_API_ENTRY
#define CL_API_ENTRY __declspec(dllexport)
#endif
#ifndef CL_TARGET_OPENCL_VERSION
#define CL_TARGET_OPENCL_VERSION 300
#endif
#include <CL/opencl.h>
#include "config.h"
#include "pocl_export.h"
#include "pocl_context.h"
#include "vccompat.hpp"
/* The maximum file, directory and path name lengths.
NOTE: GDB seems to fail to load symbols from .so files which have
longer pathnames than 511, thus the quite small dir/filename length
limiter. */
#define POCL_MAX_DIRNAME_LENGTH 64
#define POCL_MAX_FILENAME_LENGTH (POCL_MAX_DIRNAME_LENGTH)
/* -5: Leave space for .so and for additional .o if temp file debugging
is enabled. */
#define POCL_HASH_FILENAME_LENGTH (POCL_MAX_DIRNAME_LENGTH - 5)
#define POCL_MAX_PATHNAME_LENGTH 4096
/* official Khronos ID */
#ifndef CL_KHRONOS_VENDOR_ID_POCL
#define CL_KHRONOS_VENDOR_ID_POCL 0x10006
#endif
#if __cplusplus >= 201103L
# define POCL_ALIGNAS(x) alignas(x)
#elif __STDC_VERSION__ >= 201112L /* C11 */
# define POCL_ALIGNAS(x) _Alignas(x)
#elif defined(_MSC_VER)
# define POCL_ALIGNAS(x) __declspec(align(x))
#elif defined(__clang__) || defined(__GNUC__)
# define POCL_ALIGNAS(x) __attribute__ ((aligned (x)))
#else
/* Emit an error for potential correctness issues if alignas() is ignored. */
# error "Don't know alignas() counterpart for this compiler!"
#endif
#ifndef _WIN32
#define POCL_PATH_SEPARATOR "/"
#else
#define POCL_PATH_SEPARATOR "\\"
#endif
/* Always align buffer allocations, context data, and kernel arguments to
sizeof(double16) to enforce easier vectorization for targets with naturally
aligned loads/stores. */
#define MAX_EXTENDED_ALIGNMENT 128
/* represents a single buffer to host memory mapping */
typedef struct mem_mapping
{
void *host_ptr; /* the location of the mapped buffer chunk in the host memory
*/
size_t offset; /* offset to the beginning of the buffer */
size_t size;
struct mem_mapping *prev, *next;
/* This is required, because two clEnqueueMap() with the same
buffer+size+offset, will create two identical mappings in the
buffer->mappings LL. Without this flag, both corresponding clEnqUnmap()s
will find the same mapping (the first one in mappings LL), which will lead
to memory double-free corruption later. */
int unmap_requested;
cl_map_flags map_flags;
/* image mapping data */
size_t origin[3];
size_t region[3];
size_t row_pitch;
size_t slice_pitch;
} mem_mapping_t;
/* A single buffer can and will be allocated in
multiple different devices which might or might not share the global memory.
This struct captures a single buffer to global memory allocation's
properties.
*/
typedef struct pocl_mem_identifier
{
/* Global-memory-specific pointer to hardware resource that represents
memory. This may be anything, but must be non-NULL while the memory is
actually allocated, and NULL when it's not */
void *mem_ptr;
/* If mem_ptr represents an address in the device global memory which
is pinned for the lifetime of the buffer. */
int is_pinned;
/* The device-side memory address (if known). If is_pinned is true, this
must be a valid value (note: 0 means invalid address). */
void *device_addr;
/* Content version tracking. Every write use (clEnqWriteBuffer,
* clMapBuffer(CL_MAP_WRITE), write_only image, read/write buffers as kernel
* args etc) increases the version; read uses do not. This version is used in
* implicit migrations to find the latest copy of the data.
*
* In theory, a simple bool of "valid/invalid" could be used;
* the only difference is that a version numbering scheme saves history,
* which could be useful for LRU global memory garbage collection among other
* things in the future.
*/
uint64_t version;
/* Extra pointer for drivers to use for anything.
*
* Currently CUDA uses it to track ALLOC_HOST_PTR allocations.
* Vulkan uses it to store host-mapped staging memory.
*/
void *extra_ptr;
/* Extra integer for drivers to use for anything.
*
* Currently Vulkan uses it to track vulkan memory requirements.
*/
uint64_t extra;
} pocl_mem_identifier;
typedef char pixel_t[16];
typedef struct _mem_destructor_callback mem_destructor_callback_t;
/* represents a memory object destructor callback */
struct _mem_destructor_callback
{
void (CL_CALLBACK * pfn_notify) (cl_mem, void*); /* callback function */
void *user_data; /* user supplied data passed to callback function */
mem_destructor_callback_t *next;
};
typedef struct _build_program_callback build_program_callback_t;
struct _build_program_callback
{
void (CL_CALLBACK * callback_function) (cl_program, void*); /* callback function */
void *user_data; /* user supplied data passed to callback function */
};
// same as SHA1_DIGEST_SIZE
#define POCL_KERNEL_DIGEST_SIZE 20
typedef uint8_t pocl_kernel_hash_t[POCL_KERNEL_DIGEST_SIZE];
/* For clEnqueueNDRangeKernel(). */
typedef struct
{
void *hash;
void *wg; /* The work group function ptr. Device specific. */
cl_kernel kernel;
/* The launch data that can be passed to the kernel execution environment. */
struct pocl_context pc;
struct pocl_argument *arguments;
/* Can be used to store/cache arbitrary device-specific data. */
void *device_data;
/* If set to 1, disallow any work-group function specialization. */
int force_generic_wg_func;
/* If set to 1, disallow "small grid" WG function specialization. */
int force_large_grid_wg_func;
} _cl_command_run;
/* For clEnqueueCommandBufferKHR(). */
typedef struct
{
cl_command_buffer_khr buffer;
} _cl_command_replay;
/* For clEnqueueNativeKernel(). */
typedef struct
{
void *args;
size_t cb_args;
/* The argument buffers are stored in _cl_command's migr_info list. */
void **arg_locs;
void(CL_CALLBACK *user_func) (void *);
} _cl_command_native;
/* For clEnqueueReadBuffer(). */
typedef struct
{
void *__restrict__ dst_host_ptr;
pocl_mem_identifier *src_content_size_mem_id;
size_t offset;
size_t size;
size_t *content_size;
cl_mem src;
} _cl_command_read;
/* For clEnqueueWriteBuffer(). */
typedef struct
{
const void *__restrict__ src_host_ptr;
size_t offset;
size_t size;
cl_mem dst;
} _cl_command_write;
/* For clEnqueueCopyBuffer(). */
typedef struct
{
pocl_mem_identifier *src_content_size_mem_id;
cl_mem src;
cl_mem src_content_size;
cl_mem dst;
size_t src_offset;
size_t dst_offset;
size_t size;
} _cl_command_copy;
/* For clEnqueueReadBufferRect(). */
typedef struct
{
void *__restrict__ dst_host_ptr;
size_t buffer_origin[3];
size_t host_origin[3];
size_t region[3];
size_t buffer_row_pitch;
size_t buffer_slice_pitch;
size_t host_row_pitch;
size_t host_slice_pitch;
cl_mem src;
} _cl_command_read_rect;
/* For clEnqueueWriteBufferRect(). */
typedef struct
{
const void *__restrict__ src_host_ptr;
size_t buffer_origin[3];
size_t host_origin[3];
size_t region[3];
size_t buffer_row_pitch;
size_t buffer_slice_pitch;
size_t host_row_pitch;
size_t host_slice_pitch;
cl_mem dst;
} _cl_command_write_rect;
/* For clEnqueueCopyBufferRect(). */
typedef struct
{
cl_mem src;
cl_mem dst;
size_t dst_origin[3];
size_t src_origin[3];
size_t region[3];
size_t src_row_pitch;
size_t src_slice_pitch;
size_t dst_row_pitch;
size_t dst_slice_pitch;
} _cl_command_copy_rect;
/* For clEnqueueMapBuffer(). */
typedef struct
{
mem_mapping_t *mapping;
cl_mem buffer;
} _cl_command_map;
/* For clEnqueueUnMapMemObject(). */
typedef struct
{
mem_mapping_t *mapping;
cl_mem buffer;
} _cl_command_unmap;
/* For clEnqueueFillBuffer(). */
typedef struct
{
size_t size;
size_t offset;
void *__restrict__ pattern;
size_t pattern_size;
cl_mem dst;
} _cl_command_fill_mem;
/* For clEnqueue(Write/Read)Image(). */
typedef struct
{
void *__restrict__ dst_host_ptr;
cl_mem src;
cl_mem dst;
size_t dst_offset;
size_t origin[3];
size_t region[3];
size_t dst_row_pitch;
size_t dst_slice_pitch;
} _cl_command_read_image;
typedef struct
{
const void *__restrict__ src_host_ptr;
cl_mem src;
cl_mem dst;
size_t src_offset;
size_t origin[3];
size_t region[3];
size_t src_row_pitch;
size_t src_slice_pitch;
} _cl_command_write_image;
typedef struct
{
cl_mem src;
cl_mem dst;
size_t dst_origin[3];
size_t src_origin[3];
size_t region[3];
} _cl_command_copy_image;
/* For clEnqueueFillImage(). */
typedef struct
{
pixel_t fill_pixel;
cl_uint4 orig_pixel;
size_t pixel_size;
size_t origin[3];
size_t region[3];
cl_mem dst;
} _cl_command_fill_image;
/* For clEnqueueMarkerWithWaitlist(). */
typedef struct
{
void *data;
int has_wait_list;
} _cl_command_marker;
/* For clEnqueueBarrierWithWaitlist(). */
typedef _cl_command_marker _cl_command_barrier;
typedef enum pocl_migration_type_e {
ENQUEUE_MIGRATE_TYPE_NOP,
ENQUEUE_MIGRATE_TYPE_D2H,
ENQUEUE_MIGRATE_TYPE_H2D,
ENQUEUE_MIGRATE_TYPE_D2D
} pocl_migration_type_t;
/* For clEnqueueMigrateMemObjects(). */
typedef struct
{
pocl_migration_type_t type;
cl_device_id src_device;
/** For migrating a buffer that has a size buffer as per
* cl_pocl_content_size */
uint64_t migration_size;
pocl_mem_identifier *src_content_size_mem_id;
/* Number of buffers to migrate. The actual buffers are in migr_info in
the _cl_command_node. */
size_t num_buffers;
/* Set to 1 if this is a migration command generated by the runtime to
ensure coherence for the input buffers of commands. */
char implicit;
} _cl_command_migrate;
typedef struct
{
void* data;
cl_command_queue queue;
unsigned num_svm_pointers;
void **svm_pointers;
void (CL_CALLBACK *pfn_free_func) ( cl_command_queue queue,
unsigned num_svm_pointers,
void *svm_pointers[],
void *user_data);
} _cl_command_svm_free;
typedef struct
{
unsigned num_svm_pointers;
size_t *sizes;
void **svm_pointers;
} _cl_command_svm_migrate;
typedef struct
{
void* svm_ptr;
size_t size;
cl_map_flags flags;
} _cl_command_svm_map;
typedef struct
{
void* svm_ptr;
size_t size;
} _cl_command_svm_unmap;
typedef struct
{
const void *__restrict__ src;
void *__restrict__ dst;
size_t size;
} _cl_command_svm_cpy;
typedef struct
{
const void *__restrict__ src;
void *__restrict__ dst;
size_t region[3];
size_t src_origin[3];
size_t dst_origin[3];
size_t src_row_pitch;
size_t src_slice_pitch;
size_t dst_row_pitch;
size_t dst_slice_pitch;
} _cl_command_svm_cpy_rect;
typedef struct
{
void *__restrict__ svm_ptr;
size_t size;
void *__restrict__ pattern;
size_t pattern_size;
} _cl_command_svm_fill;
typedef struct
{
void *__restrict__ svm_ptr;
void *__restrict__ pattern;
size_t region[3];
size_t origin[3];
size_t row_pitch;
size_t slice_pitch;
size_t pattern_size;
} _cl_command_svm_fill_rect;
typedef struct
{
const void *ptr;
size_t size;
cl_mem_advice_intel advice;
} _cl_command_svm_memadvise;
typedef union
{
_cl_command_run run;
_cl_command_native native;
_cl_command_replay replay;
_cl_command_read read;
_cl_command_write write;
_cl_command_copy copy;
_cl_command_read_rect read_rect;
_cl_command_write_rect write_rect;
_cl_command_copy_rect copy_rect;
_cl_command_fill_mem memfill;
_cl_command_read_image read_image;
_cl_command_write_image write_image;
_cl_command_copy_image copy_image;
_cl_command_fill_image fill_image;
_cl_command_map map;
_cl_command_unmap unmap;
_cl_command_marker marker;
_cl_command_barrier barrier;
_cl_command_migrate migrate;
_cl_command_svm_free svm_free;
_cl_command_svm_map svm_map;
_cl_command_svm_unmap svm_unmap;
_cl_command_svm_cpy svm_memcpy;
_cl_command_svm_cpy_rect svm_memcpy_rect;
_cl_command_svm_fill svm_fill;
_cl_command_svm_fill_rect svm_fill_rect;
_cl_command_svm_migrate svm_migrate;
_cl_command_svm_memadvise mem_advise;
} _cl_command_t;
typedef struct _pocl_buffer_migration_info
pocl_buffer_migration_info;
/**
* Enum that represents different possible states of a command node.
*/
typedef enum
{
POCL_COMMAND_FAILED = -1,
/* Default unless set to ready. */
POCL_COMMAND_NOT_READY = 0,
POCL_COMMAND_READY = 1,
} pocl_command_state;
/* One item in a command queue or a command buffer. */
typedef struct _cl_command_node _cl_command_node;
struct _cl_command_node
{
_cl_command_t command;
cl_command_type type;
_cl_command_node *next; // for linked-list storage
_cl_command_node *prev;
/***
* Command buffers use sync points as a template for synchronizing commands
* within the buffer. Commands outside the buffer can't depend on sync points
* and individual commands in the buffer can't depend on events. Because this
* struct is used both for recorded and immediately enqueued commands, the
* two synchronization mechanisms are made mutually exclusive here.
* */
union
{
struct
{
cl_event event;
} event;
struct
{
cl_sync_point_khr sync_point;
cl_uint num_sync_points_in_wait_list;
cl_sync_point_khr *sync_point_wait_list;
} syncpoint;
} sync;
cl_device_id device;
/* The index of the targeted device in the **program** device list. */
unsigned program_device_i;
/* Used to indicate if a node is ready to be executed or failed. */
pocl_command_state state;
/* true if the node belongs to a cl_command_buffer_khr */
cl_int buffered;
/* Which of the command queues in the command buffer's queue list
* this command was recorded for. */
cl_uint queue_idx;
/* List of buffers this command accesses, used for inserting migrations */
pocl_buffer_migration_info *migr_infos;
/* back pointer to the command buffer, when this command is recorded
* in a cl_khr_command_buffer. NULL for other (non-buffered) commands */
cl_command_buffer_khr cmd_buffer;
};
/**
* Enumeration for different modes of converting automatic locals
*/
typedef enum
{
POCL_AUTOLOCALS_TO_ARGS_NEVER = 0,
POCL_AUTOLOCALS_TO_ARGS_ALWAYS = 1,
/* convert autolocals to args only if there are dynamic local memory
* function arguments in the kernel. */
POCL_AUTOLOCALS_TO_ARGS_ONLY_IF_DYNAMIC_LOCALS_PRESENT = 2,
} pocl_autolocals_to_args_strategy;
typedef struct pocl_version_t
{
unsigned major;
unsigned minor;
#ifdef __cplusplus
pocl_version_t () : major{0}, minor{0} {}
pocl_version_t (unsigned the_major, unsigned the_minor)
: major{the_major}, minor{the_minor}
{
}
#endif
} pocl_version_t;
#endif /* POCL_H */
drivers/pocl-gen7/share/pocl/include/pocl_device.h
+78
View File
@@ -0,0 +1,78 @@
/* pocl_device.h - global pocl declarations to be used in the device binaries in
case applicable by the target
Copyright (c) 2012-2018 Pekka Jääskeläinen / Tampere University of Technology
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef POCL_DEVICE_H
#define POCL_DEVICE_H
#include "pocl_types.h"
#define MAX_KERNEL_ARGS 64
#define MAX_KERNEL_NAME_LENGTH 64
/* Metadata of a single kernel stored in the device.*/
typedef struct {
const uchar name[MAX_KERNEL_NAME_LENGTH];
ushort num_args;
ushort num_locals;
void *work_group_func;
} __kernel_metadata;
#ifdef _MSC_VER
#define ALIGN4(x) __declspec(align(4)) x
#define ALIGN8(x) __declspec(align(4)) x
#else
#define ALIGN4(x) x __attribute__ ((aligned (4)))
#define ALIGN8(x) x __attribute__ ((aligned (8)))
#endif
/* A kernel invocation command. */
typedef struct {
/* The kernel to execute. Points to the metadata in the device global
memory. It will be casted to a __kernel_metadata* */
ALIGN8 (uint kernel_meta);
/* Pointer to the kernel arguments in the local memory. Will be
casted to 32 bit void* */
ALIGN8 (uint args);
ALIGN8 (uint args_size);
/* pocl context */
ALIGN8 (uint ctx);
ALIGN8 (uint ctx_size);
/* The execution status of this queue slot. */
ALIGN8 (uint status);
} __kernel_exec_cmd;
/* Kernel execution statuses. */
/* The invocation entry is free to use. */
#define POCL_KST_FREE 1
/* The kernel structure has been populated and is waiting to be
executed. */
#define POCL_KST_READY 2
/* The kernel is currently running in the device. */
#define POCL_KST_RUNNING 3
/* The kernel has finished execution. The results can be collected and the
execution entry be freed (by writing POCL_KST_FREE to the status). */
#define POCL_KST_FINISHED 4
#endif
drivers/pocl-gen7/share/pocl/include/pocl_image_types.h
+52
View File
@@ -0,0 +1,52 @@
/* pocl_image_types.h - image data structure used by device implementations
Copyright (c) 2013 Ville Korhonen
Copyright (c) 2017 Michal Babej / Tampere University of Technology
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef __X86_IMAGE_H__
#define __X86_IMAGE_H__
#ifdef __CBUILD__
#define INTTYPE cl_int
#else
#define INTTYPE int
#endif
typedef uintptr_t dev_sampler_t;
typedef struct dev_image_t {
void *_data;
INTTYPE _width;
INTTYPE _height;
INTTYPE _depth;
INTTYPE _image_array_size;
INTTYPE _row_pitch;
INTTYPE _slice_pitch;
INTTYPE _num_mip_levels; /* maybe not needed */
INTTYPE _num_samples; /* maybe not needed */
INTTYPE _order;
INTTYPE _data_type;
INTTYPE _num_channels;
INTTYPE _elem_size;
} dev_image_t;
#endif
drivers/pocl-gen7/share/pocl/include/pocl_spir.h
+33
View File
@@ -0,0 +1,33 @@
/* pocl-spir.h - global pocl declarations for the SPIR(-V) support.
Copyright (c) 2018-2022 Pekka Jääskeläinen
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to
deal in the Software without restriction, including without limitation the
rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
sell copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
IN THE SOFTWARE.
*/
#ifndef POCL_SPIR_H
#define POCL_SPIR_H
#define SPIR_ADDRESS_SPACE_PRIVATE 0
#define SPIR_ADDRESS_SPACE_GLOBAL 1
#define SPIR_ADDRESS_SPACE_CONSTANT 2
#define SPIR_ADDRESS_SPACE_LOCAL 3
#define SPIR_ADDRESS_SPACE_GENERIC 4
#endif
drivers/pocl-gen7/share/pocl/include/pocl_types.h
+179
View File
@@ -0,0 +1,179 @@
/* pocl_types.h - The basic OpenCL C device side scalar data types.
Copyright (c) 2018 Pekka Jääskeläinen / Tampere University of Technology
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* This header is designed to be included both from the device and the host.
In case compiling OpenCL C sources, __OPENCL_VERSION__ should be set.
In case compiling in the host, all but the device-specific types are
defined (size_t and others). Devices should avoid including the C
stdint.h instead of this one as OpenCL C size_t et al. is allowed to
be of different width than when targeting C.
TODO: replace this header (partially) with Clang's opencl-c.h
*/
#ifndef POCL_DEVICE_TYPES_H
#define POCL_DEVICE_TYPES_H
#ifdef __OPENCL_VERSION__
/* required for LLVM 13+ */
typedef unsigned reserve_id_t;
#ifdef __USE_CLANG_OPENCL_C_H
/* Minimal definitions, only the target specific macro overrides,
just in case Clang export the C ones which might differ for
OpenCL C. */
#ifdef __INTPTR_TYPE__
#undef __INTPTR_TYPE__
#endif
#ifdef __UINTPTR_TYPE__
#undef __UINTPTR_TYPE__
#endif
#ifdef __SIZE_TYPE__
#undef __SIZE_TYPE__
#endif
#ifdef __SIZE_MAX__
#undef __SIZE_MAX__
#endif
#if defined(POCL_DEVICE_ADDRESS_BITS) && POCL_DEVICE_ADDRESS_BITS == 32
#define __SIZE_TYPE__ uint
#define __SIZE_MAX__ UINT_MAX
#define __INTPTR_TYPE__ int
#define __UINTPTR_TYPE__ uint
#else
#define __SIZE_TYPE__ ulong
#define __SIZE_MAX__ ULONG_MAX
#define __INTPTR_TYPE__ long
#define __UINTPTR_TYPE__ ulong
#endif
#else
/* Compiling Device-specific OpenCL C or builtin library C. */
#if defined cl_khr_fp64 && !defined cl_khr_int64
#error "cl_khr_fp64 requires cl_khr_int64"
#endif
/* TODO FIXME We should not use these in OpenCL library's C code at all.
* The problem is that 1) these are predefined by glibc, 2) while we can
* re-define "ulong", we cannot control the size of "long" at all.
* which can lead to "ulong" being 64bit and "long" 32bit, resulting in
* mysterious errors and bugs. Therefore OpenCL library's C code should
* use the fixed size C types where integer size matters. */
#ifdef __CBUILD__
/* Builtin library C code definitions. */
#define size_t csize_t
#define uintptr_t cuintptr_t
#define intptr_t cintptr_t
#include <stdint.h>
#undef size_t
#undef uintptr_t
#undef intptr_t
typedef uint8_t uchar;
typedef uint16_t ushort;
typedef uint32_t uint;
#ifdef cl_khr_int64
#if __APPLE__
typedef unsigned long ulong;
#else
typedef uint64_t ulong;
#endif
#else
typedef uint32_t ulong;
#endif
#endif
/* The definitions below intentionally lead to errors if these types
are used when they are not available in the language. This prevents
accidentally using them if the compiler does not disable these
types, but only e.g. defines them with an incorrect size.*/
#ifndef cl_khr_fp64
typedef struct error_undefined_type_double error_undefined_type_double;
#define double error_undefined_type_double
#endif
#ifdef __SIZE_TYPE__
#undef __SIZE_TYPE__
#endif
#ifdef __SIZE_MAX__
#undef __SIZE_MAX__
#endif
#if defined(POCL_DEVICE_ADDRESS_BITS) && POCL_DEVICE_ADDRESS_BITS == 32
#define __SIZE_TYPE__ uint
#define __SIZE_MAX__ UINT_MAX
#else
#define __SIZE_TYPE__ ulong
#define __SIZE_MAX__ ULONG_MAX
#endif
typedef __SIZE_TYPE__ size_t;
typedef __PTRDIFF_TYPE__ ptrdiff_t;
typedef ptrdiff_t intptr_t;
typedef size_t uintptr_t;
#endif /* #ifdef __USE_CLANG_OPENCL_C_H */
#else /* #ifdef __OPENCL_VERSION__ */
/* Including from a host source (runtime API implementation). Introduce
the fixed width datatypes, but do not override C's size_t and other
target specific datatypes. */
typedef unsigned char uchar;
/* FIXME see the above TODO about these types. */
#if !(defined(_SYS_TYPES_H) && defined(__USE_MISC))
/* glibc, when including sys/types.h, typedefs these. */
typedef unsigned long int ulong;
typedef unsigned short int ushort;
typedef unsigned int uint;
#endif
#include <stdint.h>
#endif
#endif
drivers/pocl-gen7/share/pocl/kernel-x86_64-unknown-linux-gnu-ivybridge.bc
BIN
View File
Binary file not shown.