From 339c932c579f1829d8156f64a2281a571b253211 Mon Sep 17 00:00:00 2001 From: John Kesapides <46718829+JohnKesapidesARM@users.noreply.github.com> Date: Wed, 19 Jul 2023 14:18:49 +0100 Subject: [PATCH] Use the CTS typewrappers in image_r8 (#1539) Signed-off-by: John Kesapides --- test_conformance/basic/test_image_r8.cpp | 208 +++++++++-------------- 1 file changed, 78 insertions(+), 130 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 #include +#include +#include #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 &v) +{ + RandomSeed seed(gRandomSeed); + + auto random_generator = [&seed]() { + return static_cast(genrand_int32(seed)); + }; + + std::generate(v.begin(), v.end(), random_generator); } -static int -verify_8bit_image(unsigned char *image, unsigned char *outptr, int w, int h) -{ - int i; - - for (i=0; i input(length); + std::vector 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]) + clCreateBuffer(context, CL_MEM_READ_WRITE, length, nullptr, &err); + test_error(err, "clCreateBuffer failed."); + + 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"); + test_error(err, "create_single_kernel_helper 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 = 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("clCreateBuffer failed\n"); - return -1; + log_error("READ_IMAGE_R_UNSIGNED_INT8 test failed\n"); + err = -1; } - - 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) + else { - log_error("clWriteImage failed: %d\n", err); - return -1; + log_info("READ_IMAGE_R_UNSIGNED_INT8 test passed\n"); } - 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; - } - - cl_sampler 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); - 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) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -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) - { - log_error("clEnqueueReadBuffer failed\n"); - return -1; - } - - 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; } - - - - -