Test infrastructure for built-in functions, test Samplers.

10 files changed, 229 insertions, 7 deletions

diff --git a/src/api/api_event.cpp b/src/api/api_event.cpp
index 029506c..14f6ac0 100644
--- a/src/api/api_event.cpp
+++ b/src/api/api_event.cpp
@@ -91,7 +91,10 @@ clReleaseEvent(cl_event event)
         return CL_INVALID_EVENT;
 
     if (event->dereference())
+    {
+        event->freeDeviceData();
         delete event;
+    }
 
     return CL_SUCCESS;
 }
diff --git a/src/core/commandqueue.cpp b/src/core/commandqueue.cpp
index 32858a5..9e6aeca 100644
--- a/src/core/commandqueue.cpp
+++ b/src/core/commandqueue.cpp
@@ -403,7 +403,7 @@ Event::Event(CommandQueue *parent,
     }
 }
 
-Event::~Event()
+void Event::freeDeviceData()
 {
     if (parent() && p_device_data)
     {
@@ -412,7 +412,10 @@ Event::~Event()
 
         device->freeEventDeviceData(this);
     }
+}
 
+Event::~Event()
+{
     for (int i=0; i<p_num_events_in_wait_list; ++i)
         clReleaseEvent((cl_event)p_event_wait_list[i]);
 
diff --git a/src/core/commandqueue.h b/src/core/commandqueue.h
index 0ea10bd..a64c506 100644
--- a/src/core/commandqueue.h
+++ b/src/core/commandqueue.h
@@ -112,6 +112,7 @@ class Event : public Object
               const Event **event_wait_list,
               cl_int *errcode_ret);
 
+        void freeDeviceData();
         virtual ~Event();
 
         virtual Type type() const = 0;
diff --git a/src/core/sampler.cpp b/src/core/sampler.cpp
index aae8b89..bdf0f14 100644
--- a/src/core/sampler.cpp
+++ b/src/core/sampler.cpp
@@ -13,7 +13,7 @@ Sampler::Sampler(Context *ctx,
                  cl_addressing_mode addressing_mode,
                  cl_filter_mode filter_mode,
                  cl_int *errcode_ret)
