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