xref: /aosp_15_r20/external/OpenCL-CTS/test_conformance/basic/test_image_r8.cpp (revision 6467f958c7de8070b317fc65bcb0f6472e388d82)
1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 //    http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "harness/compat.h"
17 
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23 
24 #include <algorithm>
25 #include <vector>
26 
27 #include "procs.h"
28 
29 namespace {
30 const char *r_uint8_kernel_code = R"(
31 __kernel void test_r_uint8(read_only image2d_t srcimg, __global unsigned char *dst, sampler_t sampler)
32 {
33     int    tid_x = get_global_id(0);
34     int    tid_y = get_global_id(1);
35     int    indx = tid_y * get_image_width(srcimg) + tid_x;
36     uint4  color;
37 
38     color = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));
39     dst[indx] = (unsigned char)(color.x);
40 })";
41 
42 
generate_random_inputs(std::vector<cl_uchar> & v)43 void generate_random_inputs(std::vector<cl_uchar> &v)
44 {
45     RandomSeed seed(gRandomSeed);
46 
47     auto random_generator = [&seed]() {
48         return static_cast<cl_uchar>(genrand_int32(seed));
49     };
50 
51     std::generate(v.begin(), v.end(), random_generator);
52 }
53 
54 }
test_image_r8(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)55 int test_image_r8(cl_device_id device, cl_context context,
56                   cl_command_queue queue, int num_elements)
57 {
58     clMemWrapper streams[2];
59     clProgramWrapper program;
60     clKernelWrapper kernel;
61     const size_t img_width = 512;
62     const size_t img_height = 512;
63     const size_t length = img_width * img_height;
64     int err;
65 
66     PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
67 
68     const cl_image_format img_format = { CL_R, CL_UNSIGNED_INT8 };
69 
70     // early out if this image type is not supported
71     if (!is_image_format_supported(context, CL_MEM_READ_ONLY,
72                                    CL_MEM_OBJECT_IMAGE2D, &img_format))
73     {
74         log_info("WARNING: Image type not supported; skipping test.\n");
75         return TEST_SKIPPED_ITSELF;
76     }
77 
78     std::vector<cl_uchar> input(length);
79     std::vector<cl_uchar> output(length);
80 
81     generate_random_inputs(input);
82 
83     streams[0] = create_image_2d(context, CL_MEM_READ_ONLY, &img_format,
84                                  img_width, img_height, 0, nullptr, &err);
85     test_error(err, "create_image_2d failed.");
86 
87     streams[1] =
88         clCreateBuffer(context, CL_MEM_READ_WRITE, length, nullptr, &err);
89     test_error(err, "clCreateBuffer failed.");
90 
91     const size_t origin[3] = { 0, 0, 0 },
92                  region[3] = { img_width, img_height, 1 };
93     err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0,
94                               input.data(), 0, nullptr, nullptr);
95     test_error(err, "clEnqueueWriteImage failed.");
96 
97     err = create_single_kernel_helper(context, &program, &kernel, 1,
98                                       &r_uint8_kernel_code, "test_r_uint8");
99     test_error(err, "create_single_kernel_helper failed.");
100 
101     clSamplerWrapper sampler = clCreateSampler(
102         context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
103     test_error(err, "clCreateSampler failed");
104 
105     err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
106     err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
107     err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
108     test_error(err, "clSetKernelArgs failed\n");
109 
110     size_t threads[] = { img_width, img_height };
111     err = clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, threads, nullptr, 0,
112                                  nullptr, nullptr);
113     test_error(err, "clEnqueueNDRangeKernel failed\n");
114 
115 
116     err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length,
117                               output.data(), 0, nullptr, nullptr);
118     test_error(err, "clEnqueueReadBuffer failed\n");
119 
120     if (0 != memcmp(input.data(), output.data(), length))
121     {
122         log_error("READ_IMAGE_R_UNSIGNED_INT8 test failed\n");
123         err = -1;
124     }
125     else
126     {
127         log_info("READ_IMAGE_R_UNSIGNED_INT8 test passed\n");
128     }
129 
130     return err;
131 }
132