[Pkg-opencl-devel] [beignet] 38/66: Imported Upstream version 0.1+git20130626+41005e0
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:06 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 a60dfce39e8300ec3d0dacaf210869ec73f9bbcd
Author: Simon Richter <sjr at debian.org>
Date: Wed Jun 26 10:06:29 2013 +0200
Imported Upstream version 0.1+git20130626+41005e0
---
backend/src/backend/gen/gen_mesa_disasm.c | 2 +
backend/src/backend/gen_context.cpp | 2 +
backend/src/backend/gen_defs.hpp | 2 +
backend/src/backend/gen_encoder.cpp | 2 +
backend/src/backend/gen_encoder.hpp | 2 +
backend/src/backend/gen_insn_selection.cpp | 14 ++-
backend/src/backend/gen_insn_selection.hxx | 2 +
backend/src/ir/instruction.cpp | 2 +
backend/src/ir/instruction.hpp | 4 +
backend/src/ir/instruction.hxx | 2 +
backend/src/llvm/llvm_gen_backend.cpp | 4 +
backend/src/llvm/llvm_gen_ocl_function.hxx | 4 +
backend/src/ocl_stdlib.h | 66 ++++++++++-
kernels/compiler_clz_int.cl | 5 +
kernels/compiler_clz_short.cl | 5 +
kernels/get_program_info.cl | 10 --
src/cl_api.c | 63 ++++++-----
src/cl_command_queue.h | 15 +--
src/cl_context.c | 1 +
utests/CMakeLists.txt | 4 +-
utests/compiler_clz_int.cpp | 31 ++++++
utests/compiler_clz_short.cpp | 31 ++++++
utests/compiler_shader_toy.cpp | 2 +-
utests/{get_program_info.cpp => get_cl_info.cpp} | 133 +++++++++++++++++++++--
utests/utest_helper.cpp | 6 +-
utests/utest_helper.hpp | 7 +-
26 files changed, 359 insertions(+), 62 deletions(-)
diff --git a/backend/src/backend/gen/gen_mesa_disasm.c b/backend/src/backend/gen/gen_mesa_disasm.c
index 17fc845..9a4e283 100644
--- a/backend/src/backend/gen/gen_mesa_disasm.c
+++ b/backend/src/backend/gen/gen_mesa_disasm.c
@@ -63,6 +63,8 @@ static const struct {
[GEN_OPCODE_RNDZ] = { .name = "rndz", .nsrc = 1, .ndst = 1 },
[GEN_OPCODE_NOT] = { .name = "not", .nsrc = 1, .ndst = 1 },
[GEN_OPCODE_LZD] = { .name = "lzd", .nsrc = 1, .ndst = 1 },
+ [GEN_OPCODE_FBH] = { .name = "fbh", .nsrc = 1, .ndst = 1 },
+ [GEN_OPCODE_FBL] = { .name = "fbl", .nsrc = 1, .ndst = 1 },
[GEN_OPCODE_MUL] = { .name = "mul", .nsrc = 2, .ndst = 1 },
[GEN_OPCODE_MAC] = { .name = "mac", .nsrc = 2, .ndst = 1 },
diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
index 53ba73c..93d3932 100644
--- a/backend/src/backend/gen_context.cpp
+++ b/backend/src/backend/gen_context.cpp
@@ -139,6 +139,8 @@ namespace gbe
const GenRegister src = ra->genReg(insn.src(0));
switch (insn.opcode) {
case SEL_OP_MOV: p->MOV(dst, src); break;
+ case SEL_OP_FBH: p->FBH(dst, src); break;
+ case SEL_OP_FBL: p->FBL(dst, src); break;
case SEL_OP_NOT: p->NOT(dst, src); break;
case SEL_OP_RNDD: p->RNDD(dst, src); break;
case SEL_OP_RNDU: p->RNDU(dst, src); break;
diff --git a/backend/src/backend/gen_defs.hpp b/backend/src/backend/gen_defs.hpp
index 9d8db5b..d1ce6b2 100644
--- a/backend/src/backend/gen_defs.hpp
+++ b/backend/src/backend/gen_defs.hpp
@@ -154,6 +154,8 @@ enum opcode {
GEN_OPCODE_MAC = 72,
GEN_OPCODE_MACH = 73,
GEN_OPCODE_LZD = 74,
+ GEN_OPCODE_FBH = 75,
+ GEN_OPCODE_FBL = 76,
GEN_OPCODE_SAD2 = 80,
GEN_OPCODE_SADA2 = 81,
GEN_OPCODE_DP4 = 84,
diff --git a/backend/src/backend/gen_encoder.cpp b/backend/src/backend/gen_encoder.cpp
index ae981b2..e96678b 100644
--- a/backend/src/backend/gen_encoder.cpp
+++ b/backend/src/backend/gen_encoder.cpp
@@ -824,6 +824,8 @@ namespace gbe
ALU1(RNDE)
ALU1(RNDD)
ALU1(RNDU)
+ ALU1(FBH)
+ ALU1(FBL)
ALU2(SEL)
ALU1(NOT)
ALU2(AND)
diff --git a/backend/src/backend/gen_encoder.hpp b/backend/src/backend/gen_encoder.hpp
index 1a5dcf9..88a3e77 100644
--- a/backend/src/backend/gen_encoder.hpp
+++ b/backend/src/backend/gen_encoder.hpp
@@ -90,6 +90,8 @@ namespace gbe
#define ALU2(OP) void OP(GenRegister dest, GenRegister src0, GenRegister src1);
#define ALU3(OP) void OP(GenRegister dest, GenRegister src0, GenRegister src1, GenRegister src2);
ALU1(MOV)
+ ALU1(FBH)
+ ALU1(FBL)
ALU1(RNDZ)
ALU1(RNDE)
ALU1(RNDD)
diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
index 1e5f514..8fb2a80 100644
--- a/backend/src/backend/gen_insn_selection.cpp
+++ b/backend/src/backend/gen_insn_selection.cpp
@@ -431,6 +431,8 @@ namespace gbe
ALU2(MACH)
ALU1(LZD)
ALU3(MAD)
+ ALU1(FBH)
+ ALU1(FBL)
#undef ALU1
#undef ALU2
#undef ALU3
@@ -1211,10 +1213,16 @@ namespace gbe
/*! Unary instruction patterns */
DECL_PATTERN(UnaryInstruction)
{
+ static ir::Type getType(const ir::Opcode opcode) {
+ if (opcode == ir::OP_FBH || opcode == ir::OP_FBL)
+ return ir::TYPE_U32;
+ return ir::TYPE_FLOAT;
+ }
+
INLINE bool emitOne(Selection::Opaque &sel, const ir::UnaryInstruction &insn) const {
const ir::Opcode opcode = insn.getOpcode();
- const GenRegister dst = sel.selReg(insn.getDst(0));
- const GenRegister src = sel.selReg(insn.getSrc(0));
+ const GenRegister dst = sel.selReg(insn.getDst(0), getType(opcode));
+ const GenRegister src = sel.selReg(insn.getSrc(0), getType(opcode));
switch (opcode) {
case ir::OP_ABS: sel.MOV(dst, GenRegister::abs(src)); break;
case ir::OP_MOV:
@@ -1228,6 +1236,8 @@ namespace gbe
case ir::OP_RNDE: sel.RNDE(dst, src); break;
case ir::OP_RNDU: sel.RNDU(dst, src); break;
case ir::OP_RNDZ: sel.RNDZ(dst, src); break;
+ case ir::OP_FBH: sel.FBH(dst, src); break;
+ case ir::OP_FBL: sel.FBL(dst, src); break;
case ir::OP_COS: sel.MATH(dst, GEN_MATH_FUNCTION_COS, src); break;
case ir::OP_SIN: sel.MATH(dst, GEN_MATH_FUNCTION_SIN, src); break;
case ir::OP_LOG: sel.MATH(dst, GEN_MATH_FUNCTION_LOG, src); break;
diff --git a/backend/src/backend/gen_insn_selection.hxx b/backend/src/backend/gen_insn_selection.hxx
index 4b5525b..cc2be08 100644
--- a/backend/src/backend/gen_insn_selection.hxx
+++ b/backend/src/backend/gen_insn_selection.hxx
@@ -41,3 +41,5 @@ DECL_SELECTION_IR(BYTE_SCATTER, ByteScatterInstruction)
DECL_SELECTION_IR(SAMPLE, SampleInstruction)
DECL_SELECTION_IR(TYPED_WRITE, TypedWriteInstruction)
DECL_SELECTION_IR(GET_IMAGE_INFO, GetImageInfoInstruction)
+DECL_SELECTION_IR(FBH, UnaryInstruction)
+DECL_SELECTION_IR(FBL, UnaryInstruction)
diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp
index a57c204..67a4c12 100644
--- a/backend/src/ir/instruction.cpp
+++ b/backend/src/ir/instruction.cpp
@@ -1239,6 +1239,8 @@ DECL_MEM_FN(GetImageInfoInstruction, uint32_t, getInfoType(void), getInfoType())
}
DECL_EMIT_FUNCTION(MOV)
+ DECL_EMIT_FUNCTION(FBH)
+ DECL_EMIT_FUNCTION(FBL)
DECL_EMIT_FUNCTION(COS)
DECL_EMIT_FUNCTION(SIN)
DECL_EMIT_FUNCTION(LOG)
diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp
index 7662b6a..0f3bd34 100644
--- a/backend/src/ir/instruction.hpp
+++ b/backend/src/ir/instruction.hpp
@@ -485,6 +485,10 @@ namespace ir {
Instruction COS(Type type, Register dst, Register src);
/*! sin.type dst src */
Instruction SIN(Type type, Register dst, Register src);
+ /*! fbh.type dst src */
+ Instruction FBH(Type type, Register dst, Register src);
+ /*! fbl.type dst src */
+ Instruction FBL(Type type, Register dst, Register src);
/*! tan.type dst src */
Instruction RCP(Type type, Register dst, Register src);
/*! abs.type dst src */
diff --git a/backend/src/ir/instruction.hxx b/backend/src/ir/instruction.hxx
index 5cf37d2..acfb45a 100644
--- a/backend/src/ir/instruction.hxx
+++ b/backend/src/ir/instruction.hxx
@@ -71,3 +71,5 @@ DECL_INSN(SAMPLE, SampleInstruction)
DECL_INSN(SYNC, SyncInstruction)
DECL_INSN(LABEL, LabelInstruction)
DECL_INSN(GET_IMAGE_INFO, GetImageInfoInstruction)
+DECL_INSN(FBH, UnaryInstruction)
+DECL_INSN(FBL, UnaryInstruction)
diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 5b7754c..08500ba 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -1678,6 +1678,8 @@ namespace gbe
regTranslator.newScalarProxy(ir::ocl::goffset2, dst); break;
case GEN_OCL_GET_WORK_DIM:
regTranslator.newScalarProxy(ir::ocl::workdim, dst); break;
+ case GEN_OCL_FBH:
+ case GEN_OCL_FBL:
case GEN_OCL_COS:
case GEN_OCL_SIN:
case GEN_OCL_SQR:
@@ -1842,6 +1844,8 @@ namespace gbe
ctx.POW(ir::TYPE_FLOAT, dst, src0, src1);
break;
}
+ case GEN_OCL_FBH: this->emitUnaryCallInst(I,CS,ir::OP_FBH); break;
+ case GEN_OCL_FBL: this->emitUnaryCallInst(I,CS,ir::OP_FBL); break;
case GEN_OCL_COS: this->emitUnaryCallInst(I,CS,ir::OP_COS); break;
case GEN_OCL_SIN: this->emitUnaryCallInst(I,CS,ir::OP_SIN); break;
case GEN_OCL_LOG: this->emitUnaryCallInst(I,CS,ir::OP_LOG); break;
diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
index 6cd7298..fe19844 100644
--- a/backend/src/llvm/llvm_gen_ocl_function.hxx
+++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
@@ -96,3 +96,7 @@ DECL_LLVM_GEN_FUNCTION(USUB_SAT_CHAR, _Z12ocl_usub_sathh)
DECL_LLVM_GEN_FUNCTION(USUB_SAT_SHORT, _Z12ocl_usub_sattt)
DECL_LLVM_GEN_FUNCTION(USUB_SAT_INT, _Z12ocl_usub_satjj)
DECL_LLVM_GEN_FUNCTION(USUB_SAT_LONG, _Z12ocl_usub_satmm)
+
+// integer built-in functions
+DECL_LLVM_GEN_FUNCTION(FBH, __gen_ocl_fbh)
+DECL_LLVM_GEN_FUNCTION(FBL, __gen_ocl_fbl)
diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
index 81a0193..3b191ab 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
@@ -4290,7 +4290,71 @@ DEC(16);
#undef DEC4
#undef DEC8
#undef DEC16
-
+/////////////////////////////////////////////////////////////////////////////
+// Integer built-in functions
+/////////////////////////////////////////////////////////////////////////////
+PURE CONST uint __gen_ocl_fbh(uint);
+PURE CONST uint __gen_ocl_fbl(uint);
+
+INLINE_OVERLOADABLE char clz(char x) {
+ if (x < 0)
+ return 0;
+ if (x == 0)
+ return 8;
+ return __gen_ocl_fbl(x) - 24;
+}
+
+INLINE_OVERLOADABLE uchar clz(uchar x) {
+ if (x == 0)
+ return 8;
+ return __gen_ocl_fbl(x) - 24;
+}
+
+INLINE_OVERLOADABLE short clz(short x) {
+ if (x < 0)
+ return 0;
+ if (x == 0)
+ return 16;
+ return __gen_ocl_fbh(x) - 16;
+}
+
+INLINE_OVERLOADABLE ushort clz(ushort x) {
+ if (x == 0)
+ return 16;
+ return __gen_ocl_fbh(x) - 16;
+}
+
+INLINE_OVERLOADABLE int clz(int x) {
+ if (x < 0)
+ return 0;
+ if (x == 0)
+ return 32;
+ return __gen_ocl_fbh(x);
+}
+
+INLINE_OVERLOADABLE uint clz(uint x) {
+ if (x == 0)
+ return 32;
+ return __gen_ocl_fbh(x);
+}
+
+#define DEC2(type) INLINE_OVERLOADABLE type##2 clz(type##2 a) { return (clz(a.s0), clz(a.s1)); }
+#define DEC3(type) INLINE_OVERLOADABLE type##3 clz(type##3 a) { return (clz(a.s0), clz(a.s1), clz(a.s2)); }
+#define DEC4(type) INLINE_OVERLOADABLE type##4 clz(type##4 a) { return (clz(a.s0), clz(a.s1), clz(a.s2), clz(a.s3)); }
+#define DEC8(type) INLINE_OVERLOADABLE type##8 clz(type##8 a) { return (clz(a.s0), clz(a.s1), clz(a.s2), clz(a.s3), clz(a.s4), clz(a.s5), clz(a.s6), clz(a.s7)); }
+#define DEC16(type) INLINE_OVERLOADABLE type##16 clz(type##16 a) { return (clz(a.s0), clz(a.s1), clz(a.s2), clz(a.s3), clz(a.s4), clz(a.s5), clz(a.s6), clz(a.s7), clz(a.s8), clz(a.s9), clz(a.sa), clz(a.sb), clz(a.sc), clz(a.sd), clz(a.se), clz(a.sf)); }
+#define DEC(n) DEC##n(char); DEC##n(uchar); DEC##n(short); DEC##n(ushort); DEC##n(int); DEC##n(uint)
+DEC(2)
+DEC(3)
+DEC(4)
+DEC(8)
+DEC(16)
+#undef DEC
+#undef DEC2
+#undef DEC3
+#undef DEC4
+#undef DEC8
+#undef DEC16
/////////////////////////////////////////////////////////////////////////////
// Work Items functions (see 6.11.1 of OCL 1.1 spec)
/////////////////////////////////////////////////////////////////////////////
diff --git a/kernels/compiler_clz_int.cl b/kernels/compiler_clz_int.cl
new file mode 100644
index 0000000..0f17f86
--- /dev/null
+++ b/kernels/compiler_clz_int.cl
@@ -0,0 +1,5 @@
+kernel void compiler_clz_int(global int *src, global int *dst) {
+ int i = get_global_id(0);
+ dst[i] = clz(src[i]);
+}
+
diff --git a/kernels/compiler_clz_short.cl b/kernels/compiler_clz_short.cl
new file mode 100644
index 0000000..1ecf7a9
--- /dev/null
+++ b/kernels/compiler_clz_short.cl
@@ -0,0 +1,5 @@
+kernel void compiler_clz_short(global short *src, global short *dst) {
+ int i = get_global_id(0);
+ dst[i] = clz(src[i]);
+}
+
diff --git a/kernels/get_program_info.cl b/kernels/get_program_info.cl
deleted file mode 100644
index 8e0dd94..0000000
--- a/kernels/get_program_info.cl
+++ /dev/null
@@ -1,10 +0,0 @@
-kernel void get_program_info( __global int *ret ) {
- uint x = get_work_dim();
- size_t y = get_global_size(0);
- y = get_global_id(0);
- y = get_local_size(0);
- y = get_local_id(0);
- y = get_num_groups(0);
- y = get_group_id(0);
- y = get_global_offset(0);
-}
diff --git a/src/cl_api.c b/src/cl_api.c
index 3c78243..f7db4bc 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -46,6 +46,19 @@
typedef intptr_t cl_device_partition_property;
#endif
+#define FILL_GETINFO_RET(TYPE, ELT, VAL, RET) \
+ do { \
+ if (param_value && param_value_size < sizeof(TYPE)*ELT) \
+ return CL_INVALID_VALUE; \
+ if (param_value) { \
+ memcpy(param_value, (VAL), sizeof(TYPE)*ELT); \
+ } \
+ \
+ if (param_value_size_ret) \
+ *param_value_size_ret = sizeof(TYPE)*ELT; \
+ return RET; \
+ } while(0)
+
static cl_int
cl_check_device_type(cl_device_type device_type)
{
@@ -341,7 +354,20 @@ clGetCommandQueueInfo(cl_command_queue command_queue,
{
cl_int err = CL_SUCCESS;
CHECK_QUEUE (command_queue);
- NOT_IMPLEMENTED;
+
+ if (param_name == CL_QUEUE_CONTEXT) {
+ FILL_GETINFO_RET (cl_context, 1, &command_queue->ctx, CL_SUCCESS);
+ } else if (param_name == CL_QUEUE_DEVICE) {
+ FILL_GETINFO_RET (cl_device_id, 1, &command_queue->ctx->device, CL_SUCCESS);
+ } else if (param_name == CL_QUEUE_REFERENCE_COUNT) {
+ cl_uint ref = command_queue->ref_n;
+ FILL_GETINFO_RET (cl_uint, 1, &ref, CL_SUCCESS);
+ } else if (param_name == CL_QUEUE_PROPERTIES) {
+ FILL_GETINFO_RET (cl_command_queue_properties, 1, &command_queue->props, CL_SUCCESS);
+ } else {
+ return CL_INVALID_VALUE;
+ }
+
error:
return err;
}
@@ -734,19 +760,6 @@ clUnloadCompiler(void)
return 0;
}
-#define FILL_AND_RET(TYPE, ELT, VAL, RET) \
- do { \
- if (param_value && param_value_size < sizeof(TYPE)*ELT) \
- return CL_INVALID_VALUE; \
- if (param_value) { \
- memcpy(param_value, (VAL), sizeof(TYPE)*ELT); \
- } \
- \
- if (param_value_size_ret) \
- *param_value_size_ret = sizeof(TYPE)*ELT; \
- return RET; \
- } while(0)
-
cl_int
clGetProgramInfo(cl_program program,
cl_program_info param_name,
@@ -761,24 +774,24 @@ clGetProgramInfo(cl_program program,
if (param_name == CL_PROGRAM_REFERENCE_COUNT) {
cl_uint ref = program->ref_n;
- FILL_AND_RET (cl_uint, 1, (&ref), CL_SUCCESS);
+ FILL_GETINFO_RET (cl_uint, 1, (&ref), CL_SUCCESS);
} else if (param_name == CL_PROGRAM_CONTEXT) {
cl_context context = program->ctx;
- FILL_AND_RET (cl_context, 1, &context, CL_SUCCESS);
+ FILL_GETINFO_RET (cl_context, 1, &context, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_NUM_DEVICES) {
cl_uint num_dev = 1; // Just 1 dev now.
- FILL_AND_RET (cl_uint, 1, &num_dev, CL_SUCCESS);
+ FILL_GETINFO_RET (cl_uint, 1, &num_dev, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_DEVICES) {
cl_device_id dev_id = program->ctx->device;
- FILL_AND_RET (cl_device_id, 1, &dev_id, CL_SUCCESS);
+ FILL_GETINFO_RET (cl_device_id, 1, &dev_id, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_SOURCE) {
if (!program->source)
- FILL_AND_RET (char, 1, &ret_str, CL_SUCCESS);
- FILL_AND_RET (char, (strlen(program->source) + 1),
+ FILL_GETINFO_RET (char, 1, &ret_str, CL_SUCCESS);
+ FILL_GETINFO_RET (char, (strlen(program->source) + 1),
program->source, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_BINARY_SIZES) {
- FILL_AND_RET (size_t, 1, (&program->bin_sz), CL_SUCCESS);
+ FILL_GETINFO_RET (size_t, 1, (&program->bin_sz), CL_SUCCESS);
} else if (param_name == CL_PROGRAM_BINARIES) {
if (!param_value)
return CL_SUCCESS;
@@ -825,15 +838,15 @@ clGetProgramBuildInfo(cl_program program,
status = CL_BUILD_ERROR;
// TODO: Support CL_BUILD_IN_PROGRESS ?
- FILL_AND_RET (cl_build_status, 1, &status, CL_SUCCESS);
+ FILL_GETINFO_RET (cl_build_status, 1, &status, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_BUILD_OPTIONS) {
if (program->is_built && program->build_opts)
ret_str = program->build_opts;
- FILL_AND_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
+ FILL_GETINFO_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
} else if (param_name == CL_PROGRAM_BUILD_LOG) {
// TODO: need to add logs in backend when compiling.
- FILL_AND_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
+ FILL_GETINFO_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
} else {
return CL_INVALID_VALUE;
}
@@ -842,8 +855,6 @@ error:
return err;
}
-#undef FILL_AND_RET
-
cl_kernel
clCreateKernel(cl_program program,
const char * kernel_name,
diff --git a/src/cl_command_queue.h b/src/cl_command_queue.h
index 0e04ff3..4f6f987 100644
--- a/src/cl_command_queue.h
+++ b/src/cl_command_queue.h
@@ -30,13 +30,14 @@ 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 */
- cl_command_queue prev, next; /* We chain the command queues together */
- cl_gpgpu gpgpu; /* Setup all GEN commands */
- cl_mem perf; /* Where to put the perf counters */
- cl_mem fulsim_out; /* Fulsim will output this buffer */
+ 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 */
+ cl_command_queue_properties props; /* Queue properties */
+ cl_command_queue prev, next; /* We chain the command queues together */
+ cl_gpgpu gpgpu; /* Setup all GEN commands */
+ cl_mem perf; /* Where to put the perf counters */
+ cl_mem fulsim_out; /* Fulsim will output this buffer */
};
/* Allocate and initialize a new command queue. Also insert it in the list of
diff --git a/src/cl_context.c b/src/cl_context.c
index fa4c7e0..0331151 100644
--- a/src/cl_context.c
+++ b/src/cl_context.c
@@ -196,6 +196,7 @@ cl_context_create_queue(cl_context ctx,
/* We create the command queue and store it in the context list of queues */
TRY_ALLOC (queue, cl_command_queue_new(ctx));
+ queue->props = properties;
exit:
if (errcode_ret)
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index c009d99..df59feb 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -22,6 +22,8 @@ set (utests_sources
compiler_array3.cpp
compiler_byte_scatter.cpp
compiler_ceil.cpp
+ compiler_clz_short.cpp
+ compiler_clz_int.cpp
compiler_convert_uchar_sat.cpp
compiler_copy_buffer.cpp
compiler_copy_image.cpp
@@ -85,7 +87,7 @@ set (utests_sources
compiler_get_image_info.cpp
compiler_vector_load_store.cpp
compiler_cl_finish.cpp
- get_program_info.cpp
+ get_cl_info.cpp
buildin_work_dim.cpp
builtin_global_size.cpp
runtime_createcontext.cpp
diff --git a/utests/compiler_clz_int.cpp b/utests/compiler_clz_int.cpp
new file mode 100644
index 0000000..c12cfc6
--- /dev/null
+++ b/utests/compiler_clz_int.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+void compiler_clz_int(void)
+{
+ const int n = 32;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_clz_int");
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ globals[0] = n;
+ locals[0] = 16;
+
+ OCL_MAP_BUFFER(0);
+ ((int*)buf_data[0])[0] = 0;
+ for (int32_t i = 1; i < (int32_t) n; ++i)
+ ((int*)buf_data[0])[i] = 0xffffffffu >> i;
+ OCL_UNMAP_BUFFER(0);
+
+ OCL_NDRANGE(1);
+
+ OCL_MAP_BUFFER(1);
+ OCL_ASSERT(((int*)buf_data[1])[0] == 32);
+ for (int i = 1; i < n; ++i)
+ OCL_ASSERT(((int*)buf_data[1])[i] == i);
+ OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_clz_int);
diff --git a/utests/compiler_clz_short.cpp b/utests/compiler_clz_short.cpp
new file mode 100644
index 0000000..eb3a370
--- /dev/null
+++ b/utests/compiler_clz_short.cpp
@@ -0,0 +1,31 @@
+#include "utest_helper.hpp"
+
+void compiler_clz_short(void)
+{
+ const size_t n = 16;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_clz_short");
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(short), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(short), NULL);
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ globals[0] = n;
+ locals[0] = 16;
+
+ OCL_MAP_BUFFER(0);
+ ((short*)buf_data[0])[0] = 0;
+ for (int32_t i = 1; i < (int32_t) n; ++i)
+ ((short*)buf_data[0])[i] = 0xffffu >> i;
+ OCL_UNMAP_BUFFER(0);
+
+ OCL_NDRANGE(1);
+
+ OCL_MAP_BUFFER(1);
+ OCL_ASSERT(((short*)buf_data[1])[0] == 16);
+ for (unsigned i = 1; i < (unsigned) n; ++i)
+ OCL_ASSERT(((short*)buf_data[1])[i] == (short)i);
+ OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_clz_short);
diff --git a/utests/compiler_shader_toy.cpp b/utests/compiler_shader_toy.cpp
index 6c34003..ead9120 100644
--- a/utests/compiler_shader_toy.cpp
+++ b/utests/compiler_shader_toy.cpp
@@ -45,7 +45,7 @@ static void run_kernel(int w, int h, const char *name)
snprintf(kernel_file, sizeof(kernel_file), "%s.cl", name);
snprintf(dst_img, sizeof(dst_img), "%s.bmp", name);
snprintf(ref_img, sizeof(ref_img), "%s_ref.bmp", name);
- OCL_CALL (cl_kernel_init, kernel_file, name, SOURCE);
+ OCL_CALL (cl_kernel_init, kernel_file, name, SOURCE, NULL);
OCL_CREATE_BUFFER(buf[0], 0, sz, NULL);
OCL_CALL (clSetKernelArg, kernel, 0, sizeof(cl_mem), &buf[0]);
diff --git a/utests/get_program_info.cpp b/utests/get_cl_info.cpp
similarity index 57%
rename from utests/get_program_info.cpp
rename to utests/get_cl_info.cpp
index 20248e8..bb31032 100644
--- a/utests/get_program_info.cpp
+++ b/utests/get_cl_info.cpp
@@ -8,10 +8,9 @@
using namespace std;
-/* ********************************************** *
- * This file to test the API of: *
- * clGetProgramInfo *
- * ********************************************** */
+/* ***************************************************** *
+ * This file to test all the API like: clGetXXXXInfo *
+ * ***************************************************** */
#define NO_STANDARD_REF 0xFFFFF
template <typename T = cl_uint>
@@ -155,19 +154,24 @@ Info_Result<T>* cast_as(void *info)
}
-#define CALL_PROGINFO_AND_RET(TYPE) \
+#define CALL_INFO_AND_RET(TYPE, FUNC, ...) \
do { \
cl_int ret; \
size_t ret_size; \
\
Info_Result<TYPE>* info = cast_as<TYPE>(x.second); \
- ret = clGetProgramInfo(program, x.first, \
+ ret = FUNC (__VA_ARGS__, x.first, \
info->size, info->get_ret(), &ret_size); \
OCL_ASSERT((!ret)); \
OCL_ASSERT((info->check_result())); \
delete info; \
} while(0)
+/* ***************************************************** *
+ * clGetProgramInfo *
+ * ***************************************************** */
+#define CALL_PROGINFO_AND_RET(TYPE) CALL_INFO_AND_RET(TYPE, clGetProgramInfo, program)
+
void get_program_info(void)
{
map<cl_program_info, void *> maps;
@@ -179,7 +183,7 @@ void get_program_info(void)
string line;
string source_code;
- sprintf(ker_path, "%s/%s", kiss_path, "get_program_info.cl");
+ sprintf(ker_path, "%s/%s", kiss_path, "compiler_if_else.cl");
ifstream in(ker_path);
while (getline(in,line)) {
@@ -192,7 +196,7 @@ void get_program_info(void)
expect_source = (char *)source_code.c_str();
- OCL_CREATE_KERNEL("get_program_info");
+ OCL_CREATE_KERNEL("compiler_if_else");
/* First test for clGetProgramInfo. We just have 1 devices now */
expect_value = 2;//One program, one kernel.
@@ -245,3 +249,116 @@ void get_program_info(void)
MAKE_UTEST_FROM_FUNCTION(get_program_info);
+/* ***************************************************** *
+ * clGetCommandQueueInfo *
+ * ***************************************************** */
+#define CALL_QUEUEINFO_AND_RET(TYPE) CALL_INFO_AND_RET(TYPE, clGetCommandQueueInfo, queue)
+
+void get_queue_info(void)
+{
+ /* use the compiler_fabs case to test us. */
+ const size_t n = 16;
+ map<cl_program_info, void *> maps;
+ int expect_ref;
+ cl_command_queue_properties prop;
+
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
+ OCL_CREATE_KERNEL("compiler_fabs");
+
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+ globals[0] = 16;
+ locals[0] = 16;
+
+ OCL_MAP_BUFFER(0);
+ for (int32_t i = 0; i < (int32_t) n; ++i)
+ ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
+ OCL_UNMAP_BUFFER(0);
+
+ // Run the kernel on GPU
+ OCL_NDRANGE(1);
+
+ /* Do our test.*/
+ maps.insert(make_pair(CL_QUEUE_CONTEXT,
+ (void *)(new Info_Result<cl_context>(ctx))));
+ maps.insert(make_pair(CL_QUEUE_DEVICE,
+ (void *)(new Info_Result<cl_device_id>(device))));
+
+ expect_ref = 1;
+ maps.insert(make_pair(CL_QUEUE_REFERENCE_COUNT,
+ (void *)(new Info_Result<>(((cl_uint)expect_ref)))));
+
+ prop = 0;
+ maps.insert(make_pair(CL_QUEUE_PROPERTIES,
+ (void *)(new Info_Result<cl_command_queue_properties>(
+ ((cl_command_queue_properties)prop)))));
+
+ std::for_each(maps.begin(), maps.end(), [](pair<cl_program_info, void *> x) {
+ switch (x.first) {
+ case CL_QUEUE_CONTEXT:
+ CALL_QUEUEINFO_AND_RET(cl_context);
+ break;
+ case CL_QUEUE_DEVICE:
+ CALL_QUEUEINFO_AND_RET(cl_device_id);
+ break;
+ case CL_QUEUE_REFERENCE_COUNT:
+ CALL_QUEUEINFO_AND_RET(cl_uint);
+ break;
+ case CL_QUEUE_PROPERTIES:
+ CALL_QUEUEINFO_AND_RET(cl_command_queue_properties);
+ break;
+ default:
+ break;
+ }
+ });
+}
+
+MAKE_UTEST_FROM_FUNCTION(get_queue_info);
+
+/* ***************************************************** *
+ * clGetProgramBuildInfo *
+ * ***************************************************** */
+#define CALL_PROG_BUILD_INFO_AND_RET(TYPE) CALL_INFO_AND_RET(TYPE, \
+ clGetProgramBuildInfo, program, device)
+
+void get_program_build_info(void)
+{
+ map<cl_program_info, void *> maps;
+ cl_build_status expect_status;
+ char build_opt[] = "-emit-llvm";
+ char log[] = "";
+ int sz;
+
+ OCL_CALL (cl_kernel_init, "compiler_if_else.cl", "compiler_if_else", SOURCE, build_opt);
+
+ /* Do our test.*/
+ expect_status = CL_BUILD_SUCCESS;
+ maps.insert(make_pair(CL_PROGRAM_BUILD_STATUS,
+ (void *)(new Info_Result<cl_build_status>(expect_status))));
+ sz = strlen(build_opt) + 1;
+ maps.insert(make_pair(CL_PROGRAM_BUILD_OPTIONS,
+ (void *)(new Info_Result<char *>(build_opt, sz))));
+ sz = strlen(log) + 1;
+ maps.insert(make_pair(CL_PROGRAM_BUILD_LOG, /* not supported now, just "" */
+ (void *)(new Info_Result<char *>(log, sz))));
+
+ std::for_each(maps.begin(), maps.end(), [](pair<cl_program_info, void *> x) {
+ switch (x.first) {
+ case CL_PROGRAM_BUILD_STATUS:
+ CALL_PROG_BUILD_INFO_AND_RET(cl_build_status);
+ break;
+ case CL_PROGRAM_BUILD_OPTIONS:
+ CALL_PROG_BUILD_INFO_AND_RET(char *);
+ break;
+ case CL_PROGRAM_BUILD_LOG:
+ CALL_PROG_BUILD_INFO_AND_RET(char *);
+ break;
+ default:
+ break;
+ }
+ });
+}
+
+MAKE_UTEST_FROM_FUNCTION(get_program_build_info);
diff --git a/utests/utest_helper.cpp b/utests/utest_helper.cpp
index 3e73db3..941b5f9 100644
--- a/utests/utest_helper.cpp
+++ b/utests/utest_helper.cpp
@@ -231,7 +231,7 @@ do_kiss_path(const char *file, cl_device_id device)
}
int
-cl_kernel_init(const char *file_name, const char *kernel_name, int format)
+cl_kernel_init(const char *file_name, const char *kernel_name, int format, const char * build_opt)
{
cl_file_map_t *fm = NULL;
char *ker_path = NULL;
@@ -259,7 +259,7 @@ cl_kernel_init(const char *file_name, const char *kernel_name, int format)
}
/* OCL requires to build the program even if it is created from a binary */
- OCL_CALL (clBuildProgram, program, 1, &device, NULL, NULL, NULL);
+ OCL_CALL (clBuildProgram, program, 1, &device, build_opt, NULL, NULL);
/* Create a kernel from the program */
kernel = clCreateKernel(program, kernel_name, &status);
@@ -394,7 +394,7 @@ cl_test_init(const char *file_name, const char *kernel_name, int format)
goto error;
/* Load the kernel */
- if ((status = cl_kernel_init(file_name, kernel_name, format)) != CL_SUCCESS)
+ if ((status = cl_kernel_init(file_name, kernel_name, format, NULL)) != CL_SUCCESS)
goto error;
error:
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index d642351..0365040 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -62,12 +62,12 @@ extern EGLSurface eglSurface;
#define OCL_CREATE_KERNEL(NAME) \
do { \
- OCL_CALL (cl_kernel_init, NAME".cl", NAME, SOURCE); \
+ OCL_CALL (cl_kernel_init, NAME".cl", NAME, SOURCE, NULL); \
} while (0)
#define OCL_CREATE_KERNEL_FROM_FILE(FILE_NAME, KERNEL_NAME) \
do { \
- OCL_CALL(cl_kernel_init, FILE_NAME".cl", KERNEL_NAME, SOURCE); \
+ OCL_CALL(cl_kernel_init, FILE_NAME".cl", KERNEL_NAME, SOURCE, NULL); \
} while (0)
#define OCL_FLUSH() \
@@ -177,7 +177,8 @@ enum {
extern int cl_ocl_init(void);
/* Init program and kernel for the test */
-extern int cl_kernel_init(const char *file_name, const char *kernel_name, int format);
+extern int cl_kernel_init(const char *file_name,
+ const char *kernel_name, int format, const char * build_opt);
/* init the bunch of global varaibles here */
extern int cl_test_init(const char *file_name, const char *kernel_name, int format);
--
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