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 <vector>
25
26 #include "procs.h"
27
28 namespace {
29 const char *loop_kernel_code = R"(
30 __kernel void test_loop(__global int *src, __global int *loopindx, __global int *loopcnt, __global int *dst)
31 {
32 int tid = get_global_id(0);
33 int n = get_global_size(0);
34 int i, j;
35
36 dst[tid] = 0;
37 for (i=0, j=loopindx[tid]; i<loopcnt[tid]; i++, j++)
38 {
39 if (j >= n)
40 j = 0;
41 dst[tid] += src[j];
42 }
43 }
44 )";
45
46
verify_loop(std::vector<cl_int> inptr,std::vector<cl_int> loopindx,std::vector<cl_int> loopcnt,std::vector<cl_int> outptr,int n)47 int verify_loop(std::vector<cl_int> inptr, std::vector<cl_int> loopindx,
48 std::vector<cl_int> loopcnt, std::vector<cl_int> outptr, int n)
49 {
50 for (int i = 0; i < n; i++)
51 {
52 int r = 0;
53 for (int j = 0, k = loopindx[i]; j < loopcnt[i]; j++, k++)
54 {
55 if (k >= n) k = 0;
56 r += inptr[k];
57 }
58
59 if (r != outptr[i])
60 {
61 log_error("LOOP test failed: %d found, expected %d\n", outptr[i],
62 r);
63 return -1;
64 }
65 }
66
67 log_info("LOOP test passed\n");
68 return 0;
69 }
70 }
test_loop(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)71 int test_loop(cl_device_id device, cl_context context, cl_command_queue queue,
72 int num_elements)
73 {
74 clMemWrapper streams[4];
75 clProgramWrapper program;
76 clKernelWrapper kernel;
77 int err;
78
79 size_t length = sizeof(cl_int) * num_elements;
80 std::vector<cl_int> input(length);
81 std::vector<cl_int> loop_indx(length);
82 std::vector<cl_int> loop_cnt(length);
83 std::vector<cl_int> output(length);
84
85 for (auto &stream : streams)
86 {
87 stream =
88 clCreateBuffer(context, CL_MEM_READ_WRITE, length, nullptr, &err);
89 test_error(err, "clCreateBuffer failed.");
90 }
91
92 RandomSeed seed(gRandomSeed);
93 for (int i = 0; i < num_elements; i++)
94 {
95 input[i] = static_cast<int>(genrand_int32(seed));
96 loop_indx[i] =
97 static_cast<int>(get_random_float(0, num_elements - 1, seed));
98 loop_cnt[i] =
99 static_cast<int>(get_random_float(0, num_elements / 32, seed));
100 };
101
102 err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length,
103 input.data(), 0, nullptr, nullptr);
104 test_error(err, "clEnqueueWriteBuffer failed.");
105 err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length,
106 loop_indx.data(), 0, nullptr, nullptr);
107 test_error(err, "clEnqueueWriteBuffer failed.");
108 err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length,
109 loop_cnt.data(), 0, nullptr, nullptr);
110 test_error(err, "clEnqueueWriteBuffer failed.");
111
112 err = create_single_kernel_helper(context, &program, &kernel, 1,
113 &loop_kernel_code, "test_loop");
114 test_error(err, "create_single_kernel_helper failed.");
115
116 for (int i = 0; i < ARRAY_SIZE(streams); i++)
117 {
118 err = clSetKernelArg(kernel, i, sizeof streams[i], &streams[i]);
119 test_error(err, "clSetKernelArgs failed\n");
120 }
121
122 size_t threads[] = { (size_t)num_elements };
123 err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, threads, nullptr, 0,
124 nullptr, nullptr);
125 test_error(err, "clEnqueueNDRangeKernel failed\n");
126
127 err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length,
128 output.data(), 0, nullptr, nullptr);
129 test_error(err, "clEnqueueReadBuffer failed\n");
130
131 err = verify_loop(input, loop_indx, loop_cnt, output, num_elements);
132
133
134 return err;
135 }
136