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
25 #include "procs.h"
26
27 const char *conditional_kernel_code =
28 "__kernel void test_if(__global int *src, __global int *dst)\n"
29 "{\n"
30 " int tid = get_global_id(0);\n"
31 "\n"
32 " if (src[tid] == 0)\n"
33 " dst[tid] = 0x12345678;\n"
34 " else if (src[tid] == 1)\n"
35 " dst[tid] = 0x23456781;\n"
36 " else if (src[tid] == 2)\n"
37 " dst[tid] = 0x34567812;\n"
38 " else if (src[tid] == 3)\n"
39 " dst[tid] = 0x45678123;\n"
40 " else if (src[tid] == 4)\n"
41 " dst[tid] = 0x56781234;\n"
42 " else if (src[tid] == 5)\n"
43 " dst[tid] = 0x67812345;\n"
44 " else if (src[tid] == 6)\n"
45 " dst[tid] = 0x78123456;\n"
46 " else if (src[tid] == 7)\n"
47 " dst[tid] = 0x81234567;\n"
48 " else\n"
49 " dst[tid] = 0x7FFFFFFF;\n"
50 "\n"
51 "}\n";
52
53 const int results[] = {
54 0x12345678,
55 0x23456781,
56 0x34567812,
57 0x45678123,
58 0x56781234,
59 0x67812345,
60 0x78123456,
61 0x81234567,
62 };
63
64 int
verify_if(int * inptr,int * outptr,int n)65 verify_if(int *inptr, int *outptr, int n)
66 {
67 int r, i;
68
69 for (i=0; i<n; i++)
70 {
71 if (inptr[i] <= 7)
72 r = results[inptr[i]];
73 else
74 r = 0x7FFFFFFF;
75
76 if (r != outptr[i])
77 {
78 log_error("IF test failed\n");
79 return -1;
80 }
81 }
82
83 log_info("IF test passed\n");
84 return 0;
85 }
86
test_if(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)87 int test_if(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
88 {
89 cl_mem streams[2];
90 cl_int *input_ptr, *output_ptr;
91 cl_program program;
92 cl_kernel kernel;
93 size_t threads[1];
94 int err, i;
95 MTdata d = init_genrand( gRandomSeed );
96
97 size_t length = sizeof(cl_int) * num_elements;
98 input_ptr = (cl_int*)malloc(length);
99 output_ptr = (cl_int*)malloc(length);
100
101 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
102 if (!streams[0])
103 {
104 log_error("clCreateBuffer failed\n");
105 return -1;
106 }
107 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
108 if (!streams[1])
109 {
110 log_error("clCreateBuffer failed\n");
111 return -1;
112 }
113
114 for (i=0; i<num_elements; i++)
115 input_ptr[i] = (int)get_random_float(0, 32, d);
116
117 free_mtdata(d); d = NULL;
118
119 err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr, 0, NULL, NULL);
120 if (err != CL_SUCCESS)
121 {
122 log_error("clEnqueueWriteBuffer failed\n");
123 return -1;
124 }
125
126 err = create_single_kernel_helper(context, &program, &kernel, 1, &conditional_kernel_code, "test_if" );
127 if (err)
128 return -1;
129
130 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
131 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
132 if (err != CL_SUCCESS)
133 {
134 log_error("clSetKernelArgs failed\n");
135 return -1;
136 }
137
138 threads[0] = (unsigned int)num_elements;
139 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
140 if (err != CL_SUCCESS)
141 {
142 log_error("clEnqueueNDRangeKernel failed\n");
143 return -1;
144 }
145
146 err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
147 if (err != CL_SUCCESS)
148 {
149 log_error("clReadArray failed\n");
150 return -1;
151 }
152
153 err = verify_if(input_ptr, output_ptr, num_elements);
154
155 // cleanup
156 clReleaseMemObject(streams[0]);
157 clReleaseMemObject(streams[1]);
158 clReleaseKernel(kernel);
159 clReleaseProgram(program);
160 free(input_ptr);
161 free(output_ptr);
162
163 return err;
164 }
165
166
167