1 /*
2 * Copyright (c) 2025, Oracle and/or its affiliates. All rights reserved.
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
4 *
5 * This code is free software; you can redistribute it and/or modify it
6 * under the terms of the GNU General Public License version 2 only, as
7 * published by the Free Software Foundation. Oracle designates this
8 * particular file as subject to the "Classpath" exception as provided
9 * by Oracle in the LICENSE file that accompanied this code.
10 *
11 * This code is distributed in the hope that it will be useful, but WITHOUT
12 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
13 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
14 * version 2 for more details (a copy is included in the LICENSE file that
15 * accompanied this code).
16 *
17 * You should have received a copy of the GNU General Public License version
18 * 2 along with this work; if not, write to the Free Software Foundation,
19 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
20 *
21 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
22 * or visit www.oracle.com if you need additional information or have any
23 * questions.
24 */
25 package hat.test;
26
27 import hat.Accelerator;
28 import hat.ComputeContext;
29 import hat.NDRange;
30
31 import static hat.KernelContext.*;
32 import hat.backend.Backend;
33 import hat.buffer.*;
34 import hat.device.DeviceSchema;
35 import hat.device.NonMappableIface;
36 import optkl.ifacemapper.BoundSchema;
37 import optkl.ifacemapper.Buffer;
38 import optkl.ifacemapper.Schema;
39 import jdk.incubator.code.Reflect;
40 import hat.test.annotation.HatTest;
41 import hat.test.exceptions.HATAsserts;
42
43 import java.lang.foreign.ValueLayout;
44 import java.lang.invoke.MethodHandles;
45 import java.util.Random;
46
47 import static java.lang.foreign.ValueLayout.JAVA_BYTE;
48
49 public class TestArrayView {
50
51 /*
52 * simple square kernel example using S32Array's ArrayView
53 */
54 @Reflect
55 public static void squareKernel(S32Array s32Array) {
56 if (GIX() < GSX()){
57 int[] arr = s32Array.arrayView();
58 arr[GIX()] *= arr[GIX()];
59 }
60 }
61
62 @Reflect
63 public static void square(ComputeContext cc, S32Array s32Array) {
64 cc.dispatchKernel(NDRange.of1D(s32Array.length()), () -> squareKernel( s32Array));
65 }
66
67 @HatTest
68 @Reflect
69 public static void testSquare() {
70
71 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);
72 var arr = S32Array.create(accelerator, 32);
73 for (int i = 0; i < arr.length(); i++) {
74 arr.array(i, i);
75 }
76 accelerator.compute(
77 cc -> square(cc, arr)
78 );
79 for (int i = 0; i < arr.length(); i++) {
80 HATAsserts.assertEquals(i * i, arr.array(i));
81 }
82 }
83
84 /*
85 * making sure arrayviews aren't reliant on varOps
86 */
87 @Reflect
88 public static void squareKernelNoVarOp( S32Array s32Array) {
89 if (GIX()<GSX()){
90 s32Array.arrayView()[GIX()] *= s32Array.arrayView()[GIX()];
91 }
92 }
93
94 @Reflect
95 public static void squareNoVarOp(ComputeContext cc, S32Array s32Array) {
96 cc.dispatchKernel(NDRange.of1D(s32Array.length()),
97 ()-> squareKernelNoVarOp( s32Array)
98 );
99 }
100
101 @HatTest
102 @Reflect
103 public static void testSquareNoVarOp() {
104 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);
105 var arr = S32Array.create(accelerator, 32);
106 for (int i = 0; i < arr.length(); i++) {
107 arr.array(i, i);
108 }
109 accelerator.compute(
110 cc -> squareNoVarOp(cc, arr)
111 );
112 for (int i = 0; i < arr.length(); i++) {
113 HATAsserts.assertEquals(i * i, arr.array(i));
114 }
115 }
116
117 @Reflect
118 public static void square2DKernel( S32Array2D s32Array2D) {
119 if (GIX() < GSX()){
120 int[][] arr = s32Array2D.arrayView();
121 arr[GIX()][GIY()] *= arr[GIX()][GIY()];
122 }
123 }
124
125 @Reflect
126 public static void square2D(ComputeContext cc, S32Array2D s32Array2D) {
127 cc.dispatchKernel(NDRange.of1D(s32Array2D.width() * s32Array2D.height()),
128 ()-> square2DKernel( s32Array2D)
129 );
130 }
131
132 @HatTest
133 @Reflect
134 public static void testSquare2D() {
135
136 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
137 var arr = S32Array2D.create(accelerator, 5, 5);
138 for (int i = 0; i < arr.height(); i++) {
139 for (int j = 0; j < arr.width(); j++) {
140 arr.set(i, j, i * 5 + j);
141 }
142 }
143 accelerator.compute(
144 cc -> square2D(cc, arr)
145 );
146 for (int i = 0; i < arr.height(); i++) {
147 for (int j = 0; j < arr.width(); j++) {
148 HATAsserts.assertEquals((i * 5 + j) * (i * 5 + j), arr.get(i, j));
149 }
150 }
151 }
152
153 /*
154 * simplified version of Game of Life using ArrayView
155 */
156 public static final byte ALIVE = (byte) 0xff;
157 public static final byte DEAD = 0x00;
158
159 public interface CellGrid extends Buffer {
160 /*
161 * struct CellGrid{
162 * int width;
163 * int height;
164 * byte[width*height*2] cellArray;
165 * }
166 */
167 int width();
168
169 int height();
170
171 byte array(long idx);
172
173 void array(long idx, byte b);
174
175 Schema<CellGrid> schema = Schema.of(CellGrid.class, lifeData -> lifeData
176 .arrayLen("width", "height").stride(2).array("array")
177 );
178
179 static CellGrid create(Accelerator accelerator, int width, int height) {
180 return BoundSchema.of(accelerator ,schema, width, height).allocate();
181 }
182
183 ValueLayout valueLayout = JAVA_BYTE;
184
185 default byte[][] arrayView() {
186 return null;
187 }
188 }
189
190 public interface Control extends Buffer {
191 /*
192 * struct Control{
193 * int from;
194 * int to;
195 * }
196 */
197 int from();
198
199 void from(int from);
200
201 int to();
202
203 void to(int to);
204
205 Schema<Control> schema = Schema.of(
206 Control.class, control ->
207 control.fields("from", "to"));//, "generation", "requiredFrameRate", "maxGenerations"));
208
209 static Control create(Accelerator accelerator, CellGrid cellGrid) {
210 var instance = BoundSchema.of(accelerator ,schema).allocate();
211 instance.from(cellGrid.width() * cellGrid.height());
212 instance.to(0);
213 return instance;
214 }
215 }
216
217 public static byte[][] lifeCheck(CellGrid cellGrid) {
218 int w = cellGrid.width();
219 int h = cellGrid.height();
220
221 byte[][] res = new byte[h][w];
222
223 for (int y = 0; y < h; y++) {
224 for (int x = 0; x < w; x++) {
225 int idx = y * w + x;
226 byte cell = cellGrid.array(idx);
227 if (x > 0 && x < (w - 1) && y > 0 && y < (h - 1)) { // passports please
228 int count =
229 (cellGrid.array((y - 1) * w + (x - 1)) & 1)
230 + (cellGrid.array((y + 0) * w + (x - 1)) & 1)
231 + (cellGrid.array((y + 1) * w + (x - 1)) & 1)
232 + (cellGrid.array((y - 1) * w + (x + 0)) & 1)
233 + (cellGrid.array((y + 1) * w + (x + 0)) & 1)
234 + (cellGrid.array((y - 1) * w + (x + 1)) & 1)
235 + (cellGrid.array((y + 0) * w + (x + 1)) & 1)
236 + (cellGrid.array((y + 1) * w + (x + 1))& 1);
237 cell = ((count == 3) || ((count == 2) && (cell == ALIVE))) ? ALIVE : DEAD;// B3/S23.
238 }
239 res[x][y] = cell;
240 }
241 }
242 return res;
243 }
244
245 public static class Compute {
246 @Reflect
247 // TODO: switch cellGridRes to WO
248 public static void lifePerIdx(int idx, CellGrid cellGrid, CellGrid cellGridRes) {
249 int w = cellGrid.width();
250 int h = cellGrid.height();
251 int x = idx % w;
252 int y = idx / w;
253
254 // byte[] bytes = cellGrid.arrayView();
255 // byte cell = bytes[idx + from];
256 // byte[] lookup = new byte[]{};
257 // if (x > 0 && x < (w - 1) && y > 0 && y < (h - 1)) { // passports please
258 // int lookupIdx =
259 // (bytes[(y - 1) * w + x - 1 + from]&1 <<0)
260 // |(bytes[(y + 0) * w + x - 1 + from]&1 <<1)
261 // |(bytes[(y + 1) * w + x - 1 + from]&1 <<2)
262 // |(bytes[(y - 1) * w + x + 0 + from]&1 <<3)
263 // |(bytes[(y - 0) * w + x + 0 + from]&1 <<4) // current cell added
264 // |(bytes[(y + 1) * w + x + 0 + from]&1 <<5)
265 // |(bytes[(y + 0) * w + x + 1 + from]&1 <<6)
266 // |(bytes[(y - 1) * w + x + 1 + from]&1 <<7)
267 // |(bytes[(y + 1) * w + x + 1 + from]&1 <<8) ;
268 // // conditional removed!
269 // bytes[idx + to] = lookup[lookupIdx];
270 // }
271
272 byte[][] bytes = cellGrid.arrayView();
273 byte cell = bytes[x][y];
274 if (x > 0 && x < (w - 1) && y > 0 && y < (h - 1)) { // passports please
275 int count =
276 (bytes[x - 1][y - 1] & 1)
277 + (bytes[x - 1][y + 0] & 1)
278 + (bytes[x - 1][y + 1] & 1)
279 + (bytes[x + 0][y - 1] & 1)
280 + (bytes[x + 0][y + 1] & 1)
281 + (bytes[x + 1][y - 1] & 1)
282 + (bytes[x + 1][y + 0] & 1)
283 + (bytes[x + 1][y + 1] & 1);
284 cell = ((count == 3) || ((count == 2) && (cell == ALIVE))) ? ALIVE : DEAD;// B3/S23.
285 }
286 byte[][] res = cellGridRes.arrayView();
287 res[x][y] = cell;
288 }
289
290 @Reflect
291 public static void life( CellGrid cellGrid, CellGrid cellGridRes) {
292 if (GIX() < GSX()) {
293 Compute.lifePerIdx(GIX(), cellGrid, cellGridRes);
294 }
295 }
296
297 @Reflect
298 public static void compute(final ComputeContext cc, CellGrid grid, CellGrid gridRes) {
299 int range = grid.width() * grid.height();
300 cc.dispatchKernel(NDRange.of1D(range),()-> Compute.life( grid, gridRes));
301 }
302 }
303
304 @HatTest
305 @Reflect
306 public static void testLife() {
307 Accelerator accelerator = new Accelerator(MethodHandles.lookup());
308
309 // int w = 20;
310 // int h = 20;
311 // // We oversize the grid by adding 1 to n,e,w and s
312 // CellGrid cellGrid = CellGrid.create(accelerator, w, h);
313 // CellGrid cellGridRes = CellGrid.create(accelerator, w, h);
314 //
315 // Random rand = new Random();
316 // byte[][] actualGrid = new byte[w][h];
317 // for (int y = 0; y < h; y++) {
318 // for (int x = 0; x < w; x++) {
319 // actualGrid[x][y] = rand.nextBoolean() ? ALIVE : DEAD;
320 // }
321 // }
322
323 // We oversize the grid by adding 1 to n,e,w and s
324 CellGrid cellGrid = CellGrid.create(accelerator, 17, 17);
325 CellGrid cellGridRes = CellGrid.create(accelerator, 17, 17);
326
327 byte[][] actualGrid = new byte[][]{
328 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
329 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
330 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
331 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
332 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
333 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
334 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
335 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
336 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
337 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
338 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
339 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
340 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
341 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
342 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
343 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
344 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
345 };
346
347 // By shifting all cells +1,+1 so we only need to scan 1..width-1, 1..height-1
348 // we don't worry about possibly finding cells in 0,n width,n or n,0 height,n
349 for (int i = 0; i < cellGrid.height(); i++) {
350 for (int j = 0; j < cellGrid.width(); j++) {
351 cellGrid.array(((long) i * cellGrid.width()) + j, actualGrid[i][j]);
352 }
353 }
354
355 // Control control = Control.create(accelerator, cellGrid);
356
357 accelerator.compute(cc -> Compute.compute(cc, cellGrid, cellGridRes));
358
359 byte[][] resultGrid = lifeCheck(cellGrid);
360
361 for (int i = 0; i < cellGrid.height(); i++) {
362 for (int j = 0; j < cellGrid.width(); j++) {
363 HATAsserts.assertEquals(resultGrid[i][j], cellGridRes.array(((long) i * cellGrid.width()) + j));
364 }
365 }
366 }
367
368 /*
369 * simplified version of mandel using ArrayView
370 */
371 @Reflect
372 public static int mandelCheck(int i, int j, float width, float height, int[] pallette, float offsetx, float offsety, float scale) {
373 float x = (i * scale - (scale / 2f * width)) / width + offsetx;
374 float y = (j * scale - (scale / 2f * height)) / height + offsety;
375 float zx = x;
376 float zy = y;
377 float new_zx;
378 int colorIdx = 0;
379 while ((colorIdx < pallette.length) && (((zx * zx) + (zy * zy)) < 4f)) {
380 new_zx = ((zx * zx) - (zy * zy)) + x;
381 zy = (2f * zx * zy) + y;
382 zx = new_zx;
383 colorIdx++;
384 }
385 return colorIdx < pallette.length ? pallette[colorIdx] : 0;
386 }
387
388 @Reflect
389 public static void mandel( S32Array2D s32Array2D, S32Array pallette, float offsetx, float offsety, float scale) {
390 if (GIX() < GSX()) {
391 int[] pal = pallette.arrayView();
392 int[][] s32 = s32Array2D.arrayView();
393 float width = s32Array2D.width();
394 float height = s32Array2D.height();
395 float x = ((GIX() % s32Array2D.width()) * scale - (scale / 2f * width)) / width + offsetx;
396 float y = ((GIX() / s32Array2D.width()) * scale - (scale / 2f * height)) / height + offsety;
397 float zx = x;
398 float zy = y;
399 float new_zx;
400 int colorIdx = 0;
401 while ((colorIdx < pal.length) && (((zx * zx) + (zy * zy)) < 4f)) {
402 new_zx = ((zx * zx) - (zy * zy)) + x;
403 zy = (2f * zx * zy) + y;
404 zx = new_zx;
405 colorIdx++;
406 }
407 int color = colorIdx < pal.length ? pal[colorIdx] : 0;
408 s32[GIX() % s32Array2D.width()][GIX() / s32Array2D.width()] = color;
409 }
410 }
411
412
413 @Reflect
414 static public void compute(final ComputeContext computeContext, S32Array pallete, S32Array2D s32Array2D, float x, float y, float scale) {
415
416 computeContext.dispatchKernel(
417 NDRange.of1D(s32Array2D.width()*s32Array2D.height()), //0..S32Array2D.size()
418 ()-> mandel( s32Array2D, pallete, x, y, scale));
419 }
420
421 @HatTest
422 @Reflect
423 public static void testMandel() {
424 final int width = 1024;
425 final int height = 1024;
426 final float defaultScale = 3f;
427 final float originX = -1f;
428 final float originY = 0;
429 final int maxIterations = 64;
430
431 Accelerator accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);
432
433 S32Array2D s32Array2D = S32Array2D.create(accelerator, width, height);
434
435 int[] palletteArray = new int[maxIterations];
436
437 for (int i = 1; i < maxIterations; i++) {
438 palletteArray[i]=(i/8+1);
439 }
440 palletteArray[0]=0;
441 S32Array pallette = S32Array.createFrom(accelerator, palletteArray);
442
443 accelerator.compute(cc -> compute(cc, pallette, s32Array2D, originX, originY, defaultScale));
444
445 int subsample = 16;
446 char[] charPallette9 = new char []{' ', '.', ',',':', '-', '+','*', '#', '@', '%'};
447 for (int y = 0; y<height/subsample; y++) {
448 for (int x = 0; x<width/subsample; x++) {
449 int palletteValue = s32Array2D.get(x*subsample,y*subsample); // so 0->8
450 int paletteCheck = mandelCheck(x*subsample, y*subsample, width, height, palletteArray, originX, originY, defaultScale);
451 HATAsserts.assertEquals(paletteCheck, palletteValue);
452 }
453 }
454 }
455
456 /*
457 * simplified version of BlackScholes using ArrayView
458 */
459 @Reflect
460 public static void blackScholesKernel(
461 F32Array call,
462 F32Array put,
463 F32Array sArray,
464 F32Array xArray,
465 F32Array tArray,
466 float r,
467 float v) {
468 if (GIX()<GSX()){
469 float[] callArr = call.arrayView();
470 float[] putArr = put.arrayView();
471 float[] sArr = sArray.arrayView();
472 float[] xArr = xArray.arrayView();
473 float[] tArr = tArray.arrayView();
474
475 float expNegRt = (float) Math.exp(-r * tArr[GIX()]);
476 float d1 = (float) ((Math.log(sArr[GIX()] / xArr[GIX()]) + (r + v * v * .5f) * tArr[GIX()]) / (v * Math.sqrt(tArr[GIX()])));
477 float d2 = (float) (d1 - v * Math.sqrt(tArr[GIX()]));
478 float cnd1 = CND(d1);
479 float cnd2 = CND(d2);
480 float value = sArr[GIX()] * cnd1 - expNegRt * xArr[GIX()] * cnd2;
481 callArr[GIX()] = value;
482 putArr[GIX()] = expNegRt * xArr[GIX()] * (1 - cnd2) - sArr[GIX()] * (1 - cnd1);
483 }
484 }
485
486 @Reflect
487 public static float CND(float input) {
488 float x = input;
489 if (input < 0f) { // input = Math.abs(input)?
490 x = -input;
491 }
492
493 float term = 1f / (1f + (0.2316419f * x));
494 float term_pow2 = term * term;
495 float term_pow3 = term_pow2 * term;
496 float term_pow4 = term_pow2 * term_pow2;
497 float term_pow5 = term_pow2 * term_pow3;
498
499 float part1 = (1f / (float)Math.sqrt(2f * 3.1415926535f)) * (float)Math.exp((-x * x) * 0.5f);
500
501 float part2 = (0.31938153f * term) +
502 (-0.356563782f * term_pow2) +
503 (1.781477937f * term_pow3) +
504 (-1.821255978f * term_pow4) +
505 (1.330274429f * term_pow5);
506
507 if (input >= 0f) {
508 return 1f - part1 * part2;
509 }
510 return part1 * part2;
511
512 }
513
514 @Reflect
515 public static void blackScholes(ComputeContext cc, F32Array call, F32Array put, F32Array S, F32Array X, F32Array T, float r, float v) {
516 cc.dispatchKernel(NDRange.of1D(call.length()),
517 ()-> blackScholesKernel(call, put, S, X, T, r, v)
518 );
519 }
520
521 static F32Array floatArray(Accelerator accelerator, int size, float low, float high, Random rand) {
522 F32Array array = F32Array.create(accelerator, size);
523 for (int i = 0; i <size; i++) {
524 array.array(i, rand.nextFloat() * (high - low) + low);
525 }
526 return array;
527 }
528
529 public static void blackScholesKernelSeq(F32Array call, F32Array put, F32Array sArray, F32Array xArray, F32Array tArray, float r, float v) {
530 for (int i = 0; i <call.length() ; i++) {
531 float S = sArray.array(i);
532 float X = xArray.array(i);
533 float T = tArray.array(i);
534 float expNegRt = (float) Math.exp(-r * T);
535 float d1 = (float) ((Math.log(S / X) + (r + v * v * .5f) * T) / (v * Math.sqrt(T)));
536 float d2 = (float) (d1 - v * Math.sqrt(T));
537 float cnd1 = CND(d1);
538 float cnd2 = CND(d2);
539 float value = S * cnd1 - expNegRt * X * cnd2;
540 call.array(i, value);
541 put.array(i, expNegRt * X * (1 - cnd2) - S * (1 - cnd1));
542 }
543 }
544
545 @HatTest
546 @Reflect
547 public static void testBlackScholes() {
548 int size = 1024;
549 Random rand = new Random();
550 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);
551 var call = F32Array.create(accelerator, size);
552 var put = F32Array.create(accelerator, size);
553 for (int i = 0; i < size; i++) {
554 call.array(i, i);
555 put.array(i, i);
556 }
557
558 var S = floatArray(accelerator, size,1f, 100f, rand);
559 var X = floatArray(accelerator, size,1f, 100f, rand);
560 var T = floatArray(accelerator,size, 0.25f, 10f, rand);
561 float r = 0.02f;
562 float v = 0.30f;
563
564 accelerator.compute(cc -> blackScholes(cc, call, put, S, X, T, r, v));
565
566 var seqCall = F32Array.create(accelerator, size);
567 var seqPut = F32Array.create(accelerator, size);
568 for (int i = 0; i < seqCall.length(); i++) {
569 seqCall.array(i, i);
570 seqPut.array(i, i);
571 }
572
573 blackScholesKernelSeq(seqCall, seqPut, S, X, T, r, v);
574
575 for (int i = 0; i < call.length(); i++) {
576 HATAsserts.assertEquals(seqCall.array(i), call.array(i), 0.01f);
577 HATAsserts.assertEquals(seqPut.array(i), put.array(i), 0.01f);
578 }
579 }
580
581 /*
582 * basic test of local and private buffer ArrayViews
583 */
584 private interface SharedMemory extends NonMappableIface {
585 void array(long index, int value);
586 int array(long index);
587 DeviceSchema<SharedMemory> deviceSchema = DeviceSchema.of(SharedMemory.class,
588 arr -> arr.array("array", 1024));
589
590 static SharedMemory createLocal() { return null; }
591
592 default int[] localArrayView() {
593 int[] view = new int[1024];
594 for (int i = 0; i < 1024; i++) {
595 view[i] = this.array(i);
596 }
597 return view;
598 }
599 }
600
601 public interface PrivateArray extends NonMappableIface {
602 void array(long index, int value);
603 int array(long index);
604 DeviceSchema<PrivateArray> deviceSchema = DeviceSchema.of(PrivateArray.class,
605 arr -> arr.array("array", 16));
606
607 static PrivateArray createPrivate() { return null; }
608
609 default int[] privateArrayView() {
610 int[] view = new int[16];
611 for (int i = 0; i < 16; i++) {
612 view[i] = this.array(i);
613 }
614 return view;
615 }
616 }
617
618 @Reflect
619 public static void squareKernelWithPrivateAndLocal( S32Array s32Array) {
620 SharedMemory shared = SharedMemory.createLocal();
621 if (GIX() < GSX()){
622 int[] arr = s32Array.arrayView();
623 arr[GIX()] += arr[GIX()];
624
625 PrivateArray priv = PrivateArray.createPrivate();
626 int[] privView = priv.privateArrayView();
627 privView[0] = 1;
628 arr[GIX()] += privView[0];
629
630 int[] sharedView = shared.localArrayView();
631 sharedView[0] = 16;
632 barrier();
633 arr[GIX()] += sharedView[0];
634 }
635 }
636
637 @Reflect
638 public static void privateAndLocal(ComputeContext cc, S32Array s32Array) {
639 cc.dispatchKernel(NDRange.of1D(s32Array.length()),
640 ()-> squareKernelWithPrivateAndLocal( s32Array)
641 );
642 }
643
644 @HatTest
645 @Reflect
646 public static void testPrivateAndLocal() {
647
648 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
649 var arr = S32Array.create(accelerator, 32);
650 for (int i = 0; i < arr.length(); i++) {
651 arr.array(i, i);
652 }
653 accelerator.compute(
654 cc -> privateAndLocal(cc, arr)
655 );
656 for (int i = 0; i < arr.length(); i++) {
657 HATAsserts.assertEquals(2 * i + 17, arr.array(i));
658 }
659 }
660
661 /*
662 * testing basic DeviceTypes
663 */
664
665 public interface SharedNonMappableIface extends NonMappableIface {
666 void array(long index, int value);
667 int array(long index);
668 DeviceSchema<SharedNonMappableIface> deviceSchema = DeviceSchema.of(SharedNonMappableIface.class,
669 arr -> arr.array("array", 1024));
670 static SharedNonMappableIface create(Accelerator accelerator) {
671 return null;
672 }
673 static SharedNonMappableIface createLocal() {
674 return null;
675 }
676
677 default int[] localArrayView() {
678 return null;
679 }
680 }
681
682 public interface PrivateNonMappableIface extends NonMappableIface {
683 void array(long index, int value);
684 int array(long index);
685 DeviceSchema<PrivateNonMappableIface> deviceSchema = DeviceSchema.of(PrivateNonMappableIface.class,
686 arr -> arr.array("array", 32));
687 static PrivateNonMappableIface create(Accelerator accelerator) {
688 return null;
689 }
690 static PrivateNonMappableIface createPrivate() {
691 return null;
692 }
693
694 default int[] privateArrayView() {
695 return null;
696 }
697 }
698
699 @Reflect
700 public static void kernelBasicDeviceType( S32Array s32Array) {
701 SharedNonMappableIface shared = SharedNonMappableIface.createLocal();
702 if (GIX() < GSX()){
703 PrivateNonMappableIface priv = PrivateNonMappableIface.createPrivate();
704
705 int[] arr = s32Array.arrayView();
706 int[] privView = priv.privateArrayView();
707 int[] sharedView = shared.localArrayView();
708
709 privView[GIX()] = arr[GIX()];
710 sharedView[GIX()] = arr[GIX()];
711 barrier();
712 arr[GIX()] = privView[GIX()] + sharedView[GIX()];
713 }
714 }
715
716 @Reflect
717 public static void basicDeviceType(ComputeContext cc, S32Array s32Array) {
718 cc.dispatchKernel(NDRange.of1D(s32Array.length()),
719 ()-> kernelBasicDeviceType( s32Array)
720 );
721 }
722
723 @HatTest
724 @Reflect
725 public static void testBasicDeviceType() {
726 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
727 var arr = S32Array.create(accelerator, 32);
728 for (int i = 0; i < arr.length(); i++) {
729 arr.array(i, i);
730 }
731 accelerator.compute(cc -> basicDeviceType(cc, arr));
732 for (int i = 0; i < arr.length(); i++) {
733 HATAsserts.assertEquals(2 * i, arr.array(i));
734 }
735 }
736
737 @Reflect
738 public static void squareKernelDeviceType( S32Array s32Array) {
739 SharedNonMappableIface shared = SharedNonMappableIface.createLocal();
740 if (GIX() < GSX()){
741 PrivateNonMappableIface priv = PrivateNonMappableIface.createPrivate();
742
743 int[] arr = s32Array.arrayView();
744 int[] privView = priv.privateArrayView();
745 int[] sharedView = shared.localArrayView();
746
747 privView[GIX()] = arr[GIX()];
748 sharedView[privView[GIX()]] = 16 * privView[GIX()];
749 barrier();
750 arr[GIX()] += privView[GIX()] + sharedView[GIX()];
751 }
752 }
753
754 @Reflect
755 public static void deviceType(ComputeContext cc, S32Array s32Array) {
756 cc.dispatchKernel(NDRange.of1D(s32Array.length()),
757 ()-> squareKernelDeviceType( s32Array)
758 );
759 }
760
761 @HatTest
762 @Reflect
763 public static void testDeviceType() {
764 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST); //new JavaMultiThreadedBackend());
765 var arr = S32Array.create(accelerator, 32);
766 for (int i = 0; i < arr.length(); i++) {
767 arr.array(i, i);
768 }
769 accelerator.compute(cc -> deviceType(cc, arr));
770 for (int i = 0; i < arr.length(); i++) {
771 HATAsserts.assertEquals(18 * i, arr.array(i));
772 }
773 }
774 }