diff options
Diffstat (limited to 'tests/test_builtins.cpp')
| -rw-r--r-- | tests/test_builtins.cpp | 419 |
1 files changed, 419 insertions, 0 deletions
diff --git a/tests/test_builtins.cpp b/tests/test_builtins.cpp new file mode 100644 index 0000000..9a6d651 --- /dev/null +++ b/tests/test_builtins.cpp @@ -0,0 +1,419 @@ +/* + * Copyright (c) 2011, Denis Steckelmacher <steckdenis@yahoo.fr> + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the copyright holder nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE CONTRIBUTORS BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#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" + "}\n"; + +const char barrier_source[] = + "__kernel void test_case(__global uint *rs) {\n" + " *rs = 0;\n" + " int i; for (i=0; i<3; i++) barrier(0);\n" + " *rs += 1;\n" + "}\n"; + +const char image_source[] = + "__kernel void test_case(__global uint *rs, __write_only image2d_t image1,\n" + " __write_only image2d_t image2,\n" + " __read_only image2d_t image3) {\n" + " float4 fcolor;\n" + " int4 scolor;\n" + " int2 coord;\n" + " sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |\n" + " CLK_ADDRESS_MIRRORED_REPEAT |\n" + " CLK_FILTER_NEAREST;\n" + "\n" + " if (get_image_width(image1) != 4) *rs = 1;\n" + " if (get_image_height(image1) != 4) *rs = 2;\n" + " if (get_image_channel_data_type(image2) != CLK_SIGNED_INT16) *rs = 3;\n" + " if (get_image_channel_order(image2) != CLK_RGBA) *rs = 4;\n" + "\n" + " if (*rs != 0) return;\n" + "\n" + " fcolor.x = 1.0f;\n" + " fcolor.y = 0.5f;\n" + " fcolor.z = 0.0f;\n" + " fcolor.w = 1.0f;\n" + "\n" + " scolor.x = -3057;\n" + " scolor.y = 65;\n" + " scolor.z = 0;\n" + " scolor.w = 32767;\n" + "\n" + " coord.x = 3;\n" + " coord.y = 1;\n" + "\n" + " write_imagef(image1, coord, fcolor);\n" + " write_imagei(image2, coord, scolor);\n" + "\n" + " coord.x = 1;\n" + " coord.y = 1;\n" + " fcolor = read_imagef(image3, 0, coord);\n" + " if (fcolor.x < 0.99f || fcolor.y < 0.99f || fcolor.z > 0.01f ||\n" + " fcolor.w > 0.01f) { *rs = 5; return; }\n" + "\n" + " float2 fcoords;\n" + " fcoords.x = 0.31f;\n" + " fcoords.y = 0.1415f;\n" + " fcolor = read_imagef(image3, sampler, fcoords);\n" + "}\n"; + +const char builtins_source[] = + "__kernel void test_case(__global uint *rs) {\n" + " float2 f2;\n" + " float2 f2b;\n" + "\n" + " f2.x = 1.0f;\n" + " f2.y = 0.0f;\n" + " f2b.x = -0.5f;\n" + " f2b.y = (float)M_PI;\n" + "\n" + " if (cos(f2).y != 1.0f) { *rs = 1; return; }\n" + " if (cos(0.0f) != 1.0f) { *rs = 2; return; }\n" + " if (copysign(1.0f, -0.5f) != -1.0f) { *rs = 3; return; }\n" + " if (copysign(f2, f2b).x != -1.0f) { *rs = 4; return; }\n" + " if (exp2(3.0f) != 8.0f) { *rs = 5; return; }\n" + "}\n"; + +enum TestCaseKind +{ + NormalKind, + SamplerKind, + BarrierKind, + ImageKind +}; + +/* + * 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; + cl_mem mem1, mem2, mem3; + cl_image_format fmt; + + unsigned char image2d_data[3*3*4] = { + 255, 0, 0, 0, 0, 255, 0, 0, 128, 128, 128, 0, + 0, 0, 255, 0, 255, 255, 0, 0, 0, 128, 0, 0, + 255, 128, 0, 0, 128, 0, 255, 0, 0, 0, 0, 0 + }; + + 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; + + case ImageKind: + fmt.image_channel_data_type = CL_UNORM_INT8; + fmt.image_channel_order = CL_RGBA; + + mem1 = clCreateImage2D(ctx, CL_MEM_WRITE_ONLY, &fmt, 4, 4, 0, 0, &result); + if (result != CL_SUCCESS) return 65548; + + mem3 = clCreateImage2D(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &fmt, 3, 3, 0, image2d_data, &result); + if (result != CL_SUCCESS) return 65548; + + fmt.image_channel_data_type = CL_SIGNED_INT16; + + mem2 = clCreateImage2D(ctx, CL_MEM_WRITE_ONLY, &fmt, 4, 4, 0, 0, &result); + if (result != CL_SUCCESS) return 65548; + + result = clSetKernelArg(kernel, 1, sizeof(cl_mem), &mem1); + if (result != CL_SUCCESS) return 65549; + + result = clSetKernelArg(kernel, 2, sizeof(cl_mem), &mem2); + if (result != CL_SUCCESS) return 65549; + + result = clSetKernelArg(kernel, 3, sizeof(cl_mem), &mem3); + if (result != CL_SUCCESS) return 65549; + break; + + default: + break; + } + + if (kind == BarrierKind) + { + size_t local_size = 64; + size_t global_size = 64; + + result = clEnqueueNDRangeKernel(queue, kernel, 1, 0, &global_size, + &local_size, 0, 0, &event); + if (result != CL_SUCCESS) return 65544; + } + else + { + 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); + if (kind == ImageKind) + { + clReleaseMemObject(mem1); + clReleaseMemObject(mem2); + clReleaseMemObject(mem3); + } + 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 the 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"; + case 65548: + return "Cannot create an Image2D object"; + case 65549: + return "Cannot set image 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 + +START_TEST (test_barrier) +{ + uint32_t rs = run_kernel(barrier_source, BarrierKind); + + fail_if( + rs != 0x40, + default_error(rs) + ); +} +END_TEST + +START_TEST (test_image) +{ + uint32_t rs = run_kernel(image_source, ImageKind); + const char *errstr = 0; + + switch (rs) + { + case 1: + errstr = "Image1 must have width of 4"; + break; + case 2: + errstr = "Image1 must have width of 4"; + break; + case 3: + errstr = "Image2 must have type SIGNED_FLOAT16"; + break; + case 4: + errstr = "Image2 must have channel order RGBA"; + break; + case 5: + errstr = "The value read from the image is not good"; + break; + default: + errstr = default_error(rs); + } + + fail_if( + errstr != 0, + errstr + ); +} +END_TEST + +START_TEST (test_builtins) +{ + uint32_t rs = run_kernel(builtins_source, NormalKind); + const char *errstr = 0; + + switch (rs) + { + case 1: + errstr = "float2 cos(float2) doesn't behave correctly"; + break; + case 2: + errstr = "float cos(float) doesn't behave correctly"; + break; + case 3: + errstr = "float copysign(float) doesn't behave correctly"; + break; + case 4: + errstr = "float2 copysign(float2) doesn't behave correctly"; + break; + case 5: + errstr = "exp2() doesn't behave correctly"; + 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); + tcase_add_test(tc, test_barrier); + //tcase_add_test(tc, test_image); + tcase_add_test(tc, test_builtins); + return tc; +} |