openharmony/third_party_mesa3d
1
0
Fork 0
mirror of https://gitee.com/openharmony/third_party_mesa3d synced 2024-11-23 23:41:13 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

clover: Import OpenCL state tracker.

Browse Source
This commit is contained in:
Francisco Jerez 2012-04-20 16:56:19 +02:00
parent 309a186987
commit c6db1b3396
51 changed files with 14729 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

31
configure.ac
View File

@ -616,7 +616,11 @@ AC_ARG_ENABLE([va],
[enable va library @<:@default=auto@:>@])],
[enable_va="$enableval"],
[enable_va=auto])
AC_ARG_ENABLE([opencl],
[AS_HELP_STRING([--enable-opencl],
[enable OpenCL library @<:@default=no@:>@])],
[enable_opencl="$enableval"],
[enable_opencl=no])
AC_ARG_ENABLE([xlib_glx],
[AS_HELP_STRING([--enable-xlib-glx],
[make GLX library Xlib-based instead of DRI-based @<:@default=disable@:>@])],
@ -676,7 +680,8 @@ if test "x$enable_opengl" = xno -a \
"x$enable_d3d1x" = xno -a \
"x$enable_xvmc" = xno -a \
"x$enable_vdpau" = xno -a \
"x$enable_va" = xno; then
"x$enable_va" = xno -a \
"x$enable_opencl" = xno; then
AC_MSG_ERROR([at least one API should be enabled])
fi
@ -1603,6 +1608,18 @@ if test "x$enable_va" = xyes; then
HAVE_ST_VA="yes"
fi
dnl
dnl OpenCL configuration
dnl
if test "x$enable_opencl" = xyes; then
if test "x$with_gallium_drivers" = x; then
AC_MSG_ERROR([cannot enable OpenCL without Gallium])
fi
GALLIUM_STATE_TRACKERS_DIRS="$GALLIUM_STATE_TRACKERS_DIRS clover"
GALLIUM_TARGET_DIRS="$GALLIUM_TARGET_DIRS opencl"
fi
dnl
dnl GLU configuration
dnl
@ -1851,6 +1868,14 @@ AC_ARG_WITH([va-libdir],
[VA_LIB_INSTALL_DIR='${libdir}/va'])
AC_SUBST([VA_LIB_INSTALL_DIR])
dnl Directory for OpenCL libs
AC_ARG_WITH([opencl-libdir],
[AS_HELP_STRING([--with-opencl-libdir=DIR],
[directory for the OpenCL libraries @<:@default=${libdir}/opencl@:>@])],
[OPENCL_LIB_INSTALL_DIR="$withval"],
[OPENCL_LIB_INSTALL_DIR='${libdir}/opencl'])
AC_SUBST([OPENCL_LIB_INSTALL_DIR])
dnl
dnl Gallium helper functions
dnl
@ -2039,9 +2064,11 @@ CXXFLAGS="$CXXFLAGS $USER_CXXFLAGS"
dnl Substitute the config
AC_CONFIG_FILES([configs/autoconf
src/gallium/auxiliary/pipe-loader/Makefile
src/gallium/state_trackers/clover/Makefile
src/gallium/drivers/Makefile
src/gallium/drivers/r300/Makefile
src/gallium/drivers/r600/Makefile
src/gallium/targets/opencl/Makefile
src/gbm/Makefile
src/gbm/main/gbm.pc
src/egl/drivers/Makefile

998
include/CL/cl.h Normal file
View File

@ -0,0 +1,998 @@
/*******************************************************************************
* Copyright (c) 2008-2010 The Khronos Group Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and/or associated documentation files (the
* "Materials"), to deal in the Materials without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Materials, and to
* permit persons to whom the Materials are 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 Materials.
*
* THE MATERIALS ARE 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
* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
******************************************************************************/
/* $Revision: 11985 $ on $Date: 2010-07-15 11:16:06 -0700 (Thu, 15 Jul 2010) $ */
#ifndef __OPENCL_CL_H
#define __OPENCL_CL_H
#ifdef __APPLE__
#include <OpenCL/cl_platform.h>
#else
#include <CL/cl_platform.h>
#endif
#ifdef __cplusplus
extern "C" {
#endif
/******************************************************************************/
typedef struct _cl_platform_id * cl_platform_id;
typedef struct _cl_device_id * cl_device_id;
typedef struct _cl_context * cl_context;
typedef struct _cl_command_queue * cl_command_queue;
typedef struct _cl_mem * cl_mem;
typedef struct _cl_program * cl_program;
typedef struct _cl_kernel * cl_kernel;
typedef struct _cl_event * cl_event;
typedef struct _cl_sampler * cl_sampler;
typedef cl_uint cl_bool; /* WARNING! Unlike cl_ types in cl_platform.h, cl_bool is not guaranteed to be the same size as the bool in kernels. */
typedef cl_ulong cl_bitfield;
typedef cl_bitfield cl_device_type;
typedef cl_uint cl_platform_info;
typedef cl_uint cl_device_info;
typedef cl_bitfield cl_device_fp_config;
typedef cl_uint cl_device_mem_cache_type;
typedef cl_uint cl_device_local_mem_type;
typedef cl_bitfield cl_device_exec_capabilities;
typedef cl_bitfield cl_command_queue_properties;
typedef intptr_t cl_context_properties;
typedef cl_uint cl_context_info;
typedef cl_uint cl_command_queue_info;
typedef cl_uint cl_channel_order;
typedef cl_uint cl_channel_type;
typedef cl_bitfield cl_mem_flags;
typedef cl_uint cl_mem_object_type;
typedef cl_uint cl_mem_info;
typedef cl_uint cl_image_info;
typedef cl_uint cl_buffer_create_type;
typedef cl_uint cl_addressing_mode;
typedef cl_uint cl_filter_mode;
typedef cl_uint cl_sampler_info;
typedef cl_bitfield cl_map_flags;
typedef cl_uint cl_program_info;
typedef cl_uint cl_program_build_info;
typedef cl_int cl_build_status;
typedef cl_uint cl_kernel_info;
typedef cl_uint cl_kernel_work_group_info;
typedef cl_uint cl_event_info;
typedef cl_uint cl_command_type;
typedef cl_uint cl_profiling_info;
typedef struct _cl_image_format {
cl_channel_order image_channel_order;
cl_channel_type image_channel_data_type;
} cl_image_format;
typedef struct _cl_buffer_region {
size_t origin;
size_t size;
} cl_buffer_region;
/******************************************************************************/
/* Error Codes */
#define CL_SUCCESS 0
#define CL_DEVICE_NOT_FOUND -1
#define CL_DEVICE_NOT_AVAILABLE -2
#define CL_COMPILER_NOT_AVAILABLE -3
#define CL_MEM_OBJECT_ALLOCATION_FAILURE -4
#define CL_OUT_OF_RESOURCES -5
#define CL_OUT_OF_HOST_MEMORY -6
#define CL_PROFILING_INFO_NOT_AVAILABLE -7
#define CL_MEM_COPY_OVERLAP -8
#define CL_IMAGE_FORMAT_MISMATCH -9
#define CL_IMAGE_FORMAT_NOT_SUPPORTED -10
#define CL_BUILD_PROGRAM_FAILURE -11
#define CL_MAP_FAILURE -12
#define CL_MISALIGNED_SUB_BUFFER_OFFSET -13
#define CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST -14
#define CL_INVALID_VALUE -30
#define CL_INVALID_DEVICE_TYPE -31
#define CL_INVALID_PLATFORM -32
#define CL_INVALID_DEVICE -33
#define CL_INVALID_CONTEXT -34
#define CL_INVALID_QUEUE_PROPERTIES -35
#define CL_INVALID_COMMAND_QUEUE -36
#define CL_INVALID_HOST_PTR -37
#define CL_INVALID_MEM_OBJECT -38
#define CL_INVALID_IMAGE_FORMAT_DESCRIPTOR -39
#define CL_INVALID_IMAGE_SIZE -40
#define CL_INVALID_SAMPLER -41
#define CL_INVALID_BINARY -42
#define CL_INVALID_BUILD_OPTIONS -43
#define CL_INVALID_PROGRAM -44
#define CL_INVALID_PROGRAM_EXECUTABLE -45
#define CL_INVALID_KERNEL_NAME -46
#define CL_INVALID_KERNEL_DEFINITION -47
#define CL_INVALID_KERNEL -48
#define CL_INVALID_ARG_INDEX -49
#define CL_INVALID_ARG_VALUE -50
#define CL_INVALID_ARG_SIZE -51
#define CL_INVALID_KERNEL_ARGS -52
#define CL_INVALID_WORK_DIMENSION -53
#define CL_INVALID_WORK_GROUP_SIZE -54
#define CL_INVALID_WORK_ITEM_SIZE -55
#define CL_INVALID_GLOBAL_OFFSET -56
#define CL_INVALID_EVENT_WAIT_LIST -57
#define CL_INVALID_EVENT -58
#define CL_INVALID_OPERATION -59
#define CL_INVALID_GL_OBJECT -60
#define CL_INVALID_BUFFER_SIZE -61
#define CL_INVALID_MIP_LEVEL -62
#define CL_INVALID_GLOBAL_WORK_SIZE -63
#define CL_INVALID_PROPERTY -64
/* OpenCL Version */
#define CL_VERSION_1_0 1
#define CL_VERSION_1_1 1
/* cl_bool */
#define CL_FALSE 0
#define CL_TRUE 1
/* cl_platform_info */
#define CL_PLATFORM_PROFILE 0x0900
#define CL_PLATFORM_VERSION 0x0901
#define CL_PLATFORM_NAME 0x0902
#define CL_PLATFORM_VENDOR 0x0903
#define CL_PLATFORM_EXTENSIONS 0x0904
/* cl_device_type - bitfield */
#define CL_DEVICE_TYPE_DEFAULT (1 << 0)
#define CL_DEVICE_TYPE_CPU (1 << 1)
#define CL_DEVICE_TYPE_GPU (1 << 2)
#define CL_DEVICE_TYPE_ACCELERATOR (1 << 3)
#define CL_DEVICE_TYPE_ALL 0xFFFFFFFF
/* cl_device_info */
#define CL_DEVICE_TYPE 0x1000
#define CL_DEVICE_VENDOR_ID 0x1001
#define CL_DEVICE_MAX_COMPUTE_UNITS 0x1002
#define CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS 0x1003
#define CL_DEVICE_MAX_WORK_GROUP_SIZE 0x1004
#define CL_DEVICE_MAX_WORK_ITEM_SIZES 0x1005
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR 0x1006
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT 0x1007
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT 0x1008
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG 0x1009
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT 0x100A
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE 0x100B
#define CL_DEVICE_MAX_CLOCK_FREQUENCY 0x100C
#define CL_DEVICE_ADDRESS_BITS 0x100D
#define CL_DEVICE_MAX_READ_IMAGE_ARGS 0x100E
#define CL_DEVICE_MAX_WRITE_IMAGE_ARGS 0x100F
#define CL_DEVICE_MAX_MEM_ALLOC_SIZE 0x1010
#define CL_DEVICE_IMAGE2D_MAX_WIDTH 0x1011
#define CL_DEVICE_IMAGE2D_MAX_HEIGHT 0x1012
#define CL_DEVICE_IMAGE3D_MAX_WIDTH 0x1013
#define CL_DEVICE_IMAGE3D_MAX_HEIGHT 0x1014
#define CL_DEVICE_IMAGE3D_MAX_DEPTH 0x1015
#define CL_DEVICE_IMAGE_SUPPORT 0x1016
#define CL_DEVICE_MAX_PARAMETER_SIZE 0x1017
#define CL_DEVICE_MAX_SAMPLERS 0x1018
#define CL_DEVICE_MEM_BASE_ADDR_ALIGN 0x1019
#define CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE 0x101A
#define CL_DEVICE_SINGLE_FP_CONFIG 0x101B
#define CL_DEVICE_GLOBAL_MEM_CACHE_TYPE 0x101C
#define CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE 0x101D
#define CL_DEVICE_GLOBAL_MEM_CACHE_SIZE 0x101E
#define CL_DEVICE_GLOBAL_MEM_SIZE 0x101F
#define CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE 0x1020
#define CL_DEVICE_MAX_CONSTANT_ARGS 0x1021
#define CL_DEVICE_LOCAL_MEM_TYPE 0x1022
#define CL_DEVICE_LOCAL_MEM_SIZE 0x1023
#define CL_DEVICE_ERROR_CORRECTION_SUPPORT 0x1024
#define CL_DEVICE_PROFILING_TIMER_RESOLUTION 0x1025
#define CL_DEVICE_ENDIAN_LITTLE 0x1026
#define CL_DEVICE_AVAILABLE 0x1027
#define CL_DEVICE_COMPILER_AVAILABLE 0x1028
#define CL_DEVICE_EXECUTION_CAPABILITIES 0x1029
#define CL_DEVICE_QUEUE_PROPERTIES 0x102A
#define CL_DEVICE_NAME 0x102B
#define CL_DEVICE_VENDOR 0x102C
#define CL_DRIVER_VERSION 0x102D
#define CL_DEVICE_PROFILE 0x102E
#define CL_DEVICE_VERSION 0x102F
#define CL_DEVICE_EXTENSIONS 0x1030
#define CL_DEVICE_PLATFORM 0x1031
/* 0x1032 reserved for CL_DEVICE_DOUBLE_FP_CONFIG */
/* 0x1033 reserved for CL_DEVICE_HALF_FP_CONFIG */
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF 0x1034
#define CL_DEVICE_HOST_UNIFIED_MEMORY 0x1035
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR 0x1036
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT 0x1037
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_INT 0x1038
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG 0x1039
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT 0x103A
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE 0x103B
#define CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF 0x103C
#define CL_DEVICE_OPENCL_C_VERSION 0x103D
/* cl_device_fp_config - bitfield */
#define CL_FP_DENORM (1 << 0)
#define CL_FP_INF_NAN (1 << 1)
#define CL_FP_ROUND_TO_NEAREST (1 << 2)
#define CL_FP_ROUND_TO_ZERO (1 << 3)
#define CL_FP_ROUND_TO_INF (1 << 4)
#define CL_FP_FMA (1 << 5)
#define CL_FP_SOFT_FLOAT (1 << 6)
/* cl_device_mem_cache_type */
#define CL_NONE 0x0
#define CL_READ_ONLY_CACHE 0x1
#define CL_READ_WRITE_CACHE 0x2
/* cl_device_local_mem_type */
#define CL_LOCAL 0x1
#define CL_GLOBAL 0x2
/* cl_device_exec_capabilities - bitfield */
#define CL_EXEC_KERNEL (1 << 0)
#define CL_EXEC_NATIVE_KERNEL (1 << 1)
/* cl_command_queue_properties - bitfield */
#define CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE (1 << 0)
#define CL_QUEUE_PROFILING_ENABLE (1 << 1)
/* cl_context_info */
#define CL_CONTEXT_REFERENCE_COUNT 0x1080
#define CL_CONTEXT_DEVICES 0x1081
#define CL_CONTEXT_PROPERTIES 0x1082
#define CL_CONTEXT_NUM_DEVICES 0x1083
/* cl_context_info + cl_context_properties */
#define CL_CONTEXT_PLATFORM 0x1084
/* cl_command_queue_info */
#define CL_QUEUE_CONTEXT 0x1090
#define CL_QUEUE_DEVICE 0x1091
#define CL_QUEUE_REFERENCE_COUNT 0x1092
#define CL_QUEUE_PROPERTIES 0x1093
/* cl_mem_flags - bitfield */
#define CL_MEM_READ_WRITE (1 << 0)
#define CL_MEM_WRITE_ONLY (1 << 1)
#define CL_MEM_READ_ONLY (1 << 2)
#define CL_MEM_USE_HOST_PTR (1 << 3)
#define CL_MEM_ALLOC_HOST_PTR (1 << 4)
#define CL_MEM_COPY_HOST_PTR (1 << 5)
/* cl_channel_order */
#define CL_R 0x10B0
#define CL_A 0x10B1
#define CL_RG 0x10B2
#define CL_RA 0x10B3
#define CL_RGB 0x10B4
#define CL_RGBA 0x10B5
#define CL_BGRA 0x10B6
#define CL_ARGB 0x10B7
#define CL_INTENSITY 0x10B8
#define CL_LUMINANCE 0x10B9
#define CL_Rx 0x10BA
#define CL_RGx 0x10BB
#define CL_RGBx 0x10BC
/* cl_channel_type */
#define CL_SNORM_INT8 0x10D0
#define CL_SNORM_INT16 0x10D1
#define CL_UNORM_INT8 0x10D2
#define CL_UNORM_INT16 0x10D3
#define CL_UNORM_SHORT_565 0x10D4
#define CL_UNORM_SHORT_555 0x10D5
#define CL_UNORM_INT_101010 0x10D6
#define CL_SIGNED_INT8 0x10D7
#define CL_SIGNED_INT16 0x10D8
#define CL_SIGNED_INT32 0x10D9
#define CL_UNSIGNED_INT8 0x10DA
#define CL_UNSIGNED_INT16 0x10DB
#define CL_UNSIGNED_INT32 0x10DC
#define CL_HALF_FLOAT 0x10DD
#define CL_FLOAT 0x10DE
/* cl_mem_object_type */
#define CL_MEM_OBJECT_BUFFER 0x10F0
#define CL_MEM_OBJECT_IMAGE2D 0x10F1
#define CL_MEM_OBJECT_IMAGE3D 0x10F2
/* cl_mem_info */
#define CL_MEM_TYPE 0x1100
#define CL_MEM_FLAGS 0x1101
#define CL_MEM_SIZE 0x1102
#define CL_MEM_HOST_PTR 0x1103
#define CL_MEM_MAP_COUNT 0x1104
#define CL_MEM_REFERENCE_COUNT 0x1105
#define CL_MEM_CONTEXT 0x1106
#define CL_MEM_ASSOCIATED_MEMOBJECT 0x1107
#define CL_MEM_OFFSET 0x1108
/* cl_image_info */
#define CL_IMAGE_FORMAT 0x1110
#define CL_IMAGE_ELEMENT_SIZE 0x1111
#define CL_IMAGE_ROW_PITCH 0x1112
#define CL_IMAGE_SLICE_PITCH 0x1113
#define CL_IMAGE_WIDTH 0x1114
#define CL_IMAGE_HEIGHT 0x1115
#define CL_IMAGE_DEPTH 0x1116
/* cl_addressing_mode */
#define CL_ADDRESS_NONE 0x1130
#define CL_ADDRESS_CLAMP_TO_EDGE 0x1131
#define CL_ADDRESS_CLAMP 0x1132
#define CL_ADDRESS_REPEAT 0x1133
#define CL_ADDRESS_MIRRORED_REPEAT 0x1134
/* cl_filter_mode */
#define CL_FILTER_NEAREST 0x1140
#define CL_FILTER_LINEAR 0x1141
/* cl_sampler_info */
#define CL_SAMPLER_REFERENCE_COUNT 0x1150
#define CL_SAMPLER_CONTEXT 0x1151
#define CL_SAMPLER_NORMALIZED_COORDS 0x1152
#define CL_SAMPLER_ADDRESSING_MODE 0x1153
#define CL_SAMPLER_FILTER_MODE 0x1154
/* cl_map_flags - bitfield */
#define CL_MAP_READ (1 << 0)
#define CL_MAP_WRITE (1 << 1)
/* cl_program_info */
#define CL_PROGRAM_REFERENCE_COUNT 0x1160
#define CL_PROGRAM_CONTEXT 0x1161
#define CL_PROGRAM_NUM_DEVICES 0x1162
#define CL_PROGRAM_DEVICES 0x1163
#define CL_PROGRAM_SOURCE 0x1164
#define CL_PROGRAM_BINARY_SIZES 0x1165
#define CL_PROGRAM_BINARIES 0x1166
/* cl_program_build_info */
#define CL_PROGRAM_BUILD_STATUS 0x1181
#define CL_PROGRAM_BUILD_OPTIONS 0x1182
#define CL_PROGRAM_BUILD_LOG 0x1183
/* cl_build_status */
#define CL_BUILD_SUCCESS 0
#define CL_BUILD_NONE -1
#define CL_BUILD_ERROR -2
#define CL_BUILD_IN_PROGRESS -3
/* cl_kernel_info */
#define CL_KERNEL_FUNCTION_NAME 0x1190
#define CL_KERNEL_NUM_ARGS 0x1191
#define CL_KERNEL_REFERENCE_COUNT 0x1192
#define CL_KERNEL_CONTEXT 0x1193
#define CL_KERNEL_PROGRAM 0x1194
/* cl_kernel_work_group_info */
#define CL_KERNEL_WORK_GROUP_SIZE 0x11B0
#define CL_KERNEL_COMPILE_WORK_GROUP_SIZE 0x11B1
#define CL_KERNEL_LOCAL_MEM_SIZE 0x11B2
#define CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 0x11B3
#define CL_KERNEL_PRIVATE_MEM_SIZE 0x11B4
/* cl_event_info */
#define CL_EVENT_COMMAND_QUEUE 0x11D0
#define CL_EVENT_COMMAND_TYPE 0x11D1
#define CL_EVENT_REFERENCE_COUNT 0x11D2
#define CL_EVENT_COMMAND_EXECUTION_STATUS 0x11D3
#define CL_EVENT_CONTEXT 0x11D4
/* cl_command_type */
#define CL_COMMAND_NDRANGE_KERNEL 0x11F0
#define CL_COMMAND_TASK 0x11F1
#define CL_COMMAND_NATIVE_KERNEL 0x11F2
#define CL_COMMAND_READ_BUFFER 0x11F3
#define CL_COMMAND_WRITE_BUFFER 0x11F4
#define CL_COMMAND_COPY_BUFFER 0x11F5
#define CL_COMMAND_READ_IMAGE 0x11F6
#define CL_COMMAND_WRITE_IMAGE 0x11F7
#define CL_COMMAND_COPY_IMAGE 0x11F8
#define CL_COMMAND_COPY_IMAGE_TO_BUFFER 0x11F9
#define CL_COMMAND_COPY_BUFFER_TO_IMAGE 0x11FA
#define CL_COMMAND_MAP_BUFFER 0x11FB
#define CL_COMMAND_MAP_IMAGE 0x11FC
#define CL_COMMAND_UNMAP_MEM_OBJECT 0x11FD
#define CL_COMMAND_MARKER 0x11FE
#define CL_COMMAND_ACQUIRE_GL_OBJECTS 0x11FF
#define CL_COMMAND_RELEASE_GL_OBJECTS 0x1200
#define CL_COMMAND_READ_BUFFER_RECT 0x1201
#define CL_COMMAND_WRITE_BUFFER_RECT 0x1202
#define CL_COMMAND_COPY_BUFFER_RECT 0x1203
#define CL_COMMAND_USER 0x1204
/* command execution status */
#define CL_COMPLETE 0x0
#define CL_RUNNING 0x1
#define CL_SUBMITTED 0x2
#define CL_QUEUED 0x3
/* cl_buffer_create_type */
#define CL_BUFFER_CREATE_TYPE_REGION 0x1220
/* cl_profiling_info */
#define CL_PROFILING_COMMAND_QUEUED 0x1280
#define CL_PROFILING_COMMAND_SUBMIT 0x1281
#define CL_PROFILING_COMMAND_START 0x1282
#define CL_PROFILING_COMMAND_END 0x1283
/********************************************************************************************************/
/* Platform API */
extern CL_API_ENTRY cl_int CL_API_CALL
clGetPlatformIDs(cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetPlatformInfo(cl_platform_id /* platform */,
cl_platform_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Device APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clGetDeviceIDs(cl_platform_id /* platform */,
cl_device_type /* device_type */,
cl_uint /* num_entries */,
cl_device_id * /* devices */,
cl_uint * /* num_devices */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetDeviceInfo(cl_device_id /* device */,
cl_device_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Context APIs */
extern CL_API_ENTRY cl_context CL_API_CALL
clCreateContext(const cl_context_properties * /* properties */,
cl_uint /* num_devices */,
const cl_device_id * /* devices */,
void (CL_CALLBACK * /* pfn_notify */)(const char *, const void *, size_t, void *),
void * /* user_data */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_context CL_API_CALL
clCreateContextFromType(const cl_context_properties * /* properties */,
cl_device_type /* device_type */,
void (CL_CALLBACK * /* pfn_notify*/ )(const char *, const void *, size_t, void *),
void * /* user_data */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetContextInfo(cl_context /* context */,
cl_context_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Command Queue APIs */
extern CL_API_ENTRY cl_command_queue CL_API_CALL
clCreateCommandQueue(cl_context /* context */,
cl_device_id /* device */,
cl_command_queue_properties /* properties */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetCommandQueueInfo(cl_command_queue /* command_queue */,
cl_command_queue_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
#ifdef CL_USE_DEPRECATED_OPENCL_1_0_APIS
#warning CL_USE_DEPRECATED_OPENCL_1_0_APIS is defined. These APIs are unsupported and untested in OpenCL 1.1!
/*
* WARNING:
* This API introduces mutable state into the OpenCL implementation. It has been REMOVED
* to better facilitate thread safety. The 1.0 API is not thread safe. It is not tested by the
* OpenCL 1.1 conformance test, and consequently may not work or may not work dependably.
* It is likely to be non-performant. Use of this API is not advised. Use at your own risk.
*
* Software developers previously relying on this API are instructed to set the command queue
* properties when creating the queue, instead.
*/
extern CL_API_ENTRY cl_int CL_API_CALL
clSetCommandQueueProperty(cl_command_queue /* command_queue */,
cl_command_queue_properties /* properties */,
cl_bool /* enable */,
cl_command_queue_properties * /* old_properties */) CL_EXT_SUFFIX__VERSION_1_0_DEPRECATED;
#endif /* CL_USE_DEPRECATED_OPENCL_1_0_APIS */
/* Memory Object APIs */
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateBuffer(cl_context /* context */,
cl_mem_flags /* flags */,
size_t /* size */,
void * /* host_ptr */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateSubBuffer(cl_mem /* buffer */,
cl_mem_flags /* flags */,
cl_buffer_create_type /* buffer_create_type */,
const void * /* buffer_create_info */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateImage2D(cl_context /* context */,
cl_mem_flags /* flags */,
const cl_image_format * /* image_format */,
size_t /* image_width */,
size_t /* image_height */,
size_t /* image_row_pitch */,
void * /* host_ptr */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateImage3D(cl_context /* context */,
cl_mem_flags /* flags */,
const cl_image_format * /* image_format */,
size_t /* image_width */,
size_t /* image_height */,
size_t /* image_depth */,
size_t /* image_row_pitch */,
size_t /* image_slice_pitch */,
void * /* host_ptr */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetSupportedImageFormats(cl_context /* context */,
cl_mem_flags /* flags */,
cl_mem_object_type /* image_type */,
cl_uint /* num_entries */,
cl_image_format * /* image_formats */,
cl_uint * /* num_image_formats */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetMemObjectInfo(cl_mem /* memobj */,
cl_mem_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetImageInfo(cl_mem /* image */,
cl_image_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clSetMemObjectDestructorCallback( cl_mem /* memobj */,
void (CL_CALLBACK * /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/),
void * /*user_data */ ) CL_API_SUFFIX__VERSION_1_1;
/* Sampler APIs */
extern CL_API_ENTRY cl_sampler CL_API_CALL
clCreateSampler(cl_context /* context */,
cl_bool /* normalized_coords */,
cl_addressing_mode /* addressing_mode */,
cl_filter_mode /* filter_mode */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetSamplerInfo(cl_sampler /* sampler */,
cl_sampler_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Program Object APIs */
extern CL_API_ENTRY cl_program CL_API_CALL
clCreateProgramWithSource(cl_context /* context */,
cl_uint /* count */,
const char ** /* strings */,
const size_t * /* lengths */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_program CL_API_CALL
clCreateProgramWithBinary(cl_context /* context */,
cl_uint /* num_devices */,
const cl_device_id * /* device_list */,
const size_t * /* lengths */,
const unsigned char ** /* binaries */,
cl_int * /* binary_status */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clBuildProgram(cl_program /* program */,
cl_uint /* num_devices */,
const cl_device_id * /* device_list */,
const char * /* options */,
void (CL_CALLBACK * /* pfn_notify */)(cl_program /* program */, void * /* user_data */),
void * /* user_data */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clUnloadCompiler(void) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetProgramInfo(cl_program /* program */,
cl_program_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetProgramBuildInfo(cl_program /* program */,
cl_device_id /* device */,
cl_program_build_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Kernel Object APIs */
extern CL_API_ENTRY cl_kernel CL_API_CALL
clCreateKernel(cl_program /* program */,
const char * /* kernel_name */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clCreateKernelsInProgram(cl_program /* program */,
cl_uint /* num_kernels */,
cl_kernel * /* kernels */,
cl_uint * /* num_kernels_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainKernel(cl_kernel /* kernel */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseKernel(cl_kernel /* kernel */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clSetKernelArg(cl_kernel /* kernel */,
cl_uint /* arg_index */,
size_t /* arg_size */,
const void * /* arg_value */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetKernelInfo(cl_kernel /* kernel */,
cl_kernel_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetKernelWorkGroupInfo(cl_kernel /* kernel */,
cl_device_id /* device */,
cl_kernel_work_group_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Event Object APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clWaitForEvents(cl_uint /* num_events */,
const cl_event * /* event_list */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetEventInfo(cl_event /* event */,
cl_event_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_event CL_API_CALL
clCreateUserEvent(cl_context /* context */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clSetUserEventStatus(cl_event /* event */,
cl_int /* execution_status */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clSetEventCallback( cl_event /* event */,
cl_int /* command_exec_callback_type */,
void (CL_CALLBACK * /* pfn_notify */)(cl_event, cl_int, void *),
void * /* user_data */) CL_API_SUFFIX__VERSION_1_1;
/* Profiling APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clGetEventProfilingInfo(cl_event /* event */,
cl_profiling_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
/* Flush and Finish APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clFlush(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clFinish(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
/* Enqueued Commands APIs */
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueReadBuffer(cl_command_queue /* command_queue */,
cl_mem /* buffer */,
cl_bool /* blocking_read */,
size_t /* offset */,
size_t /* cb */,
void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueReadBufferRect(cl_command_queue /* command_queue */,
cl_mem /* buffer */,
cl_bool /* blocking_read */,
const size_t * /* buffer_origin */,
const size_t * /* host_origin */,
const size_t * /* region */,
size_t /* buffer_row_pitch */,
size_t /* buffer_slice_pitch */,
size_t /* host_row_pitch */,
size_t /* host_slice_pitch */,
void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueWriteBuffer(cl_command_queue /* command_queue */,
cl_mem /* buffer */,
cl_bool /* blocking_write */,
size_t /* offset */,
size_t /* cb */,
const void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueWriteBufferRect(cl_command_queue /* command_queue */,
cl_mem /* buffer */,
cl_bool /* blocking_write */,
const size_t * /* buffer_origin */,
const size_t * /* host_origin */,
const size_t * /* region */,
size_t /* buffer_row_pitch */,
size_t /* buffer_slice_pitch */,
size_t /* host_row_pitch */,
size_t /* host_slice_pitch */,
const void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyBuffer(cl_command_queue /* command_queue */,
cl_mem /* src_buffer */,
cl_mem /* dst_buffer */,
size_t /* src_offset */,
size_t /* dst_offset */,
size_t /* cb */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyBufferRect(cl_command_queue /* command_queue */,
cl_mem /* src_buffer */,
cl_mem /* dst_buffer */,
const size_t * /* src_origin */,
const size_t * /* dst_origin */,
const size_t * /* region */,
size_t /* src_row_pitch */,
size_t /* src_slice_pitch */,
size_t /* dst_row_pitch */,
size_t /* dst_slice_pitch */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueReadImage(cl_command_queue /* command_queue */,
cl_mem /* image */,
cl_bool /* blocking_read */,
const size_t * /* origin[3] */,
const size_t * /* region[3] */,
size_t /* row_pitch */,
size_t /* slice_pitch */,
void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueWriteImage(cl_command_queue /* command_queue */,
cl_mem /* image */,
cl_bool /* blocking_write */,
const size_t * /* origin[3] */,
const size_t * /* region[3] */,
size_t /* input_row_pitch */,
size_t /* input_slice_pitch */,
const void * /* ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyImage(cl_command_queue /* command_queue */,
cl_mem /* src_image */,
cl_mem /* dst_image */,
const size_t * /* src_origin[3] */,
const size_t * /* dst_origin[3] */,
const size_t * /* region[3] */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyImageToBuffer(cl_command_queue /* command_queue */,
cl_mem /* src_image */,
cl_mem /* dst_buffer */,
const size_t * /* src_origin[3] */,
const size_t * /* region[3] */,
size_t /* dst_offset */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueCopyBufferToImage(cl_command_queue /* command_queue */,
cl_mem /* src_buffer */,
cl_mem /* dst_image */,
size_t /* src_offset */,
const size_t * /* dst_origin[3] */,
const size_t * /* region[3] */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY void * CL_API_CALL
clEnqueueMapBuffer(cl_command_queue /* command_queue */,
cl_mem /* buffer */,
cl_bool /* blocking_map */,
cl_map_flags /* map_flags */,
size_t /* offset */,
size_t /* cb */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY void * CL_API_CALL
clEnqueueMapImage(cl_command_queue /* command_queue */,
cl_mem /* image */,
cl_bool /* blocking_map */,
cl_map_flags /* map_flags */,
const size_t * /* origin[3] */,
const size_t * /* region[3] */,
size_t * /* image_row_pitch */,
size_t * /* image_slice_pitch */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueUnmapMemObject(cl_command_queue /* command_queue */,
cl_mem /* memobj */,
void * /* mapped_ptr */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueNDRangeKernel(cl_command_queue /* command_queue */,
cl_kernel /* kernel */,
cl_uint /* work_dim */,
const size_t * /* global_work_offset */,
const size_t * /* global_work_size */,
const size_t * /* local_work_size */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueTask(cl_command_queue /* command_queue */,
cl_kernel /* kernel */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueNativeKernel(cl_command_queue /* command_queue */,
void (*user_func)(void *),
void * /* args */,
size_t /* cb_args */,
cl_uint /* num_mem_objects */,
const cl_mem * /* mem_list */,
const void ** /* args_mem_loc */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueMarker(cl_command_queue /* command_queue */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueWaitForEvents(cl_command_queue /* command_queue */,
cl_uint /* num_events */,
const cl_event * /* event_list */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueBarrier(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
/* Extension function access
*
* Returns the extension function address for the given function name,
* or NULL if a valid function can not be found. The client must
* check to make sure the address is not NULL, before using or
* calling the returned function address.
*/
extern CL_API_ENTRY void * CL_API_CALL clGetExtensionFunctionAddress(const char * /* func_name */) CL_API_SUFFIX__VERSION_1_0;
#ifdef __cplusplus
}
#endif
#endif /* __OPENCL_CL_H */

4011
include/CL/cl.hpp Normal file
View File

File diff suppressed because it is too large Load Diff

213
include/CL/cl_ext.h Normal file
View File

@ -0,0 +1,213 @@
/*******************************************************************************
* Copyright (c) 2008-2010 The Khronos Group Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and/or associated documentation files (the
* "Materials"), to deal in the Materials without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Materials, and to
* permit persons to whom the Materials are 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 Materials.
*
* THE MATERIALS ARE 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
* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
******************************************************************************/
/* $Revision: 11928 $ on $Date: 2010-07-13 09:04:56 -0700 (Tue, 13 Jul 2010) $ */
/* cl_ext.h contains OpenCL extensions which don't have external */
/* (OpenGL, D3D) dependencies. */
#ifndef __CL_EXT_H
#define __CL_EXT_H
#ifdef __cplusplus
extern "C" {
#endif
#ifdef __APPLE__
#include <OpenCL/cl.h>
#include <AvailabilityMacros.h>
#else
#include <CL/cl.h>
#endif
/* cl_khr_fp64 extension - no extension #define since it has no functions */
#define CL_DEVICE_DOUBLE_FP_CONFIG 0x1032
/* cl_khr_fp16 extension - no extension #define since it has no functions */
#define CL_DEVICE_HALF_FP_CONFIG 0x1033
/* Memory object destruction
*
* Apple extension for use to manage externally allocated buffers used with cl_mem objects with CL_MEM_USE_HOST_PTR
*
* Registers a user callback function that will be called when the memory object is deleted and its resources
* freed. Each call to clSetMemObjectCallbackFn registers the specified user callback function on a callback
* stack associated with memobj. The registered user callback functions are called in the reverse order in
* which they were registered. The user callback functions are called and then the memory object is deleted
* and its resources freed. This provides a mechanism for the application (and libraries) using memobj to be
* notified when the memory referenced by host_ptr, specified when the memory object is created and used as
* the storage bits for the memory object, can be reused or freed.
*
* The application may not call CL api's with the cl_mem object passed to the pfn_notify.
*
* Please check for the "cl_APPLE_SetMemObjectDestructor" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
* before using.
*/
#define cl_APPLE_SetMemObjectDestructor 1
cl_int CL_API_ENTRY clSetMemObjectDestructorAPPLE( cl_mem /* memobj */,
void (* /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/),
void * /*user_data */ ) CL_EXT_SUFFIX__VERSION_1_0;
/* Context Logging Functions
*
* The next three convenience functions are intended to be used as the pfn_notify parameter to clCreateContext().
* Please check for the "cl_APPLE_ContextLoggingFunctions" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
* before using.
*
* clLogMessagesToSystemLog fowards on all log messages to the Apple System Logger
*/
#define cl_APPLE_ContextLoggingFunctions 1
extern void CL_API_ENTRY clLogMessagesToSystemLogAPPLE( const char * /* errstr */,
const void * /* private_info */,
size_t /* cb */,
void * /* user_data */ ) CL_EXT_SUFFIX__VERSION_1_0;
/* clLogMessagesToStdout sends all log messages to the file descriptor stdout */
extern void CL_API_ENTRY clLogMessagesToStdoutAPPLE( const char * /* errstr */,
const void * /* private_info */,
size_t /* cb */,
void * /* user_data */ ) CL_EXT_SUFFIX__VERSION_1_0;
/* clLogMessagesToStderr sends all log messages to the file descriptor stderr */
extern void CL_API_ENTRY clLogMessagesToStderrAPPLE( const char * /* errstr */,
const void * /* private_info */,
size_t /* cb */,
void * /* user_data */ ) CL_EXT_SUFFIX__VERSION_1_0;
/************************
* cl_khr_icd extension *
************************/
#define cl_khr_icd 1
/* cl_platform_info */
#define CL_PLATFORM_ICD_SUFFIX_KHR 0x0920
/* Additional Error Codes */
#define CL_PLATFORM_NOT_FOUND_KHR -1001
extern CL_API_ENTRY cl_int CL_API_CALL
clIcdGetPlatformIDsKHR(cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */);
typedef CL_API_ENTRY cl_int (CL_API_CALL *clIcdGetPlatformIDsKHR_fn)(
cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */);
/******************************************
* cl_nv_device_attribute_query extension *
******************************************/
/* cl_nv_device_attribute_query extension - no extension #define since it has no functions */
#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000
#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001
#define CL_DEVICE_REGISTERS_PER_BLOCK_NV 0x4002
#define CL_DEVICE_WARP_SIZE_NV 0x4003
#define CL_DEVICE_GPU_OVERLAP_NV 0x4004
#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005
#define CL_DEVICE_INTEGRATED_MEMORY_NV 0x4006
/*********************************
* cl_amd_device_attribute_query *
*********************************/
#define CL_DEVICE_PROFILING_TIMER_OFFSET_AMD 0x4036
#ifdef CL_VERSION_1_1
/***********************************
* cl_ext_device_fission extension *
***********************************/
#define cl_ext_device_fission 1
extern CL_API_ENTRY cl_int CL_API_CALL
clReleaseDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *clReleaseDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
extern CL_API_ENTRY cl_int CL_API_CALL
clRetainDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *clRetainDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
typedef cl_ulong cl_device_partition_property_ext;
extern CL_API_ENTRY cl_int CL_API_CALL
clCreateSubDevicesEXT( cl_device_id /*in_device*/,
const cl_device_partition_property_ext * /* properties */,
cl_uint /*num_entries*/,
cl_device_id * /*out_devices*/,
cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
typedef CL_API_ENTRY cl_int
( CL_API_CALL * clCreateSubDevicesEXT_fn)( cl_device_id /*in_device*/,
const cl_device_partition_property_ext * /* properties */,
cl_uint /*num_entries*/,
cl_device_id * /*out_devices*/,
cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
/* cl_device_partition_property_ext */
#define CL_DEVICE_PARTITION_EQUALLY_EXT 0x4050
#define CL_DEVICE_PARTITION_BY_COUNTS_EXT 0x4051
#define CL_DEVICE_PARTITION_BY_NAMES_EXT 0x4052
#define CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT 0x4053
/* clDeviceGetInfo selectors */
#define CL_DEVICE_PARENT_DEVICE_EXT 0x4054
#define CL_DEVICE_PARTITION_TYPES_EXT 0x4055
#define CL_DEVICE_AFFINITY_DOMAINS_EXT 0x4056
#define CL_DEVICE_REFERENCE_COUNT_EXT 0x4057
#define CL_DEVICE_PARTITION_STYLE_EXT 0x4058
/* error codes */
#define CL_DEVICE_PARTITION_FAILED_EXT -1057
#define CL_INVALID_PARTITION_COUNT_EXT -1058
#define CL_INVALID_PARTITION_NAME_EXT -1059
/* CL_AFFINITY_DOMAINs */
#define CL_AFFINITY_DOMAIN_L1_CACHE_EXT 0x1
#define CL_AFFINITY_DOMAIN_L2_CACHE_EXT 0x2
#define CL_AFFINITY_DOMAIN_L3_CACHE_EXT 0x3
#define CL_AFFINITY_DOMAIN_L4_CACHE_EXT 0x4
#define CL_AFFINITY_DOMAIN_NUMA_EXT 0x10
#define CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT 0x100
/* cl_device_partition_property_ext list terminators */
#define CL_PROPERTIES_LIST_END_EXT ((cl_device_partition_property_ext) 0)
#define CL_PARTITION_BY_COUNTS_LIST_END_EXT ((cl_device_partition_property_ext) 0)
#define CL_PARTITION_BY_NAMES_LIST_END_EXT ((cl_device_partition_property_ext) 0 - 1)
#endif /* CL_VERSION_1_1 */
#ifdef __cplusplus
}
#endif
#endif /* __CL_EXT_H */

155
include/CL/cl_gl.h Normal file
View File

@ -0,0 +1,155 @@
/**********************************************************************************
* Copyright (c) 2008-2010 The Khronos Group Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and/or associated documentation files (the
* "Materials"), to deal in the Materials without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Materials, and to
* permit persons to whom the Materials are 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 Materials.
*
* THE MATERIALS ARE 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
* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
**********************************************************************************/
/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
/*
* cl_gl.h contains Khronos-approved (KHR) OpenCL extensions which have
* OpenGL dependencies. The application is responsible for #including
* OpenGL or OpenGL ES headers before #including cl_gl.h.
*/
#ifndef __OPENCL_CL_GL_H
#define __OPENCL_CL_GL_H
#ifdef __APPLE__
#include <OpenCL/cl.h>
#include <OpenGL/CGLDevice.h>
#else
#include <CL/cl.h>
#endif
#ifdef __cplusplus
extern "C" {
#endif
typedef cl_uint cl_gl_object_type;
typedef cl_uint cl_gl_texture_info;
typedef cl_uint cl_gl_platform_info;
typedef struct __GLsync *cl_GLsync;
/* cl_gl_object_type */
#define CL_GL_OBJECT_BUFFER 0x2000
#define CL_GL_OBJECT_TEXTURE2D 0x2001
#define CL_GL_OBJECT_TEXTURE3D 0x2002
#define CL_GL_OBJECT_RENDERBUFFER 0x2003
/* cl_gl_texture_info */
#define CL_GL_TEXTURE_TARGET 0x2004
#define CL_GL_MIPMAP_LEVEL 0x2005
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateFromGLBuffer(cl_context /* context */,
cl_mem_flags /* flags */,
cl_GLuint /* bufobj */,
int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateFromGLTexture2D(cl_context /* context */,
cl_mem_flags /* flags */,
cl_GLenum /* target */,
cl_GLint /* miplevel */,
cl_GLuint /* texture */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateFromGLTexture3D(cl_context /* context */,
cl_mem_flags /* flags */,
cl_GLenum /* target */,
cl_GLint /* miplevel */,
cl_GLuint /* texture */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_mem CL_API_CALL
clCreateFromGLRenderbuffer(cl_context /* context */,
cl_mem_flags /* flags */,
cl_GLuint /* renderbuffer */,
cl_int * /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetGLObjectInfo(cl_mem /* memobj */,
cl_gl_object_type * /* gl_object_type */,
cl_GLuint * /* gl_object_name */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clGetGLTextureInfo(cl_mem /* memobj */,
cl_gl_texture_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueAcquireGLObjects(cl_command_queue /* command_queue */,
cl_uint /* num_objects */,
const cl_mem * /* mem_objects */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
extern CL_API_ENTRY cl_int CL_API_CALL
clEnqueueReleaseGLObjects(cl_command_queue /* command_queue */,
cl_uint /* num_objects */,
const cl_mem * /* mem_objects */,
cl_uint /* num_events_in_wait_list */,
const cl_event * /* event_wait_list */,
cl_event * /* event */) CL_API_SUFFIX__VERSION_1_0;
/* cl_khr_gl_sharing extension */
#define cl_khr_gl_sharing 1
typedef cl_uint cl_gl_context_info;
/* Additional Error Codes */
#define CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR -1000
/* cl_gl_context_info */
#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR 0x2006
#define CL_DEVICES_FOR_GL_CONTEXT_KHR 0x2007
/* Additional cl_context_properties */
#define CL_GL_CONTEXT_KHR 0x2008
#define CL_EGL_DISPLAY_KHR 0x2009
#define CL_GLX_DISPLAY_KHR 0x200A
#define CL_WGL_HDC_KHR 0x200B
#define CL_CGL_SHAREGROUP_KHR 0x200C
extern CL_API_ENTRY cl_int CL_API_CALL
clGetGLContextInfoKHR(const cl_context_properties * /* properties */,
cl_gl_context_info /* param_name */,
size_t /* param_value_size */,
void * /* param_value */,
size_t * /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetGLContextInfoKHR_fn)(
const cl_context_properties * properties,
cl_gl_context_info param_name,
size_t param_value_size,
void * param_value,
size_t * param_value_size_ret);
#ifdef __cplusplus
}
#endif
#endif /* __OPENCL_CL_GL_H */

69
include/CL/cl_gl_ext.h Normal file
View File

@ -0,0 +1,69 @@
/**********************************************************************************
* Copyright (c) 2008-2010 The Khronos Group Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and/or associated documentation files (the
* "Materials"), to deal in the Materials without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Materials, and to
* permit persons to whom the Materials are 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 Materials.
*
* THE MATERIALS ARE 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
* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
**********************************************************************************/
/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
/* cl_gl_ext.h contains vendor (non-KHR) OpenCL extensions which have */
/* OpenGL dependencies. */
#ifndef __OPENCL_CL_GL_EXT_H
#define __OPENCL_CL_GL_EXT_H
#ifdef __cplusplus
extern "C" {
#endif
#ifdef __APPLE__
#include <OpenCL/cl_gl.h>
#else
#include <CL/cl_gl.h>
#endif
/*
* For each extension, follow this template
* /* cl_VEN_extname extension */
/* #define cl_VEN_extname 1
* ... define new types, if any
* ... define new tokens, if any
* ... define new APIs, if any
*
* If you need GLtypes here, mirror them with a cl_GLtype, rather than including a GL header
* This allows us to avoid having to decide whether to include GL headers or GLES here.
*/
/*
* cl_khr_gl_event extension
* See section 9.9 in the OpenCL 1.1 spec for more information
*/
#define CL_COMMAND_GL_FENCE_SYNC_OBJECT_KHR 0x200D
extern CL_API_ENTRY cl_event CL_API_CALL
clCreateEventFromGLsyncKHR(cl_context /* context */,
cl_GLsync /* cl_GLsync */,
cl_int * /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1;
#ifdef __cplusplus
}
#endif
#endif /* __OPENCL_CL_GL_EXT_H */

1198
include/CL/cl_platform.h Normal file
View File

File diff suppressed because it is too large Load Diff

54
include/CL/opencl.h Normal file
View File

@ -0,0 +1,54 @@
/*******************************************************************************
* Copyright (c) 2008-2010 The Khronos Group Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and/or associated documentation files (the
* "Materials"), to deal in the Materials without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Materials, and to
* permit persons to whom the Materials are 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 Materials.
*
* THE MATERIALS ARE 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
* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
******************************************************************************/
/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
#ifndef __OPENCL_H
#define __OPENCL_H
#ifdef __cplusplus
extern "C" {
#endif
#ifdef __APPLE__
#include <OpenCL/cl.h>
#include <OpenCL/cl_gl.h>
#include <OpenCL/cl_gl_ext.h>
#include <OpenCL/cl_ext.h>
#else
#include <CL/cl.h>
#include <CL/cl_gl.h>
#include <CL/cl_gl_ext.h>
#include <CL/cl_ext.h>
#endif
#ifdef __cplusplus
}
#endif
#endif /* __OPENCL_H */

2
src/gallium/state_trackers/Makefile
View File

@ -17,7 +17,7 @@ subdirs:
clean:
rm -f `find . -name \*.[oa]`
rm -f `find . -regex '.*\.l?[oa]'`
rm -f `find . -name depend`

1716
src/gallium/state_trackers/clover/Doxyfile Normal file
View File

File diff suppressed because it is too large Load Diff

71
src/gallium/state_trackers/clover/Makefile.am Normal file
View File

@ -0,0 +1,71 @@
AUTOMAKE_OPTIONS = subdir-objects
AM_CPPFLAGS = \
$(GALLIUM_PIPE_LOADER_DEFINES) \
-DMESA_VERSION=\"$(MESA_VERSION)\" \
-DPIPE_SEARCH_DIR=\"$(OPENCL_LIB_INSTALL_DIR)\" \
-I$(top_srcdir)/include \
-I$(top_srcdir)/src/gallium/include \
-I$(top_srcdir)/src/gallium/drivers \
-I$(top_srcdir)/src/gallium/auxiliary \
-I$(top_srcdir)/src/gallium/winsys \
-I$(srcdir)
noinst_LTLIBRARIES = libclover.la libcltgsi.la libclllvm.la
libcltgsi_la_CXXFLAGS = \
-std=c++0x
libcltgsi_la_SOURCES = \
tgsi/compiler.cpp
libclllvm_la_CXXFLAGS = \
-std=c++98
libclllvm_la_SOURCES = \
llvm/invocation.cpp
libclover_la_CXXFLAGS = \
-std=c++0x
libclover_la_LIBADD = \
libcltgsi.la libclllvm.la
libclover_la_SOURCES = \
core/base.hpp \
core/compat.hpp \
core/compiler.hpp \
core/geometry.hpp \
core/device.hpp \
core/device.cpp \
core/context.hpp \
core/context.cpp \
core/queue.hpp \
core/queue.cpp \
core/format.hpp \
core/format.cpp \
core/memory.hpp \
core/memory.cpp \
core/resource.hpp \
core/resource.cpp \
core/sampler.hpp \
core/sampler.cpp \
core/event.hpp \
core/event.cpp \
core/program.hpp \
core/program.cpp \
core/kernel.hpp \
core/kernel.cpp \
core/module.hpp \
core/module.cpp \
api/util.hpp \
api/platform.cpp \
api/device.cpp \
api/context.cpp \
api/queue.cpp \
api/memory.cpp \
api/transfer.cpp \
api/sampler.cpp \
api/event.cpp \
api/program.cpp \
api/kernel.cpp

120
src/gallium/state_trackers/clover/api/context.cpp Normal file
View File

@ -0,0 +1,120 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/context.hpp"
using namespace clover;
PUBLIC cl_context
clCreateContext(const cl_context_properties *props, cl_uint num_devs,
const cl_device_id *devs,
void (CL_CALLBACK *pfn_notify)(const char *, const void *,
size_t, void *),
void *user_data, cl_int *errcode_ret) try {
auto mprops = property_map(props);
if (!devs || !num_devs ||
(!pfn_notify && user_data))
throw error(CL_INVALID_VALUE);
if (any_of(is_zero<cl_device_id>(), devs, devs + num_devs))
throw error(CL_INVALID_DEVICE);
for (auto p : mprops) {
if (!(p.first == CL_CONTEXT_PLATFORM &&
(cl_platform_id)p.second == NULL))
throw error(CL_INVALID_PROPERTY);
}
ret_error(errcode_ret, CL_SUCCESS);
return new context(
property_vector(mprops),
std::vector<cl_device_id>(devs, devs + num_devs));
} catch(error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_context
clCreateContextFromType(const cl_context_properties *props,
cl_device_type type,
void (CL_CALLBACK *pfn_notify)(
const char *, const void *, size_t, void *),
void *user_data, cl_int *errcode_ret) {
cl_device_id dev;
cl_int ret;
ret = clGetDeviceIDs(0, type, 1, &dev, 0);
if (ret) {
ret_error(errcode_ret, ret);
return NULL;
}
return clCreateContext(props, 1, &dev, pfn_notify, user_data, errcode_ret);
}
PUBLIC cl_int
clRetainContext(cl_context ctx) {
if (!ctx)
return CL_INVALID_CONTEXT;
ctx->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseContext(cl_context ctx) {
if (!ctx)
return CL_INVALID_CONTEXT;
if (ctx->release())
delete ctx;
return CL_SUCCESS;
}
PUBLIC cl_int
clGetContextInfo(cl_context ctx, cl_context_info param,
size_t size, void *buf, size_t *size_ret) {
if (!ctx)
return CL_INVALID_CONTEXT;
switch (param) {
case CL_CONTEXT_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, ctx->ref_count());
case CL_CONTEXT_NUM_DEVICES:
return scalar_property<cl_uint>(buf, size, size_ret, ctx->devs.size());
case CL_CONTEXT_DEVICES:
return vector_property<cl_device_id>(buf, size, size_ret, ctx->devs);
case CL_CONTEXT_PROPERTIES:
return vector_property<cl_context_properties>(buf, size, size_ret,
ctx->props());
default:
return CL_INVALID_VALUE;
}
}

262
src/gallium/state_trackers/clover/api/device.cpp Normal file
View File

@ -0,0 +1,262 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/device.hpp"
using namespace clover;
static device_registry registry;
PUBLIC cl_int
clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
cl_uint num_entries, cl_device_id *devices,
cl_uint *num_devices) {
std::vector<cl_device_id> devs;
if (platform != NULL)
return CL_INVALID_PLATFORM;
if ((!num_entries && devices) ||
(!num_devices && !devices))
return CL_INVALID_VALUE;
// Collect matching devices
for (device &dev : registry) {
if (((device_type & CL_DEVICE_TYPE_DEFAULT) &&
&dev == &registry.front()) ||
(device_type & dev.type()))
devs.push_back(&dev);
}
if (devs.empty())
return CL_DEVICE_NOT_FOUND;
// ...and return the requested data.
if (num_devices)
*num_devices = devs.size();
if (devices)
std::copy_n(devs.begin(),
std::min((cl_uint)devs.size(), num_entries),
devices);
return CL_SUCCESS;
}
PUBLIC cl_int
clGetDeviceInfo(cl_device_id dev, cl_device_info param,
size_t size, void *buf, size_t *size_ret) {
if (!dev)
return CL_INVALID_DEVICE;
switch (param) {
case CL_DEVICE_TYPE:
return scalar_property<cl_device_type>(buf, size, size_ret, dev->type());
case CL_DEVICE_VENDOR_ID:
return scalar_property<cl_uint>(buf, size, size_ret, dev->vendor_id());
case CL_DEVICE_MAX_COMPUTE_UNITS:
return scalar_property<cl_uint>(buf, size, size_ret, 1);
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
return scalar_property<cl_uint>(buf, size, size_ret,
dev->max_block_size().size());
case CL_DEVICE_MAX_WORK_ITEM_SIZES:
return vector_property<size_t>(buf, size, size_ret,
dev->max_block_size());
case CL_DEVICE_MAX_WORK_GROUP_SIZE:
return scalar_property<size_t>(buf, size, size_ret, SIZE_MAX);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
return scalar_property<cl_uint>(buf, size, size_ret, 16);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
return scalar_property<cl_uint>(buf, size, size_ret, 8);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
return scalar_property<cl_uint>(buf, size, size_ret, 4);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
return scalar_property<cl_uint>(buf, size, size_ret, 2);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
return scalar_property<cl_uint>(buf, size, size_ret, 4);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
return scalar_property<cl_uint>(buf, size, size_ret, 2);
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
return scalar_property<cl_uint>(buf, size, size_ret, 0);
case CL_DEVICE_MAX_CLOCK_FREQUENCY:
return scalar_property<cl_uint>(buf, size, size_ret, 0);
case CL_DEVICE_ADDRESS_BITS:
return scalar_property<cl_uint>(buf, size, size_ret, 32);
case CL_DEVICE_MAX_READ_IMAGE_ARGS:
return scalar_property<cl_uint>(buf, size, size_ret,
dev->max_images_read());
case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
return scalar_property<cl_uint>(buf, size, size_ret,
dev->max_images_write());
case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret, 0);
case CL_DEVICE_IMAGE2D_MAX_WIDTH:
case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
return scalar_property<size_t>(buf, size, size_ret,
1 << dev->max_image_levels_2d());
case CL_DEVICE_IMAGE3D_MAX_WIDTH:
case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
case CL_DEVICE_IMAGE3D_MAX_DEPTH:
return scalar_property<size_t>(buf, size, size_ret,
1 << dev->max_image_levels_3d());
case CL_DEVICE_IMAGE_SUPPORT:
return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
case CL_DEVICE_MAX_PARAMETER_SIZE:
return scalar_property<size_t>(buf, size, size_ret,
dev->max_mem_input());
case CL_DEVICE_MAX_SAMPLERS:
return scalar_property<cl_uint>(buf, size, size_ret,
dev->max_samplers());
case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
return scalar_property<cl_uint>(buf, size, size_ret, 128);
case CL_DEVICE_SINGLE_FP_CONFIG:
return scalar_property<cl_device_fp_config>(buf, size, size_ret,
CL_FP_DENORM | CL_FP_INF_NAN | CL_FP_ROUND_TO_NEAREST);
case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
return scalar_property<cl_device_mem_cache_type>(buf, size, size_ret,
CL_NONE);
case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
return scalar_property<cl_uint>(buf, size, size_ret, 0);
case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret, 0);
case CL_DEVICE_GLOBAL_MEM_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret,
dev->max_mem_global());
case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret,
dev->max_const_buffer_size());
case CL_DEVICE_MAX_CONSTANT_ARGS:
return scalar_property<cl_uint>(buf, size, size_ret,
dev->max_const_buffers());
case CL_DEVICE_LOCAL_MEM_TYPE:
return scalar_property<cl_device_local_mem_type>(buf, size, size_ret,
CL_LOCAL);
case CL_DEVICE_LOCAL_MEM_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret,
dev->max_mem_local());
case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
return scalar_property<cl_bool>(buf, size, size_ret, CL_FALSE);
case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
return scalar_property<size_t>(buf, size, size_ret, 0);
case CL_DEVICE_ENDIAN_LITTLE:
return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
case CL_DEVICE_AVAILABLE:
case CL_DEVICE_COMPILER_AVAILABLE:
return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
case CL_DEVICE_EXECUTION_CAPABILITIES:
return scalar_property<cl_device_exec_capabilities>(buf, size, size_ret,
CL_EXEC_KERNEL);
case CL_DEVICE_QUEUE_PROPERTIES:
return scalar_property<cl_command_queue_properties>(buf, size, size_ret,
CL_QUEUE_PROFILING_ENABLE);
case CL_DEVICE_NAME:
return string_property(buf, size, size_ret, dev->device_name());
case CL_DEVICE_VENDOR:
return string_property(buf, size, size_ret, dev->vendor_name());
case CL_DRIVER_VERSION:
return string_property(buf, size, size_ret, MESA_VERSION);
case CL_DEVICE_PROFILE:
return string_property(buf, size, size_ret, "FULL_PROFILE");
case CL_DEVICE_VERSION:
return string_property(buf, size, size_ret, "OpenCL 1.1 MESA " MESA_VERSION);
case CL_DEVICE_EXTENSIONS:
return string_property(buf, size, size_ret, "");
case CL_DEVICE_PLATFORM:
return scalar_property<cl_platform_id>(buf, size, size_ret, NULL);
case CL_DEVICE_HOST_UNIFIED_MEMORY:
return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
return scalar_property<cl_uint>(buf, size, size_ret, 16);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
return scalar_property<cl_uint>(buf, size, size_ret, 8);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
return scalar_property<cl_uint>(buf, size, size_ret, 4);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
return scalar_property<cl_uint>(buf, size, size_ret, 2);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
return scalar_property<cl_uint>(buf, size, size_ret, 4);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
return scalar_property<cl_uint>(buf, size, size_ret, 2);
case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
return scalar_property<cl_uint>(buf, size, size_ret, 0);
case CL_DEVICE_OPENCL_C_VERSION:
return string_property(buf, size, size_ret, "OpenCL C 1.1");
default:
return CL_INVALID_VALUE;
}
}

239
src/gallium/state_trackers/clover/api/event.cpp Normal file
View File

@ -0,0 +1,239 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/event.hpp"
using namespace clover;
PUBLIC cl_event
clCreateUserEvent(cl_context ctx, cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
ret_error(errcode_ret, CL_SUCCESS);
return new soft_event(*ctx, {}, false);
} catch(error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clSetUserEventStatus(cl_event ev, cl_int status) {
if (!dynamic_cast<soft_event *>(ev))
return CL_INVALID_EVENT;
if (status > 0)
return CL_INVALID_VALUE;
if (ev->status() <= 0)
return CL_INVALID_OPERATION;
if (status)
ev->abort(status);
else
ev->trigger();
return CL_SUCCESS;
}
PUBLIC cl_int
clWaitForEvents(cl_uint num_evs, const cl_event *evs) try {
if (!num_evs || !evs)
throw error(CL_INVALID_VALUE);
std::for_each(evs, evs + num_evs, [&](const cl_event ev) {
if (!ev)
throw error(CL_INVALID_EVENT);
if (&ev->ctx != &evs[0]->ctx)
throw error(CL_INVALID_CONTEXT);
if (ev->status() < 0)
throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
});
// Create a temporary soft event that depends on all the events in
// the wait list
ref_ptr<soft_event> sev = transfer(
new soft_event(evs[0]->ctx, { evs, evs + num_evs }, true));
// ...and wait on it.
sev->wait();
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clGetEventInfo(cl_event ev, cl_event_info param,
size_t size, void *buf, size_t *size_ret) {
if (!ev)
return CL_INVALID_EVENT;
switch (param) {
case CL_EVENT_COMMAND_QUEUE:
return scalar_property<cl_command_queue>(buf, size, size_ret, ev->queue());
case CL_EVENT_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret, &ev->ctx);
case CL_EVENT_COMMAND_TYPE:
return scalar_property<cl_command_type>(buf, size, size_ret, ev->command());
case CL_EVENT_COMMAND_EXECUTION_STATUS:
return scalar_property<cl_int>(buf, size, size_ret, ev->status());
case CL_EVENT_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, ev->ref_count());
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clSetEventCallback(cl_event ev, cl_int type,
void (CL_CALLBACK *pfn_event_notify)(cl_event, cl_int,
void *),
void *user_data) try {
if (!ev)
throw error(CL_INVALID_EVENT);
if (!pfn_event_notify || type != CL_COMPLETE)
throw error(CL_INVALID_VALUE);
// Create a temporary soft event that depends on ev, with
// pfn_event_notify as completion action.
ref_ptr<soft_event> sev = transfer(
new soft_event(ev->ctx, { ev }, true,
[=](event &) {
ev->wait();
pfn_event_notify(ev, ev->status(), user_data);
}));
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clRetainEvent(cl_event ev) {
if (!ev)
return CL_INVALID_EVENT;
ev->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseEvent(cl_event ev) {
if (!ev)
return CL_INVALID_EVENT;
if (ev->release())
delete ev;
return CL_SUCCESS;
}
PUBLIC cl_int
clEnqueueMarker(cl_command_queue q, cl_event *ev) try {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (!ev)
throw error(CL_INVALID_VALUE);
*ev = new hard_event(*q, CL_COMMAND_MARKER, {});
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueBarrier(cl_command_queue q) {
if (!q)
return CL_INVALID_COMMAND_QUEUE;
// No need to do anything, q preserves data ordering strictly.
return CL_SUCCESS;
}
PUBLIC cl_int
clEnqueueWaitForEvents(cl_command_queue q, cl_uint num_evs,
const cl_event *evs) try {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (!num_evs || !evs)
throw error(CL_INVALID_VALUE);
std::for_each(evs, evs + num_evs, [&](const cl_event ev) {
if (!ev)
throw error(CL_INVALID_EVENT);
if (&ev->ctx != &q->ctx)
throw error(CL_INVALID_CONTEXT);
});
// Create a hard event that depends on the events in the wait list:
// subsequent commands in the same queue will be implicitly
// serialized with respect to it -- hard events always are.
ref_ptr<hard_event> hev = transfer(
new hard_event(*q, 0, { evs, evs + num_evs }));
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clGetEventProfilingInfo(cl_event ev, cl_profiling_info param,
size_t size, void *buf, size_t *size_ret) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
PUBLIC cl_int
clFinish(cl_command_queue q) try {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
// Create a temporary hard event -- it implicitly depends on all
// the previously queued hard events.
ref_ptr<hard_event> hev = transfer(new hard_event(*q, 0, { }));
// And wait on it.
hev->wait();
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}

318
src/gallium/state_trackers/clover/api/kernel.cpp Normal file
View File

@ -0,0 +1,318 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/kernel.hpp"
#include "core/event.hpp"
using namespace clover;
PUBLIC cl_kernel
clCreateKernel(cl_program prog, const char *name,
cl_int *errcode_ret) try {
if (!prog)
throw error(CL_INVALID_PROGRAM);
if (!name)
throw error(CL_INVALID_VALUE);
if (prog->binaries().empty())
throw error(CL_INVALID_PROGRAM_EXECUTABLE);
auto sym = prog->binaries().begin()->second.sym(name);
ret_error(errcode_ret, CL_SUCCESS);
return new kernel(*prog, name, { sym.args.begin(), sym.args.end() });
} catch (module::noent_error &e) {
ret_error(errcode_ret, CL_INVALID_KERNEL_NAME);
return NULL;
} catch(error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clCreateKernelsInProgram(cl_program prog, cl_uint count,
cl_kernel *kerns, cl_uint *count_ret) {
if (!prog)
throw error(CL_INVALID_PROGRAM);
if (prog->binaries().empty())
throw error(CL_INVALID_PROGRAM_EXECUTABLE);
auto &syms = prog->binaries().begin()->second.syms;
if (kerns && count < syms.size())
throw error(CL_INVALID_VALUE);
if (kerns)
std::transform(syms.begin(), syms.end(), kerns,
[=](const module::symbol &sym) {
return new kernel(*prog, compat::string(sym.name),
{ sym.args.begin(), sym.args.end() });
});
if (count_ret)
*count_ret = syms.size();
return CL_SUCCESS;
}
PUBLIC cl_int
clRetainKernel(cl_kernel kern) {
if (!kern)
return CL_INVALID_KERNEL;
kern->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseKernel(cl_kernel kern) {
if (!kern)
return CL_INVALID_KERNEL;
if (kern->release())
delete kern;
return CL_SUCCESS;
}
PUBLIC cl_int
clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size,
const void *value) try {
if (!kern)
throw error(CL_INVALID_KERNEL);
if (idx >= kern->args.size())
throw error(CL_INVALID_ARG_INDEX);
kern->args[idx]->set(size, value);
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
size_t size, void *buf, size_t *size_ret) {
if (!kern)
return CL_INVALID_KERNEL;
switch (param) {
case CL_KERNEL_FUNCTION_NAME:
return string_property(buf, size, size_ret, kern->name());
case CL_KERNEL_NUM_ARGS:
return scalar_property<cl_uint>(buf, size, size_ret,
kern->args.size());
case CL_KERNEL_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret,
kern->ref_count());
case CL_KERNEL_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret,
&kern->prog.ctx);
case CL_KERNEL_PROGRAM:
return scalar_property<cl_program>(buf, size, size_ret,
&kern->prog);
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev,
cl_kernel_work_group_info param,
size_t size, void *buf, size_t *size_ret) {
if (!kern)
return CL_INVALID_KERNEL;
if ((!dev && kern->prog.binaries().size() != 1) ||
(dev && !kern->prog.binaries().count(dev)))
return CL_INVALID_DEVICE;
switch (param) {
case CL_KERNEL_WORK_GROUP_SIZE:
return scalar_property<size_t>(buf, size, size_ret,
kern->max_block_size());
case CL_KERNEL_COMPILE_WORK_GROUP_SIZE:
return vector_property<size_t>(buf, size, size_ret,
kern->block_size());
case CL_KERNEL_LOCAL_MEM_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret,
kern->mem_local());
case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
return scalar_property<size_t>(buf, size, size_ret, 1);
case CL_KERNEL_PRIVATE_MEM_SIZE:
return scalar_property<cl_ulong>(buf, size, size_ret,
kern->mem_private());
default:
return CL_INVALID_VALUE;
}
}
namespace {
///
/// Common argument checking shared by kernel invocation commands.
///
void
kernel_validate(cl_command_queue q, cl_kernel kern,
cl_uint dims, const size_t *grid_offset,
const size_t *grid_size, const size_t *block_size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (!kern)
throw error(CL_INVALID_KERNEL);
if (&kern->prog.ctx != &q->ctx ||
any_of([&](const cl_event ev) {
return &ev->ctx != &q->ctx;
}, deps, deps + num_deps))
throw error(CL_INVALID_CONTEXT);
if (bool(num_deps) != bool(deps) ||
any_of(is_zero<cl_event>(), deps, deps + num_deps))
throw error(CL_INVALID_EVENT_WAIT_LIST);
if (any_of([](std::unique_ptr<kernel::argument> &arg) {
return !arg->set();
}, kern->args.begin(), kern->args.end()))
throw error(CL_INVALID_KERNEL_ARGS);
if (!kern->prog.binaries().count(&q->dev))
throw error(CL_INVALID_PROGRAM_EXECUTABLE);
if (dims < 1 || dims > q->dev.max_block_size().size())
throw error(CL_INVALID_WORK_DIMENSION);
if (!grid_size || any_of(is_zero<size_t>(), grid_size, grid_size + dims))
throw error(CL_INVALID_GLOBAL_WORK_SIZE);
if (block_size && any_of([](size_t b, size_t max) {
return b == 0 || b > max;
}, block_size, block_size + dims,
q->dev.max_block_size().begin()))
throw error(CL_INVALID_WORK_ITEM_SIZE);
if (block_size && any_of([](size_t b, size_t g) {
return g % b;
}, block_size, block_size + dims, grid_size))
throw error(CL_INVALID_WORK_GROUP_SIZE);
}
///
/// Common event action shared by kernel invocation commands.
///
std::function<void (event &)>
kernel_op(cl_command_queue q, cl_kernel kern,
const std::vector<size_t> &grid_offset,
const std::vector<size_t> &grid_size,
const std::vector<size_t> &block_size) {
const std::vector<size_t> reduced_grid_size = map(
std::divides<size_t>(), grid_size.begin(), grid_size.end(),
block_size.begin());
return [=](event &) {
kern->launch(*q, grid_offset, reduced_grid_size, block_size);
};
}
template<typename T, typename S>
std::vector<T>
opt_vector(const T *p, S n) {
if (p)
return { p, p + n };
else
return { n };
}
}
PUBLIC cl_int
clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern,
cl_uint dims, const size_t *pgrid_offset,
const size_t *pgrid_size, const size_t *pblock_size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
const std::vector<size_t> grid_offset = opt_vector(pgrid_offset, dims);
const std::vector<size_t> grid_size = opt_vector(pgrid_size, dims);
const std::vector<size_t> block_size = opt_vector(pblock_size, dims);
kernel_validate(q, kern, dims, pgrid_offset, pgrid_size, pblock_size,
num_deps, deps, ev);
hard_event *hev = new hard_event(
*q, CL_COMMAND_NDRANGE_KERNEL, { deps, deps + num_deps },
kernel_op(q, kern, grid_offset, grid_size, block_size));
ret_object(ev, hev);
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueTask(cl_command_queue q, cl_kernel kern,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
const std::vector<size_t> grid_offset = { 0 };
const std::vector<size_t> grid_size = { 1 };
const std::vector<size_t> block_size = { 1 };
kernel_validate(q, kern, 1, grid_offset.data(), grid_size.data(),
block_size.data(), num_deps, deps, ev);
hard_event *hev = new hard_event(
*q, CL_COMMAND_TASK, { deps, deps + num_deps },
kernel_op(q, kern, grid_offset, grid_size, block_size));
ret_object(ev, hev);
return CL_SUCCESS;
} catch(error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueNativeKernel(cl_command_queue q, void (*func)(void *),
void *args, size_t args_size,
cl_uint obj_count, const cl_mem *obj_list,
const void **obj_args, cl_uint num_deps,
const cl_event *deps, cl_event *ev) {
return CL_INVALID_OPERATION;
}

305
src/gallium/state_trackers/clover/api/memory.cpp Normal file
View File

@ -0,0 +1,305 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/memory.hpp"
#include "core/format.hpp"
using namespace clover;
PUBLIC cl_mem
clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size,
void *host_ptr, cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
CL_MEM_COPY_HOST_PTR)))
throw error(CL_INVALID_HOST_PTR);
if (!size)
throw error(CL_INVALID_BUFFER_SIZE);
if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
CL_MEM_COPY_HOST_PTR))
throw error(CL_INVALID_VALUE);
ret_error(errcode_ret, CL_SUCCESS);
return new root_buffer(*ctx, flags, size, host_ptr);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_mem
clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op,
const void *op_info, cl_int *errcode_ret) try {
root_buffer *parent = dynamic_cast<root_buffer *>(obj);
if (!parent)
throw error(CL_INVALID_MEM_OBJECT);
if ((flags & (CL_MEM_USE_HOST_PTR |
CL_MEM_ALLOC_HOST_PTR |
CL_MEM_COPY_HOST_PTR)) ||
(~flags & parent->flags() & (CL_MEM_READ_ONLY |
CL_MEM_WRITE_ONLY)))
throw error(CL_INVALID_VALUE);
if (op == CL_BUFFER_CREATE_TYPE_REGION) {
const cl_buffer_region *reg = (const cl_buffer_region *)op_info;
if (!reg ||
reg->origin > parent->size() ||
reg->origin + reg->size > parent->size())
throw error(CL_INVALID_VALUE);
if (!reg->size)
throw error(CL_INVALID_BUFFER_SIZE);
ret_error(errcode_ret, CL_SUCCESS);
return new sub_buffer(*parent, flags, reg->origin, reg->size);
} else {
throw error(CL_INVALID_VALUE);
}
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_mem
clCreateImage2D(cl_context ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t row_pitch,
void *host_ptr, cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
CL_MEM_COPY_HOST_PTR))
throw error(CL_INVALID_VALUE);
if (!format)
throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
if (width < 1 || height < 1)
throw error(CL_INVALID_IMAGE_SIZE);
if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
CL_MEM_COPY_HOST_PTR)))
throw error(CL_INVALID_HOST_PTR);
if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE2D).count(*format))
throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED);
ret_error(errcode_ret, CL_SUCCESS);
return new image2d(*ctx, flags, format, width, height,
row_pitch, host_ptr);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_mem
clCreateImage3D(cl_context ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch,
void *host_ptr, cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
CL_MEM_COPY_HOST_PTR))
throw error(CL_INVALID_VALUE);
if (!format)
throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
if (width < 1 || height < 1 || depth < 2)
throw error(CL_INVALID_IMAGE_SIZE);
if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
CL_MEM_COPY_HOST_PTR)))
throw error(CL_INVALID_HOST_PTR);
if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE3D).count(*format))
throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED);
ret_error(errcode_ret, CL_SUCCESS);
return new image3d(*ctx, flags, format, width, height, depth,
row_pitch, slice_pitch, host_ptr);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clGetSupportedImageFormats(cl_context ctx, cl_mem_flags flags,
cl_mem_object_type type, cl_uint count,
cl_image_format *buf, cl_uint *count_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
CL_MEM_COPY_HOST_PTR))
throw error(CL_INVALID_VALUE);
if (!count && buf)
throw error(CL_INVALID_VALUE);
auto formats = supported_formats(ctx, type);
if (buf)
std::copy_n(formats.begin(), std::min((cl_uint)formats.size(), count),
buf);
if (count_ret)
*count_ret = formats.size();
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clGetMemObjectInfo(cl_mem obj, cl_mem_info param,
size_t size, void *buf, size_t *size_ret) {
if (!obj)
return CL_INVALID_MEM_OBJECT;
switch (param) {
case CL_MEM_TYPE:
return scalar_property<cl_mem_object_type>(buf, size, size_ret,
obj->type());
case CL_MEM_FLAGS:
return scalar_property<cl_mem_flags>(buf, size, size_ret, obj->flags());
case CL_MEM_SIZE:
return scalar_property<size_t>(buf, size, size_ret, obj->size());
case CL_MEM_HOST_PTR:
return scalar_property<void *>(buf, size, size_ret, obj->host_ptr());
case CL_MEM_MAP_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, 0);
case CL_MEM_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, obj->ref_count());
case CL_MEM_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret, &obj->ctx);
case CL_MEM_ASSOCIATED_MEMOBJECT: {
sub_buffer *sub = dynamic_cast<sub_buffer *>(obj);
return scalar_property<cl_mem>(buf, size, size_ret,
(sub ? &sub->parent : NULL));
}
case CL_MEM_OFFSET: {
sub_buffer *sub = dynamic_cast<sub_buffer *>(obj);
return scalar_property<size_t>(buf, size, size_ret,
(sub ? sub->offset() : 0));
}
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clGetImageInfo(cl_mem obj, cl_image_info param,
size_t size, void *buf, size_t *size_ret) {
image *img = dynamic_cast<image *>(obj);
if (!img)
return CL_INVALID_MEM_OBJECT;
switch (param) {
case CL_IMAGE_FORMAT:
return scalar_property<cl_image_format>(buf, size, size_ret,
img->format());
case CL_IMAGE_ELEMENT_SIZE:
return scalar_property<size_t>(buf, size, size_ret, 0);
case CL_IMAGE_ROW_PITCH:
return scalar_property<size_t>(buf, size, size_ret, img->row_pitch());
case CL_IMAGE_SLICE_PITCH:
return scalar_property<size_t>(buf, size, size_ret, img->slice_pitch());
case CL_IMAGE_WIDTH:
return scalar_property<size_t>(buf, size, size_ret, img->width());
case CL_IMAGE_HEIGHT:
return scalar_property<size_t>(buf, size, size_ret, img->height());
case CL_IMAGE_DEPTH:
return scalar_property<size_t>(buf, size, size_ret, img->depth());
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clRetainMemObject(cl_mem obj) {
if (!obj)
return CL_INVALID_MEM_OBJECT;
obj->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseMemObject(cl_mem obj) {
if (!obj)
return CL_INVALID_MEM_OBJECT;
if (obj->release())
delete obj;
return CL_SUCCESS;
}
PUBLIC cl_int
clSetMemObjectDestructorCallback(cl_mem obj,
void (CL_CALLBACK *pfn_notify)(cl_mem, void *),
void *user_data) {
if (!obj)
return CL_INVALID_MEM_OBJECT;
if (!pfn_notify)
return CL_INVALID_VALUE;
obj->destroy_notify([=]{ pfn_notify(obj, user_data); });
return CL_SUCCESS;
}

68
src/gallium/state_trackers/clover/api/platform.cpp Normal file
View File

@ -0,0 +1,68 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
using namespace clover;
PUBLIC cl_int
clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms,
cl_uint *num_platforms) {
if ((!num_entries && platforms) ||
(!num_platforms && !platforms))
return CL_INVALID_VALUE;
if (num_platforms)
*num_platforms = 1;
if (platforms)
*platforms = NULL;
return CL_SUCCESS;
}
PUBLIC cl_int
clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
size_t size, void *buf, size_t *size_ret) {
if (platform != NULL)
return CL_INVALID_PLATFORM;
switch (param_name) {
case CL_PLATFORM_PROFILE:
return string_property(buf, size, size_ret, "FULL_PROFILE");
case CL_PLATFORM_VERSION:
return string_property(buf, size, size_ret,
"OpenCL 1.1 MESA " MESA_VERSION);
case CL_PLATFORM_NAME:
return string_property(buf, size, size_ret, "Default");
case CL_PLATFORM_VENDOR:
return string_property(buf, size, size_ret, "Mesa");
case CL_PLATFORM_EXTENSIONS:
return string_property(buf, size, size_ret, "");
default:
return CL_INVALID_VALUE;
}
}

241
src/gallium/state_trackers/clover/api/program.cpp Normal file
View File

@ -0,0 +1,241 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/program.hpp"
using namespace clover;
PUBLIC cl_program
clCreateProgramWithSource(cl_context ctx, cl_uint count,
const char **strings, const size_t *lengths,
cl_int *errcode_ret) try {
std::string source;
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (!count || !strings ||
any_of(is_zero<const char *>(), strings, strings + count))
throw error(CL_INVALID_VALUE);
// Concatenate all the provided fragments together
for (unsigned i = 0; i < count; ++i)
source += (lengths && lengths[i] ?
std::string(strings[i], strings[i] + lengths[i]) :
std::string(strings[i]));
// ...and create a program object for them.
ret_error(errcode_ret, CL_SUCCESS);
return new program(*ctx, source);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_program
clCreateProgramWithBinary(cl_context ctx, cl_uint count,
const cl_device_id *devs, const size_t *lengths,
const unsigned char **binaries, cl_int *status_ret,
cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (!count || !devs || !lengths || !binaries)
throw error(CL_INVALID_VALUE);
if (any_of([&](const cl_device_id dev) {
return !ctx->has_device(dev);
}, devs, devs + count))
throw error(CL_INVALID_DEVICE);
// Deserialize the provided binaries,
auto modules = map(
[](const unsigned char *p, size_t l) -> std::pair<cl_int, module> {
if (!p || !l)
return { CL_INVALID_VALUE, {} };
try {
compat::istream::buffer_t bin(p, l);
compat::istream s(bin);
return { CL_SUCCESS, module::deserialize(s) };
} catch (compat::istream::error &e) {
return { CL_INVALID_BINARY, {} };
}
},
binaries, binaries + count, lengths);
// update the status array,
if (status_ret)
std::transform(modules.begin(), modules.end(), status_ret,
keys<cl_int, module>);
if (any_of(key_equals<cl_int, module>(CL_INVALID_VALUE),
modules.begin(), modules.end()))
throw error(CL_INVALID_VALUE);
if (any_of(key_equals<cl_int, module>(CL_INVALID_BINARY),
modules.begin(), modules.end()))
throw error(CL_INVALID_BINARY);
// initialize a program object with them.
ret_error(errcode_ret, CL_SUCCESS);
return new program(*ctx, { devs, devs + count },
map(values<cl_int, module>,
modules.begin(), modules.end()));
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clRetainProgram(cl_program prog) {
if (!prog)
return CL_INVALID_PROGRAM;
prog->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseProgram(cl_program prog) {
if (!prog)
return CL_INVALID_PROGRAM;
if (prog->release())
delete prog;
return CL_SUCCESS;
}
PUBLIC cl_int
clBuildProgram(cl_program prog, cl_uint count, const cl_device_id *devs,
const char *opts, void (*pfn_notify)(cl_program, void *),
void *user_data) try {
if (!prog)
throw error(CL_INVALID_PROGRAM);
if (bool(count) != bool(devs) ||
(!pfn_notify && user_data))
throw error(CL_INVALID_VALUE);
if (any_of([&](const cl_device_id dev) {
return !prog->ctx.has_device(dev);
}, devs, devs + count))
throw error(CL_INVALID_DEVICE);
prog->build({ devs, devs + count });
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clUnloadCompiler() {
return CL_SUCCESS;
}
PUBLIC cl_int
clGetProgramInfo(cl_program prog, cl_program_info param,
size_t size, void *buf, size_t *size_ret) {
if (!prog)
return CL_INVALID_PROGRAM;
switch (param) {
case CL_PROGRAM_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret,
prog->ref_count());
case CL_PROGRAM_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret,
&prog->ctx);
case CL_PROGRAM_NUM_DEVICES:
return scalar_property<cl_uint>(buf, size, size_ret,
prog->binaries().size());
case CL_PROGRAM_DEVICES:
return vector_property<cl_device_id>(
buf, size, size_ret,
map(keys<device *, module>,
prog->binaries().begin(), prog->binaries().end()));
case CL_PROGRAM_SOURCE:
return string_property(buf, size, size_ret, prog->source());
case CL_PROGRAM_BINARY_SIZES:
return vector_property<size_t>(
buf, size, size_ret,
map([](const std::pair<device *, module> &ent) {
compat::ostream::buffer_t bin;
compat::ostream s(bin);
ent.second.serialize(s);
return bin.size();
},
prog->binaries().begin(), prog->binaries().end()));
case CL_PROGRAM_BINARIES:
return matrix_property<unsigned char>(
buf, size, size_ret,
map([](const std::pair<device *, module> &ent) {
compat::ostream::buffer_t bin;
compat::ostream s(bin);
ent.second.serialize(s);
return bin;
},
prog->binaries().begin(), prog->binaries().end()));
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clGetProgramBuildInfo(cl_program prog, cl_device_id dev,
cl_program_build_info param,
size_t size, void *buf, size_t *size_ret) {
if (!prog)
return CL_INVALID_PROGRAM;
if (!prog->ctx.has_device(dev))
return CL_INVALID_DEVICE;
switch (param) {
case CL_PROGRAM_BUILD_STATUS:
return scalar_property<cl_build_status>(buf, size, size_ret,
prog->build_status(dev));
case CL_PROGRAM_BUILD_OPTIONS:
return string_property(buf, size, size_ret, prog->build_opts(dev));
case CL_PROGRAM_BUILD_LOG:
return string_property(buf, size, size_ret, prog->build_log(dev));
default:
return CL_INVALID_VALUE;
}
}

102
src/gallium/state_trackers/clover/api/queue.cpp Normal file
View File

@ -0,0 +1,102 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/queue.hpp"
using namespace clover;
PUBLIC cl_command_queue
clCreateCommandQueue(cl_context ctx, cl_device_id dev,
cl_command_queue_properties props,
cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (!ctx->has_device(dev))
throw error(CL_INVALID_DEVICE);
if (props & ~(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE |
CL_QUEUE_PROFILING_ENABLE))
throw error(CL_INVALID_VALUE);
ret_error(errcode_ret, CL_SUCCESS);
return new command_queue(*ctx, *dev, props);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clRetainCommandQueue(cl_command_queue q) {
if (!q)
return CL_INVALID_COMMAND_QUEUE;
q->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseCommandQueue(cl_command_queue q) {
if (!q)
return CL_INVALID_COMMAND_QUEUE;
if (q->release())
delete q;
return CL_SUCCESS;
}
PUBLIC cl_int
clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param,
size_t size, void *buf, size_t *size_ret) {
if (!q)
return CL_INVALID_COMMAND_QUEUE;
switch (param) {
case CL_QUEUE_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret, &q->ctx);
case CL_QUEUE_DEVICE:
return scalar_property<cl_device_id>(buf, size, size_ret, &q->dev);
case CL_QUEUE_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, q->ref_count());
case CL_QUEUE_PROPERTIES:
return scalar_property<cl_command_queue_properties>(buf, size, size_ret,
q->props());
default:
return CL_INVALID_VALUE;
}
}
PUBLIC cl_int
clFlush(cl_command_queue q) {
if (!q)
return CL_INVALID_COMMAND_QUEUE;
q->flush();
return CL_SUCCESS;
}

90
src/gallium/state_trackers/clover/api/sampler.cpp Normal file
View File

@ -0,0 +1,90 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "api/util.hpp"
#include "core/sampler.hpp"
using namespace clover;
PUBLIC cl_sampler
clCreateSampler(cl_context ctx, cl_bool norm_mode,
cl_addressing_mode addr_mode, cl_filter_mode filter_mode,
cl_int *errcode_ret) try {
if (!ctx)
throw error(CL_INVALID_CONTEXT);
ret_error(errcode_ret, CL_SUCCESS);
return new sampler(*ctx, norm_mode, addr_mode, filter_mode);
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clRetainSampler(cl_sampler s) {
if (!s)
throw error(CL_INVALID_SAMPLER);
s->retain();
return CL_SUCCESS;
}
PUBLIC cl_int
clReleaseSampler(cl_sampler s) {
if (!s)
throw error(CL_INVALID_SAMPLER);
if (s->release())
delete s;
return CL_SUCCESS;
}
PUBLIC cl_int
clGetSamplerInfo(cl_sampler s, cl_sampler_info param,
size_t size, void *buf, size_t *size_ret) {
if (!s)
throw error(CL_INVALID_SAMPLER);
switch (param) {
case CL_SAMPLER_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, s->ref_count());
case CL_SAMPLER_CONTEXT:
return scalar_property<cl_context>(buf, size, size_ret, &s->ctx);
case CL_SAMPLER_NORMALIZED_COORDS:
return scalar_property<cl_bool>(buf, size, size_ret, s->norm_mode());
case CL_SAMPLER_ADDRESSING_MODE:
return scalar_property<cl_addressing_mode>(buf, size, size_ret,
s->addr_mode());
case CL_SAMPLER_FILTER_MODE:
return scalar_property<cl_filter_mode>(buf, size, size_ret,
s->filter_mode());
default:
return CL_INVALID_VALUE;
}
}

506
src/gallium/state_trackers/clover/api/transfer.cpp Normal file
View File

@ -0,0 +1,506 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <cstring>
#include "api/util.hpp"
#include "core/event.hpp"
#include "core/resource.hpp"
using namespace clover;
namespace {
typedef resource::point point;
///
/// Common argument checking shared by memory transfer commands.
///
void
validate_base(cl_command_queue q, cl_uint num_deps, const cl_event *deps) {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (bool(num_deps) != bool(deps) ||
any_of(is_zero<cl_event>(), deps, deps + num_deps))
throw error(CL_INVALID_EVENT_WAIT_LIST);
if (any_of([&](const cl_event ev) {
return &ev->ctx != &q->ctx;
}, deps, deps + num_deps))
throw error(CL_INVALID_CONTEXT);
}
///
/// Memory object-specific argument checking shared by most memory
/// transfer commands.
///
void
validate_obj(cl_command_queue q, cl_mem obj) {
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
if (&obj->ctx != &q->ctx)
throw error(CL_INVALID_CONTEXT);
}
///
/// Class that encapsulates the task of mapping an object of type
/// \a T. The return value of get() should be implicitly
/// convertible to \a void *.
///
template<typename T> struct __map;
template<> struct __map<void *> {
static void *
get(cl_command_queue q, void *obj, cl_map_flags flags,
size_t offset, size_t size) {
return (char *)obj + offset;
}
};
template<> struct __map<const void *> {
static const void *
get(cl_command_queue q, const void *obj, cl_map_flags flags,
size_t offset, size_t size) {
return (const char *)obj + offset;
}
};
template<> struct __map<memory_obj *> {
static mapping
get(cl_command_queue q, memory_obj *obj, cl_map_flags flags,
size_t offset, size_t size) {
return { *q, obj->resource(q), flags, true, { offset }, { size }};
}
};
///
/// Software copy from \a src_obj to \a dst_obj. They can be
/// either pointers or memory objects.
///
template<typename T, typename S>
std::function<void (event &)>
soft_copy_op(cl_command_queue q,
T dst_obj, const point &dst_orig, const point &dst_pitch,
S src_obj, const point &src_orig, const point &src_pitch,
const point &region) {
return [=](event &) {
auto dst = __map<T>::get(q, dst_obj, CL_MAP_WRITE,
dst_pitch(dst_orig), dst_pitch(region));
auto src = __map<S>::get(q, src_obj, CL_MAP_READ,
src_pitch(src_orig), src_pitch(region));
point p;
for (p[2] = 0; p[2] < region[2]; ++p[2]) {
for (p[1] = 0; p[1] < region[1]; ++p[1]) {
std::memcpy(static_cast<char *>(dst) + dst_pitch(p),
static_cast<const char *>(src) + src_pitch(p),
src_pitch[0] * region[0]);
}
}
};
}
///
/// Hardware copy from \a src_obj to \a dst_obj.
///
template<typename T, typename S>
std::function<void (event &)>
hard_copy_op(cl_command_queue q, T dst_obj, const point &dst_orig,
S src_obj, const point &src_orig, const point &region) {
return [=](event &) {
dst_obj->resource(q).copy(*q, dst_orig, region,
src_obj->resource(q), src_orig);
};
}
}
PUBLIC cl_int
clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
size_t offset, size_t size, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
if (!ptr || offset > obj->size() || offset + size > obj->size())
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_READ_BUFFER, { deps, deps + num_deps },
soft_copy_op(q,
ptr, { 0 }, { 1 },
obj, { offset }, { 1 },
{ size, 1, 1 }));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
size_t offset, size_t size, const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
if (!ptr || offset > obj->size() || offset + size > obj->size())
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_WRITE_BUFFER, { deps, deps + num_deps },
soft_copy_op(q,
obj, { offset }, { 1 },
ptr, { 0 }, { 1 },
{ size, 1, 1 }));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
const size_t *obj_origin, const size_t *host_origin,
const size_t *region,
size_t obj_row_pitch, size_t obj_slice_pitch,
size_t host_row_pitch, size_t host_slice_pitch,
void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
if (!ptr)
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_READ_BUFFER_RECT, { deps, deps + num_deps },
soft_copy_op(q,
ptr, host_origin,
{ 1, host_row_pitch, host_slice_pitch },
obj, obj_origin,
{ 1, obj_row_pitch, obj_slice_pitch },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
const size_t *obj_origin, const size_t *host_origin,
const size_t *region,
size_t obj_row_pitch, size_t obj_slice_pitch,
size_t host_row_pitch, size_t host_slice_pitch,
const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
if (!ptr)
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_WRITE_BUFFER_RECT, { deps, deps + num_deps },
soft_copy_op(q,
obj, obj_origin,
{ 1, obj_row_pitch, obj_slice_pitch },
ptr, host_origin,
{ 1, host_row_pitch, host_slice_pitch },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
size_t src_offset, size_t dst_offset, size_t size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, src_obj);
validate_obj(q, dst_obj);
hard_event *hev = new hard_event(
*q, CL_COMMAND_COPY_BUFFER, { deps, deps + num_deps },
hard_copy_op(q, dst_obj, { dst_offset },
src_obj, { src_offset },
{ size, 1, 1 }));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
const size_t *src_origin, const size_t *dst_origin,
const size_t *region,
size_t src_row_pitch, size_t src_slice_pitch,
size_t dst_row_pitch, size_t dst_slice_pitch,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, src_obj);
validate_obj(q, dst_obj);
hard_event *hev = new hard_event(
*q, CL_COMMAND_COPY_BUFFER_RECT, { deps, deps + num_deps },
soft_copy_op(q,
dst_obj, dst_origin,
{ 1, dst_row_pitch, dst_slice_pitch },
src_obj, src_origin,
{ 1, src_row_pitch, src_slice_pitch },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
const size_t *origin, const size_t *region,
size_t row_pitch, size_t slice_pitch, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
validate_obj(q, img);
if (!ptr)
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_READ_IMAGE, { deps, deps + num_deps },
soft_copy_op(q,
ptr, {},
{ 1, row_pitch, slice_pitch },
obj, origin,
{ 1, img->row_pitch(), img->slice_pitch() },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
const size_t *origin, const size_t *region,
size_t row_pitch, size_t slice_pitch, const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
validate_obj(q, img);
if (!ptr)
throw error(CL_INVALID_VALUE);
hard_event *hev = new hard_event(
*q, CL_COMMAND_WRITE_IMAGE, { deps, deps + num_deps },
soft_copy_op(q,
obj, origin,
{ 1, img->row_pitch(), img->slice_pitch() },
ptr, {},
{ 1, row_pitch, slice_pitch },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
const size_t *src_origin, const size_t *dst_origin,
const size_t *region,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
image *src_img = dynamic_cast<image *>(src_obj);
image *dst_img = dynamic_cast<image *>(dst_obj);
validate_base(q, num_deps, deps);
validate_obj(q, src_img);
validate_obj(q, dst_img);
hard_event *hev = new hard_event(
*q, CL_COMMAND_COPY_IMAGE, { deps, deps + num_deps },
hard_copy_op(q, dst_obj, dst_origin, src_obj, src_origin, region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
const size_t *src_origin, const size_t *region,
size_t dst_offset,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
image *src_img = dynamic_cast<image *>(src_obj);
validate_base(q, num_deps, deps);
validate_obj(q, src_img);
validate_obj(q, dst_obj);
hard_event *hev = new hard_event(
*q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, { deps, deps + num_deps },
soft_copy_op(q,
dst_obj, { dst_offset },
{ 0, 0, 0 },
src_obj, src_origin,
{ 1, src_img->row_pitch(), src_img->slice_pitch() },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC cl_int
clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
size_t src_offset,
const size_t *dst_origin, const size_t *region,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
image *dst_img = dynamic_cast<image *>(src_obj);
validate_base(q, num_deps, deps);
validate_obj(q, src_obj);
validate_obj(q, dst_img);
hard_event *hev = new hard_event(
*q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, { deps, deps + num_deps },
soft_copy_op(q,
dst_obj, dst_origin,
{ 1, dst_img->row_pitch(), dst_img->slice_pitch() },
src_obj, { src_offset },
{ 0, 0, 0 },
region));
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}
PUBLIC void *
clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
cl_map_flags flags, size_t offset, size_t size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev, cl_int *errcode_ret) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
if (offset > obj->size() || offset + size > obj->size())
throw error(CL_INVALID_VALUE);
void *map = obj->resource(q).add_map(
*q, flags, blocking, { offset }, { size });
ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER,
{ deps, deps + num_deps }));
ret_error(errcode_ret, CL_SUCCESS);
return map;
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC void *
clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
cl_map_flags flags,
const size_t *origin, const size_t *region,
size_t *row_pitch, size_t *slice_pitch,
cl_uint num_deps, const cl_event *deps,
cl_event *ev, cl_int *errcode_ret) try {
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
validate_obj(q, img);
void *map = obj->resource(q).add_map(
*q, flags, blocking, origin, region);
ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE,
{ deps, deps + num_deps }));
ret_error(errcode_ret, CL_SUCCESS);
return map;
} catch (error &e) {
ret_error(errcode_ret, e);
return NULL;
}
PUBLIC cl_int
clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
validate_base(q, num_deps, deps);
validate_obj(q, obj);
hard_event *hev = new hard_event(
*q, CL_COMMAND_UNMAP_MEM_OBJECT, { deps, deps + num_deps },
[=](event &) {
obj->resource(q).del_map(ptr);
});
ret_object(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
return e.get();
}

166
src/gallium/state_trackers/clover/api/util.hpp Normal file
View File

@ -0,0 +1,166 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CL_UTIL_HPP__
#define __CL_UTIL_HPP__
#include <cstdint>
#include <cstring>
#include <algorithm>
#include <map>
#include "core/base.hpp"
#include "pipe/p_compiler.h"
namespace clover {
///
/// Return a matrix (a container of containers) in \a buf with
/// argument and bounds checking. Intended to be used by
/// implementations of \a clGetXXXInfo().
///
template<typename T, typename V>
cl_int
matrix_property(void *buf, size_t size, size_t *size_ret, const V& v) {
if (buf && size < sizeof(T *) * v.size())
return CL_INVALID_VALUE;
if (size_ret)
*size_ret = sizeof(T *) * v.size();
if (buf)
for_each([](typename V::value_type src, T *dst) {
if (dst)
std::copy(src.begin(), src.end(), dst);
},
v.begin(), v.end(), (T **)buf);
return CL_SUCCESS;
}
///
/// Return a vector in \a buf with argument and bounds checking.
/// Intended to be used by implementations of \a clGetXXXInfo().
///
template<typename T, typename V>
cl_int
vector_property(void *buf, size_t size, size_t *size_ret, const V& v) {
if (buf && size < sizeof(T) * v.size())
return CL_INVALID_VALUE;
if (size_ret)
*size_ret = sizeof(T) * v.size();
if (buf)
std::copy(v.begin(), v.end(), (T *)buf);
return CL_SUCCESS;
}
///
/// Return a scalar in \a buf with argument and bounds checking.
/// Intended to be used by implementations of \a clGetXXXInfo().
///
template<typename T>
cl_int
scalar_property(void *buf, size_t size, size_t *size_ret, T v) {
return vector_property<T>(buf, size, size_ret, std::vector<T>(1, v));
}
///
/// Return a string in \a buf with argument and bounds checking.
/// Intended to be used by implementations of \a clGetXXXInfo().
///
inline cl_int
string_property(void *buf, size_t size, size_t *size_ret,
const std::string &v) {
if (buf && size < v.size() + 1)
return CL_INVALID_VALUE;
if (size_ret)
*size_ret = v.size() + 1;
if (buf)
std::strcpy((char *)buf, v.c_str());
return CL_SUCCESS;
}
///
/// Convert a NULL-terminated property list into an std::map.
///
template<typename T>
std::map<T, T>
property_map(const T *props) {
std::map<T, T> m;
while (props && *props) {
T key = *props++;
T value = *props++;
if (m.count(key))
throw clover::error(CL_INVALID_PROPERTY);
m.insert({ key, value });
}
return m;
}
///
/// Convert an std::map into a NULL-terminated property list.
///
template<typename T>
std::vector<T>
property_vector(const std::map<T, T> &m) {
std::vector<T> v;
for (auto &p : m) {
v.push_back(p.first);
v.push_back(p.second);
}
v.push_back(0);
return v;
}
///
/// Return an error code in \a p if non-zero.
///
inline void
ret_error(cl_int *p, const clover::error &e) {
if (p)
*p = e.get();
}
///
/// Return a reference-counted object in \a p if non-zero.
/// Otherwise release object ownership.
///
template<typename T, typename S>
void
ret_object(T p, S v) {
if (p)
*p = v;
else
v->release();
}
}
#endif

285
src/gallium/state_trackers/clover/core/base.hpp Normal file
View File

@ -0,0 +1,285 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_BASE_HPP__
#define __CORE_BASE_HPP__
#include <stdexcept>
#include <atomic>
#include <cassert>
#include <tuple>
#include <vector>
#include <functional>
#include "CL/cl.h"
///
/// Main namespace of the CL state tracker.
///
namespace clover {
///
/// Class that represents an error that can be converted to an
/// OpenCL status code.
///
class error : public std::runtime_error {
public:
error(cl_int code, std::string what = "") :
std::runtime_error(what), code(code) {
}
cl_int get() const {
return code;
}
protected:
cl_int code;
};
///
/// Base class for objects that support reference counting.
///
class ref_counter {
public:
ref_counter() : __ref_count(1) {}
unsigned ref_count() {
return __ref_count;
}
void retain() {
__ref_count++;
}
bool release() {
return (--__ref_count) == 0;
}
private:
std::atomic<unsigned> __ref_count;
};
///
/// Intrusive smart pointer for objects that implement the
/// clover::ref_counter interface.
///
template<typename T>
class ref_ptr {
public:
ref_ptr(T *q = NULL) : p(NULL) {
reset(q);
}
template<typename S>
ref_ptr(const ref_ptr<S> &ref) : p(NULL) {
reset(ref.p);
}
~ref_ptr() {
reset(NULL);
}
void reset(T *q = NULL) {
if (q)
q->retain();
if (p && p->release())
delete p;
p = q;
}
ref_ptr &operator=(const ref_ptr &ref) {
reset(ref.p);
return *this;
}
T *operator*() const {
return p;
}
T *operator->() const {
return p;
}
operator bool() const {
return p;
}
private:
T *p;
};
///
/// Transfer the caller's ownership of a reference-counted object
/// to a clover::ref_ptr smart pointer.
///
template<typename T>
inline ref_ptr<T>
transfer(T *p) {
ref_ptr<T> ref { p };
p->release();
return ref;
}
template<typename T, typename S, int N>
struct __iter_helper {
template<typename F, typename Its, typename... Args>
static T
step(F op, S state, Its its, Args... args) {
return __iter_helper<T, S, N - 1>::step(
op, state, its, *(std::get<N>(its)++), args...);
}
};
template<typename T, typename S>
struct __iter_helper<T, S, 0> {
template<typename F, typename Its, typename... Args>
static T
step(F op, S state, Its its, Args... args) {
return op(state, *(std::get<0>(its)++), args...);
}
};
struct __empty {};
template<typename T>
struct __iter_helper<T, __empty, 0> {
template<typename F, typename Its, typename... Args>
static T
step(F op, __empty state, Its its, Args... args) {
return op(*(std::get<0>(its)++), args...);
}
};
template<typename F, typename... Its>
struct __result_helper {
typedef typename std::remove_const<
typename std::result_of<
F (typename std::iterator_traits<Its>::value_type...)
>::type
>::type type;
};
///
/// Iterate \a op on the result of zipping all the specified
/// iterators together.
///
/// Similar to std::for_each, but it accepts functions of an
/// arbitrary number of arguments.
///
template<typename F, typename It0, typename... Its>
F
for_each(F op, It0 it0, It0 end0, Its... its) {
while (it0 != end0)
__iter_helper<void, __empty, sizeof...(Its)>::step(
op, {}, std::tie(it0, its...));
return op;
}
///
/// Iterate \a op on the result of zipping all the specified
/// iterators together, storing return values in a new container.
///
/// Similar to std::transform, but it accepts functions of an
/// arbitrary number of arguments and it doesn't have to be
/// provided with an output iterator.
///
template<typename F, typename It0, typename... Its,
typename C = std::vector<
typename __result_helper<F, It0, Its...>::type>>
C
map(F op, It0 it0, It0 end0, Its... its) {
C c;
while (it0 != end0)
c.push_back(
__iter_helper<typename C::value_type, __empty, sizeof...(Its)>
::step(op, {}, std::tie(it0, its...)));
return c;
}
///
/// Reduce the result of zipping all the specified iterators
/// together, using iterative application of \a op from left to
/// right.
///
/// Similar to std::accumulate, but it accepts functions of an
/// arbitrary number of arguments.
///
template<typename F, typename T, typename It0, typename... Its>
T
fold(F op, T a, It0 it0, It0 end0, Its... its) {
while (it0 != end0)
a = __iter_helper<T, T, sizeof...(Its)>::step(
op, a, std::tie(it0, its...));
return a;
}
///
/// Iterate \a op on the result of zipping the specified iterators
/// together, checking if any of the evaluations returns \a true.
///
/// Similar to std::any_of, but it accepts functions of an
/// arbitrary number of arguments.
///
template<typename F, typename It0, typename... Its>
bool
any_of(F op, It0 it0, It0 end0, Its... its) {
while (it0 != end0)
if (__iter_helper<bool, __empty, sizeof...(Its)>::step(
op, {}, std::tie(it0, its...)))
return true;
return false;
}
template<typename T, typename S>
T
keys(const std::pair<T, S> &ent) {
return ent.first;
}
template<typename T, typename S>
std::function<bool (const std::pair<T, S> &)>
key_equals(const T &x) {
return [=](const std::pair<T, S> &ent) {
return ent.first == x;
};
}
template<typename T, typename S>
S
values(const std::pair<T, S> &ent) {
return ent.second;
}
template<typename T>
std::function<bool (const T &)>
is_zero() {
return [](const T &x) {
return x == 0;
};
}
}
#endif

290
src/gallium/state_trackers/clover/core/compat.hpp Normal file
View File

@ -0,0 +1,290 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_COMPAT_HPP__
#define __CORE_COMPAT_HPP__
#include <new>
#include <cstring>
#include <cstdlib>
#include <string>
#include <stdint.h>
namespace clover {
namespace compat {
// XXX - For cases where we can't rely on STL... I.e. the
// interface between code compiled as C++98 and C++11
// source. Get rid of this as soon as everything can be
// compiled as C++11.
template<typename T>
class vector {
protected:
static T *
alloc(int n, const T *q, int m) {
T *p = reinterpret_cast<T *>(std::malloc(n * sizeof(T)));
for (int i = 0; i < m; ++i)
new(&p[i]) T(q[i]);
return p;
}
static void
free(int n, T *p) {
for (int i = 0; i < n; ++i)
p[i].~T();
std::free(p);
}
public:
vector() : p(NULL), n(0) {
}
vector(const vector &v) : p(alloc(v.n, v.p, v.n)), n(v.n) {
}
vector(T *p, size_t n) : p(alloc(n, p, n)), n(n) {
}
template<typename C>
vector(const C &v) :
p(alloc(v.size(), &*v.begin(), v.size())), n(v.size()) {
}
~vector() {
free(n, p);
}
vector &
operator=(const vector &v) {
free(n, p);
p = alloc(v.n, v.p, v.n);
n = v.n;
return *this;
}
void
reserve(size_t m) {
if (n < m) {
T *q = alloc(m, p, n);
free(n, p);
p = q;
n = m;
}
}
void
resize(size_t m, T x = T()) {
size_t n = size();
reserve(m);
for (size_t i = n; i < m; ++i)
new(&p[i]) T(x);
}
void
push_back(const T &x) {
size_t n = size();
reserve(n + 1);
new(&p[n]) T(x);
}
size_t
size() const {
return n;
}
T *
begin() {
return p;
}
const T *
begin() const {
return p;
}
T *
end() {
return p + n;
}
const T *
end() const {
return p + n;
}
T &
operator[](int i) {
return p[i];
}
const T &
operator[](int i) const {
return p[i];
}
private:
T *p;
size_t n;
};
template<typename T>
class vector_ref {
public:
vector_ref(T *p, size_t n) : p(p), n(n) {
}
template<typename C>
vector_ref(C &v) : p(&*v.begin()), n(v.size()) {
}
size_t
size() const {
return n;
}
T *
begin() {
return p;
}
const T *
begin() const {
return p;
}
T *
end() {
return p + n;
}
const T *
end() const {
return p + n;
}
T &
operator[](int i) {
return p[i];
}
const T &
operator[](int i) const {
return p[i];
}
private:
T *p;
size_t n;
};
class istream {
public:
typedef vector_ref<const unsigned char> buffer_t;
class error {
public:
virtual ~error() {}
};
istream(const buffer_t &buf) : buf(buf), offset(0) {}
void
read(char *p, size_t n) {
if (offset + n > buf.size())
throw error();
std::memcpy(p, buf.begin() + offset, n);
offset += n;
}
private:
const buffer_t &buf;
size_t offset;
};
class ostream {
public:
typedef vector<unsigned char> buffer_t;
ostream(buffer_t &buf) : buf(buf), offset(buf.size()) {}
void
write(const char *p, size_t n) {
buf.resize(offset + n);
std::memcpy(buf.begin() + offset, p, n);
offset += n;
}
private:
buffer_t &buf;
size_t offset;
};
class string : public vector_ref<const char> {
public:
string(const char *p) : vector_ref(p, std::strlen(p)) {
}
template<typename C>
string(const C &v) : vector_ref(v) {
}
operator std::string() const {
return std::string(begin(), end());
}
const char *
find(const string &s) const {
for (size_t i = 0; i + s.size() < size(); ++i) {
if (!std::memcmp(begin() + i, s.begin(), s.size()))
return begin() + i;
}
return end();
}
};
template<typename T>
bool
operator==(const vector_ref<T> &a, const vector_ref<T> &b) {
if (a.size() != b.size())
return false;
for (size_t i = 0; i < a.size(); ++i)
if (a[i] != b[i])
return false;
return true;
}
}
}
#endif

53
src/gallium/state_trackers/clover/core/compiler.hpp Normal file
View File

@ -0,0 +1,53 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_COMPILER_HPP__
#define __CORE_COMPILER_HPP__
#include "core/compat.hpp"
#include "core/module.hpp"
namespace clover {
class build_error {
public:
build_error(const compat::string &log) : log(log) {
}
virtual ~build_error() {
}
compat::string what() {
return log;
}
private:
compat::vector<char> log;
};
module compile_program_llvm(const compat::string &source,
const compat::string &target);
module compile_program_tgsi(const compat::string &source,
const compat::string &target);
}
#endif

37
src/gallium/state_trackers/clover/core/context.cpp Normal file
View File

@ -0,0 +1,37 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <algorithm>
#include "core/context.hpp"
using namespace clover;
_cl_context::_cl_context(const std::vector<cl_context_properties> &props,
const std::vector<device *> &devs) :
devs(devs), __props(props) {
}
bool
_cl_context::has_device(clover::device *dev) const {
return std::count(devs.begin(), devs.end(), dev);
}

51
src/gallium/state_trackers/clover/core/context.hpp Normal file
View File

@ -0,0 +1,51 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_CONTEXT_HPP__
#define __CORE_CONTEXT_HPP__
#include "core/base.hpp"
#include "core/device.hpp"
namespace clover {
typedef struct _cl_context context;
}
struct _cl_context : public clover::ref_counter {
public:
_cl_context(const std::vector<cl_context_properties> &props,
const std::vector<clover::device *> &devs);
_cl_context(const _cl_context &ctx) = delete;
bool has_device(clover::device *dev) const;
const std::vector<cl_context_properties> &props() const {
return __props;
}
const std::vector<clover::device *> devs;
private:
std::vector<cl_context_properties> __props;
};
#endif

179
src/gallium/state_trackers/clover/core/device.cpp Normal file
View File

@ -0,0 +1,179 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/device.hpp"
#include "pipe/p_screen.h"
#include "pipe/p_state.h"
using namespace clover;
namespace {
template<typename T>
std::vector<T>
get_compute_param(pipe_screen *pipe, pipe_compute_cap cap) {
int sz = pipe->get_compute_param(pipe, cap, NULL);
std::vector<T> v(sz / sizeof(T));
pipe->get_compute_param(pipe, cap, &v.front());
return v;
}
}
_cl_device_id::_cl_device_id(pipe_loader_device *ldev) : ldev(ldev) {
pipe = pipe_loader_create_screen(ldev, PIPE_SEARCH_DIR);
if (!pipe || !pipe->get_param(pipe, PIPE_CAP_COMPUTE))
throw error(CL_INVALID_DEVICE);
}
_cl_device_id::_cl_device_id(_cl_device_id &&dev) : pipe(dev.pipe), ldev(dev.ldev) {
dev.ldev = NULL;
dev.pipe = NULL;
}
_cl_device_id::~_cl_device_id() {
if (pipe)
pipe->destroy(pipe);
if (ldev)
pipe_loader_release(&ldev, 1);
}
cl_device_type
_cl_device_id::type() const {
switch (ldev->type) {
case PIPE_LOADER_DEVICE_SOFTWARE:
return CL_DEVICE_TYPE_CPU;
case PIPE_LOADER_DEVICE_PCI:
return CL_DEVICE_TYPE_GPU;
default:
assert(0);
return 0;
}
}
cl_uint
_cl_device_id::vendor_id() const {
switch (ldev->type) {
case PIPE_LOADER_DEVICE_SOFTWARE:
return 0;
case PIPE_LOADER_DEVICE_PCI:
return ldev->pci.vendor_id;
default:
assert(0);
return 0;
}
}
size_t
_cl_device_id::max_images_read() const {
return PIPE_MAX_SHADER_RESOURCES;
}
size_t
_cl_device_id::max_images_write() const {
return PIPE_MAX_SHADER_RESOURCES;
}
cl_uint
_cl_device_id::max_image_levels_2d() const {
return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_2D_LEVELS);
}
cl_uint
_cl_device_id::max_image_levels_3d() const {
return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_3D_LEVELS);
}
cl_uint
_cl_device_id::max_samplers() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_TEXTURE_SAMPLERS);
}
cl_ulong
_cl_device_id::max_mem_global() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE)[0];
}
cl_ulong
_cl_device_id::max_mem_local() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE)[0];
}
cl_ulong
_cl_device_id::max_mem_input() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_INPUT_SIZE)[0];
}
cl_ulong
_cl_device_id::max_const_buffer_size() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_CONSTS) * 16;
}
cl_uint
_cl_device_id::max_const_buffers() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_CONST_BUFFERS);
}
std::vector<size_t>
_cl_device_id::max_block_size() const {
return get_compute_param<uint64_t>(pipe, PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE);
}
std::string
_cl_device_id::device_name() const {
return pipe->get_name(pipe);
}
std::string
_cl_device_id::vendor_name() const {
return pipe->get_vendor(pipe);
}
std::string
_cl_device_id::ir_target() const {
switch (pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_PREFERRED_IR)) {
case PIPE_SHADER_IR_TGSI:
return "tgsi";
default:
assert(0);
return "";
}
}
device_registry::device_registry() {
int n = pipe_loader_probe(NULL, 0);
std::vector<pipe_loader_device *> ldevs(n);
pipe_loader_probe(&ldevs.front(), n);
for (pipe_loader_device *ldev : ldevs) {
try {
devs.emplace_back(ldev);
} catch (error &) {}
}
}

107
src/gallium/state_trackers/clover/core/device.hpp Normal file
View File

@ -0,0 +1,107 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_DEVICE_HPP__
#define __CORE_DEVICE_HPP__
#include <set>
#include <vector>
#include "core/base.hpp"
#include "core/format.hpp"
#include "pipe-loader/pipe_loader.h"
namespace clover {
typedef struct _cl_device_id device;
class root_resource;
class hard_event;
}
struct _cl_device_id {
public:
_cl_device_id(pipe_loader_device *ldev);
_cl_device_id(_cl_device_id &&dev);
_cl_device_id(const _cl_device_id &dev) = delete;
~_cl_device_id();
cl_device_type type() const;
cl_uint vendor_id() const;
size_t max_images_read() const;
size_t max_images_write() const;
cl_uint max_image_levels_2d() const;
cl_uint max_image_levels_3d() const;
cl_uint max_samplers() const;
cl_ulong max_mem_global() const;
cl_ulong max_mem_local() const;
cl_ulong max_mem_input() const;
cl_ulong max_const_buffer_size() const;
cl_uint max_const_buffers() const;
std::vector<size_t> max_block_size() const;
std::string device_name() const;
std::string vendor_name() const;
std::string ir_target() const;
friend struct _cl_command_queue;
friend class clover::root_resource;
friend class clover::hard_event;
friend std::set<cl_image_format>
clover::supported_formats(cl_context, cl_mem_object_type);
private:
pipe_screen *pipe;
pipe_loader_device *ldev;
};
namespace clover {
///
/// Container of all the compute devices that are available in the
/// system.
///
class device_registry {
public:
typedef std::vector<device>::iterator iterator;
device_registry();
iterator begin() {
return devs.begin();
}
iterator end() {
return devs.end();
}
device &front() {
return devs.front();
}
device &back() {
return devs.back();
}
protected:
std::vector<device> devs;
};
}
#endif

175
src/gallium/state_trackers/clover/core/event.cpp Normal file
View File

@ -0,0 +1,175 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/event.hpp"
#include "pipe/p_screen.h"
using namespace clover;
_cl_event::_cl_event(clover::context &ctx,
std::vector<clover::event *> deps,
action action_ok, action action_fail) :
ctx(ctx), __status(0), wait_count(1),
action_ok(action_ok), action_fail(action_fail) {
for (auto ev : deps)
ev->chain(this);
}
_cl_event::~_cl_event() {
}
void
_cl_event::trigger() {
if (!--wait_count) {
action_ok(*this);
while (!__chain.empty()) {
__chain.back()->trigger();
__chain.pop_back();
}
}
}
void
_cl_event::abort(cl_int status) {
__status = status;
action_fail(*this);
while (!__chain.empty()) {
__chain.back()->abort(status);
__chain.pop_back();
}
}
bool
_cl_event::signalled() const {
return !wait_count;
}
void
_cl_event::chain(clover::event *ev) {
if (wait_count) {
ev->wait_count++;
__chain.push_back(ev);
ev->deps.push_back(this);
}
}
hard_event::hard_event(clover::command_queue &q, cl_command_type command,
std::vector<clover::event *> deps, action action) :
_cl_event(q.ctx, deps, action, [](event &ev){}),
__queue(q), __command(command), __fence(NULL) {
q.sequence(this);
trigger();
}
hard_event::~hard_event() {
pipe_screen *screen = queue()->dev.pipe;
screen->fence_reference(screen, &__fence, NULL);
}
cl_int
hard_event::status() const {
pipe_screen *screen = queue()->dev.pipe;
if (__status < 0)
return __status;
else if (!__fence)
return CL_QUEUED;
else if (!screen->fence_signalled(screen, __fence))
return CL_SUBMITTED;
else
return CL_COMPLETE;
}
cl_command_queue
hard_event::queue() const {
return &__queue;
}
cl_command_type
hard_event::command() const {
return __command;
}
void
hard_event::wait() const {
pipe_screen *screen = queue()->dev.pipe;
if (status() == CL_QUEUED)
queue()->flush();
if (!__fence ||
!screen->fence_finish(screen, __fence, PIPE_TIMEOUT_INFINITE))
throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
}
void
hard_event::fence(pipe_fence_handle *fence) {
pipe_screen *screen = queue()->dev.pipe;
screen->fence_reference(screen, &__fence, fence);
}
soft_event::soft_event(clover::context &ctx,
std::vector<clover::event *> deps,
bool __trigger, action action) :
_cl_event(ctx, deps, action, action) {
if (__trigger)
trigger();
}
cl_int
soft_event::status() const {
if (__status < 0)
return __status;
else if (!signalled() ||
any_of([](const ref_ptr<event> &ev) {
return ev->status() != CL_COMPLETE;
}, deps.begin(), deps.end()))
return CL_SUBMITTED;
else
return CL_COMPLETE;
}
cl_command_queue
soft_event::queue() const {
return NULL;
}
cl_command_type
soft_event::command() const {
return CL_COMMAND_USER;
}
void
soft_event::wait() const {
for (auto ev : deps)
ev->wait();
if (status() != CL_COMPLETE)
throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
}

138
src/gallium/state_trackers/clover/core/event.hpp Normal file
View File

@ -0,0 +1,138 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_EVENT_HPP__
#define __CORE_EVENT_HPP__
#include <functional>
#include "core/base.hpp"
#include "core/queue.hpp"
namespace clover {
typedef struct _cl_event event;
}
///
/// Class that represents a task that might be executed asynchronously
/// at some point in the future.
///
/// An event consists of a list of dependencies, a boolean signalled()
/// flag, and an associated task. An event is considered signalled as
/// soon as all its dependencies (if any) are signalled as well, and
/// the trigger() method is called; at that point the associated task
/// will be started through the specified \a action_ok. If the
/// abort() method is called instead, the specified \a action_fail is
/// executed and the associated task will never be started. Dependent
/// events will be aborted recursively.
///
/// The execution status of the associated task can be queried using
/// the status() method, and it can be waited for completion using the
/// wait() method.
///
struct _cl_event : public clover::ref_counter {
public:
typedef std::function<void (clover::event &)> action;
_cl_event(clover::context &ctx, std::vector<clover::event *> deps,
action action_ok, action action_fail);
virtual ~_cl_event();
void trigger();
void abort(cl_int status);
bool signalled() const;
virtual cl_int status() const = 0;
virtual cl_command_queue queue() const = 0;
virtual cl_command_type command() const = 0;
virtual void wait() const = 0;
clover::context &ctx;
protected:
void chain(clover::event *ev);
cl_int __status;
std::vector<clover::ref_ptr<clover::event>> deps;
private:
unsigned wait_count;
action action_ok;
action action_fail;
std::vector<clover::ref_ptr<clover::event>> __chain;
};
namespace clover {
///
/// Class that represents a task executed by a command queue.
///
/// Similar to a normal clover::event. In addition it's associated
/// with a given command queue \a q and a given OpenCL \a command.
/// hard_event instances created for the same queue are implicitly
/// ordered with respect to each other, and they are implicitly
/// triggered on construction.
///
/// A hard_event is considered complete when the associated
/// hardware task finishes execution.
///
class hard_event : public event {
public:
hard_event(clover::command_queue &q, cl_command_type command,
std::vector<clover::event *> deps,
action action = [](event &){});
~hard_event();
virtual cl_int status() const;
virtual cl_command_queue queue() const;
virtual cl_command_type command() const;
virtual void wait() const;
friend class ::_cl_command_queue;
private:
virtual void fence(pipe_fence_handle *fence);
clover::command_queue &__queue;
cl_command_type __command;
pipe_fence_handle *__fence;
};
///
/// Class that represents a software event.
///
/// A soft_event is not associated with any specific hardware task
/// or command queue. It's considered complete as soon as all its
/// dependencies finish execution.
///
class soft_event : public event {
public:
soft_event(clover::context &ctx, std::vector<clover::event *> deps,
bool trigger, action action = [](event &){});
virtual cl_int status() const;
virtual cl_command_queue queue() const;
virtual cl_command_type command() const;
virtual void wait() const;
};
}
#endif

167
src/gallium/state_trackers/clover/core/format.cpp Normal file
View File

@ -0,0 +1,167 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <algorithm>
#include "core/format.hpp"
#include "core/memory.hpp"
#include "pipe/p_screen.h"
#include "pipe/p_context.h"
namespace clover {
static const std::map<cl_image_format, pipe_format> formats {
{ { CL_BGRA, CL_UNORM_INT8 }, PIPE_FORMAT_B8G8R8A8_UNORM },
{ { CL_ARGB, CL_UNORM_INT8 }, PIPE_FORMAT_A8R8G8B8_UNORM },
{ { CL_RGB, CL_UNORM_SHORT_565 }, PIPE_FORMAT_B5G6R5_UNORM },
{ { CL_LUMINANCE, CL_UNORM_INT8 }, PIPE_FORMAT_L8_UNORM },
{ { CL_A, CL_UNORM_INT8 }, PIPE_FORMAT_A8_UNORM },
{ { CL_INTENSITY, CL_UNORM_INT8 }, PIPE_FORMAT_I8_UNORM },
{ { CL_LUMINANCE, CL_UNORM_INT16 }, PIPE_FORMAT_L16_UNORM },
{ { CL_R, CL_FLOAT }, PIPE_FORMAT_R32_FLOAT },
{ { CL_RG, CL_FLOAT }, PIPE_FORMAT_R32G32_FLOAT },
{ { CL_RGB, CL_FLOAT }, PIPE_FORMAT_R32G32B32_FLOAT },
{ { CL_RGBA, CL_FLOAT }, PIPE_FORMAT_R32G32B32A32_FLOAT },
{ { CL_R, CL_UNORM_INT16 }, PIPE_FORMAT_R16_UNORM },
{ { CL_RG, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16_UNORM },
{ { CL_RGB, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16B16_UNORM },
{ { CL_RGBA, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16B16A16_UNORM },
{ { CL_R, CL_SNORM_INT16 }, PIPE_FORMAT_R16_SNORM },
{ { CL_RG, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16_SNORM },
{ { CL_RGB, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16B16_SNORM },
{ { CL_RGBA, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16B16A16_SNORM },
{ { CL_R, CL_UNORM_INT8 }, PIPE_FORMAT_R8_UNORM },
{ { CL_RG, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8_UNORM },
{ { CL_RGB, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8_UNORM },
{ { CL_RGBA, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8A8_UNORM },
{ { CL_R, CL_SNORM_INT8 }, PIPE_FORMAT_R8_SNORM },
{ { CL_RG, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8_SNORM },
{ { CL_RGB, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8B8_SNORM },
{ { CL_RGBA, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8B8A8_SNORM },
{ { CL_R, CL_HALF_FLOAT }, PIPE_FORMAT_R16_FLOAT },
{ { CL_RG, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16_FLOAT },
{ { CL_RGB, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16B16_FLOAT },
{ { CL_RGBA, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16B16A16_FLOAT },
{ { CL_RGBx, CL_UNORM_SHORT_555 }, PIPE_FORMAT_B5G5R5X1_UNORM },
{ { CL_RGBx, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8X8_UNORM },
{ { CL_A, CL_UNORM_INT16 }, PIPE_FORMAT_A16_UNORM },
{ { CL_INTENSITY, CL_UNORM_INT16 }, PIPE_FORMAT_I16_UNORM },
{ { CL_LUMINANCE, CL_SNORM_INT8 }, PIPE_FORMAT_L8_SNORM },
{ { CL_INTENSITY, CL_SNORM_INT8 }, PIPE_FORMAT_I8_SNORM },
{ { CL_A, CL_SNORM_INT16 }, PIPE_FORMAT_A16_SNORM },
{ { CL_LUMINANCE, CL_SNORM_INT16 }, PIPE_FORMAT_L16_SNORM },
{ { CL_INTENSITY, CL_SNORM_INT16 }, PIPE_FORMAT_I16_SNORM },
{ { CL_A, CL_HALF_FLOAT }, PIPE_FORMAT_A16_FLOAT },
{ { CL_LUMINANCE, CL_HALF_FLOAT }, PIPE_FORMAT_L16_FLOAT },
{ { CL_INTENSITY, CL_HALF_FLOAT }, PIPE_FORMAT_I16_FLOAT },
{ { CL_A, CL_FLOAT }, PIPE_FORMAT_A32_FLOAT },
{ { CL_LUMINANCE, CL_FLOAT }, PIPE_FORMAT_L32_FLOAT },
{ { CL_INTENSITY, CL_FLOAT }, PIPE_FORMAT_I32_FLOAT },
{ { CL_RA, CL_UNORM_INT8 }, PIPE_FORMAT_R8A8_UNORM },
{ { CL_R, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8_UINT },
{ { CL_RG, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8_UINT },
{ { CL_RGB, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8B8_UINT },
{ { CL_RGBA, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8B8A8_UINT },
{ { CL_R, CL_SIGNED_INT8 }, PIPE_FORMAT_R8_SINT },
{ { CL_RG, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8_SINT },
{ { CL_RGB, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8B8_SINT },
{ { CL_RGBA, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8B8A8_SINT },
{ { CL_R, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16_UINT },
{ { CL_RG, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16_UINT },
{ { CL_RGB, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16B16_UINT },
{ { CL_RGBA, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16B16A16_UINT },
{ { CL_R, CL_SIGNED_INT16 }, PIPE_FORMAT_R16_SINT },
{ { CL_RG, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16_SINT },
{ { CL_RGB, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16B16_SINT },
{ { CL_RGBA, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16B16A16_SINT },
{ { CL_R, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32_UINT },
{ { CL_RG, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32_UINT },
{ { CL_RGB, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32B32_UINT },
{ { CL_RGBA, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32B32A32_UINT },
{ { CL_R, CL_SIGNED_INT32 }, PIPE_FORMAT_R32_SINT },
{ { CL_RG, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32_SINT },
{ { CL_RGB, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32B32_SINT },
{ { CL_RGBA, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32B32A32_SINT },
{ { CL_A, CL_UNSIGNED_INT8 }, PIPE_FORMAT_A8_UINT },
{ { CL_INTENSITY, CL_UNSIGNED_INT8 }, PIPE_FORMAT_I8_UINT },
{ { CL_LUMINANCE, CL_UNSIGNED_INT8 }, PIPE_FORMAT_L8_UINT },
{ { CL_A, CL_SIGNED_INT8 }, PIPE_FORMAT_A8_SINT },
{ { CL_INTENSITY, CL_SIGNED_INT8 }, PIPE_FORMAT_I8_SINT },
{ { CL_LUMINANCE, CL_SIGNED_INT8 }, PIPE_FORMAT_L8_SINT },
{ { CL_A, CL_UNSIGNED_INT16 }, PIPE_FORMAT_A16_UINT },
{ { CL_INTENSITY, CL_UNSIGNED_INT16 }, PIPE_FORMAT_I16_UINT },
{ { CL_LUMINANCE, CL_UNSIGNED_INT16 }, PIPE_FORMAT_L16_UINT },
{ { CL_A, CL_SIGNED_INT16 }, PIPE_FORMAT_A16_SINT },
{ { CL_INTENSITY, CL_SIGNED_INT16 }, PIPE_FORMAT_I16_SINT },
{ { CL_LUMINANCE, CL_SIGNED_INT16 }, PIPE_FORMAT_L16_SINT },
{ { CL_A, CL_UNSIGNED_INT32 }, PIPE_FORMAT_A32_UINT },
{ { CL_INTENSITY, CL_UNSIGNED_INT32 }, PIPE_FORMAT_I32_UINT },
{ { CL_LUMINANCE, CL_UNSIGNED_INT32 }, PIPE_FORMAT_L32_UINT },
{ { CL_A, CL_SIGNED_INT32 }, PIPE_FORMAT_A32_SINT },
{ { CL_INTENSITY, CL_SIGNED_INT32 }, PIPE_FORMAT_I32_SINT },
{ { CL_LUMINANCE, CL_SIGNED_INT32 }, PIPE_FORMAT_L32_SINT }
};
pipe_texture_target
translate_target(cl_mem_object_type type) {
switch (type) {
case CL_MEM_OBJECT_BUFFER:
return PIPE_BUFFER;
case CL_MEM_OBJECT_IMAGE2D:
return PIPE_TEXTURE_2D;
case CL_MEM_OBJECT_IMAGE3D:
return PIPE_TEXTURE_3D;
default:
throw error(CL_INVALID_VALUE);
}
}
pipe_format
translate_format(const cl_image_format &format) {
auto it = formats.find(format);
if (it == formats.end())
throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED);
return it->second;
}
std::set<cl_image_format>
supported_formats(cl_context ctx, cl_mem_object_type type) {
std::set<cl_image_format> s;
pipe_texture_target target = translate_target(type);
unsigned bindings = (PIPE_BIND_SAMPLER_VIEW |
PIPE_BIND_COMPUTE_RESOURCE |
PIPE_BIND_TRANSFER_READ |
PIPE_BIND_TRANSFER_WRITE);
for (auto f : formats) {
if (std::all_of(ctx->devs.begin(), ctx->devs.end(),
[=](const device *dev) {
return dev->pipe->is_format_supported(
dev->pipe, f.second, target, 1, bindings);
}))
s.insert(f.first);
}
return s;
}
}

51
src/gallium/state_trackers/clover/core/format.hpp Normal file
View File

@ -0,0 +1,51 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_FORMAT_HPP__
#define __CORE_FORMAT_HPP__
#include <set>
#include "core/base.hpp"
#include "pipe/p_defines.h"
#include "pipe/p_format.h"
namespace clover {
pipe_texture_target translate_target(cl_mem_object_type type);
pipe_format translate_format(const cl_image_format &format);
///
/// Return all the image formats supported by a given context for
/// the given memory object type.
///
std::set<cl_image_format> supported_formats(cl_context ctx,
cl_mem_object_type type);
}
static inline bool
operator<(const cl_image_format &a, const cl_image_format &b) {
return (a.image_channel_order != b.image_channel_order ?
a.image_channel_order < b.image_channel_order :
a.image_channel_data_type < b.image_channel_data_type);
}
#endif

72
src/gallium/state_trackers/clover/core/geometry.hpp Normal file
View File

@ -0,0 +1,72 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_GEOMETRY_HPP__
#define __CORE_GEOMETRY_HPP__
#include <array>
#include <algorithm>
namespace clover {
///
/// N-dimensional coordinate array.
///
template<typename T, int N>
class point {
public:
point() : a() {
}
point(std::initializer_list<T> v) {
auto it = std::copy(v.begin(), v.end(), a.begin());
std::fill(it, a.end(), 0);
}
point(const T *v) {
std::copy(v, v + N, a.begin());
}
T &operator[](int i) {
return a[i];
}
const T &operator[](int i) const {
return a[i];
}
point operator+(const point &p) const {
point q;
std::transform(a.begin(), a.end(), p.a.begin(),
q.a.begin(), std::plus<T>());
return q;
}
T operator()(const point &p) const {
return std::inner_product(p.a.begin(), p.a.end(), a.begin(), 0);
}
protected:
std::array<T, N> a;
};
}
#endif

393
src/gallium/state_trackers/clover/core/kernel.cpp Normal file
View File

@ -0,0 +1,393 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/kernel.hpp"
#include "core/resource.hpp"
#include "pipe/p_context.h"
using namespace clover;
_cl_kernel::_cl_kernel(clover::program &prog,
const std::string &name,
const std::vector<clover::module::argument> &args) :
prog(prog), __name(name), exec(*this) {
for (auto arg : args) {
if (arg.type == module::argument::scalar)
this->args.emplace_back(new scalar_argument(arg.size));
else if (arg.type == module::argument::global)
this->args.emplace_back(new global_argument(arg.size));
else if (arg.type == module::argument::local)
this->args.emplace_back(new local_argument());
else if (arg.type == module::argument::constant)
this->args.emplace_back(new constant_argument());
else if (arg.type == module::argument::image2d_rd ||
arg.type == module::argument::image3d_rd)
this->args.emplace_back(new image_rd_argument());
else if (arg.type == module::argument::image2d_wr ||
arg.type == module::argument::image3d_wr)
this->args.emplace_back(new image_wr_argument());
else if (arg.type == module::argument::sampler)
this->args.emplace_back(new sampler_argument());
else
throw error(CL_INVALID_KERNEL_DEFINITION);
}
}
template<typename T, typename V>
static inline std::vector<T>
pad_vector(clover::command_queue &q, const V &v, T x) {
std::vector<T> w { v.begin(), v.end() };
w.resize(q.dev.max_block_size().size(), x);
return w;
}
void
_cl_kernel::launch(clover::command_queue &q,
const std::vector<size_t> &grid_offset,
const std::vector<size_t> &grid_size,
const std::vector<size_t> &block_size) {
void *st = exec.bind(&q);
auto g_handles = map([&](size_t h) { return (uint32_t *)&exec.input[h]; },
exec.g_handles.begin(), exec.g_handles.end());
q.pipe->bind_compute_state(q.pipe, st);
q.pipe->bind_compute_sampler_states(q.pipe, 0, exec.samplers.size(),
exec.samplers.data());
q.pipe->set_compute_sampler_views(q.pipe, 0, exec.sviews.size(),
exec.sviews.data());
q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(),
exec.resources.data());
q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(),
exec.g_buffers.data(), g_handles.data());
q.pipe->launch_grid(q.pipe,
pad_vector<uint>(q, block_size, 1).data(),
pad_vector<uint>(q, grid_size, 1).data(),
module(q).sym(__name).offset,
exec.input.data());
q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(), NULL, NULL);
q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(), NULL);
q.pipe->set_compute_sampler_views(q.pipe, 0, exec.sviews.size(), NULL);
q.pipe->bind_compute_sampler_states(q.pipe, 0, exec.samplers.size(), NULL);
exec.unbind();
}
size_t
_cl_kernel::mem_local() const {
size_t sz = 0;
for (auto &arg : args) {
if (dynamic_cast<local_argument *>(arg.get()))
sz += arg->storage();
}
return sz;
}
size_t
_cl_kernel::mem_private() const {
return 0;
}
size_t
_cl_kernel::max_block_size() const {
return SIZE_MAX;
}
const std::string &
_cl_kernel::name() const {
return __name;
}
std::vector<size_t>
_cl_kernel::block_size() const {
return { 0, 0, 0 };
}
const clover::module &
_cl_kernel::module(const clover::command_queue &q) const {
return prog.binaries().find(&q.dev)->second;
}
_cl_kernel::exec_context::exec_context(clover::kernel &kern) :
kern(kern), q(NULL), mem_local(0), st(NULL) {
}
_cl_kernel::exec_context::~exec_context() {
if (st)
q->pipe->delete_compute_state(q->pipe, st);
}
void *
_cl_kernel::exec_context::bind(clover::command_queue *__q) {
std::swap(q, __q);
for (auto &arg : kern.args)
arg->bind(*this);
// Create a new compute state if anything changed.
if (!st || q != __q ||
cs.req_local_mem != mem_local ||
cs.req_input_mem != input.size()) {
if (st)
__q->pipe->delete_compute_state(__q->pipe, st);
cs.prog = kern.module(*q).sec(module::section::text).data.begin();
cs.req_local_mem = mem_local;
cs.req_input_mem = input.size();
st = q->pipe->create_compute_state(q->pipe, &cs);
}
return st;
}
void
_cl_kernel::exec_context::unbind() {
for (auto &arg : kern.args)
arg->unbind(*this);
input.clear();
samplers.clear();
sviews.clear();
resources.clear();
g_buffers.clear();
g_handles.clear();
mem_local = 0;
}
_cl_kernel::argument::argument(size_t size) :
__size(size), __set(false) {
}
bool
_cl_kernel::argument::set() const {
return __set;
}
size_t
_cl_kernel::argument::storage() const {
return 0;
}
_cl_kernel::scalar_argument::scalar_argument(size_t size) :
argument(size) {
}
void
_cl_kernel::scalar_argument::set(size_t size, const void *value) {
if (size != __size)
throw error(CL_INVALID_ARG_SIZE);
v = { (uint8_t *)value, (uint8_t *)value + size };
__set = true;
}
void
_cl_kernel::scalar_argument::bind(exec_context &ctx) {
ctx.input.insert(ctx.input.end(), v.begin(), v.end());
}
void
_cl_kernel::scalar_argument::unbind(exec_context &ctx) {
}
_cl_kernel::global_argument::global_argument(size_t size) :
argument(size) {
}
void
_cl_kernel::global_argument::set(size_t size, const void *value) {
if (size != sizeof(cl_mem))
throw error(CL_INVALID_ARG_SIZE);
obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
__set = true;
}
void
_cl_kernel::global_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.g_buffers.size();
ctx.input.resize(offset + __size);
ctx.g_buffers.resize(idx + 1);
ctx.g_buffers[idx] = obj->resource(ctx.q).pipe;
ctx.g_handles.resize(idx + 1);
ctx.g_handles[idx] = offset;
}
void
_cl_kernel::global_argument::unbind(exec_context &ctx) {
}
_cl_kernel::local_argument::local_argument() :
argument(sizeof(uint32_t)) {
}
size_t
_cl_kernel::local_argument::storage() const {
return __storage;
}
void
_cl_kernel::local_argument::set(size_t size, const void *value) {
if (value)
throw error(CL_INVALID_ARG_VALUE);
__storage = size;
__set = true;
}
void
_cl_kernel::local_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t ptr = ctx.mem_local;
ctx.input.resize(offset + sizeof(uint32_t));
*(uint32_t *)&ctx.input[offset] = ptr;
ctx.mem_local += __storage;
}
void
_cl_kernel::local_argument::unbind(exec_context &ctx) {
}
_cl_kernel::constant_argument::constant_argument() :
argument(sizeof(uint32_t)) {
}
void
_cl_kernel::constant_argument::set(size_t size, const void *value) {
if (size != sizeof(cl_mem))
throw error(CL_INVALID_ARG_SIZE);
obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
__set = true;
}
void
_cl_kernel::constant_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.resources.size();
ctx.input.resize(offset + sizeof(uint32_t));
*(uint32_t *)&ctx.input[offset] = idx << 24;
ctx.resources.resize(idx + 1);
ctx.resources[idx] = st = obj->resource(ctx.q).bind_surface(*ctx.q, false);
}
void
_cl_kernel::constant_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_surface(*ctx.q, st);
}
_cl_kernel::image_rd_argument::image_rd_argument() :
argument(sizeof(uint32_t)) {
}
void
_cl_kernel::image_rd_argument::set(size_t size, const void *value) {
if (size != sizeof(cl_mem))
throw error(CL_INVALID_ARG_SIZE);
obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
__set = true;
}
void
_cl_kernel::image_rd_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.sviews.size();
ctx.input.resize(offset + sizeof(uint32_t));
*(uint32_t *)&ctx.input[offset] = idx;
ctx.sviews.resize(idx + 1);
ctx.sviews[idx] = st = obj->resource(ctx.q).bind_sampler_view(*ctx.q);
}
void
_cl_kernel::image_rd_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_sampler_view(*ctx.q, st);
}
_cl_kernel::image_wr_argument::image_wr_argument() :
argument(sizeof(uint32_t)) {
}
void
_cl_kernel::image_wr_argument::set(size_t size, const void *value) {
if (size != sizeof(cl_mem))
throw error(CL_INVALID_ARG_SIZE);
obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
__set = true;
}
void
_cl_kernel::image_wr_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.resources.size();
ctx.input.resize(offset + sizeof(uint32_t));
*(uint32_t *)&ctx.input[offset] = idx;
ctx.resources.resize(idx + 1);
ctx.resources[idx] = st = obj->resource(ctx.q).bind_surface(*ctx.q, true);
}
void
_cl_kernel::image_wr_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_surface(*ctx.q, st);
}
_cl_kernel::sampler_argument::sampler_argument() :
argument(0) {
}
void
_cl_kernel::sampler_argument::set(size_t size, const void *value) {
if (size != sizeof(cl_sampler))
throw error(CL_INVALID_ARG_SIZE);
obj = *(cl_sampler *)value;
__set = true;
}
void
_cl_kernel::sampler_argument::bind(exec_context &ctx) {
size_t idx = ctx.samplers.size();
ctx.samplers.resize(idx + 1);
ctx.samplers[idx] = st = obj->bind(*ctx.q);
}
void
_cl_kernel::sampler_argument::unbind(exec_context &ctx) {
obj->unbind(*ctx.q, st);
}

214
src/gallium/state_trackers/clover/core/kernel.hpp Normal file
View File

@ -0,0 +1,214 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_KERNEL_HPP__
#define __CORE_KERNEL_HPP__
#include <memory>
#include "core/base.hpp"
#include "core/program.hpp"
#include "core/memory.hpp"
#include "core/sampler.hpp"
#include "pipe/p_state.h"
namespace clover {
typedef struct _cl_kernel kernel;
class argument;
}
struct _cl_kernel : public clover::ref_counter {
private:
///
/// Class containing all the state required to execute a compute
/// kernel.
///
struct exec_context {
exec_context(clover::kernel &kern);
~exec_context();
void *bind(clover::command_queue *q);
void unbind();
clover::kernel &kern;
clover::command_queue *q;
std::vector<uint8_t> input;
std::vector<void *> samplers;
std::vector<pipe_sampler_view *> sviews;
std::vector<pipe_surface *> resources;
std::vector<pipe_resource *> g_buffers;
std::vector<size_t> g_handles;
size_t mem_local;
private:
void *st;
pipe_compute_state cs;
};
public:
class argument {
public:
argument(size_t size);
/// \a true if the argument has been set.
bool set() const;
/// Argument size in the input buffer.
size_t size() const;
/// Storage space required for the referenced object.
virtual size_t storage() const;
/// Set this argument to some object.
virtual void set(size_t size, const void *value) = 0;
/// Allocate the necessary resources to bind the specified
/// object to this argument, and update \a ctx accordingly.
virtual void bind(exec_context &ctx) = 0;
/// Free any resources that were allocated in bind().
virtual void unbind(exec_context &ctx) = 0;
protected:
size_t __size;
bool __set;
};
_cl_kernel(clover::program &prog,
const std::string &name,
const std::vector<clover::module::argument> &args);
void launch(clover::command_queue &q,
const std::vector<size_t> &grid_offset,
const std::vector<size_t> &grid_size,
const std::vector<size_t> &block_size);
size_t mem_local() const;
size_t mem_private() const;
size_t max_block_size() const;
const std::string &name() const;
std::vector<size_t> block_size() const;
clover::program &prog;
std::vector<std::unique_ptr<argument>> args;
private:
const clover::module &
module(const clover::command_queue &q) const;
class scalar_argument : public argument {
public:
scalar_argument(size_t size);
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
std::vector<uint8_t> v;
};
class global_argument : public argument {
public:
global_argument(size_t size);
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
clover::buffer *obj;
};
class local_argument : public argument {
public:
local_argument();
virtual size_t storage() const;
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
size_t __storage;
};
class constant_argument : public argument {
public:
constant_argument();
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
clover::buffer *obj;
pipe_surface *st;
};
class image_rd_argument : public argument {
public:
image_rd_argument();
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
clover::image *obj;
pipe_sampler_view *st;
};
class image_wr_argument : public argument {
public:
image_wr_argument();
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
clover::image *obj;
pipe_surface *st;
};
class sampler_argument : public argument {
public:
sampler_argument();
virtual void set(size_t size, const void *value);
virtual void bind(exec_context &ctx);
virtual void unbind(exec_context &ctx);
private:
clover::sampler *obj;
void *st;
};
std::string __name;
exec_context exec;
};
#endif

198
src/gallium/state_trackers/clover/core/memory.cpp Normal file
View File

@ -0,0 +1,198 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/memory.hpp"
#include "core/resource.hpp"
using namespace clover;
_cl_mem::_cl_mem(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr) :
ctx(ctx), __flags(flags),
__size(size), __host_ptr(host_ptr),
__destroy_notify([]{}),
data((char *)host_ptr, (host_ptr ? size : 0)) {
}
_cl_mem::~_cl_mem() {
__destroy_notify();
}
void
_cl_mem::destroy_notify(std::function<void ()> f) {
__destroy_notify = f;
}
cl_mem_flags
_cl_mem::flags() const {
return __flags;
}
size_t
_cl_mem::size() const {
return __size;
}
void *
_cl_mem::host_ptr() const {
return __host_ptr;
}
buffer::buffer(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr) :
memory_obj(ctx, flags, size, host_ptr) {
}
cl_mem_object_type
buffer::type() const {
return CL_MEM_OBJECT_BUFFER;
}
root_buffer::root_buffer(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr) :
buffer(ctx, flags, size, host_ptr) {
}
clover::resource &
root_buffer::resource(cl_command_queue q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = (!resources.empty() ?
new root_resource(q->dev, *this, *resources.begin()->second) :
new root_resource(q->dev, *this, data));
resources.insert(std::make_pair(&q->dev,
std::unique_ptr<root_resource>(r)));
data.clear();
}
return *resources.find(&q->dev)->second;
}
sub_buffer::sub_buffer(clover::root_buffer &parent, cl_mem_flags flags,
size_t offset, size_t size) :
buffer(parent.ctx, flags, size,
(char *)parent.host_ptr() + offset),
parent(parent), __offset(offset) {
}
clover::resource &
sub_buffer::resource(cl_command_queue q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = new sub_resource(parent.resource(q), { offset() });
resources.insert(std::make_pair(&q->dev,
std::unique_ptr<sub_resource>(r)));
}
return *resources.find(&q->dev)->second;
}
size_t
sub_buffer::offset() const {
return __offset;
}
image::image(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch, size_t size,
void *host_ptr) :
memory_obj(ctx, flags, size, host_ptr),
__format(*format), __width(width), __height(height), __depth(depth),
__row_pitch(row_pitch), __slice_pitch(slice_pitch) {
}
clover::resource &
image::resource(cl_command_queue q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = (!resources.empty() ?
new root_resource(q->dev, *this, *resources.begin()->second) :
new root_resource(q->dev, *this, data));
resources.insert(std::make_pair(&q->dev,
std::unique_ptr<root_resource>(r)));
data.clear();
}
return *resources.find(&q->dev)->second;
}
cl_image_format
image::format() const {
return __format;
}
size_t
image::width() const {
return __width;
}
size_t
image::height() const {
return __height;
}
size_t
image::depth() const {
return __depth;
}
size_t
image::row_pitch() const {
return __row_pitch;
}
size_t
image::slice_pitch() const {
return __slice_pitch;
}
image2d::image2d(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format, size_t width,
size_t height, size_t row_pitch,
void *host_ptr) :
image(ctx, flags, format, width, height, 0,
row_pitch, 0, height * row_pitch, host_ptr) {
}
cl_mem_object_type
image2d::type() const {
return CL_MEM_OBJECT_IMAGE2D;
}
image3d::image3d(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch,
void *host_ptr) :
image(ctx, flags, format, width, height, depth,
row_pitch, slice_pitch, depth * slice_pitch,
host_ptr) {
}
cl_mem_object_type
image3d::type() const {
return CL_MEM_OBJECT_IMAGE3D;
}

157
src/gallium/state_trackers/clover/core/memory.hpp Normal file
View File

@ -0,0 +1,157 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_MEMORY_HPP__
#define __CORE_MEMORY_HPP__
#include <functional>
#include <map>
#include <memory>
#include "core/base.hpp"
#include "core/queue.hpp"
namespace clover {
typedef struct _cl_mem memory_obj;
class resource;
class sub_resource;
}
struct _cl_mem : public clover::ref_counter {
protected:
_cl_mem(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr);
_cl_mem(const _cl_mem &obj) = delete;
public:
virtual ~_cl_mem();
virtual cl_mem_object_type type() const = 0;
virtual clover::resource &resource(cl_command_queue q) = 0;
void destroy_notify(std::function<void ()> f);
cl_mem_flags flags() const;
size_t size() const;
void *host_ptr() const;
clover::context &ctx;
private:
cl_mem_flags __flags;
size_t __size;
void *__host_ptr;
std::function<void ()> __destroy_notify;
protected:
std::string data;
};
namespace clover {
struct buffer : public memory_obj {
protected:
buffer(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr);
public:
virtual cl_mem_object_type type() const;
};
struct root_buffer : public buffer {
public:
root_buffer(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr);
virtual clover::resource &resource(cl_command_queue q);
private:
std::map<clover::device *,
std::unique_ptr<clover::root_resource>> resources;
};
struct sub_buffer : public buffer {
public:
sub_buffer(clover::root_buffer &parent, cl_mem_flags flags,
size_t offset, size_t size);
virtual clover::resource &resource(cl_command_queue q);
size_t offset() const;
clover::root_buffer &parent;
private:
size_t __offset;
std::map<clover::device *,
std::unique_ptr<clover::sub_resource>> resources;
};
struct image : public memory_obj {
protected:
image(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch, size_t size,
void *host_ptr);
public:
virtual clover::resource &resource(cl_command_queue q);
cl_image_format format() const;
size_t width() const;
size_t height() const;
size_t depth() const;
size_t row_pitch() const;
size_t slice_pitch() const;
private:
cl_image_format __format;
size_t __width;
size_t __height;
size_t __depth;
size_t __row_pitch;
size_t __slice_pitch;
std::map<clover::device *,
std::unique_ptr<clover::root_resource>> resources;
};
struct image2d : public image {
public:
image2d(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format, size_t width,
size_t height, size_t row_pitch,
void *host_ptr);
virtual cl_mem_object_type type() const;
};
struct image3d : public image {
public:
image3d(clover::context &ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch,
void *host_ptr);
virtual cl_mem_object_type type() const;
};
}
#endif

172
src/gallium/state_trackers/clover/core/module.cpp Normal file
View File

@ -0,0 +1,172 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <type_traits>
#include <algorithm>
#include "core/module.hpp"
using namespace clover;
namespace {
template<typename T, typename = void>
struct __serializer;
/// Serialize the specified object.
template<typename T>
void
__proc(compat::ostream &os, const T &x) {
__serializer<T>::proc(os, x);
}
/// Deserialize the specified object.
template<typename T>
void
__proc(compat::istream &is, T &x) {
__serializer<T>::proc(is, x);
}
template<typename T>
T
__proc(compat::istream &is) {
T x;
__serializer<T>::proc(is, x);
return x;
}
/// (De)serialize a scalar value.
template<typename T>
struct __serializer<T, typename std::enable_if<
std::is_scalar<T>::value>::type> {
static void
proc(compat::ostream &os, const T &x) {
os.write(reinterpret_cast<const char *>(&x), sizeof(x));
}
static void
proc(compat::istream &is, T &x) {
is.read(reinterpret_cast<char *>(&x), sizeof(x));
}
};
/// (De)serialize a vector.
template<typename T>
struct __serializer<compat::vector<T>> {
static void
proc(compat::ostream &os, const compat::vector<T> &v) {
__proc<uint32_t>(os, v.size());
for (size_t i = 0; i < v.size(); i++)
__proc<T>(os, v[i]);
}
static void
proc(compat::istream &is, compat::vector<T> &v) {
v.reserve(__proc<uint32_t>(is));
for (size_t i = 0; i < v.size(); i++)
new(&v[i]) T(__proc<T>(is));
}
};
/// (De)serialize a module::section.
template<>
struct __serializer<module::section> {
template<typename S, typename QT>
static void
proc(S &s, QT &x) {
__proc(s, x.type);
__proc(s, x.size);
__proc(s, x.data);
}
};
/// (De)serialize a module::argument.
template<>
struct __serializer<module::argument> {
template<typename S, typename QT>
static void
proc(S &s, QT &x) {
__proc(s, x.type);
__proc(s, x.size);
}
};
/// (De)serialize a module::symbol.
template<>
struct __serializer<module::symbol> {
template<typename S, typename QT>
static void
proc(S &s, QT &x) {
__proc(s, x.section);
__proc(s, x.offset);
__proc(s, x.args);
}
};
/// (De)serialize a module.
template<>
struct __serializer<module> {
template<typename S, typename QT>
static void
proc(S &s, QT &x) {
__proc(s, x.syms);
__proc(s, x.secs);
}
};
};
namespace clover {
void
module::serialize(compat::ostream &os) const {
__proc(os, *this);
}
module
module::deserialize(compat::istream &is) {
return __proc<module>(is);
}
const module::symbol &
module::sym(compat::string name) const {
auto it = std::find_if(syms.begin(), syms.end(), [&](const symbol &x) {
return compat::string(x.name) == name;
});
if (it == syms.end())
throw noent_error();
return *it;
}
const module::section &
module::sec(typename section::type type) const {
auto it = std::find_if(secs.begin(), secs.end(), [&](const section &x) {
return x.type == type;
});
if (it == secs.end())
throw noent_error();
return *it;
}
}

93
src/gallium/state_trackers/clover/core/module.hpp Normal file
View File

@ -0,0 +1,93 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_MODULE_HPP__
#define __CORE_MODULE_HPP__
#include "core/compat.hpp"
namespace clover {
struct module {
class noent_error {
public:
virtual ~noent_error() {}
};
typedef uint32_t resource_id;
typedef uint32_t size_t;
struct section {
enum type {
text,
data_constant,
data_global,
data_local,
data_private
};
resource_id id;
type type;
size_t size;
clover::compat::vector<char> data;
};
struct argument {
enum type {
scalar,
constant,
global,
local,
image2d_rd,
image2d_wr,
image3d_rd,
image3d_wr,
sampler
};
type type;
size_t size;
};
struct symbol {
clover::compat::vector<char> name;
resource_id section;
size_t offset;
clover::compat::vector<argument> args;
};
void serialize(compat::ostream &os) const;
static module deserialize(compat::istream &is);
/// Look up a symbol by name. Throws module::noent_error if not
/// found.
const symbol &sym(compat::string name) const;
/// Look up a section by type. Throws module::noent_error if not
/// found.
const section &sec(typename section::type type) const;
clover::compat::vector<symbol> syms;
clover::compat::vector<section> secs;
};
}
#endif

85
src/gallium/state_trackers/clover/core/program.cpp Normal file
View File

@ -0,0 +1,85 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/program.hpp"
#include "core/compiler.hpp"
using namespace clover;
_cl_program::_cl_program(clover::context &ctx,
const std::string &source) :
ctx(ctx), __source(source) {
}
_cl_program::_cl_program(clover::context &ctx,
const std::vector<clover::device *> &devs,
const std::vector<clover::module> &binaries) :
ctx(ctx) {
for_each([&](clover::device *dev, const clover::module &bin) {
__binaries.insert({ dev, bin });
},
devs.begin(), devs.end(), binaries.begin());
}
void
_cl_program::build(const std::vector<clover::device *> &devs) {
__binaries.clear();
__logs.clear();
for (auto dev : devs) {
try {
auto module = (dev->ir_target() == "tgsi" ?
compile_program_tgsi(__source, dev->ir_target()) :
compile_program_llvm(__source, dev->ir_target()));
__binaries.insert({ dev, module });
} catch (build_error &e) {
__logs.insert({ dev, e.what() });
throw error(CL_BUILD_PROGRAM_FAILURE);
}
}
}
const std::string &
_cl_program::source() const {
return __source;
}
const std::map<clover::device *, clover::module> &
_cl_program::binaries() const {
return __binaries;
}
cl_build_status
_cl_program::build_status(clover::device *dev) const {
return __binaries.count(dev) ? CL_BUILD_SUCCESS : CL_BUILD_NONE;
}
std::string
_cl_program::build_opts(clover::device *dev) const {
return {};
}
std::string
_cl_program::build_log(clover::device *dev) const {
return __logs.count(dev) ? __logs.find(dev)->second : "";
}

61
src/gallium/state_trackers/clover/core/program.hpp Normal file
View File

@ -0,0 +1,61 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_PROGRAM_HPP__
#define __CORE_PROGRAM_HPP__
#include <map>
#include "core/base.hpp"
#include "core/context.hpp"
#include "core/module.hpp"
namespace clover {
typedef struct _cl_program program;
}
struct _cl_program : public clover::ref_counter {
public:
_cl_program(clover::context &ctx,
const std::string &source);
_cl_program(clover::context &ctx,
const std::vector<clover::device *> &devs,
const std::vector<clover::module> &binaries);
void build(const std::vector<clover::device *> &devs);
const std::string &source() const;
const std::map<clover::device *, clover::module> &binaries() const;
cl_build_status build_status(clover::device *dev) const;
std::string build_opts(clover::device *dev) const;
std::string build_log(clover::device *dev) const;
clover::context &ctx;
private:
std::map<clover::device *, clover::module> __binaries;
std::map<clover::device *, std::string> __logs;
std::string __source;
};
#endif

69
src/gallium/state_trackers/clover/core/queue.cpp Normal file
View File

@ -0,0 +1,69 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <algorithm>
#include "core/queue.hpp"
#include "core/event.hpp"
#include "pipe/p_screen.h"
#include "pipe/p_context.h"
using namespace clover;
_cl_command_queue::_cl_command_queue(context &ctx, device &dev,
cl_command_queue_properties props) :
ctx(ctx), dev(dev), __props(props) {
pipe = dev.pipe->context_create(dev.pipe, NULL);
if (!pipe)
throw error(CL_INVALID_DEVICE);
}
_cl_command_queue::~_cl_command_queue() {
pipe->destroy(pipe);
}
void
_cl_command_queue::flush() {
pipe_screen *screen = dev.pipe;
pipe_fence_handle *fence = NULL;
if (!queued_events.empty()) {
// Find out which events have already been signalled.
auto first = queued_events.begin();
auto last = std::find_if(queued_events.begin(), queued_events.end(),
[](event_ptr &ev) { return !ev->signalled(); });
// Flush and fence them.
pipe->flush(pipe, &fence);
std::for_each(first, last, [&](event_ptr &ev) { ev->fence(fence); });
screen->fence_reference(screen, &fence, NULL);
queued_events.erase(first, last);
}
}
void
_cl_command_queue::sequence(clover::hard_event *ev) {
if (!queued_events.empty())
queued_events.back()->chain(ev);
queued_events.push_back(ev);
}

71
src/gallium/state_trackers/clover/core/queue.hpp Normal file
View File

@ -0,0 +1,71 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_QUEUE_HPP__
#define __CORE_QUEUE_HPP__
#include "core/base.hpp"
#include "core/context.hpp"
#include "pipe/p_context.h"
namespace clover {
typedef struct _cl_command_queue command_queue;
class resource;
class mapping;
class hard_event;
}
struct _cl_command_queue : public clover::ref_counter {
public:
_cl_command_queue(clover::context &ctx, clover::device &dev,
cl_command_queue_properties props);
_cl_command_queue(const _cl_command_queue &q) = delete;
~_cl_command_queue();
void flush();
cl_command_queue_properties props() const {
return __props;
}
clover::context &ctx;
clover::device &dev;
friend class clover::resource;
friend class clover::mapping;
friend class clover::hard_event;
friend struct _cl_sampler;
friend struct _cl_kernel;
private:
/// Serialize a hardware event with respect to the previous ones,
/// and push it to the pending list.
void sequence(clover::hard_event *ev);
cl_command_queue_properties __props;
pipe_context *pipe;
typedef clover::ref_ptr<clover::hard_event> event_ptr;
std::vector<event_ptr> queued_events;
};
#endif

192
src/gallium/state_trackers/clover/core/resource.cpp Normal file
View File

@ -0,0 +1,192 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/resource.hpp"
#include "pipe/p_screen.h"
#include "util/u_sampler.h"
using namespace clover;
namespace {
class box {
public:
box(const resource::point &origin, const resource::point &size) :
pipe({ (unsigned)origin[0], (unsigned)origin[1],
(unsigned)origin[2], (unsigned)size[0],
(unsigned)size[1], (unsigned)size[2] }) {
}
operator const pipe_box *() {
return &pipe;
}
protected:
pipe_box pipe;
};
}
resource::resource(clover::device &dev, clover::memory_obj &obj) :
dev(dev), obj(obj), pipe(NULL), offset{0} {
}
resource::~resource() {
}
void
resource::copy(command_queue &q, const point &origin, const point &region,
resource &src_res, const point &src_origin) {
point p = offset + origin;
q.pipe->resource_copy_region(q.pipe, pipe, 0, p[0], p[1], p[2],
src_res.pipe, 0,
box(src_res.offset + src_origin, region));
}
void *
resource::add_map(command_queue &q, cl_map_flags flags, bool blocking,
const point &origin, const point &region) {
maps.emplace_back(q, *this, flags, blocking, origin, region);
return maps.back();
}
void
resource::del_map(void *p) {
auto it = std::find(maps.begin(), maps.end(), p);
if (it != maps.end())
maps.erase(it);
}
unsigned
resource::map_count() const {
return maps.size();
}
pipe_sampler_view *
resource::bind_sampler_view(clover::command_queue &q) {
pipe_sampler_view info;
u_sampler_view_default_template(&info, pipe, pipe->format);
return q.pipe->create_sampler_view(q.pipe, pipe, &info);
}
void
resource::unbind_sampler_view(clover::command_queue &q,
pipe_sampler_view *st) {
q.pipe->sampler_view_destroy(q.pipe, st);
}
pipe_surface *
resource::bind_surface(clover::command_queue &q, bool rw) {
pipe_surface info {};
info.format = pipe->format;
info.usage = pipe->bind;
info.writable = rw;
if (pipe->target == PIPE_BUFFER)
info.u.buf.last_element = pipe->width0 - 1;
return q.pipe->create_surface(q.pipe, pipe, &info);
}
void
resource::unbind_surface(clover::command_queue &q, pipe_surface *st) {
q.pipe->surface_destroy(q.pipe, st);
}
root_resource::root_resource(clover::device &dev, clover::memory_obj &obj,
std::string data) :
resource(dev, obj) {
pipe_resource info {};
if (image *img = dynamic_cast<image *>(&obj)) {
info.format = translate_format(img->format());
info.width0 = img->width();
info.height0 = img->height();
info.depth0 = img->depth();
} else {
info.width0 = obj.size();
}
info.target = translate_target(obj.type());
info.bind = (PIPE_BIND_SAMPLER_VIEW |
PIPE_BIND_COMPUTE_RESOURCE |
PIPE_BIND_GLOBAL |
PIPE_BIND_TRANSFER_READ |
PIPE_BIND_TRANSFER_WRITE);
pipe = dev.pipe->resource_create(dev.pipe, &info);
if (!pipe)
throw error(CL_OUT_OF_RESOURCES);
assert(data.empty()); // XXX -- initialize it with the supplied data
}
root_resource::root_resource(clover::device &dev, clover::memory_obj &obj,
clover::root_resource &r) :
resource(dev, obj) {
assert(0); // XXX -- resource shared among dev and r.dev
}
root_resource::~root_resource() {
dev.pipe->resource_destroy(dev.pipe, pipe);
}
sub_resource::sub_resource(clover::resource &r, point offset) :
resource(r.dev, r.obj) {
pipe = r.pipe;
offset = r.offset + offset;
}
mapping::mapping(command_queue &q, resource &r,
cl_map_flags flags, bool blocking,
const resource::point &origin,
const resource::point &region) :
pctx(q.pipe) {
unsigned usage = ((flags & CL_MAP_WRITE ? PIPE_TRANSFER_WRITE : 0 ) |
(flags & CL_MAP_READ ? PIPE_TRANSFER_READ : 0 ) |
(blocking ? PIPE_TRANSFER_UNSYNCHRONIZED : 0));
pxfer = pctx->get_transfer(pctx, r.pipe, 0, usage,
box(origin + r.offset, region));
if (!pxfer)
throw error(CL_OUT_OF_RESOURCES);
p = pctx->transfer_map(pctx, pxfer);
if (!p) {
pctx->transfer_destroy(pctx, pxfer);
throw error(CL_OUT_OF_RESOURCES);
}
}
mapping::mapping(mapping &&m) :
pctx(m.pctx), pxfer(m.pxfer), p(m.p) {
m.p = NULL;
m.pxfer = NULL;
}
mapping::~mapping() {
if (pxfer) {
pctx->transfer_unmap(pctx, pxfer);
pctx->transfer_destroy(pctx, pxfer);
}
}

129
src/gallium/state_trackers/clover/core/resource.hpp Normal file
View File

@ -0,0 +1,129 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_RESOURCE_HPP__
#define __CORE_RESOURCE_HPP__
#include <list>
#include "core/base.hpp"
#include "core/memory.hpp"
#include "core/geometry.hpp"
#include "pipe/p_state.h"
namespace clover {
class mapping;
///
/// Class that represents a device-specific instance of some memory
/// object.
///
class resource {
public:
typedef clover::point<size_t, 3> point;
resource(const resource &r) = delete;
virtual ~resource();
void copy(command_queue &q, const point &origin, const point &region,
resource &src_resource, const point &src_origin);
void *add_map(command_queue &q, cl_map_flags flags, bool blocking,
const point &origin, const point &region);
void del_map(void *p);
unsigned map_count() const;
clover::device &dev;
clover::memory_obj &obj;
friend class sub_resource;
friend class mapping;
friend struct ::_cl_kernel;
protected:
resource(clover::device &dev, clover::memory_obj &obj);
pipe_sampler_view *bind_sampler_view(clover::command_queue &q);
void unbind_sampler_view(clover::command_queue &q,
pipe_sampler_view *st);
pipe_surface *bind_surface(clover::command_queue &q, bool rw);
void unbind_surface(clover::command_queue &q, pipe_surface *st);
pipe_resource *pipe;
point offset;
private:
std::list<mapping> maps;
};
///
/// Resource associated with its own top-level data storage
/// allocated in some device.
///
class root_resource : public resource {
public:
root_resource(clover::device &dev, clover::memory_obj &obj,
std::string data);
root_resource(clover::device &dev, clover::memory_obj &obj,
root_resource &r);
virtual ~root_resource();
};
///
/// Resource that reuses a portion of some other resource as data
/// storage.
///
class sub_resource : public resource {
public:
sub_resource(clover::resource &r, point offset);
};
///
/// Class that represents a mapping of some resource into the CPU
/// memory space.
///
class mapping {
public:
mapping(command_queue &q, resource &r, cl_map_flags flags,
bool blocking, const resource::point &origin,
const resource::point &region);
mapping(const mapping &m) = delete;
mapping(mapping &&m);
~mapping();
operator void *() {
return p;
}
operator char *() {
return (char *)p;
}
private:
pipe_context *pctx;
pipe_transfer *pxfer;
void *p;
};
}
#endif

73
src/gallium/state_trackers/clover/core/sampler.cpp Normal file
View File

@ -0,0 +1,73 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/sampler.hpp"
#include "pipe/p_state.h"
using namespace clover;
_cl_sampler::_cl_sampler(clover::context &ctx, bool norm_mode,
cl_addressing_mode addr_mode,
cl_filter_mode filter_mode) :
ctx(ctx), __norm_mode(norm_mode),
__addr_mode(addr_mode), __filter_mode(filter_mode) {
}
bool
_cl_sampler::norm_mode() {
return __norm_mode;
}
cl_addressing_mode
_cl_sampler::addr_mode() {
return __addr_mode;
}
cl_filter_mode
_cl_sampler::filter_mode() {
return __filter_mode;
}
void *
_cl_sampler::bind(clover::command_queue &q) {
struct pipe_sampler_state info {};
info.normalized_coords = norm_mode();
info.wrap_s = info.wrap_t = info.wrap_r =
(addr_mode() == CL_ADDRESS_CLAMP_TO_EDGE ? PIPE_TEX_WRAP_CLAMP_TO_EDGE :
addr_mode() == CL_ADDRESS_CLAMP ? PIPE_TEX_WRAP_CLAMP_TO_BORDER :
addr_mode() == CL_ADDRESS_REPEAT ? PIPE_TEX_WRAP_REPEAT :
addr_mode() == CL_ADDRESS_MIRRORED_REPEAT ? PIPE_TEX_WRAP_MIRROR_REPEAT :
PIPE_TEX_WRAP_CLAMP_TO_EDGE);
info.min_img_filter = info.mag_img_filter =
(filter_mode() == CL_FILTER_LINEAR ? PIPE_TEX_FILTER_LINEAR :
PIPE_TEX_FILTER_NEAREST);
return q.pipe->create_sampler_state(q.pipe, &info);
}
void
_cl_sampler::unbind(clover::command_queue &q, void *st) {
q.pipe->delete_sampler_state(q.pipe, st);
}

55
src/gallium/state_trackers/clover/core/sampler.hpp Normal file
View File

@ -0,0 +1,55 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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 __CORE_SAMPLER_HPP__
#define __CORE_SAMPLER_HPP__
#include "core/base.hpp"
#include "core/queue.hpp"
namespace clover {
typedef struct _cl_sampler sampler;
}
struct _cl_sampler : public clover::ref_counter {
public:
_cl_sampler(clover::context &ctx, bool norm_mode,
cl_addressing_mode addr_mode, cl_filter_mode filter_mode);
bool norm_mode();
cl_addressing_mode addr_mode();
cl_filter_mode filter_mode();
clover::context &ctx;
friend class _cl_kernel;
private:
void *bind(clover::command_queue &q);
void unbind(clover::command_queue &q, void *st);
bool __norm_mode;
cl_addressing_mode __addr_mode;
cl_filter_mode __filter_mode;
};
#endif

94
src/gallium/state_trackers/clover/llvm/invocation.cpp Normal file
View File

@ -0,0 +1,94 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include "core/compiler.hpp"
#if 0
#include <clang/Frontend/CompilerInstance.h>
#include <clang/Frontend/TextDiagnosticPrinter.h>
#include <clang/CodeGen/CodeGenAction.h>
#include <llvm/LLVMContext.h>
#include <llvm/Support/TargetSelect.h>
#include <llvm/Support/MemoryBuffer.h>
#include <iostream>
#include <iomanip>
#include <fstream>
#include <cstdio>
#endif
using namespace clover;
#if 0
namespace {
void
build_binary(const std::string &source, const std::string &target,
const std::string &name) {
clang::CompilerInstance c;
clang::EmitObjAction act(&llvm::getGlobalContext());
std::string log;
llvm::raw_string_ostream s_log(log);
LLVMInitializeTGSITarget();
LLVMInitializeTGSITargetInfo();
LLVMInitializeTGSITargetMC();
LLVMInitializeTGSIAsmPrinter();
c.getFrontendOpts().Inputs.push_back(
std::make_pair(clang::IK_OpenCL, name));
c.getHeaderSearchOpts().UseBuiltinIncludes = false;
c.getHeaderSearchOpts().UseStandardIncludes = false;
c.getLangOpts().NoBuiltin = true;
c.getTargetOpts().Triple = target;
c.getInvocation().setLangDefaults(clang::IK_OpenCL);
c.createDiagnostics(0, NULL, new clang::TextDiagnosticPrinter(
s_log, c.getDiagnosticOpts()));
c.getPreprocessorOpts().addRemappedFile(
name, llvm::MemoryBuffer::getMemBuffer(source));
if (!c.ExecuteAction(act))
throw build_error(log);
}
module
load_binary(const char *name) {
std::ifstream fs((name));
std::vector<unsigned char> str((std::istreambuf_iterator<char>(fs)),
(std::istreambuf_iterator<char>()));
compat::istream cs(str);
return module::deserialize(cs);
}
}
#endif
module
clover::compile_program_llvm(const compat::string &source,
const compat::string &target) {
#if 0
build_binary(source, target, "cl_input");
module m = load_binary("cl_input.o");
std::remove("cl_input.o");
return m;
#endif
return module();
}

100
src/gallium/state_trackers/clover/tgsi/compiler.cpp Normal file
View File

@ -0,0 +1,100 @@
//
// Copyright 2012 Francisco Jerez
//
// 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 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.
//
#include <sstream>
#include "core/compiler.hpp"
#include "tgsi/tgsi_parse.h"
#include "tgsi/tgsi_text.h"
#include "util/u_memory.h"
using namespace clover;
namespace {
void
read_header(const std::string &header, module &m) {
std::istringstream ls(header);
std::string line;
while (getline(ls, line)) {
std::istringstream ts(line);
std::string name, tok;
module::size_t offset;
compat::vector<module::argument> args;
if (!(ts >> name))
continue;
if (!(ts >> offset))
throw build_error("invalid kernel start address");
while (ts >> tok) {
if (tok == "scalar")
args.push_back({ module::argument::scalar, 4 });
else if (tok == "global")
args.push_back({ module::argument::global, 4 });
else if (tok == "local")
args.push_back({ module::argument::local, 4 });
else if (tok == "constant")
args.push_back({ module::argument::constant, 4 });
else if (tok == "image2d_rd")
args.push_back({ module::argument::image2d_rd, 4 });
else if (tok == "image2d_wr")
args.push_back({ module::argument::image2d_wr, 4 });
else if (tok == "image3d_rd")
args.push_back({ module::argument::image3d_rd, 4 });
else if (tok == "image3d_wr")
args.push_back({ module::argument::image3d_wr, 4 });
else if (tok == "sampler")
args.push_back({ module::argument::sampler, 0 });
else
throw build_error("invalid kernel argument");
}
m.syms.push_back({ name, 0, offset, args });
}
}
void
read_body(const char *source, module &m) {
tgsi_token prog[1024];
if (!tgsi_text_translate(source, prog, Elements(prog)))
throw build_error("translate failed");
unsigned sz = tgsi_num_tokens(prog) * sizeof(tgsi_token);
m.secs.push_back({ 0, module::section::text, sz, { (char *)prog, sz } });
}
}
module
clover::compile_program_tgsi(const compat::string &source,
const compat::string &target) {
const char *body = source.find("COMP\n");
module m;
read_header({ source.begin(), body }, m);
read_body(body, m);
return m;
}

36
src/gallium/targets/opencl/Makefile.am Normal file
View File

@ -0,0 +1,36 @@
AUTOMAKE_OPTIONS = subdir-objects
lib_LTLIBRARIES = libOpenCL.la
libOpenCL_la_LDFLAGS = \
-version-number 1:0
libOpenCL_la_LIBADD = \
$(top_builddir)/src/gallium/state_trackers/clover/libclover.la \
$(top_builddir)/src/gallium/auxiliary/libgallium.a \
$(GALLIUM_PIPE_LOADER_LIBS) $(LIBUDEV_LIBS) \
-ldl
libOpenCL_la_SOURCES =
# Force usage of a C++ linker
nodist_EXTRA_libOpenCL_la_SOURCES = dummy.cpp
PIPE_SRC_DIR = $(top_srcdir)/src/gallium/targets/pipe-loader
# Provide compatibility with scripts for the old Mesa build system for
# a while by putting a link to the driver into /lib of the build tree.
all-local: libOpenCL.la
@$(MAKE) -C $(PIPE_SRC_DIR)
$(MKDIR_P) $(top_builddir)/$(LIB_DIR)
ln -f .libs/libOpenCL.so* $(top_builddir)/$(LIB_DIR)/
install-exec-local:
@$(MAKE) -C $(PIPE_SRC_DIR) PIPE_INSTALL_DIR=$(OPENCL_LIB_INSTALL_DIR) install
clean-local:
@$(MAKE) -C $(PIPE_SRC_DIR) clean
# FIXME: Remove when the rest of Gallium is converted to automake.
TOP=$(top_builddir)
default: all