diff options
Diffstat (limited to 'test_conformance/basic/test_image_r8.cpp')
-rw-r--r-- | test_conformance/basic/test_image_r8.cpp | 196 |
1 files changed, 72 insertions, 124 deletions
diff --git a/test_conformance/basic/test_image_r8.cpp b/test_conformance/basic/test_image_r8.cpp index b633d6ab..2dca1611 100644 --- a/test_conformance/basic/test_image_r8.cpp +++ b/test_conformance/basic/test_image_r8.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -21,163 +21,111 @@ #include <sys/types.h> #include <sys/stat.h> +#include <algorithm> +#include <vector> #include "procs.h" -static const char *r_uint8_kernel_code = -"__kernel void test_r_uint8(read_only image2d_t srcimg, __global unsigned char *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" uint4 color;\n" -"\n" -" color = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" dst[indx] = (unsigned char)(color.x);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8bit_image(int w, int h, MTdata d) +namespace { +const char *r_uint8_kernel_code = R"( +__kernel void test_r_uint8(read_only image2d_t srcimg, __global unsigned char *dst, sampler_t sampler) { - unsigned char *ptr = (unsigned char*)malloc(w * h * sizeof(unsigned char)); - int i; + int tid_x = get_global_id(0); + int tid_y = get_global_id(1); + int indx = tid_y * get_image_width(srcimg) + tid_x; + uint4 color; - for (i=0; i<w*h; i++) - ptr[i] = (unsigned char)genrand_int32(d); + color = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y)); + dst[indx] = (unsigned char)(color.x); +})"; - return ptr; -} -static int -verify_8bit_image(unsigned char *image, unsigned char *outptr, int w, int h) +void generate_random_inputs(std::vector<cl_uchar> &v) { - int i; + RandomSeed seed(gRandomSeed); - for (i=0; i<w*h; i++) - { - if (outptr[i] != image[i]) - { - log_error("READ_IMAGE_R_UNSIGNED_INT8 test failed\n"); - return -1; - } - } + auto random_generator = [&seed]() { + return static_cast<cl_uchar>(genrand_int32(seed)); + }; - log_info("READ_IMAGE_R_UNSIGNED_INT8 test passed\n"); - return 0; + std::generate(v.begin(), v.end(), random_generator); } -int -test_image_r8(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +} +int test_image_r8(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - cl_mem streams[2]; - cl_image_format img_format; - cl_uchar *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[3]; - int img_width = 512; - int img_height = 512; - int err; - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - img_format.image_channel_order = CL_R; - img_format.image_channel_data_type = CL_UNSIGNED_INT8; + clMemWrapper streams[2]; + clProgramWrapper program; + clKernelWrapper kernel; + const size_t img_width = 512; + const size_t img_height = 512; + const size_t length = img_width * img_height; + int err; + + PASSIVE_REQUIRE_IMAGE_SUPPORT(device) + + const cl_image_format img_format = { CL_R, CL_UNSIGNED_INT8 }; // early out if this image type is not supported if (!is_image_format_supported(context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE2D, &img_format)) { log_info("WARNING: Image type not supported; skipping test.\n"); - return 0; + return TEST_SKIPPED_ITSELF; } - d = init_genrand( gRandomSeed ); - input_ptr = generate_8bit_image(img_width, img_height, d); - free_mtdata(d); d = NULL; + std::vector<cl_uchar> input(length); + std::vector<cl_uchar> output(length); + + generate_random_inputs(input); - output_ptr = (cl_uchar*)malloc(sizeof(cl_uchar) * img_width * img_height); streams[0] = create_image_2d(context, CL_MEM_READ_ONLY, &img_format, - img_width, img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed\n"); - return -1; - } + img_width, img_height, 0, nullptr, &err); + test_error(err, "create_image_2d failed."); streams[1] = - clCreateBuffer(context, CL_MEM_READ_WRITE, - sizeof(cl_uchar) * img_width * img_height, NULL, NULL); - if (!streams[1]) - { - log_error("clCreateBuffer failed\n"); - return -1; - } + clCreateBuffer(context, CL_MEM_READ_WRITE, length, nullptr, &err); + test_error(err, "clCreateBuffer failed."); - size_t origin[3] = {0,0,0}, region[3]={img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, - origin, region, 0, 0, - input_ptr, - 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed: %d\n", err); - return -1; - } + const size_t origin[3] = { 0, 0, 0 }, + region[3] = { img_width, img_height, 1 }; + err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, + input.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueWriteImage failed."); - err = create_single_kernel_helper(context, &program, &kernel, 1, &r_uint8_kernel_code, "test_r_uint8" ); - if (err) { - log_error("Failed to create kernel and program: %d\n", err); - return -1; - } + err = create_single_kernel_helper(context, &program, &kernel, 1, + &r_uint8_kernel_code, "test_r_uint8"); + test_error(err, "create_single_kernel_helper failed."); - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); + clSamplerWrapper sampler = clCreateSampler( + context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); + + err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); + err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); + err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); + test_error(err, "clSetKernelArgs failed\n"); + + size_t threads[] = { img_width, img_height }; + err = clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, threads, nullptr, 0, + nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed\n"); - err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed: %d\n", err); - return -1; - } - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) + err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, + output.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueReadBuffer failed\n"); + + if (0 != memcmp(input.data(), output.data(), length)) { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; + log_error("READ_IMAGE_R_UNSIGNED_INT8 test failed\n"); + err = -1; } - - err = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(cl_uchar)*img_width*img_height, (void *)output_ptr, 0, NULL, NULL ); - if (err != CL_SUCCESS) + else { - log_error("clEnqueueReadBuffer failed\n"); - return -1; + log_info("READ_IMAGE_R_UNSIGNED_INT8 test passed\n"); } - err = verify_8bit_image(input_ptr, output_ptr, img_width, img_height); - - - // cleanup - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - clReleaseKernel(kernel); - clReleaseProgram(program); - clReleaseSampler(sampler); - free(input_ptr); - free(output_ptr); - return err; } - - - - - |