Neko 1.99.6
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
math.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2021-2025, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34
35#ifdef __APPLE__
36#include <OpenCL/cl.h>
37#else
38#include <CL/cl.h>
39#endif
40
41#include <stdio.h>
42#include <stdlib.h>
43#include <math.h>
44#include <device/device_config.h>
45#include <device/opencl/jit.h>
46#include <device/opencl/prgm_lib.h>
47#include <device/opencl/buffer.h>
48#include <device/opencl/check.h>
49
50#include "math_kernel.cl.h"
51
55void opencl_copy(void *a, void *b, int *n, cl_command_queue cmd_queue) {
56 CL_CHECK(clEnqueueCopyBuffer(cmd_queue,
57 b, a, 0, 0, (*n) * sizeof(real),
58 0, NULL, NULL));
59}
60
65void opencl_masked_copy_0(void *a, void *b, void *mask, int *n, int *m,
66 cl_command_queue cmd_queue) {
67 cl_int err;
68
69 if (math_program == NULL)
70 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
71
72 cl_kernel kernel = clCreateKernel(math_program, "masked_copy_kernel_0", &err);
73 CL_CHECK(err);
74
75 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
76 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
77 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
78 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
79 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
80
81 const int nb = ((*m) + 256 - 1) / 256;
82 const size_t global_item_size = 256 * nb;
83 const size_t local_item_size = 256;
84
85 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
86 &global_item_size, &local_item_size,
87 0, NULL, NULL));
88 CL_CHECK(clReleaseKernel(kernel));
89
90}
91
95void opencl_masked_copy_aligned(void *a, void *b, void *mask, int *n, int *m,
96 cl_command_queue cmd_queue) {
97 cl_int err;
98
99 if (math_program == NULL)
100 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
101
102 cl_kernel kernel = clCreateKernel(math_program, "masked_copy_kernel_aligned", &err);
103 CL_CHECK(err);
104
105 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
106 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
107 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
108 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
109 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
110
111 const int nb = ((*m) + 256 - 1) / 256;
112 const size_t global_item_size = 256 * nb;
113 const size_t local_item_size = 256;
114
115 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
116 &global_item_size, &local_item_size,
117 0, NULL, NULL));
118 CL_CHECK(clReleaseKernel(kernel));
119
120}
121
125void opencl_masked_gather_copy(void *a, void *b, void *mask, int *n, int *m,
126 cl_command_queue cmd_queue) {
127 cl_int err;
128
129 if (math_program == NULL)
130 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
131
132 cl_kernel kernel = clCreateKernel(math_program, "masked_gather_copy_kernel",
133 &err);
134 CL_CHECK(err);
135
136 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
137 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
138 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
139 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
140 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
141
142 const int nb = ((*n) + 256 - 1) / 256;
143 const size_t global_item_size = 256 * nb;
144 const size_t local_item_size = 256;
145
146 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
147 &global_item_size, &local_item_size,
148 0, NULL, NULL));
149 CL_CHECK(clReleaseKernel(kernel));
150
151}
152
156void opencl_masked_gather_copy_aligned(void *a, void *b, void *mask, int *n,
157 int *m, cl_command_queue cmd_queue) {
158 cl_int err;
159
160 if (math_program == NULL)
161 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
162
163 cl_kernel kernel = clCreateKernel(math_program,
164 "masked_gather_copy_aligned_kernel", &err);
165 CL_CHECK(err);
166
167 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
168 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
169 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
170 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
171 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
172
173 const int nb = ((*n) + 256 - 1) / 256;
174 const size_t global_item_size = 256 * nb;
175 const size_t local_item_size = 256;
176
177 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
178 &global_item_size, &local_item_size,
179 0, NULL, NULL));
180 CL_CHECK(clReleaseKernel(kernel));
181
182}
183
187void opencl_face_masked_gather_copy(void *a, void *b, void *mask, void *facet,
188 int *n1, int *n2, int *lx, int *ly,
189 int *lz, int *m,
190 cl_command_queue cmd_queue) {
191 cl_int err;
192
193 if (math_program == NULL)
194 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
195
196 cl_kernel kernel = clCreateKernel(math_program,
197 "face_masked_gather_copy_kernel", &err);
198 CL_CHECK(err);
199
200 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
201 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
202 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
203 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &facet));
204 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n1));
205 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n2));
206 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), lx));
207 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), ly));
208 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), lz));
209 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), m));
210
211 const int nb = ((*m) + 256 - 1) / 256;
212 const size_t global_item_size = 256 * nb;
213 const size_t local_item_size = 256;
214
215 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
216 &global_item_size, &local_item_size,
217 0, NULL, NULL));
218 CL_CHECK(clReleaseKernel(kernel));
219
220}
221
225void opencl_masked_scatter_copy(void *a, void *b, void *mask, int *n, int *m,
226 cl_command_queue cmd_queue) {
227 cl_int err;
228
229 if (math_program == NULL)
230 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
231
232 cl_kernel kernel = clCreateKernel(math_program, "masked_scatter_copy_kernel",
233 &err);
234 CL_CHECK(err);
235
236 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
237 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
238 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
239 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
240 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
241
242 const int nb = ((*n) + 256 - 1) / 256;
243 const size_t global_item_size = 256 * nb;
244 const size_t local_item_size = 256;
245
246 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
247 &global_item_size, &local_item_size,
248 0, NULL, NULL));
249 CL_CHECK(clReleaseKernel(kernel));
250
251}
252
256void opencl_masked_scatter_copy_aligned(void *a, void *b, void *mask, int *n, int *m,
257 cl_command_queue cmd_queue) {
258 cl_int err;
259
260 if (math_program == NULL)
261 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
262
263 cl_kernel kernel = clCreateKernel(math_program, "masked_scatter_copy_aligned_kernel",
264 &err);
265 CL_CHECK(err);
266
267 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
268 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
269 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mask));
270 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
271 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), m));
272
273 const int nb = ((*n) + 256 - 1) / 256;
274 const size_t global_item_size = 256 * nb;
275 const size_t local_item_size = 256;
276
277 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
278 &global_item_size, &local_item_size,
279 0, NULL, NULL));
280 CL_CHECK(clReleaseKernel(kernel));
281
282}
283
287void opencl_cfill_mask(void* a, void* c, int* size, void* mask, int* mask_size,
288 cl_command_queue cmd_queue) {
289 cl_int err;
290
291 if (math_program == NULL)
292 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
293
294 cl_kernel kernel = clCreateKernel(math_program, "cfill_mask_kernel", &err);
295 CL_CHECK(err);
296
297 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
298 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
299 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), size));
300 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &mask));
301 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), mask_size));
302
303 const int nb = ((*mask_size) + 256 - 1) / 256;
304 const size_t global_item_size = 256 * nb;
305 const size_t local_item_size = 256;
306
307 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
308 &global_item_size, &local_item_size,
309 0, NULL, NULL));
310 CL_CHECK(clReleaseKernel(kernel));
311}
312
316void opencl_rzero(void *a, int *n, cl_command_queue cmd_queue) {
317 cl_event wait_kern;
318 real zero = 0.0;
319
320 CL_CHECK(clEnqueueFillBuffer(cmd_queue, a, &zero, sizeof(real), 0,
321 (*n) * sizeof(real), 0, NULL, &wait_kern));
322 CL_CHECK(clWaitForEvents(1, &wait_kern));
323}
324
328void opencl_rone(void *a, int *n, cl_command_queue cmd_queue) {
329 cl_event wait_kern;
330 real one = 1.0;
331
332 CL_CHECK(clEnqueueFillBuffer(cmd_queue, a, &one, sizeof(real), 0,
333 (*n) * sizeof(real), 0, NULL, &wait_kern));
334 CL_CHECK(clWaitForEvents(1, &wait_kern));
335}
336
340void opencl_cmult(void *a, real *c, int *n, cl_command_queue cmd_queue) {
341 cl_int err;
342
343 if (math_program == NULL)
344 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
345
346 cl_kernel kernel = clCreateKernel(math_program, "cmult_kernel", &err);
347 CL_CHECK(err);
348
349 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
350 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
351 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
352
353 const int nb = ((*n) + 256 - 1) / 256;
354 const size_t global_item_size = 256 * nb;
355 const size_t local_item_size = 256;
356
357 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
358 &global_item_size, &local_item_size,
359 0, NULL, NULL));
360 CL_CHECK(clReleaseKernel(kernel));
361}
362
366void opencl_cmult2(void *a, void *b, real *c, int *n,
367 cl_command_queue cmd_queue) {
368 cl_int err;
369
370 if (math_program == NULL)
371 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
372
373 cl_kernel kernel = clCreateKernel(math_program, "cmult2_kernel", &err);
374 CL_CHECK(err);
375
376 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
377 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
378 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c));
379 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
380
381 const int nb = ((*n) + 256 - 1) / 256;
382 const size_t global_item_size = 256 * nb;
383 const size_t local_item_size = 256;
384
385 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
386 &global_item_size, &local_item_size,
387 0, NULL, NULL));
388 CL_CHECK(clReleaseKernel(kernel));
389}
390
394void opencl_cdiv(void *a, real *c, int *n, cl_command_queue cmd_queue) {
395 cl_int err;
396
397 if (math_program == NULL)
398 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
399
400 cl_kernel kernel = clCreateKernel(math_program, "cdiv_kernel", &err);
401 CL_CHECK(err);
402
403 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
404 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
405 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
406
407 const int nb = ((*n) + 256 - 1) / 256;
408 const size_t global_item_size = 256 * nb;
409 const size_t local_item_size = 256;
410
411 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
412 &global_item_size, &local_item_size,
413 0, NULL, NULL));
414 CL_CHECK(clReleaseKernel(kernel));
415}
416
420void opencl_cdiv2(void *a, void *b, real *c, int *n,
421 cl_command_queue cmd_queue) {
422 cl_int err;
423
424 if (math_program == NULL)
425 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
426
427 cl_kernel kernel = clCreateKernel(math_program, "cdiv2_kernel", &err);
428 CL_CHECK(err);
429
430 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
431 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
432 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c));
433 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
434
435 const int nb = ((*n) + 256 - 1) / 256;
436 const size_t global_item_size = 256 * nb;
437 const size_t local_item_size = 256;
438
439 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
440 &global_item_size, &local_item_size,
441 0, NULL, NULL));
442 CL_CHECK(clReleaseKernel(kernel));
443}
444
448void opencl_radd(void *a, real *c, int *n, cl_command_queue cmd_queue) {
449 cl_int err;
450
451 if (math_program == NULL)
452 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
453
454 cl_kernel kernel = clCreateKernel(math_program, "radd_kernel", &err);
455 CL_CHECK(err);
456
457 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
458 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
459 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
460
461 const int nb = ((*n) + 256 - 1) / 256;
462 const size_t global_item_size = 256 * nb;
463 const size_t local_item_size = 256;
464
465 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
466 &global_item_size, &local_item_size,
467 0, NULL, NULL));
468 CL_CHECK(clReleaseKernel(kernel));
469}
470
474void opencl_cadd2(void *a, void *b, real *c, int *n,
475 cl_command_queue cmd_queue) {
476 cl_int err;
477
478 if (math_program == NULL)
479 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
480
481 cl_kernel kernel = clCreateKernel(math_program, "cadd2_kernel", &err);
482 CL_CHECK(err);
483
484 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
485 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
486 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c));
487 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
488
489 const int nb = ((*n) + 256 - 1) / 256;
490 const size_t global_item_size = 256 * nb;
491 const size_t local_item_size = 256;
492
493 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
494 &global_item_size, &local_item_size,
495 0, NULL, NULL));
496 CL_CHECK(clReleaseKernel(kernel));
497}
498
502void opencl_cwrap(void *a, real *min_val, real *max_val, int *n,
503 cl_command_queue cmd_queue) {
504 cl_int err;
505
506 if (math_program == NULL)
507 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
508
509 cl_kernel kernel = clCreateKernel(math_program, "cwrap_kernel", &err);
510 CL_CHECK(err);
511
512 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
513 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), min_val));
514 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), max_val));
515 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
516
517 const int nb = ((*n) + 256 - 1) / 256;
518 const size_t global_item_size = 256 * nb;
519 const size_t local_item_size = 256;
520
521 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
522 &global_item_size, &local_item_size,
523 0, NULL, NULL));
524 CL_CHECK(clReleaseKernel(kernel));
525}
526
530void opencl_sqrt_inplace(void *a, int *n, cl_command_queue cmd_queue) {
531 cl_int err;
532
533 if (math_program == NULL)
534 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
535
536 cl_kernel kernel = clCreateKernel(math_program, "sqrt_inplace_kernel", &err);
537 CL_CHECK(err);
538
539 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
540 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), n));
541
542 const int nb = ((*n) + 256 - 1) / 256;
543 const size_t global_item_size = 256 * nb;
544 const size_t local_item_size = 256;
545
546 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
547 &global_item_size, &local_item_size,
548 0, NULL, NULL));
549 CL_CHECK(clReleaseKernel(kernel));
550}
551
555void opencl_power(void *ap, void *a, real *p, int *n,
556 cl_command_queue cmd_queue) {
557 cl_int err;
558
559 if (math_program == NULL)
560 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
561
562 cl_kernel kernel = clCreateKernel(math_program, "power_kernel", &err);
563 CL_CHECK(err);
564
565 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &ap));
566 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &a));
567 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), p));
568 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
569
570 const int nb = ((*n) + 256 - 1) / 256;
571 const size_t global_item_size = 256 * nb;
572 const size_t local_item_size = 256;
573
574 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
575 &global_item_size, &local_item_size,
576 0, NULL, NULL));
577 CL_CHECK(clReleaseKernel(kernel));
578}
579
583void opencl_cfill(void *a, real *c, int *n, cl_command_queue cmd_queue) {
584 cl_int err;
585
586 if (math_program == NULL)
587 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
588
589 cl_kernel kernel = clCreateKernel(math_program, "cfill_kernel", &err);
590 CL_CHECK(err);
591
592 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
593 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
594 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
595
596 const int nb = ((*n) + 256 - 1) / 256;
597 const size_t global_item_size = 256 * nb;
598 const size_t local_item_size = 256;
599
600 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
601 &global_item_size, &local_item_size,
602 0, NULL, NULL));
603 CL_CHECK(clReleaseKernel(kernel));
604}
605
610void opencl_add2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
611 cl_int err;
612
613 if (math_program == NULL)
614 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
615
616 cl_kernel kernel = clCreateKernel(math_program, "add2_kernel", &err);
617 CL_CHECK(err);
618
619 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
620 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
621 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
622
623 const int nb = ((*n) + 256 - 1) / 256;
624 const size_t global_item_size = 256 * nb;
625 const size_t local_item_size = 256;
626
627 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
628 &global_item_size, &local_item_size,
629 0, NULL, NULL));
630 CL_CHECK(clReleaseKernel(kernel));
631}
632
637void opencl_add3(void *a, void *b, void *c, int *n,
638 cl_command_queue cmd_queue) {
639 cl_int err;
640
641 if (math_program == NULL)
642 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
643
644 cl_kernel kernel = clCreateKernel(math_program, "add3_kernel", &err);
645 CL_CHECK(err);
646
647 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
648 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
649 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
650 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
651
652 const int nb = ((*n) + 256 - 1) / 256;
653 const size_t global_item_size = 256 * nb;
654 const size_t local_item_size = 256;
655
656 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
657 &global_item_size, &local_item_size,
658 0, NULL, NULL));
659 CL_CHECK(clReleaseKernel(kernel));
660}
661
666void opencl_add4(void *a, void *b, void *c, void *d, int *n,
667 cl_command_queue cmd_queue) {
668 cl_int err;
669
670 if (math_program == NULL)
671 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
672
673 cl_kernel kernel = clCreateKernel(math_program, "add4_kernel", &err);
674 CL_CHECK(err);
675
676 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
677 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
678 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
679 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &d));
680 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
681
682 const int nb = ((*n) + 256 - 1) / 256;
683 const size_t global_item_size = 256 * nb;
684 const size_t local_item_size = 256;
685
686 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
687 &global_item_size, &local_item_size,
688 0, NULL, NULL));
689 CL_CHECK(clReleaseKernel(kernel));
690}
691
697void opencl_add2s1(void *a, void *b, real *c1, int *n,
698 cl_command_queue cmd_queue) {
699 cl_int err;
700
701 if (math_program == NULL)
702 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
703
704 cl_kernel kernel = clCreateKernel(math_program, "add2s1_kernel", &err);
705 CL_CHECK(err);
706
707 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
708 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
709 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c1));
710 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
711
712 const int nb = ((*n) + 256 - 1) / 256;
713 const size_t global_item_size = 256 * nb;
714 const size_t local_item_size = 256;
715
716 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
717 &global_item_size, &local_item_size,
718 0, NULL, NULL));
719 CL_CHECK(clReleaseKernel(kernel));
720}
721
727void opencl_add2s2(void *a, void *b, real *c1, int *n,
728 cl_command_queue cmd_queue) {
729 cl_int err;
730
731 if (math_program == NULL)
732 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
733
734 cl_kernel kernel = clCreateKernel(math_program, "add2s2_kernel", &err);
735 CL_CHECK(err);
736
737 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
738 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
739 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c1));
740 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
741
742 const int nb = ((*n) + 256 - 1) / 256;
743 const size_t global_item_size = 256 * nb;
744 const size_t local_item_size = 256;
745
746 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
747 &global_item_size, &local_item_size,
748 0, NULL, NULL));
749 CL_CHECK(clReleaseKernel(kernel));
750}
751
758void opencl_add2s2_many(void *x, void *p, void *alpha, int *j, int *n,
759 cl_command_queue cmd_queue) {
760 cl_int err;
761
762 if (math_program == NULL)
763 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
764
765 cl_kernel kernel = clCreateKernel(math_program, "add2s2_many_kernel", &err);
766 CL_CHECK(err);
767
768 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &x));
769 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &p));
770 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &alpha));
771 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), j));
772 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
773
774 const int nb = ((*n) + 256 - 1) / 256;
775 const size_t global_item_size = 256 * nb;
776 const size_t local_item_size = 256;
777
778 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
779 &global_item_size, &local_item_size,
780 0, NULL, NULL));
781 CL_CHECK(clReleaseKernel(kernel));
782
783}
784
790void opencl_addsqr2s2(void *a, void *b, real *c1, int *n,
791 cl_command_queue cmd_queue) {
792 cl_int err;
793
794 if (math_program == NULL)
795 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
796
797 cl_kernel kernel = clCreateKernel(math_program, "addsqr2s2_kernel", &err);
798 CL_CHECK(err);
799
800 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
801 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
802 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c1));
803 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
804
805 const int nb = ((*n) + 256 - 1) / 256;
806 const size_t global_item_size = 256 * nb;
807 const size_t local_item_size = 256;
808
809 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
810 &global_item_size, &local_item_size,
811 0, NULL, NULL));
812 CL_CHECK(clReleaseKernel(kernel));
813}
814
819void opencl_add3s2(void *a, void *b, void * c, real *c1, real *c2, int *n,
820 cl_command_queue cmd_queue) {
821 cl_int err;
822
823 if (math_program == NULL)
824 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
825
826 cl_kernel kernel = clCreateKernel(math_program, "add3s2_kernel", &err);
827 CL_CHECK(err);
828
829 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
830 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
831 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
832 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), c1));
833 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(real), c2));
834 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
835
836 const int nb = ((*n) + 256 - 1) / 256;
837 const size_t global_item_size = 256 * nb;
838 const size_t local_item_size = 256;
839
840 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
841 &global_item_size, &local_item_size,
842 0, NULL, NULL));
843 CL_CHECK(clReleaseKernel(kernel));
844}
845
850void opencl_add4s3(void *a, void *b, void * c, void * d, real *c1, real *c2,
851 real *c3, int *n, cl_command_queue cmd_queue) {
852 cl_int err;
853
854 if (math_program == NULL)
855 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
856
857 cl_kernel kernel = clCreateKernel(math_program, "add4s3_kernel", &err);
858 CL_CHECK(err);
859
860 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
861 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
862 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
863 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &d));
864 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(real), c1));
865 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(real), c2));
866 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(real), c3));
867 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), n));
868
869 const int nb = ((*n) + 256 - 1) / 256;
870 const size_t global_item_size = 256 * nb;
871 const size_t local_item_size = 256;
872
873 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
874 &global_item_size, &local_item_size,
875 0, NULL, NULL));
876 CL_CHECK(clReleaseKernel(kernel));
877}
878
883void opencl_add5s4(void *a, void *b, void * c, void * d, void * e, real *c1,
884 real *c2, real *c3, real * c4, int *n,
885 cl_command_queue cmd_queue) {
886 cl_int err;
887
888 if (math_program == NULL)
889 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
890
891 cl_kernel kernel = clCreateKernel(math_program, "add5s4_kernel", &err);
892 CL_CHECK(err);
893
894 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
895 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
896 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
897 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &d));
898 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &e));
899 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(real), c1));
900 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(real), c2));
901 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(real), c3));
902 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(real), c4));
903 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), n));
904
905 const int nb = ((*n) + 256 - 1) / 256;
906 const size_t global_item_size = 256 * nb;
907 const size_t local_item_size = 256;
908
909 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
910 &global_item_size, &local_item_size,
911 0, NULL, NULL));
912 CL_CHECK(clReleaseKernel(kernel));
913}
914
919void opencl_invcol1(void *a, int *n, cl_command_queue cmd_queue) {
920 cl_int err;
921
922 if (math_program == NULL)
923 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
924
925 cl_kernel kernel = clCreateKernel(math_program, "invcol1_kernel", &err);
926 CL_CHECK(err);
927
928 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
929 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), n));
930
931 const int nb = ((*n) + 256 - 1) / 256;
932 const size_t global_item_size = 256 * nb;
933 const size_t local_item_size = 256;
934
935 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
936 &global_item_size, &local_item_size,
937 0, NULL, NULL));
938 CL_CHECK(clReleaseKernel(kernel));
939}
940
945void opencl_invcol2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
946 cl_int err;
947
948 if (math_program == NULL)
949 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
950
951 cl_kernel kernel = clCreateKernel(math_program, "invcol2_kernel", &err);
952 CL_CHECK(err);
953
954 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
955 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
956 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
957
958 const int nb = ((*n) + 256 - 1) / 256;
959 const size_t global_item_size = 256 * nb;
960 const size_t local_item_size = 256;
961
962 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
963 &global_item_size, &local_item_size,
964 0, NULL, NULL));
965 CL_CHECK(clReleaseKernel(kernel));
966}
967
972void opencl_invcol3(void *a, void *b, void *c, int *n,
973 cl_command_queue cmd_queue) {
974 cl_int err;
975
976 if (math_program == NULL)
977 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
978
979 cl_kernel kernel = clCreateKernel(math_program, "invcol3_kernel", &err);
980 CL_CHECK(err);
981
982 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
983 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
984 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
985 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
986
987 const int nb = ((*n) + 256 - 1) / 256;
988 const size_t global_item_size = 256 * nb;
989 const size_t local_item_size = 256;
990
991 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
992 &global_item_size, &local_item_size,
993 0, NULL, NULL));
994 CL_CHECK(clReleaseKernel(kernel));
995}
996
1001void opencl_col2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1002 cl_int err;
1003
1004 if (math_program == NULL)
1005 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1006
1007 cl_kernel kernel = clCreateKernel(math_program, "col2_kernel", &err);
1008 CL_CHECK(err);
1009
1010 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1011 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1012 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1013
1014 const int nb = ((*n) + 256 - 1) / 256;
1015 const size_t global_item_size = 256 * nb;
1016 const size_t local_item_size = 256;
1017
1018 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1019 &global_item_size, &local_item_size,
1020 0, NULL, NULL));
1021 CL_CHECK(clReleaseKernel(kernel));
1022}
1023
1028void opencl_col3(void *a, void *b, void *c, int *n,
1029 cl_command_queue cmd_queue) {
1030 cl_int err;
1031
1032 if (math_program == NULL)
1033 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1034
1035 cl_kernel kernel = clCreateKernel(math_program, "col3_kernel", &err);
1036 CL_CHECK(err);
1037
1038 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1039 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1040 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1041 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1042
1043 const int nb = ((*n) + 256 - 1) / 256;
1044 const size_t global_item_size = 256 * nb;
1045 const size_t local_item_size = 256;
1046
1047 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1048 &global_item_size, &local_item_size,
1049 0, NULL, NULL));
1050 CL_CHECK(clReleaseKernel(kernel));
1051}
1052
1057void opencl_subcol3(void *a, void *b, void *c, int *n,
1058 cl_command_queue cmd_queue) {
1059 cl_int err;
1060
1061 if (math_program == NULL)
1062 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1063
1064 cl_kernel kernel = clCreateKernel(math_program, "subcol3_kernel", &err);
1065 CL_CHECK(err);
1066
1067 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1068 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1069 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1070 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1071
1072 const int nb = ((*n) + 256 - 1) / 256;
1073 const size_t global_item_size = 256 * nb;
1074 const size_t local_item_size = 256;
1075
1076 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1077 &global_item_size, &local_item_size,
1078 0, NULL, NULL));
1079 CL_CHECK(clReleaseKernel(kernel));
1080}
1081
1086void opencl_sub2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1087 cl_int err;
1088
1089 if (math_program == NULL)
1090 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1091
1092 cl_kernel kernel = clCreateKernel(math_program, "sub2_kernel", &err);
1093 CL_CHECK(err);
1094
1095 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1096 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1097 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1098
1099 const int nb = ((*n) + 256 - 1) / 256;
1100 const size_t global_item_size = 256 * nb;
1101 const size_t local_item_size = 256;
1102
1103 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1104 &global_item_size, &local_item_size,
1105 0, NULL, NULL));
1106 CL_CHECK(clReleaseKernel(kernel));
1107}
1108
1113void opencl_sub3(void *a, void *b, void *c, int *n,
1114 cl_command_queue cmd_queue) {
1115 cl_int err;
1116
1117 if (math_program == NULL)
1118 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1119
1120 cl_kernel kernel = clCreateKernel(math_program, "sub3_kernel", &err);
1121 CL_CHECK(err);
1122
1123 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1124 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1125 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1126 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1127
1128 const int nb = ((*n) + 256 - 1) / 256;
1129 const size_t global_item_size = 256 * nb;
1130 const size_t local_item_size = 256;
1131
1132 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1133 &global_item_size, &local_item_size,
1134 0, NULL, NULL));
1135 CL_CHECK(clReleaseKernel(kernel));
1136}
1137
1142void opencl_addcol3(void *a, void *b, void *c, int *n,
1143 cl_command_queue cmd_queue) {
1144 cl_int err;
1145
1146 if (math_program == NULL)
1147 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1148
1149 cl_kernel kernel = clCreateKernel(math_program, "addcol3_kernel", &err);
1150 CL_CHECK(err);
1151
1152 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1153 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1154 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1155 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1156
1157 const int nb = ((*n) + 256 - 1) / 256;
1158 const size_t global_item_size = 256 * nb;
1159 const size_t local_item_size = 256;
1160
1161 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1162 &global_item_size, &local_item_size,
1163 0, NULL, NULL));
1164 CL_CHECK(clReleaseKernel(kernel));
1165}
1166
1171void opencl_addcol4(void *a, void *b, void *c, void *d, int *n,
1172 cl_command_queue cmd_queue) {
1173 cl_int err;
1174
1175 if (math_program == NULL)
1176 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1177
1178 cl_kernel kernel = clCreateKernel(math_program, "addcol4_kernel", &err);
1179 CL_CHECK(err);
1180
1181 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1182 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1183 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1184 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &d));
1185 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
1186
1187 const int nb = ((*n) + 256 - 1) / 256;
1188 const size_t global_item_size = 256 * nb;
1189 const size_t local_item_size = 256;
1190
1191 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
1192 NULL, &global_item_size, &local_item_size,
1193 0, NULL, NULL));
1194 CL_CHECK(clReleaseKernel(kernel));
1195}
1196
1201void opencl_addcol3s2(void *a, void *b, void *c, real *s, int *n,
1202 cl_command_queue cmd_queue) {
1203 cl_int err;
1204
1205 if (math_program == NULL)
1206 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1207
1208 cl_kernel kernel = clCreateKernel(math_program, "addcol3s2_kernel", &err);
1209 CL_CHECK(err);
1210
1211 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1212 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1213 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1214 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), s));
1215 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
1216
1217 const int nb = ((*n) + 256 - 1) / 256;
1218 const size_t global_item_size = 256 * nb;
1219 const size_t local_item_size = 256;
1220
1221 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
1222 NULL, &global_item_size, &local_item_size,
1223 0, NULL, NULL));
1224 CL_CHECK(clReleaseKernel(kernel));
1225}
1226
1232void opencl_vdot3(void *dot, void *u1, void *u2, void *u3,
1233 void *v1, void *v2, void *v3, int *n,
1234 cl_command_queue cmd_queue) {
1235 cl_int err;
1236
1237 if (math_program == NULL)
1238 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1239
1240 cl_kernel kernel = clCreateKernel(math_program, "vdot3_kernel", &err);
1241 CL_CHECK(err);
1242
1243 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &dot));
1244 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &u1));
1245 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &u2));
1246 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &u3));
1247 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &v1));
1248 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &v2));
1249 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &v3));
1250 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), n));
1251
1252 const int nb = ((*n) + 256 - 1) / 256;
1253 const size_t global_item_size = 256 * nb;
1254 const size_t local_item_size = 256;
1255
1256 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1257 &global_item_size, &local_item_size,
1258 0, NULL, NULL));
1259 CL_CHECK(clReleaseKernel(kernel));
1260}
1261
1267void opencl_vcross(void *u1, void *u2, void *u3,
1268 void *v1, void *v2, void *v3,
1269 void *w1, void *w2, void *w3,
1270 int *n, cl_command_queue cmd_queue) {
1271
1272 cl_int err;
1273
1274 if (math_program == NULL)
1275 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1276
1277 cl_kernel kernel = clCreateKernel(math_program, "vcross_kernel", &err);
1278 CL_CHECK(err);
1279
1280 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &u1));
1281 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &u2));
1282 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &u3));
1283 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &v1));
1284 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &v2));
1285 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &v3));
1286 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &w1));
1287 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &w2));
1288 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &w3));
1289 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), n));
1290
1291 const int nb = ((*n) + 256 - 1) / 256;
1292 const size_t global_item_size = 256 * nb;
1293 const size_t local_item_size = 256;
1294
1295 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1296 &global_item_size, &local_item_size,
1297 0, NULL, NULL));
1298 CL_CHECK(clReleaseKernel(kernel));
1299
1300}
1301
1302/*
1303 * Reduction buffer, owned by the device layer and released
1304 * on device teardown (opencl_buffer_free_all in opencl_finalize)
1305 */
1306opencl_buffer_t redbuf_xp = OPENCL_BUFFER_INIT;
1307
1312real_xp opencl_glsc3(void *a, void *b, void *c, int *n,
1313 cl_command_queue cmd_queue) {
1314 cl_int err;
1315 cl_event kern_wait;
1316 int i;
1317
1318 if (*n <= 0) {
1319 return 0.0;
1320 }
1321
1322 if (math_program == NULL)
1323 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1324
1325 const int nb = ((*n) + 256 - 1) / 256;
1326 const size_t global_item_size = 256 * nb;
1327 const size_t local_item_size = 256;
1328
1329 opencl_buffer_reserve(&redbuf_xp, nb * sizeof(real_xp));
1330
1331 cl_kernel kernel = clCreateKernel(math_program, "glsc3_kernel", &err);
1332 CL_CHECK(err);
1333
1334 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1335 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1336 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1337 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &redbuf_xp.dev));
1338 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
1339
1340 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1341 &global_item_size, &local_item_size,
1342 0, NULL, &kern_wait));
1343
1344 CL_CHECK(clEnqueueReadBuffer(cmd_queue, redbuf_xp.dev, CL_TRUE, 0,
1345 nb * sizeof(real_xp), ((real_xp *) redbuf_xp.host), 1,
1346 &kern_wait, NULL));
1347
1348 real_xp res = 0.0;
1349 for (i = 0; i < nb; i++) {
1350 res += ((real_xp *) redbuf_xp.host)[i];
1351 }
1352
1353 CL_CHECK(clReleaseEvent(kern_wait));
1354 CL_CHECK(clReleaseKernel(kernel));
1355
1356 return res;
1357}
1358
1363void opencl_glsc3_many(real_xp *h, void * w, void *v, void *mult, int *j,
1364 int *n, cl_command_queue cmd_queue){
1365 int i, k;
1366 cl_int err;
1367 cl_event kern_wait;
1368
1369 if (*n <= 0) {
1370 for (k = 0; k < (*j); k++) {
1371 h[k] = 0.0;
1372 }
1373 return;
1374 }
1375
1376 if (math_program == NULL)
1377 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1378
1379 int pow2 = 1;
1380 while(pow2 < (*j)){
1381 pow2 = 2*pow2;
1382 }
1383
1384 const int nt = 256 / pow2;
1385 const int nb = ((*n) + nt - 1) / nt;
1386 const size_t local_item_size[2] = {nt, pow2};
1387 const size_t global_item_size[2] = {nb * nt, pow2};
1388
1389 opencl_buffer_reserve(&redbuf_xp, (*j) * nb * sizeof(real_xp));
1390
1391 cl_kernel kernel = clCreateKernel(math_program, "glsc3_many_kernel", &err);
1392 CL_CHECK(err);
1393
1394 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &w));
1395 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &v));
1396 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &mult));
1397 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &redbuf_xp.dev));
1398 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), j));
1399 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
1400
1401 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 2, NULL,
1402 global_item_size, local_item_size,
1403 0, NULL, &kern_wait));
1404
1405 CL_CHECK(clEnqueueReadBuffer(cmd_queue, redbuf_xp.dev, CL_TRUE, 0,
1406 (*j) * nb * sizeof(real_xp),
1407 ((real_xp *) redbuf_xp.host), 1, &kern_wait, NULL));
1408
1409 for (k = 0; k < (*j); k++) {
1410 h[k] = 0.0;
1411 }
1412
1413 for (i = 0; i < nb; i++) {
1414 for (k = 0; k < (*j); k++) {
1415 h[k] += ((real_xp *) redbuf_xp.host)[i*(*j)+k];
1416 }
1417 }
1418
1419 CL_CHECK(clReleaseEvent(kern_wait));
1420 CL_CHECK(clReleaseKernel(kernel));
1421}
1422
1427real_xp opencl_glsc2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1428 cl_int err;
1429 cl_event kern_wait;
1430 int i;
1431
1432 if (*n <= 0) {
1433 return 0.0;
1434 }
1435
1436 if (math_program == NULL)
1437 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1438
1439 const int nb = ((*n) + 256 - 1) / 256;
1440 const size_t global_item_size = 256 * nb;
1441 const size_t local_item_size = 256;
1442
1443 opencl_buffer_reserve(&redbuf_xp, nb * sizeof(real_xp));
1444
1445 cl_kernel kernel = clCreateKernel(math_program, "glsc2_kernel", &err);
1446 CL_CHECK(err);
1447
1448 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1449 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1450 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &redbuf_xp.dev));
1451 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1452
1453 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1454 &global_item_size, &local_item_size,
1455 0, NULL, &kern_wait));
1456
1457 CL_CHECK(clEnqueueReadBuffer(cmd_queue, redbuf_xp.dev, CL_TRUE, 0,
1458 nb * sizeof(real_xp), ((real_xp *) redbuf_xp.host), 1,
1459 &kern_wait, NULL));
1460
1461 real_xp res = 0.0;
1462 for (i = 0; i < nb; i++) {
1463 res += ((real_xp *) redbuf_xp.host)[i];
1464 }
1465
1466 CL_CHECK(clReleaseEvent(kern_wait));
1467 CL_CHECK(clReleaseKernel(kernel));
1468
1469 return res;
1470}
1471
1476real_xp opencl_glsubnorm2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1477 cl_int err;
1478 cl_event kern_wait;
1479 int i;
1480
1481 if (*n <= 0) {
1482 return 0.0;
1483 }
1484
1485 if (math_program == NULL)
1486 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1487
1488 const int nb = ((*n) + 256 - 1) / 256;
1489 const size_t global_item_size = 256 * nb;
1490 const size_t local_item_size = 256;
1491
1492 opencl_buffer_reserve(&redbuf_xp, nb * sizeof(real_xp));
1493
1494 cl_kernel kernel = clCreateKernel(math_program, "glsubnorm2_kernel", &err);
1495 CL_CHECK(err);
1496
1497 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1498 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1499 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &redbuf_xp.dev));
1500 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1501
1502 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1503 &global_item_size, &local_item_size,
1504 0, NULL, &kern_wait));
1505
1506 CL_CHECK(clEnqueueReadBuffer(cmd_queue, redbuf_xp.dev, CL_TRUE, 0,
1507 nb * sizeof(real_xp), ((real_xp *) redbuf_xp.host), 1,
1508 &kern_wait, NULL));
1509
1510 real_xp res = 0.0;
1511 for (i = 0; i < nb; i++) {
1512 res += ((real_xp *) redbuf_xp.host)[i];
1513 }
1514
1515 CL_CHECK(clReleaseEvent(kern_wait));
1516 CL_CHECK(clReleaseKernel(kernel));
1517
1518 return res;
1519}
1520
1525real_xp opencl_glsum(void *a, int *n, cl_command_queue cmd_queue) {
1526 cl_int err;
1527 cl_event kern_wait;
1528 int i;
1529
1530 if (*n <= 0) {
1531 return 0.0;
1532 }
1533
1534 if (math_program == NULL)
1535 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1536
1537 const int nb = ((*n) + 256 - 1) / 256;
1538 const size_t global_item_size = 256 * nb;
1539 const size_t local_item_size = 256;
1540
1541 opencl_buffer_reserve(&redbuf_xp, nb * sizeof(real_xp));
1542
1543 cl_kernel kernel = clCreateKernel(math_program, "glsum_kernel", &err);
1544 CL_CHECK(err);
1545
1546 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1547 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &redbuf_xp.dev));
1548 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1549
1550 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1551 &global_item_size, &local_item_size,
1552 0, NULL, &kern_wait));
1553
1554 CL_CHECK(clEnqueueReadBuffer(cmd_queue, redbuf_xp.dev, CL_TRUE, 0,
1555 nb * sizeof(real_xp), ((real_xp *) redbuf_xp.host), 1, &kern_wait, NULL));
1556
1557 real_xp res = 0.0;
1558 for (i = 0; i < nb; i++) {
1559 res += ((real_xp *) redbuf_xp.host)[i];
1560 }
1561
1562 CL_CHECK(clReleaseEvent(kern_wait));
1563 CL_CHECK(clReleaseKernel(kernel));
1564
1565 return res;
1566}
1567
1568real opencl_glmax(void *a, int *n, cl_command_queue cmd_queue) {
1569 cl_int err;
1570 cl_event kern_wait;
1571 int i;
1572
1573 if (*n <= 0) {
1574 return -((real) HUGE_VAL);
1575 }
1576
1577 if (math_program == NULL)
1578 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1579
1580 const int nb = ((*n) + 256 - 1) / 256;
1581 const size_t global_item_size = 256 * nb;
1582 const size_t local_item_size = 256;
1583
1584 real * buf = (real *) malloc(nb * sizeof(real));
1585
1586 cl_kernel kernel = clCreateKernel(math_program, "glmax_kernel", &err);
1587 CL_CHECK(err);
1588
1589 cl_mem buf_d = clCreateBuffer(glb_ctx, CL_MEM_READ_WRITE,
1590 nb * sizeof(real), NULL, &err);
1591 CL_CHECK(err);
1592
1593 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1594 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &buf_d));
1595 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1596
1597 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1598 &global_item_size, &local_item_size,
1599 0, NULL, &kern_wait));
1600
1601 CL_CHECK(clEnqueueReadBuffer(cmd_queue, buf_d, CL_TRUE, 0,
1602 nb * sizeof(real), buf, 1, &kern_wait, NULL));
1603
1604 real res = buf[0];
1605 for (i = 1; i < nb; i++) {
1606 res = fmax(res, buf[i]);
1607 }
1608
1609 free(buf);
1610 CL_CHECK(clReleaseEvent(kern_wait));
1611 CL_CHECK(clReleaseMemObject(buf_d));
1612 CL_CHECK(clReleaseKernel(kernel));
1613
1614 return res;
1615}
1616
1617real opencl_glmin(void *a, int *n, cl_command_queue cmd_queue) {
1618 cl_int err;
1619 cl_event kern_wait;
1620 int i;
1621
1622 if (*n <= 0) {
1623 return (real) HUGE_VAL;
1624 }
1625
1626 if (math_program == NULL)
1627 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1628
1629 const int nb = ((*n) + 256 - 1) / 256;
1630 const size_t global_item_size = 256 * nb;
1631 const size_t local_item_size = 256;
1632
1633 real * buf = (real *) malloc(nb * sizeof(real));
1634
1635 cl_kernel kernel = clCreateKernel(math_program, "glmin_kernel", &err);
1636 CL_CHECK(err);
1637
1638 cl_mem buf_d = clCreateBuffer(glb_ctx, CL_MEM_READ_WRITE,
1639 nb * sizeof(real), NULL, &err);
1640 CL_CHECK(err);
1641
1642 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1643 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &buf_d));
1644 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1645
1646 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1647 &global_item_size, &local_item_size,
1648 0, NULL, &kern_wait));
1649
1650 CL_CHECK(clEnqueueReadBuffer(cmd_queue, buf_d, CL_TRUE, 0,
1651 nb * sizeof(real), buf, 1, &kern_wait, NULL));
1652
1653 real res = buf[0];
1654 for (i = 1; i < nb; i++) {
1655 res = fmin(res, buf[i]);
1656 }
1657
1658 free(buf);
1659 CL_CHECK(clReleaseEvent(kern_wait));
1660 CL_CHECK(clReleaseMemObject(buf_d));
1661 CL_CHECK(clReleaseKernel(kernel));
1662
1663 return res;
1664}
1665
1666
1670void opencl_absval(void *a, int *n, cl_command_queue cmd_queue) {
1671 cl_int err;
1672
1673 if (math_program == NULL)
1674 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1675
1676 cl_kernel kernel = clCreateKernel(math_program, "absval_kernel", &err);
1677 CL_CHECK(err);
1678
1679 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1680 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), n));
1681
1682 const int nb = ((*n) + 256 - 1) / 256;
1683 const size_t global_item_size = 256 * nb;
1684 const size_t local_item_size = 256;
1685
1686 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL,
1687 &global_item_size, &local_item_size,
1688 0, NULL, NULL));
1689 CL_CHECK(clReleaseKernel(kernel));
1690}
1691
1695void opencl_iadd(void *a, int *c, int *n, cl_command_queue cmd_queue) {
1696 cl_int err;
1697
1698 if (math_program == NULL)
1699 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1700
1701 cl_kernel kernel = clCreateKernel(math_program, "iadd_kernel", &err);
1702 CL_CHECK(err);
1703
1704 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1705 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), c));
1706 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1707
1708 const int nb = ((*n) + 256 - 1) / 256;
1709 const size_t global_item_size = 256 * nb;
1710 const size_t local_item_size = 256;
1711
1712 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1713 NULL, &global_item_size, &local_item_size,
1714 0, NULL, NULL));
1715 CL_CHECK(clReleaseKernel(kernel));
1716}
1717
1722void opencl_pwmax_vec2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1723 cl_int err;
1724
1725 if (math_program == NULL)
1726 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1727
1728 cl_kernel kernel = clCreateKernel(math_program, "pwmax_vec2_kernel", &err);
1729 CL_CHECK(err);
1730
1731 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1732 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1733 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1734
1735 const int nb = ((*n) + 256 - 1) / 256;
1736 const size_t global_item_size = 256 * nb;
1737 const size_t local_item_size = 256;
1738
1739 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1740 NULL, &global_item_size, &local_item_size,
1741 0, NULL, NULL));
1742 CL_CHECK(clReleaseKernel(kernel));
1743}
1744
1749void opencl_pwmax_vec3(void *a, void *b, void *c,
1750 int *n, cl_command_queue cmd_queue) {
1751 cl_int err;
1752
1753 if (math_program == NULL)
1754 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1755
1756 cl_kernel kernel = clCreateKernel(math_program, "pwmax_vec3_kernel", &err);
1757 CL_CHECK(err);
1758
1759 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1760 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1761 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1762 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1763
1764 const int nb = ((*n) + 256 - 1) / 256;
1765 const size_t global_item_size = 256 * nb;
1766 const size_t local_item_size = 256;
1767
1768 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1769 NULL, &global_item_size, &local_item_size,
1770 0, NULL, NULL));
1771 CL_CHECK(clReleaseKernel(kernel));
1772}
1773
1778void opencl_pwmax_sca2(void *a, real *c, int *n, cl_command_queue cmd_queue) {
1779 cl_int err;
1780
1781 if (math_program == NULL)
1782 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1783
1784 cl_kernel kernel = clCreateKernel(math_program, "pwmax_sca2_kernel", &err);
1785 CL_CHECK(err);
1786
1787 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1788 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
1789 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1790
1791 const int nb = ((*n) + 256 - 1) / 256;
1792 const size_t global_item_size = 256 * nb;
1793 const size_t local_item_size = 256;
1794
1795 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1796 NULL, &global_item_size, &local_item_size,
1797 0, NULL, NULL));
1798 CL_CHECK(clReleaseKernel(kernel));
1799}
1800
1805void opencl_pwmax_sca3(void *a, void *b, real *c,
1806 int *n, cl_command_queue cmd_queue) {
1807 cl_int err;
1808
1809 if (math_program == NULL)
1810 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1811
1812 cl_kernel kernel = clCreateKernel(math_program, "pwmax_sca3_kernel", &err);
1813 CL_CHECK(err);
1814
1815 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1816 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1817 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c));
1818 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1819
1820 const int nb = ((*n) + 256 - 1) / 256;
1821 const size_t global_item_size = 256 * nb;
1822 const size_t local_item_size = 256;
1823
1824 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1825 NULL, &global_item_size, &local_item_size,
1826 0, NULL, NULL));
1827 CL_CHECK(clReleaseKernel(kernel));
1828}
1829
1834void opencl_pwmin_vec2(void *a, void *b, int *n, cl_command_queue cmd_queue) {
1835 cl_int err;
1836
1837 if (math_program == NULL)
1838 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1839
1840 cl_kernel kernel = clCreateKernel(math_program, "pwmin_vec2_kernel", &err);
1841 CL_CHECK(err);
1842
1843 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1844 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1845 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1846
1847 const int nb = ((*n) + 256 - 1) / 256;
1848 const size_t global_item_size = 256 * nb;
1849 const size_t local_item_size = 256;
1850
1851 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1852 NULL, &global_item_size, &local_item_size,
1853 0, NULL, NULL));
1854 CL_CHECK(clReleaseKernel(kernel));
1855}
1856
1861void opencl_pwmin_vec3(void *a, void *b, void *c,
1862 int *n, cl_command_queue cmd_queue) {
1863 cl_int err;
1864
1865 if (math_program == NULL)
1866 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1867
1868 cl_kernel kernel = clCreateKernel(math_program, "pwmin_vec3_kernel", &err);
1869 CL_CHECK(err);
1870
1871 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1872 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1873 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &c));
1874 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1875
1876 const int nb = ((*n) + 256 - 1) / 256;
1877 const size_t global_item_size = 256 * nb;
1878 const size_t local_item_size = 256;
1879
1880 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1881 NULL, &global_item_size, &local_item_size,
1882 0, NULL, NULL));
1883 CL_CHECK(clReleaseKernel(kernel));
1884}
1885
1890void opencl_pwmin_sca2(void *a, real *c, int *n, cl_command_queue cmd_queue) {
1891 cl_int err;
1892
1893 if (math_program == NULL)
1894 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1895
1896 cl_kernel kernel = clCreateKernel(math_program, "pwmin_sca2_kernel", &err);
1897 CL_CHECK(err);
1898
1899 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1900 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(real), c));
1901 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), n));
1902
1903 const int nb = ((*n) + 256 - 1) / 256;
1904 const size_t global_item_size = 256 * nb;
1905 const size_t local_item_size = 256;
1906
1907 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1908 NULL, &global_item_size, &local_item_size,
1909 0, NULL, NULL));
1910 CL_CHECK(clReleaseKernel(kernel));
1911}
1912
1917void opencl_pwmin_sca3(void *a, void *b, real *c,
1918 int *n, cl_command_queue cmd_queue) {
1919 cl_int err;
1920
1921 if (math_program == NULL)
1922 opencl_kernel_jit(math_kernel, (cl_program *) &math_program);
1923
1924 cl_kernel kernel = clCreateKernel(math_program, "pwmin_sca3_kernel", &err);
1925 CL_CHECK(err);
1926
1927 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a));
1928 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &b));
1929 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), c));
1930 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), n));
1931
1932 const int nb = ((*n) + 256 - 1) / 256;
1933 const size_t global_item_size = 256 * nb;
1934 const size_t local_item_size = 256;
1935
1936 CL_CHECK(clEnqueueNDRangeKernel(cmd_queue, kernel, 1,
1937 NULL, &global_item_size, &local_item_size,
1938 0, NULL, NULL));
1939 CL_CHECK(clReleaseKernel(kernel));
1940}
opencl_buffer_reserve
void opencl_buffer_reserve(opencl_buffer_t *buf, size_t size)
Definition buffer.c:46
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
i
const int i
Definition ax_helm_full_kernel.h:95
e
const int e
Definition ax_helm_full_kernel.h:93
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
j
const int j
Definition ax_helm_full_kernel.h:94
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:106
w3
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
Definition cdtp_kernel.h:113
device_config.h
real
double real
Definition device_config.h:12
glb_ctx
void * glb_ctx
real_xp
double real_xp
Definition device_config.h:16
opencl_kernel_jit
void opencl_kernel_jit(const char *kernel, cl_program *program)
Definition jit.c:50
opencl_iadd
void opencl_iadd(void *a, int *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1695
opencl_col3
void opencl_col3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1028
opencl_cdiv
void opencl_cdiv(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:394
opencl_masked_scatter_copy
void opencl_masked_scatter_copy(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:225
opencl_pwmax_sca2
void opencl_pwmax_sca2(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1778
opencl_vcross
void opencl_vcross(void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, void *w1, void *w2, void *w3, int *n, cl_command_queue cmd_queue)
Definition math.c:1267
opencl_masked_gather_copy
void opencl_masked_gather_copy(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:125
opencl_sub2
void opencl_sub2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1086
opencl_face_masked_gather_copy
void opencl_face_masked_gather_copy(void *a, void *b, void *mask, void *facet, int *n1, int *n2, int *lx, int *ly, int *lz, int *m, cl_command_queue cmd_queue)
Definition math.c:187
opencl_power
void opencl_power(void *ap, void *a, real *p, int *n, cl_command_queue cmd_queue)
Definition math.c:555
opencl_col2
void opencl_col2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1001
opencl_add2s2_many
void opencl_add2s2_many(void *x, void *p, void *alpha, int *j, int *n, cl_command_queue cmd_queue)
Definition math.c:758
opencl_sub3
void opencl_sub3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1113
opencl_add2s1
void opencl_add2s1(void *a, void *b, real *c1, int *n, cl_command_queue cmd_queue)
Definition math.c:697
opencl_addcol3s2
void opencl_addcol3s2(void *a, void *b, void *c, real *s, int *n, cl_command_queue cmd_queue)
Definition math.c:1201
opencl_invcol1
void opencl_invcol1(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:919
opencl_add3
void opencl_add3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:637
opencl_rone
void opencl_rone(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:328
opencl_invcol3
void opencl_invcol3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:972
opencl_add5s4
void opencl_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1, real *c2, real *c3, real *c4, int *n, cl_command_queue cmd_queue)
Definition math.c:883
opencl_add4s3
void opencl_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2, real *c3, int *n, cl_command_queue cmd_queue)
Definition math.c:850
opencl_cmult
void opencl_cmult(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:340
opencl_cfill_mask
void opencl_cfill_mask(void *a, void *c, int *size, void *mask, int *mask_size, cl_command_queue cmd_queue)
Definition math.c:287
opencl_cadd2
void opencl_cadd2(void *a, void *b, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:474
opencl_masked_scatter_copy_aligned
void opencl_masked_scatter_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:256
opencl_pwmin_sca3
void opencl_pwmin_sca3(void *a, void *b, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1917
opencl_pwmax_vec3
void opencl_pwmax_vec3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1749
redbuf_xp
opencl_buffer_t redbuf_xp
Definition math.c:1306
opencl_glsc3
real_xp opencl_glsc3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1312
opencl_add4
void opencl_add4(void *a, void *b, void *c, void *d, int *n, cl_command_queue cmd_queue)
Definition math.c:666
opencl_glsubnorm2
real_xp opencl_glsubnorm2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1476
opencl_pwmin_vec3
void opencl_pwmin_vec3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1861
opencl_radd
void opencl_radd(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:448
opencl_add2s2
void opencl_add2s2(void *a, void *b, real *c1, int *n, cl_command_queue cmd_queue)
Definition math.c:727
opencl_vdot3
void opencl_vdot3(void *dot, void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, int *n, cl_command_queue cmd_queue)
Definition math.c:1232
opencl_pwmin_vec2
void opencl_pwmin_vec2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1834
opencl_add3s2
void opencl_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, cl_command_queue cmd_queue)
Definition math.c:819
opencl_masked_copy_aligned
void opencl_masked_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:95
opencl_addsqr2s2
void opencl_addsqr2s2(void *a, void *b, real *c1, int *n, cl_command_queue cmd_queue)
Definition math.c:790
opencl_absval
void opencl_absval(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:1670
opencl_cwrap
void opencl_cwrap(void *a, real *min_val, real *max_val, int *n, cl_command_queue cmd_queue)
Definition math.c:502
opencl_glsc2
real_xp opencl_glsc2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1427
opencl_addcol3
void opencl_addcol3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1142
opencl_glmin
real opencl_glmin(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:1617
opencl_pwmax_vec2
void opencl_pwmax_vec2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:1722
opencl_sqrt_inplace
void opencl_sqrt_inplace(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:530
opencl_rzero
void opencl_rzero(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:316
opencl_copy
void opencl_copy(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:55
opencl_subcol3
void opencl_subcol3(void *a, void *b, void *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1057
opencl_add2
void opencl_add2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:610
opencl_addcol4
void opencl_addcol4(void *a, void *b, void *c, void *d, int *n, cl_command_queue cmd_queue)
Definition math.c:1171
opencl_pwmin_sca2
void opencl_pwmin_sca2(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1890
opencl_pwmax_sca3
void opencl_pwmax_sca3(void *a, void *b, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:1805
opencl_invcol2
void opencl_invcol2(void *a, void *b, int *n, cl_command_queue cmd_queue)
Definition math.c:945
opencl_glsc3_many
void opencl_glsc3_many(real_xp *h, void *w, void *v, void *mult, int *j, int *n, cl_command_queue cmd_queue)
Definition math.c:1363
opencl_cdiv2
void opencl_cdiv2(void *a, void *b, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:420
opencl_masked_copy_0
void opencl_masked_copy_0(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:65
opencl_cfill
void opencl_cfill(void *a, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:583
opencl_glsum
real_xp opencl_glsum(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:1525
opencl_masked_gather_copy_aligned
void opencl_masked_gather_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cl_command_queue cmd_queue)
Definition math.c:156
opencl_glmax
real opencl_glmax(void *a, int *n, cl_command_queue cmd_queue)
Definition math.c:1568
opencl_cmult2
void opencl_cmult2(void *a, void *b, real *c, int *n, cl_command_queue cmd_queue)
Definition math.c:366
mask
Object for handling masks in Neko.
Definition mask.f90:34
OPENCL_BUFFER_INIT
#define OPENCL_BUFFER_INIT
Definition buffer.h:63
CL_CHECK
#define CL_CHECK(err)
Definition check.h:12
prgm_lib.h
math_program
void * math_program
opencl_buffer
Definition buffer.h:54
opencl_buffer::host
void * host
Definition buffer.h:55
opencl_buffer::dev
cl_mem dev
Definition buffer.h:56