[Pkg-opencl-devel] [beignet] 11/66: Imported Upstream version 0.1+git20130418+0546d2e
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:02 UTC 2014
This is an automated email from the git hooks/post-receive script.
anbe pushed a commit to branch master
in repository beignet.
commit 2679e319c3de966203b18327d789945c51959998
Author: Simon Richter <sjr at debian.org>
Date: Thu Apr 18 11:51:37 2013 +0200
Imported Upstream version 0.1+git20130418+0546d2e
---
CMake/FindLLVM.cmake | 6 +-
CMake/FindOCLIcd.cmake | 23 ++++
CMakeLists.txt | 9 +-
backend/CMakeLists.txt | 4 +-
backend/src/backend/gen_context.cpp | 4 +-
backend/src/backend/gen_insn_selection.cpp | 9 +-
backend/src/backend/gen_insn_selection.hpp | 2 +-
backend/src/backend/gen_reg_allocation.hpp | 2 +-
backend/src/backend/program.cpp | 4 +-
backend/src/ir/instruction.cpp | 12 +-
backend/src/llvm/llvm_to_gen.cpp | 9 +-
backend/src/sys/alloc.cpp | 2 +-
backend/src/sys/mutex.hpp | 10 +-
src/CMakeLists.txt | 13 ++-
src/cl_api.c | 11 +-
src/cl_command_queue.c | 2 +
src/cl_command_queue.h | 1 +
src/cl_context.c | 2 +
src/cl_context.h | 2 +
src/cl_device_id.c | 4 +
src/cl_device_id.h | 1 +
src/cl_event.h | 1 +
src/cl_extensions.c | 17 ++-
src/cl_extensions.h | 60 +++++-----
src/cl_kernel.c | 2 +
src/cl_kernel.h | 1 +
src/cl_khr_icd.c | 175 +++++++++++++++++++++++++++++
src/{cl_event.h => cl_khr_icd.h} | 21 ++--
src/cl_mem.c | 2 +
src/cl_mem.h | 1 +
src/cl_mem_gl.c | 6 +-
src/cl_platform_id.c | 7 +-
src/cl_platform_id.h | 6 +-
src/cl_program.c | 2 +
src/cl_program.h | 1 +
src/cl_sampler.c | 2 +
src/cl_sampler.h | 1 +
utests/utest_helper.hpp | 8 +-
38 files changed, 354 insertions(+), 91 deletions(-)
diff --git a/CMake/FindLLVM.cmake b/CMake/FindLLVM.cmake
index 6137a9c..c06b8a4 100644
--- a/CMake/FindLLVM.cmake
+++ b/CMake/FindLLVM.cmake
@@ -7,9 +7,9 @@
# LLVM_MODULE_LIBS - list of llvm libs for working with modules.
# LLVM_FOUND - True if llvm found.
if (LLVM_INSTALL_DIR)
- find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config llvm-config-32 llvm-config-3.0 DOC "llvm-config executable" PATHS ${LLVM_INSTALL_DIR} NO_DEFAULT_PATH)
-else (LLVM_INSTALL_DIR)
- find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config llvm-config-32 llvm-config-3.0 DOC "llvm-config executable")
+ find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config-32 llvm-config-3.2 llvm-config DOC "llvm-config executable" PATHS ${LLVM_INSTALL_DIR} NO_DEFAULT_PATH)
+else (LLVM_INSTALL_DIR)
+ find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config-32 llvm-config-3.2 llvm-config DOC "llvm-config executable")
endif (LLVM_INSTALL_DIR)
if (LLVM_CONFIG_EXECUTABLE)
diff --git a/CMake/FindOCLIcd.cmake b/CMake/FindOCLIcd.cmake
new file mode 100644
index 0000000..076f00e
--- /dev/null
+++ b/CMake/FindOCLIcd.cmake
@@ -0,0 +1,23 @@
+#
+# Try to find ocl_icd library and include path.
+# Once done this will define
+#
+# OCLIcd_FOUND
+# OCLIcd_INCLUDE_PATH
+#
+
+FIND_PATH(OCLIcd_INCLUDE_PATH ocl_icd.h
+ ~/include/
+ /usr/include/
+ /usr/local/include/
+ /sw/include/
+ /opt/local/include/
+ DOC "The directory where ocl_icd.h resides")
+
+IF(OCLIcd_INCLUDE_PATH)
+ SET(OCLIcd_FOUND 1 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
+ELSE(OCLIcd_INCLUDE_PATH)
+ SET(OCLIcd_FOUND 0 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
+ENDIF(OCLIcd_INCLUDE_PATH)
+
+MARK_AS_ADVANCED(OCLIcd_FOUND)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 7f37be9..4b402ee 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -55,7 +55,7 @@ ELSE (USE_FULSIM)
ADD_DEFINITIONS(-DUSE_FULSIM=0)
ENDIF (USE_FULSIM)
-SET(CMAKE_CXX_FLAGS "-Wall -Wno-invalid-offsetof -mfpmath=sse --no-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
+SET(CMAKE_CXX_FLAGS "-Wall -Wno-invalid-offsetof -mfpmath=sse -fno-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
SET(CMAKE_C_FLAGS "-Wall -mfpmath=sse -msse2 -Wcast-align -msse2 -msse3 -mssse3 -msse4.1")
# Front end stuff we need
@@ -130,6 +130,13 @@ ELSE(EGL_FOUND)
MESSAGE(STATUS "Looking for EGL - not found")
ENDIF(EGL_FOUND)
+Find_Package(OCLIcd)
+IF(OCLIcd_FOUND)
+ MESSAGE(STATUS "Looking for OCL ICD header file - found")
+ELSE(OCLIcd_FOUND)
+ MESSAGE(STATUS "Looking for OCL ICD header file - not found")
+ENDIF(OCLIcd_FOUND)
+
ADD_SUBDIRECTORY(include)
ADD_SUBDIRECTORY(backend)
ADD_SUBDIRECTORY(src)
diff --git a/backend/CMakeLists.txt b/backend/CMakeLists.txt
index 89f3c4a..5498ff0 100644
--- a/backend/CMakeLists.txt
+++ b/backend/CMakeLists.txt
@@ -43,7 +43,7 @@ endif (GBE_DEBUG_MEMORY)
set (CMAKE_C_CXX_FLAGS "-fvisibility=hidden")
if (COMPILER STREQUAL "GCC")
- set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -funroll-loops -Wstrict-aliasing=2 -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -ffast-math -fPIC -Wall")
+ set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -funroll-loops -Wstrict-aliasing=2 -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -fPIC -Wall")
set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} ${LLVM_CFLAGS}")
set (CMAKE_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -Wno-invalid-offsetof -fno-rtti -std=c++0x")
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
@@ -70,7 +70,7 @@ elseif (COMPILER STREQUAL "CLANG")
set (CMAKE_C_FLAGS_MINSIZEREL "-Os -DNDEBUG -DGBE_DEBUG=0")
set (CMAKE_C_FLAGS_RELEASE "-O2 -DNDEBUG -DGBE_DEBUG=0")
set (CMAKE_CXX_COMPILER "clang++")
- set (CMAKE_CXX_FLAGS "-fstrict-aliasing -msse2 -ffast-math -fPIC -Wall -Wno-format-security -Wno-invalid-offsetof -std=c++0x")
+ set (CMAKE_CXX_FLAGS "-fstrict-aliasing -msse2 -fPIC -Wall -Wno-format-security -Wno-invalid-offsetof -std=c++0x")
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${VISIBILITY_FLAG}")
diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
index b4c9a65..b3d385b 100644
--- a/backend/src/backend/gen_context.cpp
+++ b/backend/src/backend/gen_context.cpp
@@ -326,9 +326,9 @@ namespace gbe
for( uint32_t quarter = 0; quarter < quarterNum; quarter++)
{
#define QUARTER_MOV0(dst_nr, src) p->MOV(GenRegister::ud8grf(dst_nr, 0), \
- GenRegister::retype(GenRegister::ud8grf(src.nr + quarter, 0), src.type))
+ GenRegister::retype(GenRegister::QnPhysical(src, quarter), src.type))
#define QUARTER_MOV1(dst_nr, src) p->MOV(GenRegister::retype(GenRegister::ud8grf(dst_nr, 0), src.type), \
- GenRegister::retype(GenRegister::ud8grf(src.nr + quarter, 0), src.type))
+ GenRegister::retype(GenRegister::QnPhysical(src,quarter), src.type))
QUARTER_MOV0(nr + 1, ucoord);
QUARTER_MOV0(nr + 2, vcoord);
QUARTER_MOV1(nr + 5, R);
diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
index ecaaeeb..e0e8920 100644
--- a/backend/src/backend/gen_insn_selection.cpp
+++ b/backend/src/backend/gen_insn_selection.cpp
@@ -1621,10 +1621,10 @@ namespace gbe
{
using namespace ir;
const uint32_t valueNum = insn.getValueNum();
- GenRegister dst[valueNum];
+ vector<GenRegister> dst(valueNum);
for (uint32_t dstID = 0; dstID < valueNum; ++dstID)
dst[dstID] = GenRegister::retype(sel.selReg(insn.getValue(dstID)), GEN_TYPE_F);
- sel.UNTYPED_READ(addr, dst, valueNum, bti);
+ sel.UNTYPED_READ(addr, dst.data(), valueNum, bti);
}
void emitByteGather(Selection::Opaque &sel,
@@ -1683,12 +1683,13 @@ namespace gbe
using namespace ir;
const uint32_t valueNum = insn.getValueNum();
const uint32_t addrID = ir::StoreInstruction::addressIndex;
- GenRegister addr, value[valueNum];
+ GenRegister addr;
+ vector<GenRegister> value(valueNum);
addr = GenRegister::retype(sel.selReg(insn.getSrc(addrID)), GEN_TYPE_F);;
for (uint32_t valueID = 0; valueID < valueNum; ++valueID)
value[valueID] = GenRegister::retype(sel.selReg(insn.getValue(valueID)), GEN_TYPE_F);
- sel.UNTYPED_WRITE(addr, value, valueNum, bti);
+ sel.UNTYPED_WRITE(addr, value.data(), valueNum, bti);
}
void emitByteScatter(Selection::Opaque &sel,
diff --git a/backend/src/backend/gen_insn_selection.hpp b/backend/src/backend/gen_insn_selection.hpp
index 0c8263d..f6735c2 100644
--- a/backend/src/backend/gen_insn_selection.hpp
+++ b/backend/src/backend/gen_insn_selection.hpp
@@ -117,7 +117,7 @@ namespace gbe
/*! To store various indices */
uint16_t index;
/*! Variable sized. Destinations and sources go here */
- GenRegister regs[];
+ GenRegister regs[0];
private:
/*! Just Selection class can create SelectionInstruction */
SelectionInstruction(SelectionOpcode, uint32_t dstNum, uint32_t srcNum);
diff --git a/backend/src/backend/gen_reg_allocation.hpp b/backend/src/backend/gen_reg_allocation.hpp
index b9859d7..5541304 100644
--- a/backend/src/backend/gen_reg_allocation.hpp
+++ b/backend/src/backend/gen_reg_allocation.hpp
@@ -32,7 +32,7 @@ namespace gbe
{
class Selection; // Pre-register allocation code generation
class GenRegister; // Pre-register allocation Gen register
- class GenRegInterval; // Liveness interval for each register
+ struct GenRegInterval; // Liveness interval for each register
class GenContext; // Gen specific context
/*! Register allocate (i.e. virtual to physical register mapping) */
diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp
index df0df28..d33c533 100644
--- a/backend/src/backend/program.cpp
+++ b/backend/src/backend/program.cpp
@@ -119,9 +119,9 @@ namespace gbe {
// Now compile the code to llvm using clang
#if LLVM_VERSION_MINOR <= 1
- std::string compileCmd = LLVM_PREFIX "/bin/clang -x cl -fno-color-diagnostics -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
+ std::string compileCmd = "clang -x cl -fno-color-diagnostics -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
#else
- std::string compileCmd = LLVM_PREFIX "/bin/clang -ffp-contract=off -emit-llvm -O3 -target nvptx -x cl -c ";
+ std::string compileCmd = "clang -ffp-contract=off -emit-llvm -O3 -target nvptx -x cl -c ";
#endif /* LLVM_VERSION_MINOR <= 1 */
compileCmd += clName;
compileCmd += " ";
diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp
index d76aa74..9fd4247 100644
--- a/backend/src/ir/instruction.cpp
+++ b/backend/src/ir/instruction.cpp
@@ -294,7 +294,7 @@ namespace ir {
LabelIndex labelIndex; //!< Index of the label the branch targets
bool hasPredicate:1; //!< Is it predicated?
bool hasLabel:1; //!< Is there any target label?
- Register dst[]; //!< No destination
+ Register dst[0]; //!< No destination
};
class ALIGNED_INSTRUCTION LoadInstruction :
@@ -334,7 +334,7 @@ namespace ir {
INLINE void out(std::ostream &out, const Function &fn) const;
INLINE bool isAligned(void) const { return !!dwAligned; }
Type type; //!< Type to store
- Register src[]; //!< Address where to load from
+ Register src[0]; //!< Address where to load from
Register offset; //!< Alias to make it similar to store
Tuple values; //!< Values to load
AddressSpace addrSpace; //!< Where to load
@@ -389,7 +389,7 @@ namespace ir {
AddressSpace addrSpace; //!< Where to store
uint8_t valueNum:7; //!< Number of values to store
uint8_t dwAligned:1; //!< DWORD aligned is what matters with GEN
- Register dst[]; //!< No destination
+ Register dst[0]; //!< No destination
};
class ALIGNED_INSTRUCTION SampleInstruction : // TODO
@@ -488,7 +488,7 @@ namespace ir {
bool wellFormed(const Function &fn, std::string &why) const;
INLINE void out(std::ostream &out, const Function &fn) const;
Register dst[1]; //!< RegisterData to store into
- Register src[]; //!< No source register
+ Register src[0]; //!< No source register
ImmediateIndex immediateIndex; //!< Index in the vector of immediates
Type type; //!< Type of the immediate
};
@@ -507,7 +507,7 @@ namespace ir {
INLINE bool wellFormed(const Function &fn, std::string &why) const;
INLINE void out(std::ostream &out, const Function &fn) const;
uint32_t parameters;
- Register dst[], src[];
+ Register dst[0], src[0];
};
class ALIGNED_INSTRUCTION LabelInstruction :
@@ -524,7 +524,7 @@ namespace ir {
INLINE bool wellFormed(const Function &fn, std::string &why) const;
INLINE void out(std::ostream &out, const Function &fn) const;
LabelIndex labelIndex; //!< Index of the label
- Register dst[], src[];
+ Register dst[0], src[0];
};
#undef ALIGNED_INSTRUCTION
diff --git a/backend/src/llvm/llvm_to_gen.cpp b/backend/src/llvm/llvm_to_gen.cpp
index 21193a5..ea3d9eb 100644
--- a/backend/src/llvm/llvm_to_gen.cpp
+++ b/backend/src/llvm/llvm_to_gen.cpp
@@ -55,7 +55,7 @@ namespace gbe
std::string errInfo;
std::unique_ptr<llvm::raw_fd_ostream> o = NULL;
if (OCL_OUTPUT_LLVM_BEFORE_EXTRA_PASS || OCL_OUTPUT_LLVM)
- o = std::unique_ptr<llvm::raw_fd_ostream>(new llvm::raw_fd_ostream("-", errInfo));
+ o = std::unique_ptr<llvm::raw_fd_ostream>(new llvm::raw_fd_ostream(fileno(stdout), false));
// Get the module from its file
SMDiagnostic Err;
@@ -83,13 +83,6 @@ namespace gbe
passes.add(createPrintModulePass(&*o));
passes.run(mod);
- // raw_fd_ostream closes stdout. We must reopen it
- if (OCL_OUTPUT_LLVM_BEFORE_EXTRA_PASS || OCL_OUTPUT_LLVM) {
- o = NULL;
- const int fd = open("/dev/tty", O_WRONLY);
- stdout = fdopen(fd, "w");
- }
-
return true;
}
} /* namespace gbe */
diff --git a/backend/src/sys/alloc.cpp b/backend/src/sys/alloc.cpp
index cc2186f..2db95c9 100644
--- a/backend/src/sys/alloc.cpp
+++ b/backend/src/sys/alloc.cpp
@@ -271,7 +271,7 @@ namespace gbe
/// Linux Platform
////////////////////////////////////////////////////////////////////////////////
-#if defined(__LINUX__)
+#if defined(__LINUX__) || defined(__GLIBC__)
#include <unistd.h>
#include <sys/mman.h>
diff --git a/backend/src/sys/mutex.hpp b/backend/src/sys/mutex.hpp
index c8e3f2f..1a462b0 100644
--- a/backend/src/sys/mutex.hpp
+++ b/backend/src/sys/mutex.hpp
@@ -42,17 +42,17 @@ namespace gbe
/*! active mutex */
class MutexActive {
public:
- INLINE MutexActive(void) : $lock(LOCK_IS_FREE) {}
+ INLINE MutexActive(void) : _lock(LOCK_IS_FREE) {}
INLINE void lock(void) {
GBE_COMPILER_READ_BARRIER;
- while (cmpxchg($lock, LOCK_IS_TAKEN, LOCK_IS_FREE) != LOCK_IS_FREE)
+ while (cmpxchg(_lock, LOCK_IS_TAKEN, LOCK_IS_FREE) != LOCK_IS_FREE)
_mm_pause();
GBE_COMPILER_READ_BARRIER;
}
- INLINE void unlock(void) { $lock.storeRelease(LOCK_IS_FREE); }
+ INLINE void unlock(void) { _lock.storeRelease(LOCK_IS_FREE); }
protected:
- enum ${ LOCK_IS_FREE = 0, LOCK_IS_TAKEN = 1 };
- Atomic $lock;
+ enum { LOCK_IS_FREE = 0, LOCK_IS_TAKEN = 1 };
+ Atomic _lock;
MutexActive(const MutexActive&); // don't implement
MutexActive& operator=(const MutexActive&); // don't implement
GBE_CLASS(MutexActive);
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 140a864..2f590c6 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -32,8 +32,19 @@ if (EGL_FOUND)
set (OPENCL_SRC ${OPENCL_SRC} cl_mem_gl.c cl_gl_api.c x11/gbm_dri2_x11_platform.c)
SET(CMAKE_CXX_FLAGS "-DHAS_EGL ${CMAKE_CXX_FLAGS}")
SET(CMAKE_C_FLAGS "-DHAS_EGL ${CMAKE_C_FLAGS}")
+SET(OPTIONAL_EGL_LIBRARY "${EGL_LIBRARY}")
+else(EGL_FOUND)
+SET(OPTIONAL_EGL_LIBRARY "")
endif (EGL_FOUND)
+if (OCLIcd_FOUND)
+set (OPENCL_SRC ${OPENCL_SRC} cl_khr_icd.c)
+SET(CMAKE_CXX_FLAGS "-DHAS_OCLIcd ${CMAKE_CXX_FLAGS}")
+SET(CMAKE_C_FLAGS "-DHAS_OCLIcd ${CMAKE_C_FLAGS}")
+endif (OCLIcd_FOUND)
+
+SET(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,-Bsymbolic")
+
link_directories (${LLVM_LIBRARY_DIR})
add_library(cl SHARED ${OPENCL_SRC})
target_link_libraries(
@@ -45,6 +56,6 @@ target_link_libraries(
${DRM_INTEL_LIBRARY}
${DRM_LIBRARY}
${OPENGL_LIBRARIES}
- ${EGL_LIBRARY}
+ ${OPTIONAL_EGL_LIBRARY}
${GBM_LIBRARY})
install (TARGETS cl LIBRARY DESTINATION lib)
diff --git a/src/cl_api.c b/src/cl_api.c
index 88e1096..c39ef83 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -30,6 +30,7 @@
#include "cl_utils.h"
#include "CL/cl.h"
+#include "CL/cl_ext.h"
#include "CL/cl_intel.h"
#include <stdio.h>
@@ -51,7 +52,7 @@ clGetPlatformInfo(cl_platform_id platform,
void * param_value,
size_t * param_value_size_ret)
{
- return cl_get_platform_into(platform,
+ return cl_get_platform_info(platform,
param_name,
param_value_size,
param_value,
@@ -1169,7 +1170,13 @@ clEnqueueBarrier(cl_command_queue command_queue)
void*
clGetExtensionFunctionAddress(const char *func_name)
{
- NOT_IMPLEMENTED;
+ if (func_name == NULL)
+ return NULL;
+#ifdef HAS_OCLIcd
+ /* cl_khr_icd */
+ if (strcmp("clIcdGetPlatformIDsKHR", func_name) == 0)
+ return (void *)clIcdGetPlatformIDsKHR;
+#endif
return NULL;
}
diff --git a/src/cl_command_queue.c b/src/cl_command_queue.c
index 37e78b4..a22884f 100644
--- a/src/cl_command_queue.c
+++ b/src/cl_command_queue.c
@@ -26,6 +26,7 @@
#include "cl_utils.h"
#include "cl_alloc.h"
#include "cl_driver.h"
+#include "cl_khr_icd.h"
#include <assert.h>
#include <stdio.h>
@@ -38,6 +39,7 @@ cl_command_queue_new(cl_context ctx)
assert(ctx);
TRY_ALLOC_NO_ERR (queue, CALLOC(struct _cl_command_queue));
+ SET_ICD(queue->dispatch)
queue->magic = CL_MAGIC_QUEUE_HEADER;
queue->ref_n = 1;
queue->ctx = ctx;
diff --git a/src/cl_command_queue.h b/src/cl_command_queue.h
index 1e2bcc1..6387ae1 100644
--- a/src/cl_command_queue.h
+++ b/src/cl_command_queue.h
@@ -29,6 +29,7 @@ struct intel_gpgpu;
/* Basically, this is a (kind-of) batch buffer */
struct _cl_command_queue {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a command queue */
volatile int ref_n; /* We reference count this object */
cl_context ctx; /* Its parent context */
diff --git a/src/cl_context.c b/src/cl_context.c
index d902537..4a1925c 100644
--- a/src/cl_context.c
+++ b/src/cl_context.c
@@ -25,6 +25,7 @@
#include "cl_alloc.h"
#include "cl_utils.h"
#include "cl_driver.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
#include "CL/cl_gl.h"
@@ -154,6 +155,7 @@ cl_context_new(struct _cl_context_prop *props)
TRY_ALLOC_NO_ERR (ctx, CALLOC(struct _cl_context));
TRY_ALLOC_NO_ERR (ctx->drv, cl_driver_new(props));
+ SET_ICD(ctx->dispatch)
ctx->props = *props;
ctx->magic = CL_MAGIC_CONTEXT_HEADER;
ctx->ref_n = 1;
diff --git a/src/cl_context.h b/src/cl_context.h
index d9f2fe4..5dff2ef 100644
--- a/src/cl_context.h
+++ b/src/cl_context.h
@@ -23,6 +23,7 @@
#include "cl_internals.h"
#include "cl_driver.h"
#include "CL/cl.h"
+#include "cl_khr_icd.h"
#include <stdint.h>
#include <pthread.h>
@@ -52,6 +53,7 @@ struct _cl_context_prop {
/* Encapsulate the whole device */
struct _cl_context {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a context */
volatile int ref_n; /* We reference count this object */
cl_driver drv; /* Handles HW or simulator */
diff --git a/src/cl_device_id.c b/src/cl_device_id.c
index 8d47aa5..9f8e6ad 100644
--- a/src/cl_device_id.c
+++ b/src/cl_device_id.c
@@ -23,6 +23,7 @@
#include "cl_utils.h"
#include "cl_driver.h"
#include "cl_device_data.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
#include <assert.h>
@@ -30,6 +31,7 @@
#include <string.h>
static struct _cl_device_id intel_ivb_gt2_device = {
+ INIT_ICD(dispatch)
.max_compute_unit = 128,
.max_thread_per_unit = 8,
.max_work_item_sizes = {512, 512, 512},
@@ -41,6 +43,7 @@ static struct _cl_device_id intel_ivb_gt2_device = {
};
static struct _cl_device_id intel_ivb_gt1_device = {
+ INIT_ICD(dispatch)
.max_compute_unit = 64,
.max_thread_per_unit = 8,
.max_work_item_sizes = {512, 512, 512},
@@ -53,6 +56,7 @@ static struct _cl_device_id intel_ivb_gt1_device = {
/* XXX we clone IVB for HSW now */
static struct _cl_device_id intel_hsw_device = {
+ INIT_ICD(dispatch)
.max_compute_unit = 64,
.max_thread_per_unit = 8,
.max_work_item_sizes = {512, 512, 512},
diff --git a/src/cl_device_id.h b/src/cl_device_id.h
index b7ba6b3..610eaf6 100644
--- a/src/cl_device_id.h
+++ b/src/cl_device_id.h
@@ -22,6 +22,7 @@
/* Store complete information about the device */
struct _cl_device_id {
+ DEFINE_ICD(dispatch)
cl_device_type device_type;
cl_uint vendor_id;
cl_uint max_compute_unit;
diff --git a/src/cl_event.h b/src/cl_event.h
index 879357c..23378e8 100644
--- a/src/cl_event.h
+++ b/src/cl_event.h
@@ -21,6 +21,7 @@
#define __CL_EVENT_H__
struct _cl_event {
+ DEFINE_ICD(dispatch)
};
#endif /* __CL_EVENT_H__ */
diff --git a/src/cl_extensions.c b/src/cl_extensions.c
index 1c87681..052b589 100644
--- a/src/cl_extensions.c
+++ b/src/cl_extensions.c
@@ -15,7 +15,7 @@ static struct cl_extensions intel_extensions =
{
{
#define DECL_EXT(name) \
- {(struct cl_extension_base){.ext_id = name##_ext_id, .ext_name = #name, .ext_enabled = 0}},
+ {(struct cl_extension_base){.ext_id = cl_##name##_ext_id, .ext_name = "cl_" #name, .ext_enabled = 0}},
DECL_ALL_EXTENSIONS
},
#undef DECL_EXT
@@ -27,7 +27,15 @@ void check_basic_extension(cl_extensions_t *extensions)
int id;
for(id = BASE_EXT_START_ID; id <= BASE_EXT_END_ID; id++)
//It seems we only support this mandatory extension.
- if (id == EXT_ID(cl_khr_byte_addressable_store))
+ if (id == EXT_ID(khr_byte_addressable_store))
+ extensions->extensions[id].base.ext_enabled = 1;
+}
+
+void check_opt1_extension(cl_extensions_t *extensions)
+{
+ int id;
+ for(id = OPT1_EXT_START_ID; id <= OPT1_EXT_END_ID; id++)
+ if (id == EXT_ID(khr_icd))
extensions->extensions[id].base.ext_enabled = 1;
}
@@ -48,9 +56,9 @@ static struct cl_gl_ext_deps egl_funcs;
&& egl_funcs.eglDestroyImageKHR_func != NULL) {
/* For now, we only support cl_khr_gl_sharing. */
for(id = GL_EXT_START_ID; id <= GL_EXT_END_ID; id++)
- if (id == EXT_ID(cl_khr_gl_sharing)) {
+ if (id == EXT_ID(khr_gl_sharing)) {
extensions->extensions[id].base.ext_enabled = 1;
- extensions->extensions[id].EXT_STRUCT_NAME(cl_khr_gl_sharing).gl_ext_deps = &egl_funcs;
+ extensions->extensions[id].EXT_STRUCT_NAME(khr_gl_sharing).gl_ext_deps = &egl_funcs;
}
}
#endif
@@ -101,6 +109,7 @@ cl_intel_platform_extension_init(cl_platform_id intel_platform)
return;
}
check_basic_extension(&intel_extensions);
+ check_opt1_extension(&intel_extensions);
check_gl_extension(&intel_extensions);
check_intel_extension(&intel_extensions);
process_extension_str(&intel_extensions);
diff --git a/src/cl_extensions.h b/src/cl_extensions.h
index 39cbff2..51eb8e0 100644
--- a/src/cl_extensions.h
+++ b/src/cl_extensions.h
@@ -2,36 +2,37 @@
* names must be returned by all device that
* support OpenCL C 1.2. */
#define DECL_BASE_EXTENSIONS \
- DECL_EXT(cl_khr_global_int32_base_atomics) \
- DECL_EXT(cl_khr_global_int32_extended_atomics) \
- DECL_EXT(cl_khr_local_int32_base_atomics) \
- DECL_EXT(cl_khr_local_int32_extended_atomics) \
- DECL_EXT(cl_khr_byte_addressable_store) \
- DECL_EXT(cl_khr_fp64)
+ DECL_EXT(khr_global_int32_base_atomics) \
+ DECL_EXT(khr_global_int32_extended_atomics) \
+ DECL_EXT(khr_local_int32_base_atomics) \
+ DECL_EXT(khr_local_int32_extended_atomics) \
+ DECL_EXT(khr_byte_addressable_store) \
+ DECL_EXT(khr_fp64)
/* The OPT1 extensions are those optional extensions
* which don't have external dependecies*/
#define DECL_OPT1_EXTENSIONS \
- DECL_EXT(cl_khr_int64_base_atomics)\
- DECL_EXT(cl_khr_int64_extended_atomics)\
- DECL_EXT(cl_khr_3d_image_writes)\
- DECL_EXT(cl_khr_fp16)\
- DECL_EXT(cl_khr_image2d_from_buffer)\
- DECL_EXT(cl_khr_initialize_memory)\
- DECL_EXT(cl_khr_context_abort)\
- DECL_EXT(cl_khr_depth_images)\
- DECL_EXT(cl_khr_spir)
+ DECL_EXT(khr_int64_base_atomics)\
+ DECL_EXT(khr_int64_extended_atomics)\
+ DECL_EXT(khr_3d_image_writes)\
+ DECL_EXT(khr_fp16)\
+ DECL_EXT(khr_image2d_from_buffer)\
+ DECL_EXT(khr_initialize_memory)\
+ DECL_EXT(khr_context_abort)\
+ DECL_EXT(khr_depth_images)\
+ DECL_EXT(khr_spir) \
+ DECL_EXT(khr_icd)
#define DECL_GL_EXTENSIONS \
- DECL_EXT(cl_khr_gl_sharing)\
- DECL_EXT(cl_khr_gl_event)\
- DECL_EXT(cl_khr_gl_depth_images)\
- DECL_EXT(cl_khr_gl_msaa_sharing)
+ DECL_EXT(khr_gl_sharing)\
+ DECL_EXT(khr_gl_event)\
+ DECL_EXT(khr_gl_depth_images)\
+ DECL_EXT(khr_gl_msaa_sharing)
#define DECL_D3D_EXTENSIONS \
- DECL_EXT(cl_khr_d3d10_sharing)\
- DECL_EXT(cl_khr_dx9_media_sharing)\
- DECL_EXT(cl_khr_d3d11_sharing)\
+ DECL_EXT(khr_d3d10_sharing)\
+ DECL_EXT(khr_dx9_media_sharing)\
+ DECL_EXT(khr_d3d11_sharing)\
#define DECL_ALL_EXTENSIONS \
DECL_BASE_EXTENSIONS \
@@ -39,8 +40,8 @@
DECL_GL_EXTENSIONS \
DECL_D3D_EXTENSIONS
-#define EXT_ID(name) name ## _ext_id
-#define EXT_STRUCT_NAME(name) name ##ext
+#define EXT_ID(name) cl_ ## name ## _ext_id
+#define EXT_STRUCT_NAME(name) cl_ ## name ## ext
/*Declare enum ids */
typedef enum {
#define DECL_EXT(name) EXT_ID(name),
@@ -49,12 +50,15 @@ DECL_ALL_EXTENSIONS
cl_khr_extension_id_max
}cl_extension_enum;
-#define BASE_EXT_START_ID EXT_ID(cl_khr_global_int32_base_atomics)
-#define BASE_EXT_END_ID EXT_ID(cl_khr_fp64)
-#define GL_EXT_START_ID EXT_ID(cl_khr_gl_sharing)
-#define GL_EXT_END_ID EXT_ID(cl_khr_gl_msaa_sharing)
+#define BASE_EXT_START_ID EXT_ID(khr_global_int32_base_atomics)
+#define BASE_EXT_END_ID EXT_ID(khr_fp64)
+#define OPT1_EXT_START_ID EXT_ID(khr_int64_base_atomics)
+#define OPT1_EXT_END_ID EXT_ID(khr_icd)
+#define GL_EXT_START_ID EXT_ID(khr_gl_sharing)
+#define GL_EXT_END_ID EXT_ID(khr_gl_msaa_sharing)
#define IS_BASE_EXTENSION(id) (id >= BASE_EXT_START_ID && id <= BASE_EXT_END_ID)
+#define IS_OPT1_EXTENSION(id) (id >= OPT1_EXT_START_ID && id <= OPT1_EXT_END_ID)
#define IS_GL_EXTENSION(id) (id >= GL_EXT_START_ID && id <= GL_EXT_END_ID)
struct cl_extension_base {
diff --git a/src/cl_kernel.c b/src/cl_kernel.c
index 356a8a7..bbd4438 100644
--- a/src/cl_kernel.c
+++ b/src/cl_kernel.c
@@ -24,6 +24,7 @@
#include "cl_mem.h"
#include "cl_alloc.h"
#include "cl_utils.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
#include "cl_sampler.h"
@@ -64,6 +65,7 @@ cl_kernel_new(cl_program p)
{
cl_kernel k = NULL;
TRY_ALLOC_NO_ERR (k, CALLOC(struct _cl_kernel));
+ SET_ICD(k->dispatch)
k->ref_n = 1;
k->magic = CL_MAGIC_KERNEL_HEADER;
k->program = p;
diff --git a/src/cl_kernel.h b/src/cl_kernel.h
index e444f3b..dd98fb3 100644
--- a/src/cl_kernel.h
+++ b/src/cl_kernel.h
@@ -43,6 +43,7 @@ typedef struct cl_argument {
/* One OCL function */
struct _cl_kernel {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a kernel */
volatile int ref_n; /* We reference count this object */
cl_buffer bo; /* The code itself */
diff --git a/src/cl_khr_icd.c b/src/cl_khr_icd.c
new file mode 100644
index 0000000..5f0180a
--- /dev/null
+++ b/src/cl_khr_icd.c
@@ -0,0 +1,175 @@
+/*
+ * Copyright © 2013 Simon Richter
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#include <ocl_icd.h>
+
+#include "cl_platform_id.h"
+
+/* The interop functions are not implemented in Beignet */
+#define CL_GL_INTEROP(x) NULL
+/* OpenCL 1.2 is not implemented in Beignet */
+#define CL_1_2_NOTYET(x) NULL
+
+/** Return platform list through ICD interface
+ * This code is used only if a client is linked directly against the library
+ * instead of using the ICD loader. In this case, no other implementations
+ * should exist in the process address space, so the call is equivalent to
+ * clGetPlatformIDs().
+ *
+ * @param[in] num_entries Number of entries allocated in return buffer
+ * @param[out] platforms Platform identifiers supported by this implementation
+ * @param[out] num_platforms Number of platform identifiers returned
+ * @return OpenCL error code
+ * @retval CL_SUCCESS Successful execution
+ * @retval CL_PLATFORM_NOT_FOUND_KHR No platforms provided
+ * @retval CL_INVALID_VALUE Invalid parameters
+ */
+cl_int
+clIcdGetPlatformIDsKHR(cl_uint num_entries,
+ cl_platform_id * platforms,
+ cl_uint * num_platforms)
+{
+ return cl_get_platform_ids(num_entries, platforms, num_platforms);
+}
+
+struct _cl_icd_dispatch const cl_khr_icd_dispatch = {
+ clGetPlatformIDs,
+ clGetPlatformInfo,
+ clGetDeviceIDs,
+ clGetDeviceInfo,
+ clCreateContext,
+ clCreateContextFromType,
+ clRetainContext,
+ clReleaseContext,
+ clGetContextInfo,
+ clCreateCommandQueue,
+ clRetainCommandQueue,
+ clReleaseCommandQueue,
+ clGetCommandQueueInfo,
+ (void *) NULL, /* clSetCommandQueueProperty */
+ clCreateBuffer,
+ clCreateImage2D,
+ clCreateImage3D,
+ clRetainMemObject,
+ clReleaseMemObject,
+ clGetSupportedImageFormats,
+ clGetMemObjectInfo,
+ clGetImageInfo,
+ clCreateSampler,
+ clRetainSampler,
+ clReleaseSampler,
+ clGetSamplerInfo,
+ clCreateProgramWithSource,
+ clCreateProgramWithBinary,
+ clRetainProgram,
+ clReleaseProgram,
+ clBuildProgram,
+ clUnloadCompiler,
+ clGetProgramInfo,
+ clGetProgramBuildInfo,
+ clCreateKernel,
+ clCreateKernelsInProgram,
+ clRetainKernel,
+ clReleaseKernel,
+ clSetKernelArg,
+ clGetKernelInfo,
+ clGetKernelWorkGroupInfo,
+ clWaitForEvents,
+ clGetEventInfo,
+ clRetainEvent,
+ clReleaseEvent,
+ clGetEventProfilingInfo,
+ clFlush,
+ clFinish,
+ clEnqueueReadBuffer,
+ clEnqueueWriteBuffer,
+ clEnqueueCopyBuffer,
+ clEnqueueReadImage,
+ clEnqueueWriteImage,
+ clEnqueueCopyImage,
+ clEnqueueCopyImageToBuffer,
+ clEnqueueCopyBufferToImage,
+ clEnqueueMapBuffer,
+ clEnqueueMapImage,
+ clEnqueueUnmapMemObject,
+ clEnqueueNDRangeKernel,
+ clEnqueueTask,
+ clEnqueueNativeKernel,
+ clEnqueueMarker,
+ clEnqueueWaitForEvents,
+ clEnqueueBarrier,
+ clGetExtensionFunctionAddress,
+ CL_GL_INTEROP(clCreateFromGLBuffer),
+ CL_GL_INTEROP(clCreateFromGLTexture2D),
+ CL_GL_INTEROP(clCreateFromGLTexture3D),
+ CL_GL_INTEROP(clCreateFromGLRenderbuffer),
+ CL_GL_INTEROP(clGetGLObjectInfo),
+ CL_GL_INTEROP(clGetGLTextureInfo),
+ CL_GL_INTEROP(clEnqueueAcquireGLObjects),
+ CL_GL_INTEROP(clEnqueueReleaseGLObjects),
+ CL_GL_INTEROP(clGetGLContextInfoKHR),
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ clSetEventCallback,
+ clCreateSubBuffer,
+ clSetMemObjectDestructorCallback,
+ clCreateUserEvent,
+ clSetUserEventStatus,
+ clEnqueueReadBufferRect,
+ clEnqueueWriteBufferRect,
+ clEnqueueCopyBufferRect,
+ CL_1_2_NOTYET(clCreateSubDevicesEXT),
+ CL_1_2_NOTYET(clRetainDeviceEXT),
+ CL_1_2_NOTYET(clReleaseDeviceEXT),
+#ifdef CL_VERSION_1_2
+ (void *) NULL,
+ CL_1_2_NOTYET(clCreateSubDevices),
+ CL_1_2_NOTYET(clRetainDevice),
+ CL_1_2_NOTYET(clReleaseDevice),
+ CL_1_2_NOTYET(clCreateImage),
+ CL_1_2_NOTYET(clCreateProgramWithBuiltInKernels),
+ CL_1_2_NOTYET(clCompileProgram),
+ CL_1_2_NOTYET(clLinkProgram),
+ CL_1_2_NOTYET(clUnloadPlatformCompiler),
+ CL_1_2_NOTYET(clGetKernelArgInfo),
+ CL_1_2_NOTYET(clEnqueueFillBuffer),
+ CL_1_2_NOTYET(clEnqueueFillImage),
+ CL_1_2_NOTYET(clEnqueueMigrateMemObjects),
+ CL_1_2_NOTYET(clEnqueueMarkerWithWaitList),
+ CL_1_2_NOTYET(clEnqueueBarrierWithWaitList),
+ CL_1_2_NOTYET(clGetExtensionFunctionAddressForPlatform),
+ CL_GL_INTEROP(clCreateFromGLTexture),
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL,
+ (void *) NULL
+#endif
+};
+
diff --git a/src/cl_event.h b/src/cl_khr_icd.h
similarity index 63%
copy from src/cl_event.h
copy to src/cl_khr_icd.h
index 879357c..6c8b9f4 100644
--- a/src/cl_event.h
+++ b/src/cl_khr_icd.h
@@ -1,5 +1,5 @@
/*
- * Copyright © 2012 Intel Corporation
+ * Copyright © 2013 Simon Richter
*
* This library is free software; you can redistribute it and/or
* modify it under the terms of the GNU Lesser General Public
@@ -13,15 +13,18 @@
*
* You should have received a copy of the GNU Lesser General Public
* License along with this library. If not, see <http://www.gnu.org/licenses/>.
- *
- * Author: Benjamin Segovia <benjamin.segovia at intel.com>
*/
-#ifndef __CL_EVENT_H__
-#define __CL_EVENT_H__
-
-struct _cl_event {
-};
+#ifdef HAS_OCLIcd
-#endif /* __CL_EVENT_H__ */
+#define SET_ICD(dispatch) \
+ dispatch = &cl_khr_icd_dispatch;
+#define INIT_ICD(member) .member = &cl_khr_icd_dispatch,
+#define DEFINE_ICD(member) struct _cl_icd_dispatch const *member;
+extern struct _cl_icd_dispatch const cl_khr_icd_dispatch;
+#else
+#define SET_ICD(dispatch)
+#define INIT_ICD(member)
+#define DEFINE_ICD(member)
+#endif
diff --git a/src/cl_mem.c b/src/cl_mem.c
index 3a8cfdd..e89aafa 100644
--- a/src/cl_mem.c
+++ b/src/cl_mem.c
@@ -24,6 +24,7 @@
#include "cl_alloc.h"
#include "cl_device_id.h"
#include "cl_driver.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
#include "CL/cl_intel.h"
@@ -63,6 +64,7 @@ cl_mem_allocate(cl_context ctx,
/* Allocate and inialize the structure itself */
TRY_ALLOC (mem, CALLOC(struct _cl_mem));
+ SET_ICD(mem->dispatch)
mem->ref_n = 1;
mem->magic = CL_MAGIC_MEM_HEADER;
mem->flags = flags;
diff --git a/src/cl_mem.h b/src/cl_mem.h
index 6992454..8e7a2dd 100644
--- a/src/cl_mem.h
+++ b/src/cl_mem.h
@@ -32,6 +32,7 @@ typedef enum cl_image_tiling {
/* Used for buffers and images */
struct _cl_mem {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a memory object */
volatile int ref_n; /* This object is reference counted */
cl_buffer bo; /* Data in GPU memory */
diff --git a/src/cl_mem_gl.c b/src/cl_mem_gl.c
index 44a97ef..3dfac00 100644
--- a/src/cl_mem_gl.c
+++ b/src/cl_mem_gl.c
@@ -139,11 +139,7 @@ EGLImageKHR cl_create_textured_egl_image(cl_context ctx,
EGLint egl_attribs[] = { EGL_GL_TEXTURE_LEVEL_KHR, miplevel, EGL_NONE};
assert(ctx->props.gl_type == CL_GL_EGL_DISPLAY);
-/* cl.h defined cl_khr_gl_sharing to 1. we have to undefine it here.*/
-#ifdef cl_khr_gl_sharing
-#undef cl_khr_gl_sharing
-#endif
- egl_funcs = CL_EXTENSION_GET_FUNCS(ctx, cl_khr_gl_sharing, gl_ext_deps);
+ egl_funcs = CL_EXTENSION_GET_FUNCS(ctx, khr_gl_sharing, gl_ext_deps);
assert(egl_funcs != NULL);
egl_display = (EGLDisplay)ctx->props.egl_display;
egl_context = (EGLDisplay)ctx->props.gl_context;
diff --git a/src/cl_platform_id.c b/src/cl_platform_id.c
index 0d46453..2f66064 100644
--- a/src/cl_platform_id.c
+++ b/src/cl_platform_id.c
@@ -21,6 +21,7 @@
#include "cl_internals.h"
#include "cl_utils.h"
#include "CL/cl.h"
+#include "CL/cl_ext.h"
#include <stdlib.h>
#include <string.h>
@@ -30,10 +31,12 @@
.JOIN(FIELD,_sz) = sizeof(STRING) + 1,
static struct _cl_platform_id intel_platform_data = {
+ INIT_ICD(dispatch)
DECL_INFO_STRING(profile, "FULL_PROFILE")
DECL_INFO_STRING(version, OCL_VERSION_STRING)
DECL_INFO_STRING(name, "Experiment Intel Gen OCL Driver")
DECL_INFO_STRING(vendor, "Intel")
+ DECL_INFO_STRING(icd_suffix_khr, "Intel")
};
#undef DECL_INFO_STRING
@@ -86,7 +89,7 @@ cl_get_platform_ids(cl_uint num_entries,
return CL_SUCCESS;
LOCAL cl_int
-cl_get_platform_into(cl_platform_id platform,
+cl_get_platform_info(cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void * param_value,
@@ -103,6 +106,7 @@ cl_get_platform_into(cl_platform_id platform,
GET_FIELD_SZ (PLATFORM_NAME, name);
GET_FIELD_SZ (PLATFORM_VENDOR, vendor);
GET_FIELD_SZ (PLATFORM_EXTENSIONS, extensions);
+ GET_FIELD_SZ (PLATFORM_ICD_SUFFIX_KHR, icd_suffix_khr);
default: return CL_INVALID_VALUE;
}
}
@@ -114,6 +118,7 @@ cl_get_platform_into(cl_platform_id platform,
DECL_FIELD (PLATFORM_NAME, name);
DECL_FIELD (PLATFORM_VENDOR, vendor);
DECL_FIELD (PLATFORM_EXTENSIONS, extensions);
+ DECL_FIELD (PLATFORM_ICD_SUFFIX_KHR, icd_suffix_khr);
default: return CL_INVALID_VALUE;
}
}
diff --git a/src/cl_platform_id.h b/src/cl_platform_id.h
index 2cbebce..edd3aae 100644
--- a/src/cl_platform_id.h
+++ b/src/cl_platform_id.h
@@ -22,19 +22,23 @@
#include "cl_internals.h"
#include "cl_extensions.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
struct _cl_platform_id {
+ DEFINE_ICD(dispatch)
const char *profile;
const char *version;
const char *name;
const char *vendor;
char *extensions;
+ const char *icd_suffix_khr;
size_t profile_sz;
size_t version_sz;
size_t name_sz;
size_t vendor_sz;
size_t extensions_sz;
+ size_t icd_suffix_khr_sz;
struct cl_extensions *internal_extensions;
};
@@ -47,7 +51,7 @@ extern cl_int cl_get_platform_ids(cl_uint num_entries,
cl_uint * num_platforms);
/* Return information for the current platform */
-extern cl_int cl_get_platform_into(cl_platform_id platform,
+extern cl_int cl_get_platform_info(cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void * param_value,
diff --git a/src/cl_program.c b/src/cl_program.c
index ecffb00..0c48ef3 100644
--- a/src/cl_program.c
+++ b/src/cl_program.c
@@ -23,6 +23,7 @@
#include "cl_context.h"
#include "cl_alloc.h"
#include "cl_utils.h"
+#include "cl_khr_icd.h"
#include "CL/cl.h"
#include "CL/cl_intel.h"
@@ -91,6 +92,7 @@ cl_program_new(cl_context ctx)
/* Allocate the structure */
TRY_ALLOC_NO_ERR (p, CALLOC(struct _cl_program));
+ SET_ICD(p->dispatch)
p->ref_n = 1;
p->magic = CL_MAGIC_PROGRAM_HEADER;
p->ctx = ctx;
diff --git a/src/cl_program.h b/src/cl_program.h
index fd00621..161d858 100644
--- a/src/cl_program.h
+++ b/src/cl_program.h
@@ -38,6 +38,7 @@ enum {
/* This maps an OCL file containing some kernels */
struct _cl_program {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a program */
volatile int ref_n; /* We reference count this object */
gbe_program opaque; /* (Opaque) program as ouput by the compiler */
diff --git a/src/cl_sampler.c b/src/cl_sampler.c
index fd88a77..d3e61da 100644
--- a/src/cl_sampler.c
+++ b/src/cl_sampler.c
@@ -21,6 +21,7 @@
#include "cl_sampler.h"
#include "cl_utils.h"
#include "cl_alloc.h"
+#include "cl_khr_icd.h"
#include <assert.h>
@@ -36,6 +37,7 @@ cl_sampler_new(cl_context ctx,
/* Allocate and inialize the structure itself */
TRY_ALLOC (sampler, CALLOC(struct _cl_sampler));
+ SET_ICD(sampler->dispatch)
sampler->ref_n = 1;
sampler->magic = CL_MAGIC_SAMPLER_HEADER;
sampler->normalized_coords = normalized_coords;
diff --git a/src/cl_sampler.h b/src/cl_sampler.h
index 800de4c..da9a488 100644
--- a/src/cl_sampler.h
+++ b/src/cl_sampler.h
@@ -25,6 +25,7 @@
/* How to access images */
struct _cl_sampler {
+ DEFINE_ICD(dispatch)
uint64_t magic; /* To identify it as a sampler object */
volatile int ref_n; /* This object is reference counted */
cl_sampler prev, next; /* We chain the samplers in the allocator */
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index b52174d..02249e1 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -41,6 +41,10 @@
#include <EGL/egl.h>
#include <EGL/eglext.h>
#include <CL/cl_gl.h>
+
+extern EGLDisplay eglDisplay;
+extern EGLContext eglContext;
+extern EGLSurface eglSurface;
#endif
#define OCL_THROW_ERROR(FN, STATUS) \
@@ -128,10 +132,6 @@ extern cl_mem buf[MAX_BUFFER_N];
extern void* buf_data[MAX_BUFFER_N];
extern size_t globals[3];
extern size_t locals[3];
-extern Display *xDisplay;
-extern EGLDisplay eglDisplay;
-extern EGLSurface eglSurface;
-
enum {
SOURCE = 0,
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/beignet.git
More information about the Pkg-opencl-devel
mailing list