1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 //    http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "harness/compat.h"
17 
18 #include <stdio.h>
19 #include <string.h>
20 #include <sys/types.h>
21 #include <sys/stat.h>
22 
23 
24 #include "procs.h"
25 
26 const char *copy_kernel_code =
27 "__kernel void test_copy(__global unsigned int *src, __global unsigned int *dst)\n"
28 "{\n"
29 "    int  tid = get_global_id(0);\n"
30 "\n"
31 "    dst[tid] = src[tid];\n"
32 "}\n";
33 
34 int
test_arraycopy(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)35 test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
36 {
37     cl_uint    *input_ptr, *output_ptr;
38     cl_mem                streams[4], results;
39     cl_program          program;
40     cl_kernel            kernel;
41     unsigned            num_elements = 128 * 1024;
42     cl_uint             num_copies = 1;
43     size_t                delta_offset;
44     unsigned            i;
45     cl_int err;
46     MTdata              d;
47 
48     int error_count = 0;
49 
50     input_ptr = (cl_uint*)malloc(sizeof(cl_uint) * num_elements);
51     output_ptr = (cl_uint*)malloc(sizeof(cl_uint) * num_elements);
52 
53     // results
54     results = clCreateBuffer(context, CL_MEM_READ_WRITE,
55                              sizeof(cl_uint) * num_elements, NULL, &err);
56     test_error(err, "clCreateBuffer failed");
57 
58 /*****************************************************************************************************************************************/
59 #pragma mark client backing
60 
61     log_info("Testing CL_MEM_USE_HOST_PTR buffer with clEnqueueCopyBuffer\n");
62     // randomize data
63     d = init_genrand( gRandomSeed );
64     for (i=0; i<num_elements; i++)
65         input_ptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
66 
67     // client backing
68     streams[0] =
69         clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
70                        sizeof(cl_uint) * num_elements, input_ptr, &err);
71     test_error(err, "clCreateBuffer failed");
72 
73     delta_offset = num_elements * sizeof(cl_uint) / num_copies;
74     for (i=0; i<num_copies; i++)
75     {
76         size_t    offset = i * delta_offset;
77         err = clEnqueueCopyBuffer(queue, streams[0], results, offset, offset, delta_offset, 0, NULL, NULL);
78         test_error(err, "clEnqueueCopyBuffer failed");
79     }
80 
81     // Try upload from client backing
82     err = clEnqueueReadBuffer( queue, results, CL_TRUE, 0, num_elements*sizeof(cl_uint), output_ptr, 0, NULL, NULL );
83     test_error(err, "clEnqueueReadBuffer failed");
84 
85     for (i=0; i<num_elements; i++)
86     {
87         if (input_ptr[i] != output_ptr[i])
88         {
89             err = -1;
90             error_count++;
91         }
92     }
93 
94     if (err)
95         log_error("\tCL_MEM_USE_HOST_PTR buffer with clEnqueueCopyBuffer FAILED\n");
96     else
97         log_info("\tCL_MEM_USE_HOST_PTR buffer with clEnqueueCopyBuffer passed\n");
98 
99 
100 
101 #pragma mark framework backing (no client data)
102 
103     log_info("Testing with clEnqueueWriteBuffer and clEnqueueCopyBuffer\n");
104     // randomize data
105     for (i=0; i<num_elements; i++)
106         input_ptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
107 
108     // no backing
109     streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
110                                 sizeof(cl_uint) * num_elements, NULL, &err);
111     test_error(err, "clCreateBuffer failed");
112 
113     for (i=0; i<num_copies; i++)
114     {
115         size_t    offset = i * delta_offset;
116 
117         // Copy the array up from host ptr
118         err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, sizeof(cl_uint)*num_elements, input_ptr, 0, NULL, NULL);
119         test_error(err, "clEnqueueWriteBuffer failed");
120 
121         err = clEnqueueCopyBuffer(queue, streams[2], results, offset, offset, delta_offset, 0, NULL, NULL);
122         test_error(err, "clEnqueueCopyBuffer failed");
123     }
124 
125     err = clEnqueueReadBuffer( queue, results, true, 0, num_elements*sizeof(cl_uint), output_ptr, 0, NULL, NULL );
126     test_error(err, "clEnqueueReadBuffer failed");
127 
128     for (i=0; i<num_elements; i++)
129     {
130         if (input_ptr[i] != output_ptr[i])
131         {
132             err = -1;
133             error_count++;
134             break;
135         }
136     }
137 
138     if (err)
139         log_error("\tclEnqueueWriteBuffer and clEnqueueCopyBuffer FAILED\n");
140     else
141         log_info("\tclEnqueueWriteBuffer and clEnqueueCopyBuffer passed\n");
142 
143 /*****************************************************************************************************************************************/
144 #pragma mark kernel copy test
145 
146     log_info("Testing CL_MEM_USE_HOST_PTR buffer with kernel copy\n");
147     // randomize data
148     for (i=0; i<num_elements; i++)
149         input_ptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
150     free_mtdata(d); d= NULL;
151 
152     // client backing
153     streams[3] =
154         clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
155                        sizeof(cl_uint) * num_elements, input_ptr, &err);
156     test_error(err, "clCreateBuffer failed");
157 
158     err = create_single_kernel_helper(context, &program, &kernel, 1,
159                                       &copy_kernel_code, "test_copy");
160     test_error(err, "create_single_kernel_helper failed");
161 
162     err = clSetKernelArg(kernel, 0, sizeof streams[3], &streams[3]);
163     err |= clSetKernelArg(kernel, 1, sizeof results, &results);
164     test_error(err, "clSetKernelArg failed");
165 
166     size_t threads[3] = { num_elements, 0, 0 };
167 
168     err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
169   test_error(err, "clEnqueueNDRangeKernel failed");
170 
171     err = clEnqueueReadBuffer( queue, results, CL_TRUE, 0, num_elements*sizeof(cl_uint), output_ptr, 0, NULL, NULL );
172     test_error(err, "clEnqueueReadBuffer failed");
173 
174     for (i=0; i<num_elements; i++)
175     {
176         if (input_ptr[i] != output_ptr[i])
177         {
178             err = -1;
179       error_count++;
180             break;
181         }
182     }
183 
184   // Keep track of multiple errors.
185   if (error_count != 0)
186     err = error_count;
187 
188     if (err)
189         log_error("\tCL_MEM_USE_HOST_PTR buffer with kernel copy FAILED\n");
190     else
191         log_info("\tCL_MEM_USE_HOST_PTR buffer with kernel copy passed\n");
192 
193 
194   clReleaseProgram(program);
195   clReleaseKernel(kernel);
196   clReleaseMemObject(results);
197   clReleaseMemObject(streams[0]);
198   clReleaseMemObject(streams[2]);
199   clReleaseMemObject(streams[3]);
200 
201   free(input_ptr);
202   free(output_ptr);
203 
204     return err;
205 }
206 
207 
208 
209