-: Object(Object::T_Sampler), p_bitfield(0)
+: Object(Object::T_Sampler, ctx), p_bitfield(0)
 {
     if (normalized_coords)
         p_bitfield |= CLK_NORMALIZED_COORDS_TRUE;
diff --git a/src/core/sampler.h b/src/core/sampler.h
index 515a0c6..6df24a2 100644
--- a/src/core/sampler.h
+++ b/src/core/sampler.h
@@ -4,6 +4,8 @@
 #include <CL/cl.h>
 #include "object.h"
 
+// WARNING: Keep in sync with stdlib.h
+
 #define CLK_NORMALIZED_COORDS_FALSE 0x00000000
 #define CLK_NORMALIZED_COORDS_TRUE  0x00000001
 #define CLK_ADDRESS_NONE            0x00000000
diff --git a/src/runtime/stdlib.h b/src/runtime/stdlib.h
index f1c7749..dae9043 100644
--- a/src/runtime/stdlib.h
+++ b/src/runtime/stdlib.h
@@ -7,13 +7,13 @@ typedef unsigned long ulong;
 typedef int *intptr_t;
 typedef uint *uintptr_t;
 
+typedef unsigned int sampler_t;
+
 /* Standard types from Clang's stddef, Copyright (C) 2008 Eli Friedman */
 typedef __typeof__(((int*)0)-((int*)0)) ptrdiff_t;
 typedef __typeof__(sizeof(int)) size_t;
 
 /* Vectors */
-
-//typedef float float4 __attribute__((ext_vector_type(4)));
 #define COAL_VECTOR(type, len)                                  \
    typedef type type##len __attribute__((ext_vector_type(len)))
 #define COAL_VECTOR_SET(type) \
@@ -55,6 +55,17 @@ COAL_VECTOR_SET(float);
 #define write_only __write_only
 #define read_only __read_only
 
+/* Defines */
+#define CLK_NORMALIZED_COORDS_FALSE 0x00000000
+#define CLK_NORMALIZED_COORDS_TRUE  0x00000001
+#define CLK_ADDRESS_NONE            0x00000000
+#define CLK_ADDRESS_MIRRORED_REPEAT 0x00000010
+#define CLK_ADDRESS_REPEAT          0x00000020
+#define CLK_ADDRESS_CLAMP_TO_EDGE   0x00000030
+#define CLK_ADDRESS_CLAMP           0x00000040
+#define CLK_FILTER_NEAREST          0x00000000
+#define CLK_FILTER_LINEAR           0x00000100
+
 /* Management functions */
 uint get_work_dim();
 size_t get_global_size(uint dimindx);
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 3cd4a58..3c21a3c 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -10,6 +10,7 @@ set(OPENCL_TESTS_SOURCE
     test_mem.cpp
     test_kernel.cpp
     test_program.cpp
+    test_builtins.cpp
     )
 
 add_executable(tests ${OPENCL_TESTS_SOURCE})
@@ -26,3 +27,4 @@ OPENCL_TEST(tests commandqueue)
 OPENCL_TEST(tests mem)
 OPENCL_TEST(tests kernel)
 OPENCL_TEST(tests program)
+OPENCL_TEST(tests builtins)
diff --git a/tests/test_builtins.cpp b/tests/test_builtins.cpp
new file mode 100644
index 0000000..1513428
--- /dev/null
+++ b/tests/test_builtins.cpp
@@ -0,0 +1,181 @@
+#include <iostream>
+#include <cstdlib>
+
+#include "test_builtins.h"
+#include "CL/cl.h"
+
+#include <stdint.h>
+
+const char sampler_source[] =
+    "__kernel void test_case(__global uint *rs, sampler_t sampler) {\n"
+    "   sampler_t good_sampler = CLK_NORMALIZED_COORDS_TRUE |\n"
+    "                            CLK_ADDRESS_MIRRORED_REPEAT |\n"
+    "                            CLK_FILTER_NEAREST;\n"
+    "\n"
+    "   if (sampler != good_sampler) *rs = 1;"
+    "}\n";
+
+enum TestCaseKind
+{
+    NormalKind,
+    SamplerKind
+};
+
+/*
+ * To ease testing, each kernel will be a Task kernel taking a pointer to an
+ * integer and running built-in functions. If an error is encountered, the
+ * integer pointed to by the arg will be set accordingly. If the kernel succeeds,
+ * this integer is set to 0.
+ */
+static uint32_t run_kernel(const char *source, TestCaseKind kind)
+{
+    cl_platform_id platform = 0;
+    cl_device_id device;
+    cl_context ctx;
+    cl_command_queue queue;
+    cl_program program;
+    cl_int result;
+    cl_kernel kernel;
+    cl_event event;
+    cl_mem rs_buf;
+
+    cl_sampler sampler;
+
+    uint32_t rs = 0;
+
+    result = clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, 0);
+    if (result != CL_SUCCESS) return 65536;
+
+    ctx = clCreateContext(0, 1, &device, 0, 0, &result);
+    if (result != CL_SUCCESS) return 65537;
+
+    queue = clCreateCommandQueue(ctx, device, 0, &result);
+    if (result != CL_SUCCESS) return 65538;
+
+    program = clCreateProgramWithSource(ctx, 1, &source, 0, &result);
+    if (result != CL_SUCCESS) return 65539;
+
+    result = clBuildProgram(program, 1, &device, "", 0, 0);
+    if (result != CL_SUCCESS)
+    {
+        // Print log
+        char *log = 0;
+        size_t len = 0;
+
+        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, 0, &len);
+        log = (char *)std::malloc(len);
+        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, len, log, 0);
+
+        std::cout << log << std::endl;
+        std::free(log);
+
+        return 65540;
+    }
+
+    kernel = clCreateKernel(program, "test_case", &result);
+    if (result != CL_SUCCESS) return 65541;
+
+    // Create the result buffer
+    rs_buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
+                            sizeof(rs), &rs, &result);
+    if (result != CL_SUCCESS) return 65542;
+
+    result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &rs_buf);
+    if (result != CL_SUCCESS) return 65543;
+
+    // Kind
+    switch (kind)
+    {
+        case NormalKind:
+            break;
+
+        case SamplerKind:
+            sampler = clCreateSampler(ctx, 1, CL_ADDRESS_MIRRORED_REPEAT, CL_FILTER_NEAREST, &result);
+            if (result != CL_SUCCESS) return 65546;
+
+            result = clSetKernelArg(kernel, 1, sizeof(cl_sampler), &sampler);
+            if (result != CL_SUCCESS) return 65547;
+            break;
+    }
+
+    result = clEnqueueTask(queue, kernel, 0, 0, &event);
+    if (result != CL_SUCCESS) return 65544;
+
+    result = clWaitForEvents(1, &event);
+    if (result != CL_SUCCESS) return 65545;
+
+    if (kind == SamplerKind) clReleaseSampler(sampler);
+    clReleaseEvent(event);
+    clReleaseMemObject(rs_buf);
+    clReleaseKernel(kernel);
+    clReleaseProgram(program);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(ctx);
+
+    return rs;
+}
+
+static const char *default_error(uint32_t errcode)
+{
+    switch (errcode)
+    {
+        case 0:
+            return 0;
+        case 65536:
+            return "Cannot get a device ID";
+        case 65537:
+            return "Cannot create a context";
+        case 65538:
+            return "Cannot create a command queue";
+        case 65539:
+            return "Cannot create a program with given source";
+        case 65540:
+            return "Cannot build the program";
+        case 65541:
+            return "Cannot create the test_case kernel";
+        case 65542:
+            return "Cannot create a buffer holding a uint32_t";
+        case 65543:
+            return "Cannot set kernel argument";
+        case 65544:
+            return "Cannot enqueue a task kernel";
+        case 65545:
+            return "Cannot wait for the event";
+        case 65546:
+            return "Cannot create a sampler";
+        case 65547:
+            return "Cannot set a sampler kernel argument";
+
+        default:
+            return "Unknown error code";
+    }
+}
+
+START_TEST (test_sampler)
+{
+    uint32_t rs = run_kernel(sampler_source, SamplerKind);
+    const char *errstr = 0;
+
+    switch (rs)
+    {
+        case 1:
+            errstr = "Sampler bitfield invalid";
+            break;
+        default:
+            errstr = default_error(rs);
+    }
+
+    fail_if(
+        errstr != 0,
+        errstr
+    );
+}
+END_TEST
+
+TCase *cl_builtins_tcase_create(void)
+{
+    TCase *tc = NULL;
+    tc = tcase_create("builtins");
+    tcase_add_test(tc, test_sampler);
+    return tc;
+}
diff --git a/tests/test_builtins.h b/tests/test_builtins.h
new file mode 100644
index 0000000..9c7151f
--- /dev/null
+++ b/tests/test_builtins.h
@@ -0,0 +1,17 @@
+#ifndef __UTEST_BUILTINS__
+#define __UTEST_BUILTINS__
+
+#include <check.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+TCase *cl_builtins_tcase_create(void);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
+
diff --git a/tests/tests.c b/tests/tests.c
index 3a37f39..5738028 100644
--- a/tests/tests.c
+++ b/tests/tests.c
@@ -5,6 +5,7 @@
 #include "test_mem.h"
 #include "test_kernel.h"
 #include "test_program.h"
+#include "test_builtins.h"
 
 #include <stdlib.h>
 #include <stdio.h>
@@ -32,17 +33,18 @@ int main(int argc, char **argv)
     TESTSUITE(mem, "mem");
     TESTSUITE(kernel, "kernel");
     TESTSUITE(program, "program");
+    TESTSUITE(builtins, "builtins");
 
     if (s == NULL) {
-        printf("test case %s does not exist", argv[1]);
+        printf("test case %s does not exist\n", argv[1]);
         return EXIT_FAILURE;
     }
 
     SRunner *sr = srunner_create(s);
-    
+
     if (argc == 3 && !strcmp("nofork", argv[2]))
         srunner_set_fork_status (sr, CK_NOFORK);
-    
+
     srunner_run_all(sr, CK_NORMAL);
 
     n_failed_tests = srunner_ntests_failed(sr);