xref: /aosp_15_r20/external/OpenCL-CTS/test_conformance/profiling/writeArray.cpp (revision 6467f958c7de8070b317fc65bcb0f6472e388d82)
1*6467f958SSadaf Ebrahimi //
2*6467f958SSadaf Ebrahimi // Copyright (c) 2017 The Khronos Group Inc.
3*6467f958SSadaf Ebrahimi //
4*6467f958SSadaf Ebrahimi // Licensed under the Apache License, Version 2.0 (the "License");
5*6467f958SSadaf Ebrahimi // you may not use this file except in compliance with the License.
6*6467f958SSadaf Ebrahimi // You may obtain a copy of the License at
7*6467f958SSadaf Ebrahimi //
8*6467f958SSadaf Ebrahimi //    http://www.apache.org/licenses/LICENSE-2.0
9*6467f958SSadaf Ebrahimi //
10*6467f958SSadaf Ebrahimi // Unless required by applicable law or agreed to in writing, software
11*6467f958SSadaf Ebrahimi // distributed under the License is distributed on an "AS IS" BASIS,
12*6467f958SSadaf Ebrahimi // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13*6467f958SSadaf Ebrahimi // See the License for the specific language governing permissions and
14*6467f958SSadaf Ebrahimi // limitations under the License.
15*6467f958SSadaf Ebrahimi //
16*6467f958SSadaf Ebrahimi #include "harness/compat.h"
17*6467f958SSadaf Ebrahimi 
18*6467f958SSadaf Ebrahimi #include <stdio.h>
19*6467f958SSadaf Ebrahimi #include <string.h>
20*6467f958SSadaf Ebrahimi #include <sys/types.h>
21*6467f958SSadaf Ebrahimi #include <sys/stat.h>
22*6467f958SSadaf Ebrahimi 
23*6467f958SSadaf Ebrahimi #include "procs.h"
24*6467f958SSadaf Ebrahimi #include "harness/testHarness.h"
25*6467f958SSadaf Ebrahimi #include "harness/errorHelpers.h"
26*6467f958SSadaf Ebrahimi #include "harness/conversions.h"
27*6467f958SSadaf Ebrahimi 
28*6467f958SSadaf Ebrahimi #ifndef uchar
29*6467f958SSadaf Ebrahimi typedef unsigned char uchar;
30*6467f958SSadaf Ebrahimi #endif
31*6467f958SSadaf Ebrahimi 
32*6467f958SSadaf Ebrahimi #ifndef TestStruct
33*6467f958SSadaf Ebrahimi typedef struct{
34*6467f958SSadaf Ebrahimi     int        a;
35*6467f958SSadaf Ebrahimi     float    b;
36*6467f958SSadaf Ebrahimi } TestStruct;
37*6467f958SSadaf Ebrahimi #endif
38*6467f958SSadaf Ebrahimi 
39*6467f958SSadaf Ebrahimi const char *stream_write_int_kernel_code[] = {
40*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_int(__global int *src, __global int *dst)\n"
41*6467f958SSadaf Ebrahimi     "{\n"
42*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
43*6467f958SSadaf Ebrahimi     "\n"
44*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
45*6467f958SSadaf Ebrahimi     "}\n",
46*6467f958SSadaf Ebrahimi 
47*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_int2(__global int2 *src, __global int2 *dst)\n"
48*6467f958SSadaf Ebrahimi     "{\n"
49*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
50*6467f958SSadaf Ebrahimi     "\n"
51*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
52*6467f958SSadaf Ebrahimi     "}\n",
53*6467f958SSadaf Ebrahimi 
54*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_int4(__global int4 *src, __global int4 *dst)\n"
55*6467f958SSadaf Ebrahimi     "{\n"
56*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
57*6467f958SSadaf Ebrahimi     "\n"
58*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
59*6467f958SSadaf Ebrahimi     "}\n",
60*6467f958SSadaf Ebrahimi 
61*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_int8(__global int8 *src, __global int8 *dst)\n"
62*6467f958SSadaf Ebrahimi     "{\n"
63*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
64*6467f958SSadaf Ebrahimi     "\n"
65*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
66*6467f958SSadaf Ebrahimi     "}\n",
67*6467f958SSadaf Ebrahimi 
68*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_int16(__global int16 *src, __global int16 *dst)\n"
69*6467f958SSadaf Ebrahimi     "{\n"
70*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
71*6467f958SSadaf Ebrahimi     "\n"
72*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
73*6467f958SSadaf Ebrahimi     "}\n" };
74*6467f958SSadaf Ebrahimi 
75*6467f958SSadaf Ebrahimi static const char *int_kernel_name[] = { "test_stream_write_int", "test_stream_write_int2", "test_stream_write_int4", "test_stream_write_int8", "test_stream_write_int16" };
76*6467f958SSadaf Ebrahimi 
77*6467f958SSadaf Ebrahimi 
78*6467f958SSadaf Ebrahimi const char *stream_write_uint_kernel_code[] = {
79*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uint(__global uint *src, __global uint *dst)\n"
80*6467f958SSadaf Ebrahimi     "{\n"
81*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
82*6467f958SSadaf Ebrahimi     "\n"
83*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
84*6467f958SSadaf Ebrahimi     "}\n",
85*6467f958SSadaf Ebrahimi 
86*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uint2(__global uint2 *src, __global uint2 *dst)\n"
87*6467f958SSadaf Ebrahimi     "{\n"
88*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
89*6467f958SSadaf Ebrahimi     "\n"
90*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
91*6467f958SSadaf Ebrahimi     "}\n",
92*6467f958SSadaf Ebrahimi 
93*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uint4(__global uint4 *src, __global uint4 *dst)\n"
94*6467f958SSadaf Ebrahimi     "{\n"
95*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
96*6467f958SSadaf Ebrahimi     "\n"
97*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
98*6467f958SSadaf Ebrahimi     "}\n",
99*6467f958SSadaf Ebrahimi 
100*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uint8(__global uint8 *src, __global uint8 *dst)\n"
101*6467f958SSadaf Ebrahimi     "{\n"
102*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
103*6467f958SSadaf Ebrahimi     "\n"
104*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
105*6467f958SSadaf Ebrahimi     "}\n",
106*6467f958SSadaf Ebrahimi 
107*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uint16(__global uint16 *src, __global uint16 *dst)\n"
108*6467f958SSadaf Ebrahimi     "{\n"
109*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
110*6467f958SSadaf Ebrahimi     "\n"
111*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
112*6467f958SSadaf Ebrahimi     "}\n" };
113*6467f958SSadaf Ebrahimi 
114*6467f958SSadaf Ebrahimi static const char *uint_kernel_name[] = { "test_stream_write_uint", "test_stream_write_uint2", "test_stream_write_uint4", "test_stream_write_uint8", "test_stream_write_uint16" };
115*6467f958SSadaf Ebrahimi 
116*6467f958SSadaf Ebrahimi 
117*6467f958SSadaf Ebrahimi const char *stream_write_ushort_kernel_code[] = {
118*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ushort(__global ushort *src, __global ushort *dst)\n"
119*6467f958SSadaf Ebrahimi     "{\n"
120*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
121*6467f958SSadaf Ebrahimi     "\n"
122*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
123*6467f958SSadaf Ebrahimi     "}\n",
124*6467f958SSadaf Ebrahimi 
125*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ushort2(__global ushort2 *src, __global ushort2 *dst)\n"
126*6467f958SSadaf Ebrahimi     "{\n"
127*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
128*6467f958SSadaf Ebrahimi     "\n"
129*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
130*6467f958SSadaf Ebrahimi     "}\n",
131*6467f958SSadaf Ebrahimi 
132*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ushort4(__global ushort4 *src, __global ushort4 *dst)\n"
133*6467f958SSadaf Ebrahimi     "{\n"
134*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
135*6467f958SSadaf Ebrahimi     "\n"
136*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
137*6467f958SSadaf Ebrahimi     "}\n",
138*6467f958SSadaf Ebrahimi 
139*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ushort8(__global ushort8 *src, __global ushort8 *dst)\n"
140*6467f958SSadaf Ebrahimi     "{\n"
141*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
142*6467f958SSadaf Ebrahimi     "\n"
143*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
144*6467f958SSadaf Ebrahimi     "}\n",
145*6467f958SSadaf Ebrahimi 
146*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ushort16(__global ushort16 *src, __global ushort16 *dst)\n"
147*6467f958SSadaf Ebrahimi     "{\n"
148*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
149*6467f958SSadaf Ebrahimi     "\n"
150*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
151*6467f958SSadaf Ebrahimi     "}\n" };
152*6467f958SSadaf Ebrahimi 
153*6467f958SSadaf Ebrahimi static const char *ushort_kernel_name[] = { "test_stream_write_ushort", "test_stream_write_ushort2", "test_stream_write_ushort4", "test_stream_write_ushort8", "test_stream_write_ushort16" };
154*6467f958SSadaf Ebrahimi 
155*6467f958SSadaf Ebrahimi 
156*6467f958SSadaf Ebrahimi 
157*6467f958SSadaf Ebrahimi const char *stream_write_short_kernel_code[] = {
158*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_short(__global short *src, __global short *dst)\n"
159*6467f958SSadaf Ebrahimi     "{\n"
160*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
161*6467f958SSadaf Ebrahimi     "\n"
162*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
163*6467f958SSadaf Ebrahimi     "}\n",
164*6467f958SSadaf Ebrahimi 
165*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_short2(__global short2 *src, __global short2 *dst)\n"
166*6467f958SSadaf Ebrahimi     "{\n"
167*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
168*6467f958SSadaf Ebrahimi     "\n"
169*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
170*6467f958SSadaf Ebrahimi     "}\n",
171*6467f958SSadaf Ebrahimi 
172*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_short4(__global short4 *src, __global short4 *dst)\n"
173*6467f958SSadaf Ebrahimi     "{\n"
174*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
175*6467f958SSadaf Ebrahimi     "\n"
176*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
177*6467f958SSadaf Ebrahimi     "}\n",
178*6467f958SSadaf Ebrahimi 
179*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_short8(__global short8 *src, __global short8 *dst)\n"
180*6467f958SSadaf Ebrahimi     "{\n"
181*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
182*6467f958SSadaf Ebrahimi     "\n"
183*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
184*6467f958SSadaf Ebrahimi     "}\n",
185*6467f958SSadaf Ebrahimi 
186*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_short16(__global short16 *src, __global short16 *dst)\n"
187*6467f958SSadaf Ebrahimi     "{\n"
188*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
189*6467f958SSadaf Ebrahimi     "\n"
190*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
191*6467f958SSadaf Ebrahimi     "}\n" };
192*6467f958SSadaf Ebrahimi 
193*6467f958SSadaf Ebrahimi static const char *short_kernel_name[] = { "test_stream_write_short", "test_stream_write_short2", "test_stream_write_short4", "test_stream_write_short8", "test_stream_write_short16" };
194*6467f958SSadaf Ebrahimi 
195*6467f958SSadaf Ebrahimi 
196*6467f958SSadaf Ebrahimi const char *stream_write_char_kernel_code[] = {
197*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_char(__global char *src, __global char *dst)\n"
198*6467f958SSadaf Ebrahimi     "{\n"
199*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
200*6467f958SSadaf Ebrahimi     "\n"
201*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
202*6467f958SSadaf Ebrahimi     "}\n",
203*6467f958SSadaf Ebrahimi 
204*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_char2(__global char2 *src, __global char2 *dst)\n"
205*6467f958SSadaf Ebrahimi     "{\n"
206*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
207*6467f958SSadaf Ebrahimi     "\n"
208*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
209*6467f958SSadaf Ebrahimi     "}\n",
210*6467f958SSadaf Ebrahimi 
211*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_char4(__global char4 *src, __global char4 *dst)\n"
212*6467f958SSadaf Ebrahimi     "{\n"
213*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
214*6467f958SSadaf Ebrahimi     "\n"
215*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
216*6467f958SSadaf Ebrahimi     "}\n",
217*6467f958SSadaf Ebrahimi 
218*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_char8(__global char8 *src, __global char8 *dst)\n"
219*6467f958SSadaf Ebrahimi     "{\n"
220*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
221*6467f958SSadaf Ebrahimi     "\n"
222*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
223*6467f958SSadaf Ebrahimi     "}\n",
224*6467f958SSadaf Ebrahimi 
225*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_char16(__global char16 *src, __global char16 *dst)\n"
226*6467f958SSadaf Ebrahimi     "{\n"
227*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
228*6467f958SSadaf Ebrahimi     "\n"
229*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
230*6467f958SSadaf Ebrahimi     "}\n" };
231*6467f958SSadaf Ebrahimi 
232*6467f958SSadaf Ebrahimi static const char *char_kernel_name[] = { "test_stream_write_char", "test_stream_write_char2", "test_stream_write_char4", "test_stream_write_char8", "test_stream_write_char16" };
233*6467f958SSadaf Ebrahimi 
234*6467f958SSadaf Ebrahimi 
235*6467f958SSadaf Ebrahimi const char *stream_write_uchar_kernel_code[] = {
236*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uchar(__global uchar *src, __global uchar *dst)\n"
237*6467f958SSadaf Ebrahimi     "{\n"
238*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
239*6467f958SSadaf Ebrahimi     "\n"
240*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
241*6467f958SSadaf Ebrahimi     "}\n",
242*6467f958SSadaf Ebrahimi 
243*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uchar2(__global uchar2 *src, __global uchar2 *dst)\n"
244*6467f958SSadaf Ebrahimi     "{\n"
245*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
246*6467f958SSadaf Ebrahimi     "\n"
247*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
248*6467f958SSadaf Ebrahimi     "}\n",
249*6467f958SSadaf Ebrahimi 
250*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uchar4(__global uchar4 *src, __global uchar4 *dst)\n"
251*6467f958SSadaf Ebrahimi     "{\n"
252*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
253*6467f958SSadaf Ebrahimi     "\n"
254*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
255*6467f958SSadaf Ebrahimi     "}\n",
256*6467f958SSadaf Ebrahimi 
257*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uchar8(__global uchar8 *src, __global uchar8 *dst)\n"
258*6467f958SSadaf Ebrahimi     "{\n"
259*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
260*6467f958SSadaf Ebrahimi     "\n"
261*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
262*6467f958SSadaf Ebrahimi     "}\n",
263*6467f958SSadaf Ebrahimi 
264*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_uchar16(__global uchar16 *src, __global uchar16 *dst)\n"
265*6467f958SSadaf Ebrahimi     "{\n"
266*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
267*6467f958SSadaf Ebrahimi     "\n"
268*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
269*6467f958SSadaf Ebrahimi     "}\n" };
270*6467f958SSadaf Ebrahimi 
271*6467f958SSadaf Ebrahimi static const char *uchar_kernel_name[] = { "test_stream_write_uchar", "test_stream_write_uchar2", "test_stream_write_uchar4", "test_stream_write_uchar8", "test_stream_write_uchar16" };
272*6467f958SSadaf Ebrahimi 
273*6467f958SSadaf Ebrahimi 
274*6467f958SSadaf Ebrahimi const char *stream_write_float_kernel_code[] = {
275*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_float(__global float *src, __global float *dst)\n"
276*6467f958SSadaf Ebrahimi     "{\n"
277*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
278*6467f958SSadaf Ebrahimi     "\n"
279*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
280*6467f958SSadaf Ebrahimi     "}\n",
281*6467f958SSadaf Ebrahimi 
282*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_float2(__global float2 *src, __global float2 *dst)\n"
283*6467f958SSadaf Ebrahimi     "{\n"
284*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
285*6467f958SSadaf Ebrahimi     "\n"
286*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
287*6467f958SSadaf Ebrahimi     "}\n",
288*6467f958SSadaf Ebrahimi 
289*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_float4(__global float4 *src, __global float4 *dst)\n"
290*6467f958SSadaf Ebrahimi     "{\n"
291*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
292*6467f958SSadaf Ebrahimi     "\n"
293*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
294*6467f958SSadaf Ebrahimi     "}\n",
295*6467f958SSadaf Ebrahimi 
296*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_float8(__global float8 *src, __global float8 *dst)\n"
297*6467f958SSadaf Ebrahimi     "{\n"
298*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
299*6467f958SSadaf Ebrahimi     "\n"
300*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
301*6467f958SSadaf Ebrahimi     "}\n",
302*6467f958SSadaf Ebrahimi 
303*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_float16(__global float16 *src, __global float16 *dst)\n"
304*6467f958SSadaf Ebrahimi     "{\n"
305*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
306*6467f958SSadaf Ebrahimi     "\n"
307*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
308*6467f958SSadaf Ebrahimi     "}\n" };
309*6467f958SSadaf Ebrahimi 
310*6467f958SSadaf Ebrahimi static const char *float_kernel_name[] = { "test_stream_write_float", "test_stream_write_float2", "test_stream_write_float4", "test_stream_write_float8", "test_stream_write_float16" };
311*6467f958SSadaf Ebrahimi 
312*6467f958SSadaf Ebrahimi 
313*6467f958SSadaf Ebrahimi const char *stream_write_half_kernel_code[] = {
314*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_half(__global half *src, __global float *dst)\n"
315*6467f958SSadaf Ebrahimi     "{\n"
316*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
317*6467f958SSadaf Ebrahimi     "\n"
318*6467f958SSadaf Ebrahimi     "    dst[tid] = vload_half( tid * 2, src );\n"
319*6467f958SSadaf Ebrahimi     "}\n",
320*6467f958SSadaf Ebrahimi 
321*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_half2(__global half2 *src, __global float2 *dst)\n"
322*6467f958SSadaf Ebrahimi     "{\n"
323*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
324*6467f958SSadaf Ebrahimi     "\n"
325*6467f958SSadaf Ebrahimi     "    dst[tid] = vload_half2( tid * 2, src );\n"
326*6467f958SSadaf Ebrahimi     "}\n",
327*6467f958SSadaf Ebrahimi 
328*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_half4(__global half4 *src, __global float4 *dst)\n"
329*6467f958SSadaf Ebrahimi     "{\n"
330*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
331*6467f958SSadaf Ebrahimi     "\n"
332*6467f958SSadaf Ebrahimi     "    dst[tid] = vload_half4( tid * 2, src );\n"
333*6467f958SSadaf Ebrahimi     "}\n",
334*6467f958SSadaf Ebrahimi 
335*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_half8(__global half8 *src, __global float8 *dst)\n"
336*6467f958SSadaf Ebrahimi     "{\n"
337*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
338*6467f958SSadaf Ebrahimi     "\n"
339*6467f958SSadaf Ebrahimi     "    dst[tid] = vload_half8( tid * 2, src );\n"
340*6467f958SSadaf Ebrahimi     "}\n",
341*6467f958SSadaf Ebrahimi 
342*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_half16(__global half16 *src, __global float16 *dst)\n"
343*6467f958SSadaf Ebrahimi     "{\n"
344*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
345*6467f958SSadaf Ebrahimi     "\n"
346*6467f958SSadaf Ebrahimi     "    dst[tid] = vload_half16( tid * 2, src );\n"
347*6467f958SSadaf Ebrahimi     "}\n" };
348*6467f958SSadaf Ebrahimi 
349*6467f958SSadaf Ebrahimi static const char *half_kernel_name[] = { "test_stream_write_half", "test_stream_write_half2", "test_stream_write_half4", "test_stream_write_half8", "test_stream_write_half16" };
350*6467f958SSadaf Ebrahimi 
351*6467f958SSadaf Ebrahimi 
352*6467f958SSadaf Ebrahimi const char *stream_write_long_kernel_code[] = {
353*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_long(__global long *src, __global long *dst)\n"
354*6467f958SSadaf Ebrahimi     "{\n"
355*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
356*6467f958SSadaf Ebrahimi     "\n"
357*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
358*6467f958SSadaf Ebrahimi     "}\n",
359*6467f958SSadaf Ebrahimi 
360*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_long2(__global long2 *src, __global long2 *dst)\n"
361*6467f958SSadaf Ebrahimi     "{\n"
362*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
363*6467f958SSadaf Ebrahimi     "\n"
364*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
365*6467f958SSadaf Ebrahimi     "}\n",
366*6467f958SSadaf Ebrahimi 
367*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_long4(__global long4 *src, __global long4 *dst)\n"
368*6467f958SSadaf Ebrahimi     "{\n"
369*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
370*6467f958SSadaf Ebrahimi     "\n"
371*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
372*6467f958SSadaf Ebrahimi     "}\n",
373*6467f958SSadaf Ebrahimi 
374*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_long8(__global long8 *src, __global long8 *dst)\n"
375*6467f958SSadaf Ebrahimi     "{\n"
376*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
377*6467f958SSadaf Ebrahimi     "\n"
378*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
379*6467f958SSadaf Ebrahimi     "}\n",
380*6467f958SSadaf Ebrahimi 
381*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_long16(__global long16 *src, __global long16 *dst)\n"
382*6467f958SSadaf Ebrahimi     "{\n"
383*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
384*6467f958SSadaf Ebrahimi     "\n"
385*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
386*6467f958SSadaf Ebrahimi     "}\n" };
387*6467f958SSadaf Ebrahimi 
388*6467f958SSadaf Ebrahimi static const char *long_kernel_name[] = { "test_stream_write_long", "test_stream_write_long2", "test_stream_write_long4", "test_stream_write_long8", "test_stream_write_long16" };
389*6467f958SSadaf Ebrahimi 
390*6467f958SSadaf Ebrahimi 
391*6467f958SSadaf Ebrahimi const char *stream_write_ulong_kernel_code[] = {
392*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ulong(__global ulong *src, __global ulong *dst)\n"
393*6467f958SSadaf Ebrahimi     "{\n"
394*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
395*6467f958SSadaf Ebrahimi     "\n"
396*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
397*6467f958SSadaf Ebrahimi     "}\n",
398*6467f958SSadaf Ebrahimi 
399*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ulong2(__global ulong2 *src, __global ulong2 *dst)\n"
400*6467f958SSadaf Ebrahimi     "{\n"
401*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
402*6467f958SSadaf Ebrahimi     "\n"
403*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
404*6467f958SSadaf Ebrahimi     "}\n",
405*6467f958SSadaf Ebrahimi 
406*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ulong4(__global ulong4 *src, __global ulong4 *dst)\n"
407*6467f958SSadaf Ebrahimi     "{\n"
408*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
409*6467f958SSadaf Ebrahimi     "\n"
410*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
411*6467f958SSadaf Ebrahimi     "}\n",
412*6467f958SSadaf Ebrahimi 
413*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ulong8(__global ulong8 *src, __global ulong8 *dst)\n"
414*6467f958SSadaf Ebrahimi     "{\n"
415*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
416*6467f958SSadaf Ebrahimi     "\n"
417*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
418*6467f958SSadaf Ebrahimi     "}\n",
419*6467f958SSadaf Ebrahimi 
420*6467f958SSadaf Ebrahimi     "__kernel void test_stream_write_ulong16(__global ulong16 *src, __global ulong16 *dst)\n"
421*6467f958SSadaf Ebrahimi     "{\n"
422*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
423*6467f958SSadaf Ebrahimi     "\n"
424*6467f958SSadaf Ebrahimi     "    dst[tid] = src[tid];\n"
425*6467f958SSadaf Ebrahimi     "}\n" };
426*6467f958SSadaf Ebrahimi 
427*6467f958SSadaf Ebrahimi static const char *ulong_kernel_name[] = { "test_stream_write_ulong", "test_stream_write_ulong2", "test_stream_write_ulong4", "test_stream_write_ulong8", "test_stream_write_ulong16" };
428*6467f958SSadaf Ebrahimi 
429*6467f958SSadaf Ebrahimi 
430*6467f958SSadaf Ebrahimi static const char *stream_write_struct_kernel_code[] = {
431*6467f958SSadaf Ebrahimi     "typedef struct{\n"
432*6467f958SSadaf Ebrahimi     "int    a;\n"
433*6467f958SSadaf Ebrahimi     "float    b;\n"
434*6467f958SSadaf Ebrahimi     "} TestStruct;\n"
435*6467f958SSadaf Ebrahimi     "__kernel void read_write_struct(__global TestStruct *src, __global TestStruct *dst)\n"
436*6467f958SSadaf Ebrahimi     "{\n"
437*6467f958SSadaf Ebrahimi     "    int  tid = get_global_id(0);\n"
438*6467f958SSadaf Ebrahimi     "\n"
439*6467f958SSadaf Ebrahimi     "    dst[tid].a = src[tid].a;\n"
440*6467f958SSadaf Ebrahimi     "     dst[tid].b = src[tid].b;\n"
441*6467f958SSadaf Ebrahimi     "}\n" };
442*6467f958SSadaf Ebrahimi 
443*6467f958SSadaf Ebrahimi static const char *struct_kernel_name[] = { "read_write_struct" };
444*6467f958SSadaf Ebrahimi 
445*6467f958SSadaf Ebrahimi 
verify_write_int(void * ptr1,void * ptr2,int n)446*6467f958SSadaf Ebrahimi static int verify_write_int( void *ptr1, void *ptr2, int n )
447*6467f958SSadaf Ebrahimi {
448*6467f958SSadaf Ebrahimi     int        i;
449*6467f958SSadaf Ebrahimi     int        *inptr = (int *)ptr1;
450*6467f958SSadaf Ebrahimi     int        *outptr = (int *)ptr2;
451*6467f958SSadaf Ebrahimi 
452*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
453*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
454*6467f958SSadaf Ebrahimi             return -1;
455*6467f958SSadaf Ebrahimi     }
456*6467f958SSadaf Ebrahimi 
457*6467f958SSadaf Ebrahimi     return 0;
458*6467f958SSadaf Ebrahimi }
459*6467f958SSadaf Ebrahimi 
460*6467f958SSadaf Ebrahimi 
verify_write_uint(void * ptr1,void * ptr2,int n)461*6467f958SSadaf Ebrahimi static int verify_write_uint( void *ptr1, void *ptr2, int n )
462*6467f958SSadaf Ebrahimi {
463*6467f958SSadaf Ebrahimi     int        i;
464*6467f958SSadaf Ebrahimi     cl_uint    *inptr = (cl_uint *)ptr1;
465*6467f958SSadaf Ebrahimi     cl_uint    *outptr = (cl_uint *)ptr2;
466*6467f958SSadaf Ebrahimi 
467*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
468*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
469*6467f958SSadaf Ebrahimi             return -1;
470*6467f958SSadaf Ebrahimi     }
471*6467f958SSadaf Ebrahimi 
472*6467f958SSadaf Ebrahimi     return 0;
473*6467f958SSadaf Ebrahimi }
474*6467f958SSadaf Ebrahimi 
475*6467f958SSadaf Ebrahimi 
verify_write_short(void * ptr1,void * ptr2,int n)476*6467f958SSadaf Ebrahimi static int verify_write_short( void *ptr1, void *ptr2, int n )
477*6467f958SSadaf Ebrahimi {
478*6467f958SSadaf Ebrahimi     int        i;
479*6467f958SSadaf Ebrahimi     short    *inptr = (short *)ptr1;
480*6467f958SSadaf Ebrahimi     short    *outptr = (short *)ptr2;
481*6467f958SSadaf Ebrahimi 
482*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
483*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
484*6467f958SSadaf Ebrahimi             return -1;
485*6467f958SSadaf Ebrahimi     }
486*6467f958SSadaf Ebrahimi 
487*6467f958SSadaf Ebrahimi     return 0;
488*6467f958SSadaf Ebrahimi }
489*6467f958SSadaf Ebrahimi 
490*6467f958SSadaf Ebrahimi 
verify_write_ushort(void * ptr1,void * ptr2,int n)491*6467f958SSadaf Ebrahimi static int verify_write_ushort( void *ptr1, void *ptr2, int n )
492*6467f958SSadaf Ebrahimi {
493*6467f958SSadaf Ebrahimi     int        i;
494*6467f958SSadaf Ebrahimi     cl_ushort    *inptr = (cl_ushort *)ptr1;
495*6467f958SSadaf Ebrahimi     cl_ushort    *outptr = (cl_ushort *)ptr2;
496*6467f958SSadaf Ebrahimi 
497*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
498*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
499*6467f958SSadaf Ebrahimi             return -1;
500*6467f958SSadaf Ebrahimi     }
501*6467f958SSadaf Ebrahimi 
502*6467f958SSadaf Ebrahimi     return 0;
503*6467f958SSadaf Ebrahimi }
504*6467f958SSadaf Ebrahimi 
505*6467f958SSadaf Ebrahimi 
verify_write_char(void * ptr1,void * ptr2,int n)506*6467f958SSadaf Ebrahimi static int verify_write_char( void *ptr1, void *ptr2, int n )
507*6467f958SSadaf Ebrahimi {
508*6467f958SSadaf Ebrahimi     int        i;
509*6467f958SSadaf Ebrahimi     char    *inptr = (char *)ptr1;
510*6467f958SSadaf Ebrahimi     char    *outptr = (char *)ptr2;
511*6467f958SSadaf Ebrahimi 
512*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
513*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
514*6467f958SSadaf Ebrahimi             return -1;
515*6467f958SSadaf Ebrahimi     }
516*6467f958SSadaf Ebrahimi 
517*6467f958SSadaf Ebrahimi     return 0;
518*6467f958SSadaf Ebrahimi }
519*6467f958SSadaf Ebrahimi 
520*6467f958SSadaf Ebrahimi 
verify_write_uchar(void * ptr1,void * ptr2,int n)521*6467f958SSadaf Ebrahimi static int verify_write_uchar( void *ptr1, void *ptr2, int n )
522*6467f958SSadaf Ebrahimi {
523*6467f958SSadaf Ebrahimi     int        i;
524*6467f958SSadaf Ebrahimi     uchar    *inptr = (uchar *)ptr1;
525*6467f958SSadaf Ebrahimi     uchar    *outptr = (uchar *)ptr2;
526*6467f958SSadaf Ebrahimi 
527*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
528*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
529*6467f958SSadaf Ebrahimi             return -1;
530*6467f958SSadaf Ebrahimi     }
531*6467f958SSadaf Ebrahimi 
532*6467f958SSadaf Ebrahimi     return 0;
533*6467f958SSadaf Ebrahimi }
534*6467f958SSadaf Ebrahimi 
535*6467f958SSadaf Ebrahimi 
verify_write_float(void * ptr1,void * ptr2,int n)536*6467f958SSadaf Ebrahimi static int verify_write_float( void *ptr1, void *ptr2, int n )
537*6467f958SSadaf Ebrahimi {
538*6467f958SSadaf Ebrahimi     int        i;
539*6467f958SSadaf Ebrahimi     float    *inptr = (float *)ptr1;
540*6467f958SSadaf Ebrahimi     float    *outptr = (float *)ptr2;
541*6467f958SSadaf Ebrahimi 
542*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
543*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
544*6467f958SSadaf Ebrahimi             return -1;
545*6467f958SSadaf Ebrahimi     }
546*6467f958SSadaf Ebrahimi 
547*6467f958SSadaf Ebrahimi     return 0;
548*6467f958SSadaf Ebrahimi }
549*6467f958SSadaf Ebrahimi 
550*6467f958SSadaf Ebrahimi 
verify_write_half(void * ptr1,void * ptr2,int n)551*6467f958SSadaf Ebrahimi static int verify_write_half( void *ptr1, void *ptr2, int n )
552*6467f958SSadaf Ebrahimi {
553*6467f958SSadaf Ebrahimi     int        i;
554*6467f958SSadaf Ebrahimi     cl_half *inptr = (cl_half *)ptr1;
555*6467f958SSadaf Ebrahimi     cl_half *outptr = (cl_half *)ptr2;
556*6467f958SSadaf Ebrahimi 
557*6467f958SSadaf Ebrahimi     for( i = 0; i < n; i++ ){
558*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
559*6467f958SSadaf Ebrahimi             return -1;
560*6467f958SSadaf Ebrahimi     }
561*6467f958SSadaf Ebrahimi 
562*6467f958SSadaf Ebrahimi     return 0;
563*6467f958SSadaf Ebrahimi }
564*6467f958SSadaf Ebrahimi 
565*6467f958SSadaf Ebrahimi 
verify_write_long(void * ptr1,void * ptr2,int n)566*6467f958SSadaf Ebrahimi static int verify_write_long( void *ptr1, void *ptr2, int n )
567*6467f958SSadaf Ebrahimi {
568*6467f958SSadaf Ebrahimi     int        i;
569*6467f958SSadaf Ebrahimi     cl_long    *inptr = (cl_long *)ptr1;
570*6467f958SSadaf Ebrahimi     cl_long    *outptr = (cl_long *)ptr2;
571*6467f958SSadaf Ebrahimi 
572*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
573*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
574*6467f958SSadaf Ebrahimi             return -1;
575*6467f958SSadaf Ebrahimi     }
576*6467f958SSadaf Ebrahimi 
577*6467f958SSadaf Ebrahimi     return 0;
578*6467f958SSadaf Ebrahimi }
579*6467f958SSadaf Ebrahimi 
580*6467f958SSadaf Ebrahimi 
verify_write_ulong(void * ptr1,void * ptr2,int n)581*6467f958SSadaf Ebrahimi static int verify_write_ulong( void *ptr1, void *ptr2, int n )
582*6467f958SSadaf Ebrahimi {
583*6467f958SSadaf Ebrahimi     int        i;
584*6467f958SSadaf Ebrahimi     cl_ulong    *inptr = (cl_ulong *)ptr1;
585*6467f958SSadaf Ebrahimi     cl_ulong    *outptr = (cl_ulong *)ptr2;
586*6467f958SSadaf Ebrahimi 
587*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
588*6467f958SSadaf Ebrahimi         if( outptr[i] != inptr[i] )
589*6467f958SSadaf Ebrahimi             return -1;
590*6467f958SSadaf Ebrahimi     }
591*6467f958SSadaf Ebrahimi 
592*6467f958SSadaf Ebrahimi     return 0;
593*6467f958SSadaf Ebrahimi }
594*6467f958SSadaf Ebrahimi 
595*6467f958SSadaf Ebrahimi 
verify_write_struct(void * ptr1,void * ptr2,int n)596*6467f958SSadaf Ebrahimi static int verify_write_struct( void *ptr1, void *ptr2, int n )
597*6467f958SSadaf Ebrahimi {
598*6467f958SSadaf Ebrahimi     int            i;
599*6467f958SSadaf Ebrahimi     TestStruct    *inptr = (TestStruct *)ptr1;
600*6467f958SSadaf Ebrahimi     TestStruct    *outptr = (TestStruct *)ptr2;
601*6467f958SSadaf Ebrahimi 
602*6467f958SSadaf Ebrahimi     for (i=0; i<n; i++){
603*6467f958SSadaf Ebrahimi         if( ( outptr[i].a != inptr[i].a ) || ( outptr[i].b != outptr[i].b ) )
604*6467f958SSadaf Ebrahimi             return -1;
605*6467f958SSadaf Ebrahimi     }
606*6467f958SSadaf Ebrahimi 
607*6467f958SSadaf Ebrahimi     return 0;
608*6467f958SSadaf Ebrahimi }
609*6467f958SSadaf Ebrahimi 
610*6467f958SSadaf Ebrahimi 
test_stream_write(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements,size_t size,const char * type,int loops,void * inptr[5],const char * kernelCode[],const char * kernelName[],int (* fn)(void *,void *,int),MTdata d)611*6467f958SSadaf Ebrahimi int test_stream_write( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements, size_t size, const char *type, int loops,
612*6467f958SSadaf Ebrahimi                       void *inptr[5], const char *kernelCode[], const char *kernelName[], int (*fn)(void *,void *,int), MTdata d )
613*6467f958SSadaf Ebrahimi {
614*6467f958SSadaf Ebrahimi     cl_mem            streams[10];
615*6467f958SSadaf Ebrahimi     void            *outptr[5];
616*6467f958SSadaf Ebrahimi     cl_program        program[5];
617*6467f958SSadaf Ebrahimi     cl_kernel        kernel[5];
618*6467f958SSadaf Ebrahimi     cl_event        writeEvent;
619*6467f958SSadaf Ebrahimi     cl_ulong    queueStart, submitStart, writeStart, writeEnd;
620*6467f958SSadaf Ebrahimi     size_t            ptrSizes[5], outPtrSizes[5];
621*6467f958SSadaf Ebrahimi     size_t            threads[1];
622*6467f958SSadaf Ebrahimi     int                err, err_count = 0;
623*6467f958SSadaf Ebrahimi     int                i, ii;
624*6467f958SSadaf Ebrahimi 
625*6467f958SSadaf Ebrahimi     threads[0] = (size_t)num_elements;
626*6467f958SSadaf Ebrahimi 
627*6467f958SSadaf Ebrahimi     ptrSizes[0] = size;
628*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
629*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
630*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
631*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
632*6467f958SSadaf Ebrahimi 
633*6467f958SSadaf Ebrahimi     loops = ( loops < 5 ? loops : 5 );
634*6467f958SSadaf Ebrahimi 
635*6467f958SSadaf Ebrahimi     for( i = 0; i < loops; i++ )
636*6467f958SSadaf Ebrahimi     {
637*6467f958SSadaf Ebrahimi         outPtrSizes[i] = ptrSizes[i];
638*6467f958SSadaf Ebrahimi     }
639*6467f958SSadaf Ebrahimi 
640*6467f958SSadaf Ebrahimi     for( i = 0; i < loops; i++ ){
641*6467f958SSadaf Ebrahimi         ii = i << 1;
642*6467f958SSadaf Ebrahimi         streams[ii] = clCreateBuffer(context, CL_MEM_READ_WRITE,
643*6467f958SSadaf Ebrahimi                                      ptrSizes[i] * num_elements, NULL, &err);
644*6467f958SSadaf Ebrahimi         if( ! streams[ii] ){
645*6467f958SSadaf Ebrahimi             free( outptr[i] );
646*6467f958SSadaf Ebrahimi             log_error( " clCreateBuffer failed\n" );
647*6467f958SSadaf Ebrahimi             return -1;
648*6467f958SSadaf Ebrahimi         }
649*6467f958SSadaf Ebrahimi         if( ! strcmp( type, "half" ) ){
650*6467f958SSadaf Ebrahimi             outptr[i] = malloc( outPtrSizes[i] * num_elements * 2 );
651*6467f958SSadaf Ebrahimi             streams[ii + 1] =
652*6467f958SSadaf Ebrahimi                 clCreateBuffer(context, CL_MEM_READ_WRITE,
653*6467f958SSadaf Ebrahimi                                outPtrSizes[i] * 2 * num_elements, NULL, &err);
654*6467f958SSadaf Ebrahimi         }
655*6467f958SSadaf Ebrahimi         else{
656*6467f958SSadaf Ebrahimi             outptr[i] = malloc( outPtrSizes[i] * num_elements );
657*6467f958SSadaf Ebrahimi             streams[ii + 1] =
658*6467f958SSadaf Ebrahimi                 clCreateBuffer(context, CL_MEM_READ_WRITE,
659*6467f958SSadaf Ebrahimi                                outPtrSizes[i] * num_elements, NULL, &err);
660*6467f958SSadaf Ebrahimi         }
661*6467f958SSadaf Ebrahimi         if( ! streams[ii+1] ){
662*6467f958SSadaf Ebrahimi             clReleaseMemObject(streams[ii]);
663*6467f958SSadaf Ebrahimi             free( outptr[i] );
664*6467f958SSadaf Ebrahimi             log_error( " clCreateBuffer failed\n" );
665*6467f958SSadaf Ebrahimi             return -1;
666*6467f958SSadaf Ebrahimi         }
667*6467f958SSadaf Ebrahimi 
668*6467f958SSadaf Ebrahimi         err = clEnqueueWriteBuffer( queue, streams[ii], false, 0, ptrSizes[i]*num_elements, inptr[i], 0, NULL, &writeEvent );
669*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
670*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
671*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
672*6467f958SSadaf Ebrahimi             free( outptr[i] );
673*6467f958SSadaf Ebrahimi             print_error( err, " clWriteArray failed" );
674*6467f958SSadaf Ebrahimi             return -1;
675*6467f958SSadaf Ebrahimi         }
676*6467f958SSadaf Ebrahimi 
677*6467f958SSadaf Ebrahimi         // This synchronization point is needed in order to assume the data is valid.
678*6467f958SSadaf Ebrahimi         // Getting profiling information is not a synchronization point.
679*6467f958SSadaf Ebrahimi         err = clWaitForEvents( 1, &writeEvent );
680*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS )
681*6467f958SSadaf Ebrahimi         {
682*6467f958SSadaf Ebrahimi             print_error( err, "Unable to wait for event completion" );
683*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
684*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
685*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
686*6467f958SSadaf Ebrahimi             free( outptr[i] );
687*6467f958SSadaf Ebrahimi             return -1;
688*6467f958SSadaf Ebrahimi         }
689*6467f958SSadaf Ebrahimi 
690*6467f958SSadaf Ebrahimi         // test profiling
691*6467f958SSadaf Ebrahimi         while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
692*6467f958SSadaf Ebrahimi               CL_PROFILING_INFO_NOT_AVAILABLE );
693*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
694*6467f958SSadaf Ebrahimi             print_error( err, "clGetEventProfilingInfo failed" );
695*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
696*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
697*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
698*6467f958SSadaf Ebrahimi             free( outptr[i] );
699*6467f958SSadaf Ebrahimi             return -1;
700*6467f958SSadaf Ebrahimi         }
701*6467f958SSadaf Ebrahimi 
702*6467f958SSadaf Ebrahimi         while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL ) ) ==
703*6467f958SSadaf Ebrahimi               CL_PROFILING_INFO_NOT_AVAILABLE );
704*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
705*6467f958SSadaf Ebrahimi             print_error( err, "clGetEventProfilingInfo failed" );
706*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
707*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
708*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
709*6467f958SSadaf Ebrahimi             free( outptr[i] );
710*6467f958SSadaf Ebrahimi             return -1;
711*6467f958SSadaf Ebrahimi         }
712*6467f958SSadaf Ebrahimi 
713*6467f958SSadaf Ebrahimi         err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
714*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
715*6467f958SSadaf Ebrahimi             print_error( err, "clGetEventProfilingInfo failed" );
716*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
717*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
718*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
719*6467f958SSadaf Ebrahimi             free( outptr[i] );
720*6467f958SSadaf Ebrahimi             return -1;
721*6467f958SSadaf Ebrahimi         }
722*6467f958SSadaf Ebrahimi 
723*6467f958SSadaf Ebrahimi         err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
724*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
725*6467f958SSadaf Ebrahimi             print_error( err, "clGetEventProfilingInfo failed" );
726*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
727*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
728*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
729*6467f958SSadaf Ebrahimi             free( outptr[i] );
730*6467f958SSadaf Ebrahimi             return -1;
731*6467f958SSadaf Ebrahimi         }
732*6467f958SSadaf Ebrahimi 
733*6467f958SSadaf Ebrahimi 
734*6467f958SSadaf Ebrahimi         err = create_single_kernel_helper( context, &program[i], &kernel[i], 1, &kernelCode[i], kernelName[i] );
735*6467f958SSadaf Ebrahimi         if( err ){
736*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
737*6467f958SSadaf Ebrahimi             clReleaseMemObject(streams[ii]);
738*6467f958SSadaf Ebrahimi             clReleaseMemObject(streams[ii+1]);
739*6467f958SSadaf Ebrahimi             free( outptr[i] );
740*6467f958SSadaf Ebrahimi             log_error( " Error creating program for %s\n", type );
741*6467f958SSadaf Ebrahimi             return -1;
742*6467f958SSadaf Ebrahimi         }
743*6467f958SSadaf Ebrahimi 
744*6467f958SSadaf Ebrahimi         err = clSetKernelArg( kernel[i], 0, sizeof( cl_mem ), (void *)&streams[ii] );
745*6467f958SSadaf Ebrahimi         err |= clSetKernelArg( kernel[i], 1, sizeof( cl_mem ), (void *)&streams[ii+1] );
746*6467f958SSadaf Ebrahimi         if (err != CL_SUCCESS){
747*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
748*6467f958SSadaf Ebrahimi             clReleaseKernel( kernel[i] );
749*6467f958SSadaf Ebrahimi             clReleaseProgram( program[i] );
750*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
751*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
752*6467f958SSadaf Ebrahimi             free( outptr[i] );
753*6467f958SSadaf Ebrahimi             print_error( err, " clSetKernelArg failed" );
754*6467f958SSadaf Ebrahimi             return -1;
755*6467f958SSadaf Ebrahimi         }
756*6467f958SSadaf Ebrahimi 
757*6467f958SSadaf Ebrahimi         err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
758*6467f958SSadaf Ebrahimi 
759*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
760*6467f958SSadaf Ebrahimi             print_error( err, " clEnqueueNDRangeKernel failed" );
761*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
762*6467f958SSadaf Ebrahimi             clReleaseKernel( kernel[i] );
763*6467f958SSadaf Ebrahimi             clReleaseProgram( program[i] );
764*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
765*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
766*6467f958SSadaf Ebrahimi             free( outptr[i] );
767*6467f958SSadaf Ebrahimi             return -1;
768*6467f958SSadaf Ebrahimi         }
769*6467f958SSadaf Ebrahimi 
770*6467f958SSadaf Ebrahimi         if( ! strcmp( type, "half" ) ){
771*6467f958SSadaf Ebrahimi             err = clEnqueueReadBuffer( queue, streams[ii+1], true, 0, outPtrSizes[i]*num_elements, outptr[i], 0, NULL, NULL );
772*6467f958SSadaf Ebrahimi         }
773*6467f958SSadaf Ebrahimi         else{
774*6467f958SSadaf Ebrahimi             err = clEnqueueReadBuffer( queue, streams[ii+1], true, 0, outPtrSizes[i]*num_elements, outptr[i], 0, NULL, NULL );
775*6467f958SSadaf Ebrahimi         }
776*6467f958SSadaf Ebrahimi         if( err != CL_SUCCESS ){
777*6467f958SSadaf Ebrahimi             clReleaseEvent(writeEvent);
778*6467f958SSadaf Ebrahimi             clReleaseKernel( kernel[i] );
779*6467f958SSadaf Ebrahimi             clReleaseProgram( program[i] );
780*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii] );
781*6467f958SSadaf Ebrahimi             clReleaseMemObject( streams[ii+1] );
782*6467f958SSadaf Ebrahimi             free( outptr[i] );
783*6467f958SSadaf Ebrahimi             print_error( err, " clEnqueueReadBuffer failed" );
784*6467f958SSadaf Ebrahimi             return -1;
785*6467f958SSadaf Ebrahimi         }
786*6467f958SSadaf Ebrahimi 
787*6467f958SSadaf Ebrahimi         char *inP = (char *)inptr[i];
788*6467f958SSadaf Ebrahimi         char *outP = (char *)outptr[i];
789*6467f958SSadaf Ebrahimi         int err2 = 0;
790*6467f958SSadaf Ebrahimi         for( size_t p = 0; p < (size_t)num_elements; p++ )
791*6467f958SSadaf Ebrahimi         {
792*6467f958SSadaf Ebrahimi             if( fn( inP, outP, (int)(ptrSizes[i] / ptrSizes[0]) ) )
793*6467f958SSadaf Ebrahimi             {
794*6467f958SSadaf Ebrahimi                 log_error( " %s%d data failed to verify\n", type, 1<<i );
795*6467f958SSadaf Ebrahimi                 err2 = -1;
796*6467f958SSadaf Ebrahimi                 err_count++;
797*6467f958SSadaf Ebrahimi             }
798*6467f958SSadaf Ebrahimi             inP += ptrSizes[i];
799*6467f958SSadaf Ebrahimi             outP += outPtrSizes[i];
800*6467f958SSadaf Ebrahimi         }
801*6467f958SSadaf Ebrahimi         if( !err2 )
802*6467f958SSadaf Ebrahimi         {
803*6467f958SSadaf Ebrahimi             log_info(" %s%d data verified\n", type, 1 << i);
804*6467f958SSadaf Ebrahimi         }
805*6467f958SSadaf Ebrahimi         err = err2;
806*6467f958SSadaf Ebrahimi 
807*6467f958SSadaf Ebrahimi         if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
808*6467f958SSadaf Ebrahimi             err_count++;
809*6467f958SSadaf Ebrahimi 
810*6467f958SSadaf Ebrahimi         // cleanup
811*6467f958SSadaf Ebrahimi         clReleaseEvent(writeEvent);
812*6467f958SSadaf Ebrahimi         clReleaseKernel( kernel[i] );
813*6467f958SSadaf Ebrahimi         clReleaseProgram( program[i] );
814*6467f958SSadaf Ebrahimi         clReleaseMemObject( streams[ii] );
815*6467f958SSadaf Ebrahimi         clReleaseMemObject( streams[ii+1] );
816*6467f958SSadaf Ebrahimi         free( outptr[i] );
817*6467f958SSadaf Ebrahimi     }
818*6467f958SSadaf Ebrahimi 
819*6467f958SSadaf Ebrahimi     return err_count;
820*6467f958SSadaf Ebrahimi 
821*6467f958SSadaf Ebrahimi }    // end test_stream_write()
822*6467f958SSadaf Ebrahimi 
823*6467f958SSadaf Ebrahimi 
test_write_array_int(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)824*6467f958SSadaf Ebrahimi int test_write_array_int( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
825*6467f958SSadaf Ebrahimi {
826*6467f958SSadaf Ebrahimi     int    *inptr[5];
827*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
828*6467f958SSadaf Ebrahimi     int        i, j, err;
829*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
830*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
831*6467f958SSadaf Ebrahimi     foo = verify_write_int;
832*6467f958SSadaf Ebrahimi 
833*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_int);
834*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
835*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
836*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
837*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
838*6467f958SSadaf Ebrahimi 
839*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
840*6467f958SSadaf Ebrahimi         inptr[i] = (int *)malloc(ptrSizes[i] * num_elements);
841*6467f958SSadaf Ebrahimi 
842*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
843*6467f958SSadaf Ebrahimi             inptr[i][j] = genrand_int32(d);
844*6467f958SSadaf Ebrahimi     }
845*6467f958SSadaf Ebrahimi 
846*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_int ), "int", 5, (void**)inptr,
847*6467f958SSadaf Ebrahimi                             stream_write_int_kernel_code, int_kernel_name, foo, d );
848*6467f958SSadaf Ebrahimi 
849*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
850*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
851*6467f958SSadaf Ebrahimi     }
852*6467f958SSadaf Ebrahimi 
853*6467f958SSadaf Ebrahimi     free_mtdata(d);
854*6467f958SSadaf Ebrahimi 
855*6467f958SSadaf Ebrahimi     return err;
856*6467f958SSadaf Ebrahimi 
857*6467f958SSadaf Ebrahimi }    // end write_int_array()
858*6467f958SSadaf Ebrahimi 
859*6467f958SSadaf Ebrahimi 
test_write_array_uint(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)860*6467f958SSadaf Ebrahimi int test_write_array_uint( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
861*6467f958SSadaf Ebrahimi {
862*6467f958SSadaf Ebrahimi     cl_uint    *inptr[5];
863*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
864*6467f958SSadaf Ebrahimi     int        i, j, err;
865*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
866*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
867*6467f958SSadaf Ebrahimi     foo = verify_write_uint;
868*6467f958SSadaf Ebrahimi 
869*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_uint);
870*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
871*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
872*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
873*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
874*6467f958SSadaf Ebrahimi 
875*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
876*6467f958SSadaf Ebrahimi         inptr[i] = (cl_uint *)malloc(ptrSizes[i] * num_elements);
877*6467f958SSadaf Ebrahimi 
878*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
879*6467f958SSadaf Ebrahimi             inptr[i][j] = genrand_int32(d);
880*6467f958SSadaf Ebrahimi     }
881*6467f958SSadaf Ebrahimi 
882*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_uint ), "uint", 5, (void **)inptr,
883*6467f958SSadaf Ebrahimi                             stream_write_uint_kernel_code, uint_kernel_name, foo, d );
884*6467f958SSadaf Ebrahimi 
885*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
886*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
887*6467f958SSadaf Ebrahimi     }
888*6467f958SSadaf Ebrahimi 
889*6467f958SSadaf Ebrahimi     free_mtdata(d);
890*6467f958SSadaf Ebrahimi     return err;
891*6467f958SSadaf Ebrahimi 
892*6467f958SSadaf Ebrahimi }    // end write_uint_array()
893*6467f958SSadaf Ebrahimi 
894*6467f958SSadaf Ebrahimi 
test_write_array_short(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)895*6467f958SSadaf Ebrahimi int test_write_array_short( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
896*6467f958SSadaf Ebrahimi {
897*6467f958SSadaf Ebrahimi     short    *inptr[5];
898*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
899*6467f958SSadaf Ebrahimi     int        i, j, err;
900*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
901*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
902*6467f958SSadaf Ebrahimi     foo = verify_write_short;
903*6467f958SSadaf Ebrahimi 
904*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_short);
905*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
906*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
907*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
908*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
909*6467f958SSadaf Ebrahimi 
910*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
911*6467f958SSadaf Ebrahimi         inptr[i] = (short *)malloc(ptrSizes[i] * num_elements);
912*6467f958SSadaf Ebrahimi 
913*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
914*6467f958SSadaf Ebrahimi             inptr[i][j] = (short)genrand_int32(d);
915*6467f958SSadaf Ebrahimi     }
916*6467f958SSadaf Ebrahimi 
917*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_short ), "short", 5, (void **)inptr,
918*6467f958SSadaf Ebrahimi                             stream_write_short_kernel_code, short_kernel_name, foo, d );
919*6467f958SSadaf Ebrahimi 
920*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
921*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
922*6467f958SSadaf Ebrahimi     }
923*6467f958SSadaf Ebrahimi 
924*6467f958SSadaf Ebrahimi     free_mtdata(d);
925*6467f958SSadaf Ebrahimi     return err;
926*6467f958SSadaf Ebrahimi 
927*6467f958SSadaf Ebrahimi }    // end write_short_array()
928*6467f958SSadaf Ebrahimi 
929*6467f958SSadaf Ebrahimi 
test_write_array_ushort(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)930*6467f958SSadaf Ebrahimi int test_write_array_ushort( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
931*6467f958SSadaf Ebrahimi {
932*6467f958SSadaf Ebrahimi     cl_ushort    *inptr[5];
933*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
934*6467f958SSadaf Ebrahimi     int        i, j, err;
935*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
936*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
937*6467f958SSadaf Ebrahimi     foo = verify_write_ushort;
938*6467f958SSadaf Ebrahimi 
939*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_ushort);
940*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
941*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
942*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
943*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
944*6467f958SSadaf Ebrahimi 
945*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
946*6467f958SSadaf Ebrahimi         inptr[i] = (cl_ushort *)malloc(ptrSizes[i] * num_elements);
947*6467f958SSadaf Ebrahimi 
948*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
949*6467f958SSadaf Ebrahimi             inptr[i][j] = (cl_ushort)genrand_int32(d);
950*6467f958SSadaf Ebrahimi     }
951*6467f958SSadaf Ebrahimi 
952*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_ushort ), "ushort", 5, (void **)inptr,
953*6467f958SSadaf Ebrahimi                             stream_write_ushort_kernel_code, ushort_kernel_name, foo, d );
954*6467f958SSadaf Ebrahimi 
955*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
956*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
957*6467f958SSadaf Ebrahimi     }
958*6467f958SSadaf Ebrahimi 
959*6467f958SSadaf Ebrahimi     free_mtdata(d);
960*6467f958SSadaf Ebrahimi     return err;
961*6467f958SSadaf Ebrahimi 
962*6467f958SSadaf Ebrahimi }    // end write_ushort_array()
963*6467f958SSadaf Ebrahimi 
964*6467f958SSadaf Ebrahimi 
test_write_array_char(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)965*6467f958SSadaf Ebrahimi int test_write_array_char( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
966*6467f958SSadaf Ebrahimi {
967*6467f958SSadaf Ebrahimi     char    *inptr[5];
968*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
969*6467f958SSadaf Ebrahimi     int        i, j, err;
970*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
971*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
972*6467f958SSadaf Ebrahimi     foo = verify_write_char;
973*6467f958SSadaf Ebrahimi 
974*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_char);
975*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
976*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
977*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
978*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
979*6467f958SSadaf Ebrahimi 
980*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
981*6467f958SSadaf Ebrahimi         inptr[i] = (char *)malloc(ptrSizes[i] * num_elements);
982*6467f958SSadaf Ebrahimi 
983*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
984*6467f958SSadaf Ebrahimi             inptr[i][j] = (char)genrand_int32(d);
985*6467f958SSadaf Ebrahimi     }
986*6467f958SSadaf Ebrahimi 
987*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_char ), "char", 5, (void **)inptr,
988*6467f958SSadaf Ebrahimi                             stream_write_char_kernel_code, char_kernel_name, foo, d );
989*6467f958SSadaf Ebrahimi 
990*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
991*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
992*6467f958SSadaf Ebrahimi     }
993*6467f958SSadaf Ebrahimi 
994*6467f958SSadaf Ebrahimi     free_mtdata(d);
995*6467f958SSadaf Ebrahimi     return err;
996*6467f958SSadaf Ebrahimi 
997*6467f958SSadaf Ebrahimi }    // end write_char_array()
998*6467f958SSadaf Ebrahimi 
999*6467f958SSadaf Ebrahimi 
test_write_array_uchar(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1000*6467f958SSadaf Ebrahimi int test_write_array_uchar( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1001*6467f958SSadaf Ebrahimi {
1002*6467f958SSadaf Ebrahimi     uchar    *inptr[5];
1003*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
1004*6467f958SSadaf Ebrahimi     int        i, j, err;
1005*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1006*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1007*6467f958SSadaf Ebrahimi     foo = verify_write_uchar;
1008*6467f958SSadaf Ebrahimi 
1009*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_uchar);
1010*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
1011*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
1012*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
1013*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
1014*6467f958SSadaf Ebrahimi 
1015*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1016*6467f958SSadaf Ebrahimi         inptr[i] = (uchar *)malloc(ptrSizes[i] * num_elements);
1017*6467f958SSadaf Ebrahimi 
1018*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1019*6467f958SSadaf Ebrahimi             inptr[i][j] = (uchar)genrand_int32(d);
1020*6467f958SSadaf Ebrahimi     }
1021*6467f958SSadaf Ebrahimi 
1022*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_uchar ), "uchar", 5, (void **)inptr,
1023*6467f958SSadaf Ebrahimi                             stream_write_uchar_kernel_code, uchar_kernel_name, foo, d );
1024*6467f958SSadaf Ebrahimi 
1025*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1026*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
1027*6467f958SSadaf Ebrahimi     }
1028*6467f958SSadaf Ebrahimi 
1029*6467f958SSadaf Ebrahimi     free_mtdata(d);
1030*6467f958SSadaf Ebrahimi     return err;
1031*6467f958SSadaf Ebrahimi 
1032*6467f958SSadaf Ebrahimi }    // end write_uchar_array()
1033*6467f958SSadaf Ebrahimi 
1034*6467f958SSadaf Ebrahimi 
test_write_array_float(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1035*6467f958SSadaf Ebrahimi int test_write_array_float( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1036*6467f958SSadaf Ebrahimi {
1037*6467f958SSadaf Ebrahimi     float    *inptr[5];
1038*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
1039*6467f958SSadaf Ebrahimi     int        i, j, err;
1040*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1041*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1042*6467f958SSadaf Ebrahimi     foo = verify_write_float;
1043*6467f958SSadaf Ebrahimi 
1044*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_float);
1045*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
1046*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
1047*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
1048*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
1049*6467f958SSadaf Ebrahimi 
1050*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1051*6467f958SSadaf Ebrahimi         inptr[i] = (float *)malloc(ptrSizes[i] * num_elements);
1052*6467f958SSadaf Ebrahimi 
1053*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1054*6467f958SSadaf Ebrahimi             inptr[i][j] = get_random_float( -FLT_MAX, FLT_MAX, d );
1055*6467f958SSadaf Ebrahimi     }
1056*6467f958SSadaf Ebrahimi 
1057*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_float ), "float", 5, (void **)inptr,
1058*6467f958SSadaf Ebrahimi                             stream_write_float_kernel_code, float_kernel_name, foo, d );
1059*6467f958SSadaf Ebrahimi 
1060*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1061*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
1062*6467f958SSadaf Ebrahimi     }
1063*6467f958SSadaf Ebrahimi 
1064*6467f958SSadaf Ebrahimi     free_mtdata(d);
1065*6467f958SSadaf Ebrahimi     return err;
1066*6467f958SSadaf Ebrahimi 
1067*6467f958SSadaf Ebrahimi }    // end write_float_array()
1068*6467f958SSadaf Ebrahimi 
1069*6467f958SSadaf Ebrahimi 
test_write_array_half(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1070*6467f958SSadaf Ebrahimi int test_write_array_half( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1071*6467f958SSadaf Ebrahimi {
1072*6467f958SSadaf Ebrahimi     float    *inptr[5];
1073*6467f958SSadaf Ebrahimi     size_t    ptrSizes[5];
1074*6467f958SSadaf Ebrahimi     int        i, j, err;
1075*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1076*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1077*6467f958SSadaf Ebrahimi     foo = verify_write_half;
1078*6467f958SSadaf Ebrahimi 
1079*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof( cl_half );
1080*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
1081*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
1082*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
1083*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
1084*6467f958SSadaf Ebrahimi 
1085*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1086*6467f958SSadaf Ebrahimi         inptr[i] = (float *)malloc(ptrSizes[i] * num_elements);
1087*6467f958SSadaf Ebrahimi 
1088*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ( ptrSizes[0] * 2 ); j++ )
1089*6467f958SSadaf Ebrahimi             inptr[i][j] = get_random_float( -FLT_MAX, FLT_MAX, d );
1090*6467f958SSadaf Ebrahimi     }
1091*6467f958SSadaf Ebrahimi 
1092*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_half ), "half", 5, (void **)inptr,
1093*6467f958SSadaf Ebrahimi                             stream_write_half_kernel_code, half_kernel_name, foo, d );
1094*6467f958SSadaf Ebrahimi 
1095*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1096*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
1097*6467f958SSadaf Ebrahimi     }
1098*6467f958SSadaf Ebrahimi 
1099*6467f958SSadaf Ebrahimi     free_mtdata(d);
1100*6467f958SSadaf Ebrahimi     return err;
1101*6467f958SSadaf Ebrahimi 
1102*6467f958SSadaf Ebrahimi }    // end write_half_array()
1103*6467f958SSadaf Ebrahimi 
1104*6467f958SSadaf Ebrahimi 
test_write_array_long(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1105*6467f958SSadaf Ebrahimi int test_write_array_long( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1106*6467f958SSadaf Ebrahimi {
1107*6467f958SSadaf Ebrahimi     cl_long    *inptr[5];
1108*6467f958SSadaf Ebrahimi     size_t        ptrSizes[5];
1109*6467f958SSadaf Ebrahimi     int            i, j, err;
1110*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1111*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1112*6467f958SSadaf Ebrahimi     foo = verify_write_long;
1113*6467f958SSadaf Ebrahimi 
1114*6467f958SSadaf Ebrahimi     if (!gHasLong)
1115*6467f958SSadaf Ebrahimi     {
1116*6467f958SSadaf Ebrahimi         log_info("write_long_array: Long types unsupported, skipping.");
1117*6467f958SSadaf Ebrahimi         return CL_SUCCESS;
1118*6467f958SSadaf Ebrahimi     }
1119*6467f958SSadaf Ebrahimi 
1120*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_long);
1121*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
1122*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
1123*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
1124*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
1125*6467f958SSadaf Ebrahimi 
1126*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1127*6467f958SSadaf Ebrahimi         inptr[i] = (cl_long *)malloc(ptrSizes[i] * num_elements);
1128*6467f958SSadaf Ebrahimi 
1129*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1130*6467f958SSadaf Ebrahimi             inptr[i][j] = (cl_long) genrand_int32(d) ^ ((cl_long) genrand_int32(d) << 32);
1131*6467f958SSadaf Ebrahimi     }
1132*6467f958SSadaf Ebrahimi 
1133*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_long ), "cl_long", 5, (void **)inptr,
1134*6467f958SSadaf Ebrahimi                             stream_write_long_kernel_code, long_kernel_name, foo, d );
1135*6467f958SSadaf Ebrahimi 
1136*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1137*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
1138*6467f958SSadaf Ebrahimi     }
1139*6467f958SSadaf Ebrahimi 
1140*6467f958SSadaf Ebrahimi     free_mtdata(d);
1141*6467f958SSadaf Ebrahimi     return err;
1142*6467f958SSadaf Ebrahimi 
1143*6467f958SSadaf Ebrahimi }    // end write_long_array()
1144*6467f958SSadaf Ebrahimi 
1145*6467f958SSadaf Ebrahimi 
test_write_array_ulong(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1146*6467f958SSadaf Ebrahimi int test_write_array_ulong( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1147*6467f958SSadaf Ebrahimi {
1148*6467f958SSadaf Ebrahimi     cl_ulong    *inptr[5];
1149*6467f958SSadaf Ebrahimi     size_t                ptrSizes[5];
1150*6467f958SSadaf Ebrahimi     int                    i, j, err;
1151*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1152*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1153*6467f958SSadaf Ebrahimi     foo = verify_write_ulong;
1154*6467f958SSadaf Ebrahimi 
1155*6467f958SSadaf Ebrahimi     if (!gHasLong)
1156*6467f958SSadaf Ebrahimi     {
1157*6467f958SSadaf Ebrahimi         log_info("write_long_array: Long types unsupported, skipping.");
1158*6467f958SSadaf Ebrahimi         return CL_SUCCESS;
1159*6467f958SSadaf Ebrahimi     }
1160*6467f958SSadaf Ebrahimi 
1161*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof(cl_ulong);
1162*6467f958SSadaf Ebrahimi     ptrSizes[1] = ptrSizes[0] << 1;
1163*6467f958SSadaf Ebrahimi     ptrSizes[2] = ptrSizes[1] << 1;
1164*6467f958SSadaf Ebrahimi     ptrSizes[3] = ptrSizes[2] << 1;
1165*6467f958SSadaf Ebrahimi     ptrSizes[4] = ptrSizes[3] << 1;
1166*6467f958SSadaf Ebrahimi 
1167*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1168*6467f958SSadaf Ebrahimi         inptr[i] = (cl_ulong *)malloc(ptrSizes[i] * num_elements);
1169*6467f958SSadaf Ebrahimi 
1170*6467f958SSadaf Ebrahimi         for( j = 0; (unsigned int)j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1171*6467f958SSadaf Ebrahimi             inptr[i][j] = (cl_ulong) genrand_int32(d) | ((cl_ulong) genrand_int32(d) << 32);
1172*6467f958SSadaf Ebrahimi     }
1173*6467f958SSadaf Ebrahimi 
1174*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( cl_ulong ), "ulong long", 5, (void **)inptr,
1175*6467f958SSadaf Ebrahimi                             stream_write_ulong_kernel_code, ulong_kernel_name, foo, d );
1176*6467f958SSadaf Ebrahimi 
1177*6467f958SSadaf Ebrahimi     for( i = 0; i < 5; i++ ){
1178*6467f958SSadaf Ebrahimi         free( (void *)inptr[i] );
1179*6467f958SSadaf Ebrahimi     }
1180*6467f958SSadaf Ebrahimi 
1181*6467f958SSadaf Ebrahimi     free_mtdata(d);
1182*6467f958SSadaf Ebrahimi     return err;
1183*6467f958SSadaf Ebrahimi 
1184*6467f958SSadaf Ebrahimi }    // end write_ulong_array()
1185*6467f958SSadaf Ebrahimi 
1186*6467f958SSadaf Ebrahimi 
test_write_array_struct(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1187*6467f958SSadaf Ebrahimi int test_write_array_struct( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
1188*6467f958SSadaf Ebrahimi {
1189*6467f958SSadaf Ebrahimi     TestStruct            *inptr[1];
1190*6467f958SSadaf Ebrahimi     size_t                ptrSizes[1];
1191*6467f958SSadaf Ebrahimi     int                    j, err;
1192*6467f958SSadaf Ebrahimi     int    (*foo)(void *,void *,int);
1193*6467f958SSadaf Ebrahimi     MTdata d = init_genrand( gRandomSeed );
1194*6467f958SSadaf Ebrahimi     foo = verify_write_struct;
1195*6467f958SSadaf Ebrahimi 
1196*6467f958SSadaf Ebrahimi     ptrSizes[0] = sizeof( TestStruct );
1197*6467f958SSadaf Ebrahimi 
1198*6467f958SSadaf Ebrahimi     inptr[0] = (TestStruct *)malloc( ptrSizes[0] * num_elements );
1199*6467f958SSadaf Ebrahimi 
1200*6467f958SSadaf Ebrahimi     for( j = 0; (unsigned int)j < ptrSizes[0] * num_elements / ptrSizes[0]; j++ ){
1201*6467f958SSadaf Ebrahimi         inptr[0][j].a = (int)genrand_int32(d);
1202*6467f958SSadaf Ebrahimi         inptr[0][j].b = get_random_float( 0.f, 1.844674407370954e+19f, d );
1203*6467f958SSadaf Ebrahimi     }
1204*6467f958SSadaf Ebrahimi 
1205*6467f958SSadaf Ebrahimi     err = test_stream_write( device, context, queue, num_elements, sizeof( TestStruct ), "struct", 1, (void **)inptr,
1206*6467f958SSadaf Ebrahimi                             stream_write_struct_kernel_code, struct_kernel_name, foo, d );
1207*6467f958SSadaf Ebrahimi 
1208*6467f958SSadaf Ebrahimi     free( (void *)inptr[0] );
1209*6467f958SSadaf Ebrahimi 
1210*6467f958SSadaf Ebrahimi     free_mtdata(d);
1211*6467f958SSadaf Ebrahimi     return err;
1212*6467f958SSadaf Ebrahimi 
1213*6467f958SSadaf Ebrahimi }    // end write_struct_array()
1214*6467f958SSadaf Ebrahimi 
1215*6467f958SSadaf Ebrahimi 
1216*6467f958SSadaf Ebrahimi 
1217*6467f958SSadaf Ebrahimi 
1218*6467f958SSadaf Ebrahimi 
1219