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