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.KernelContext;
30 import hat.backend.Backend;
31 import hat.buffer.*;
32 import hat.ifacemapper.MappableIface.*;
33 import hat.ifacemapper.Schema;
34 import jdk.incubator.code.CodeReflection;
35 import hat.test.annotation.HatTest;
36 import hat.test.engine.HatAsserts;
37
38 import java.lang.foreign.ValueLayout;
39 import java.lang.invoke.MethodHandles;
40 import java.util.Random;
41
42 import static java.lang.foreign.ValueLayout.JAVA_BYTE;
43
44 public class TestArrayView {
45
46 // simple square kernel example using S32Array's ArrayView
47 @CodeReflection
48 public static void squareKernel(@RO KernelContext kc, @RW S32Array s32Array) {
49 if (kc.x<kc.maxX){
50 int[] arr = s32Array.arrayView();
51 arr[kc.x] *= arr[kc.x];
52 }
53 }
54
55 @CodeReflection
56 public static void square(@RO ComputeContext cc, @RW S32Array s32Array) {
57 cc.dispatchKernel(s32Array.length(),
58 kc -> squareKernel(kc, s32Array)
59 );
60 }
61
62 @HatTest
63 public static void testSquare() {
64
65 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
66 var arr = S32Array.create(accelerator, 32);
67 for (int i = 0; i < arr.length(); i++) {
68 arr.array(i, i);
69 }
70 accelerator.compute(
71 cc -> square(cc, arr) //QuotableComputeContextConsumer
72 ); // extends Quotable, Consumer<ComputeContext>
73 for (int i = 0; i < arr.length(); i++) {
74 HatAsserts.assertEquals(i * i, arr.array(i));
75 }
76 }
77
78 @CodeReflection
79 public static void square2DKernel(@RO KernelContext kc, @RW S32Array2D s32Array2D) {
80 if (kc.x<kc.maxX){
81 int[][] arr = s32Array2D.arrayView();
82 arr[kc.x][kc.y] *= arr[kc.x][kc.y];
83 }
84 }
85
86 @CodeReflection
87 public static void square2D(@RO ComputeContext cc, @RW S32Array2D s32Array2D) {
88 cc.dispatchKernel(s32Array2D.width() * s32Array2D.height(),
89 kc -> square2DKernel(kc, s32Array2D)
90 );
91 }
92
93 @HatTest
94 public static void testSquare2D() {
95
96 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
97 var arr = S32Array2D.create(accelerator, 5, 5);
98 for (int i = 0; i < arr.height(); i++) {
99 for (int j = 0; j < arr.width(); j++) {
100 arr.set(i, j, i * 5 + j);
101 }
102 }
103 accelerator.compute(
104 cc -> square2D(cc, arr) //QuotableComputeContextConsumer
105 ); // extends Quotable, Consumer<ComputeContext>
106 for (int i = 0; i < arr.height(); i++) {
107 for (int j = 0; j < arr.width(); j++) {
108 HatAsserts.assertEquals((i * 5 + j) * (i * 5 + j), arr.get(i, j));
109 }
110 }
111 }
112
113 // simplified version of Game of Life using ArrayView
114 public final static byte ALIVE = (byte) 0xff;
115 public final static byte DEAD = 0x00;
116
117 public interface CellGrid extends Buffer {
118 /*
119 * struct CellGrid{
120 * int width;
121 * int height;
122 * byte[width*height*2] cellArray;
123 * }
124 */
125 int width();
126
127 int height();
128
129 byte array(long idx);
130
131 void array(long idx, byte b);
132
133 Schema<CellGrid> schema = Schema.of(CellGrid.class, lifeData -> lifeData
134 .arrayLen("width", "height").stride(2).array("array")
135 );
136
137 static CellGrid create(Accelerator accelerator, int width, int height) {
138 return schema.allocate(accelerator, width, height);
139 }
140
141 ValueLayout valueLayout = JAVA_BYTE;
142
143 default byte[] arrayView() {
144 int size = this.width() * this.height();
145 byte[] arr = new byte[size];
146 for (int i = 0; i < size; i++) {
147 arr[i] = this.array(i);
148 }
149 return arr;
150 }
151 }
152
153 public interface Control extends Buffer {
154 /*
155 * struct Control{
156 * int from;
157 * int to;
158 * }
159 */
160 int from();
161
162 void from(int from);
163
164 int to();
165
166 void to(int to);
167
168 Schema<Control> schema = Schema.of(
169 Control.class, control ->
170 control.fields("from", "to"));//, "generation", "requiredFrameRate", "maxGenerations"));
171
172 static Control create(Accelerator accelerator, CellGrid cellGrid) {
173 var instance = schema.allocate(accelerator);
174 instance.from(cellGrid.width() * cellGrid.height());
175 instance.to(0);
176 return instance;
177 }
178 }
179
180 public static class Compute {
181 @CodeReflection
182 public static void lifePerIdx(int idx, @RO Control control, @RW CellGrid cellGrid) {
183 int w = cellGrid.width();
184 int h = cellGrid.height();
185 int from = control.from();
186 int to = control.to();
187 int x = idx % w;
188 int y = idx / w;
189
190 // byte[] bytes = cellGrid.arrayView();
191 // byte cell = bytes[idx + from];
192 // byte[] lookup = new byte[]{};
193 // if (x > 0 && x < (w - 1) && y > 0 && y < (h - 1)) { // passports please
194 // int lookupIdx =
195 // (bytes[(y - 1) * w + x - 1 + from]&1 <<0)
196 // |(bytes[(y + 0) * w + x - 1 + from]&1 <<1)
197 // |(bytes[(y + 1) * w + x - 1 + from]&1 <<2)
198 // |(bytes[(y - 1) * w + x + 0 + from]&1 <<3)
199 // |(bytes[(y - 0) * w + x + 0 + from]&1 <<4) // current cell added
200 // |(bytes[(y + 1) * w + x + 0 + from]&1 <<5)
201 // |(bytes[(y + 0) * w + x + 1 + from]&1 <<6)
202 // |(bytes[(y - 1) * w + x + 1 + from]&1 <<7)
203 // |(bytes[(y + 1) * w + x + 1 + from]&1 <<8) ;
204 // // conditional removed!
205 // bytes[idx + to] = lookup[lookupIdx];
206 // }
207
208 byte[] bytes = cellGrid.arrayView();
209 byte cell = bytes[idx];
210 if (x > 0 && x < (w - 1) && y > 0 && y < (h - 1)) { // passports please
211 int count =
212 (bytes[(y - 1) * w + (x - 1)] & 1)
213 + (bytes[(y + 0) * w + (x - 1)] & 1)
214 + (bytes[(y + 1) * w + (x - 1)] & 1)
215 + (bytes[(y - 1) * w + (x + 0)] & 1)
216 + (bytes[(y + 1) * w + (x + 0)] & 1)
217 + (bytes[(y - 1) * w + (x + 1)] & 1)
218 + (bytes[(y + 0) * w + (x + 1)] & 1)
219 + (bytes[(y + 1) * w + (x + 1)] & 1);
220 cell = ((count == 3) || ((count == 2) && (cell == ALIVE))) ? ALIVE : DEAD;// B3/S23.
221 }
222 bytes[idx] = cell;
223 }
224
225 @CodeReflection
226 public static void life(@RO KernelContext kc, @RO Control control, @RW CellGrid cellGrid) {
227 if (kc.x < kc.maxX) {
228 Compute.lifePerIdx(kc.x, control, cellGrid);
229 }
230 }
231
232 @CodeReflection
233 static public void compute(final @RO ComputeContext cc, @RO Control ctrl, @RW CellGrid grid) {
234 int range = grid.width() * grid.height();
235 cc.dispatchKernel(range, kc -> Compute.life(kc, ctrl, grid));
236 }
237 }
238
239 @HatTest
240 public static void testLife() {
241 Accelerator accelerator = new Accelerator(MethodHandles.lookup());//,new OpenCLBackend("INFO,MINIMIZE_COPIES,SHOW_COMPUTE_MODEL"));
242
243 // We oversize the grid by adding 1 to n,e,w and s
244 CellGrid cellGrid = CellGrid.create(accelerator,
245 17,
246 17);
247
248 byte[][] actualGrid = new byte[][]{
249 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
250 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
251 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
252 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
253 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
254 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
255 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
256 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
257 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
258 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
259 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
260 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
261 {DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD},
262 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
263 {DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD},
264 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
265 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
266 };
267
268 // By shifting all cells +1,+1 so we only need to scan 1..width-1, 1..height-1
269 // we don't worry about possibly finding cells in 0,n width,n or n,0 height,n
270 for (int i = 0; i < cellGrid.height(); i++) {
271 for (int j = 0; j < cellGrid.width(); j++) {
272 cellGrid.array(((long) i * cellGrid.width()) + j, actualGrid[i][j]);
273 }
274 }
275
276 Control control = Control.create(accelerator, cellGrid);
277
278 accelerator.compute(cc -> Compute.compute(cc, control, cellGrid));
279
280 byte[][] resultGrid = new byte[][]{
281 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
282 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
283 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
284 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
285 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
286 {DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, ALIVE, ALIVE, DEAD, ALIVE, ALIVE, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD},
287 {DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD},
288 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
289 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
290 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
291 {DEAD, DEAD, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, ALIVE, DEAD, DEAD, DEAD},
292 {DEAD, ALIVE, ALIVE, ALIVE, DEAD, DEAD, ALIVE, ALIVE, DEAD, ALIVE, ALIVE, DEAD, DEAD, ALIVE, ALIVE, ALIVE, DEAD},
293 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
294 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, ALIVE, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
295 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
296 {DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD, ALIVE, DEAD, DEAD, DEAD, DEAD, DEAD},
297 {DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD, DEAD},
298 };
299
300 for (int i = 0; i < cellGrid.height(); i++) {
301 for (int j = 0; j < cellGrid.width(); j++) {
302 HatAsserts.assertEquals(resultGrid[i][j], cellGrid.array(((long) i * cellGrid.width()) + j));
303 }
304 }
305 }
306
307 // simplified version of mandel using ArrayView
308 @CodeReflection
309 public static int mandelCheck(int i, int j, float width, float height, int[] pallette, float offsetx, float offsety, float scale) {
310 float x = (i * scale - (scale / 2f * width)) / width + offsetx;
311 float y = (j * scale - (scale / 2f * height)) / height + offsety;
312 float zx = x;
313 float zy = y;
314 float new_zx;
315 int colorIdx = 0;
316 while ((colorIdx < pallette.length) && (((zx * zx) + (zy * zy)) < 4f)) {
317 new_zx = ((zx * zx) - (zy * zy)) + x;
318 zy = (2f * zx * zy) + y;
319 zx = new_zx;
320 colorIdx++;
321 }
322 return colorIdx < pallette.length ? pallette[colorIdx] : 0;
323 }
324
325 @CodeReflection
326 public static void mandel(@RO KernelContext kc, @RW S32Array2D s32Array2D, @RO S32Array pallette, float offsetx, float offsety, float scale) {
327 if (kc.x < kc.maxX) {
328 int[] pal = pallette.arrayView();
329 int[][] s32 = s32Array2D.arrayView();
330 float width = s32Array2D.width();
331 float height = s32Array2D.height();
332 float x = ((kc.x % s32Array2D.width()) * scale - (scale / 2f * width)) / width + offsetx;
333 float y = ((kc.x / s32Array2D.width()) * scale - (scale / 2f * height)) / height + offsety;
334 float zx = x;
335 float zy = y;
336 float new_zx;
337 int colorIdx = 0;
338 while ((colorIdx < pal.length) && (((zx * zx) + (zy * zy)) < 4f)) {
339 new_zx = ((zx * zx) - (zy * zy)) + x;
340 zy = (2f * zx * zy) + y;
341 zx = new_zx;
342 colorIdx++;
343 }
344 int color = colorIdx < pal.length ? pal[colorIdx] : 0;
345 s32[kc.x % s32Array2D.width()][kc.x / s32Array2D.width()] = color;
346 }
347 }
348
349
350 @CodeReflection
351 static public void compute(final ComputeContext computeContext, S32Array pallete, S32Array2D s32Array2D, float x, float y, float scale) {
352
353 computeContext.dispatchKernel(
354 s32Array2D.width()*s32Array2D.height(), //0..S32Array2D.size()
355 kc -> mandel(kc, s32Array2D, pallete, x, y, scale));
356 }
357
358 @HatTest
359 public static void testMandel() {
360 final int width = 1024;
361 final int height = 1024;
362 final float defaultScale = 3f;
363 final float originX = -1f;
364 final float originY = 0;
365 final int maxIterations = 64;
366
367 Accelerator accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);
368
369 S32Array2D s32Array2D = S32Array2D.create(accelerator, width, height);
370
371 int[] palletteArray = new int[maxIterations];
372
373 for (int i = 1; i < maxIterations; i++) {
374 palletteArray[i]=(i/8+1);// 0-7?
375 }
376 palletteArray[0]=0;
377 S32Array pallette = S32Array.createFrom(accelerator, palletteArray);
378
379 accelerator.compute(cc -> compute(cc, pallette, s32Array2D, originX, originY, defaultScale));
380
381 // Well take 1 in 4 samples (so 1024 -> 128 grid) of the pallette.
382 int subsample = 16;
383 char[] charPallette9 = new char []{' ', '.', ',',':', '-', '+','*', '#', '@', '%'};
384 for (int y = 0; y<height/subsample; y++) {
385 for (int x = 0; x<width/subsample; x++) {
386 int palletteValue = s32Array2D.get(x*subsample,y*subsample); // so 0->8
387 int paletteCheck = mandelCheck(x*subsample, y*subsample, width, height, palletteArray, originX, originY, defaultScale);
388 // System.out.print(charPallette9[palletteValue]);
389 HatAsserts.assertEquals(paletteCheck, palletteValue);
390 }
391 // System.out.println();
392 }
393 }
394
395 // simplified version of BlackScholes using ArrayView
396 @CodeReflection
397 public static float[] blackScholesCheck(float s, float x, float t, float r, float v) {
398 float expNegRt = (float) Math.exp(-r * t);
399 float d1 = (float) ((Math.log(s / x) + (r + v * v * .5f) * t) / (v * Math.sqrt(t)));
400 float d2 = (float) (d1 - v * Math.sqrt(t));
401 float cnd1 = CND(d1);
402 float cnd2 = CND(d2);
403 float call = s * cnd1 - expNegRt * x * cnd2;
404 float put = expNegRt * x * (1 - cnd2) - s * (1 - cnd1);
405 return new float[]{call, put};
406 }
407
408 @CodeReflection
409 public static void blackScholesKernel(@RO KernelContext kc,
410 @WO F32Array call,
411 @WO F32Array put,
412 @RO F32Array sArray,
413 @RO F32Array xArray,
414 @RO F32Array tArray,
415 float r,
416 float v) {
417 if (kc.x<kc.maxX){
418 float[] callArr = call.arrayView();
419 float[] putArr = put.arrayView();
420 float[] sArr = sArray.arrayView();
421 float[] xArr = xArray.arrayView();
422 float[] tArr = tArray.arrayView();
423
424 float expNegRt = (float) Math.exp(-r * tArr[kc.x]);
425 float d1 = (float) ((Math.log(sArr[kc.x] / xArr[kc.x]) + (r + v * v * .5f) * tArr[kc.x]) / (v * Math.sqrt(tArr[kc.x])));
426 float d2 = (float) (d1 - v * Math.sqrt(tArr[kc.x]));
427 float cnd1 = CND(d1);
428 float cnd2 = CND(d2);
429 float value = sArr[kc.x] * cnd1 - expNegRt * xArr[kc.x] * cnd2;
430 callArr[kc.x] = value;
431 putArr[kc.x] = expNegRt * xArr[kc.x] * (1 - cnd2) - sArr[kc.x] * (1 - cnd1);
432 }
433 }
434
435 @CodeReflection
436 public static float CND(float input) {
437 float x = input;
438 if (input < 0f) { // input = Math.abs(input)?
439 x = -input;
440 }
441
442 float term = 1f / (1f + (0.2316419f * x));
443 float term_pow2 = term * term;
444 float term_pow3 = term_pow2 * term;
445 float term_pow4 = term_pow2 * term_pow2;
446 float term_pow5 = term_pow2 * term_pow3;
447
448 float part1 = (1f / (float)Math.sqrt(2f * 3.1415926535f)) * (float)Math.exp((-x * x) * 0.5f);
449
450 float part2 = (0.31938153f * term) +
451 (-0.356563782f * term_pow2) +
452 (1.781477937f * term_pow3) +
453 (-1.821255978f * term_pow4) +
454 (1.330274429f * term_pow5);
455
456 if (input >= 0f) {
457 return 1f - part1 * part2;
458 }
459 return part1 * part2;
460
461 }
462
463 @CodeReflection
464 public static void blackScholes(@RO ComputeContext cc, @WO F32Array call, @WO F32Array put, @RO F32Array S, @RO F32Array X, @RO F32Array T, float r, float v) {
465 cc.dispatchKernel(call.length(),
466 kc -> blackScholesKernel(kc, call, put, S, X, T, r, v)
467 );
468 }
469
470 static F32Array floatArray(Accelerator accelerator, int size, float low, float high, Random rand) {
471 F32Array array = F32Array.create(accelerator, size);
472 for (int i = 0; i <size; i++) {
473 array.array(i, rand.nextFloat() * (high - low) + low);
474 }
475 return array;
476 }
477
478 @HatTest
479 public static void testBlackScholes() {
480 int size = 50;
481 Random rand = new Random();
482 var accelerator = new Accelerator(java.lang.invoke.MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
483 var call = F32Array.create(accelerator, size);
484 for (int i = 0; i < call.length(); i++) {
485 call.array(i, i);
486 }
487
488 var put = F32Array.create(accelerator, size);
489 for (int i = 0; i < put.length(); i++) {
490 put.array(i, i);
491 }
492
493 var S = floatArray(accelerator, size,1f, 100f, rand);
494 var X = floatArray(accelerator, size,1f, 100f, rand);
495 var T = floatArray(accelerator,size, 0.25f, 10f, rand);
496 float r = 0.02f;
497 float v = 0.30f;
498
499 accelerator.compute(cc -> blackScholes(cc, call, put, S, X, T, r, v));
500 float[] res;
501 for (int i = 0; i < call.length(); i++) {
502 res = blackScholesCheck(S.array(i), X.array(i), T.array(i), r, v);
503 HatAsserts.assertEquals(res[0], call.array(i), 0.0001);
504 HatAsserts.assertEquals(res[1], put.array(i), 0.0001);
505 }
506 }
507
508 // basic test of local and private buffer ArrayViews
509 private interface SharedMemory extends Buffer {
510 void array(long index, int value);
511 int array(long index);
512 Schema<SharedMemory> schema = Schema.of(SharedMemory.class,
513 arr -> arr.array("array", 1024));
514 static SharedMemory create(Accelerator accelerator) {
515 return schema.allocate(accelerator);
516 }
517 static SharedMemory createLocal() {
518 return schema.allocate(new Accelerator(MethodHandles.lookup(), Backend.FIRST));
519 }
520
521 default int[] localArrayView() {
522 int[] view = new int[1024];
523 for (int i = 0; i < 1024; i++) {
524 view[i] = this.array(i);
525 }
526 return view;
527 }
528 }
529
530 public interface PrivateArray extends Buffer {
531 void array(long index, int value);
532 int array(long index);
533 Schema<PrivateArray> schema = Schema.of(PrivateArray.class,
534 arr -> arr.array("array", 16));
535 static PrivateArray create(Accelerator accelerator) {
536 return schema.allocate(accelerator);
537 }
538 static PrivateArray createPrivate() {
539 return schema.allocate(new Accelerator(MethodHandles.lookup(), Backend.FIRST));
540 }
541
542 default int[] privateArrayView() {
543 int[] view = new int[16];
544 for (int i = 0; i < 16; i++) {
545 view[i] = this.array(i);
546 }
547 return view;
548 }
549 }
550
551 @CodeReflection
552 public static void squareKernelWithPrivateAndLocal(@RO KernelContext kc, @RW S32Array s32Array) {
553 SharedMemory shared = SharedMemory.createLocal();
554 if (kc.x<kc.maxX){
555 int[] arr = s32Array.arrayView();
556 arr[kc.x] += arr[kc.x];
557 // int[] a = new int[4];
558 // a[1] = 4;
559
560 PrivateArray priv = PrivateArray.createPrivate();
561 int[] privView = priv.privateArrayView();
562 privView[0] = 1;
563 arr[kc.x] += privView[0];
564
565 int[] sharedView = shared.localArrayView();
566 sharedView[0] = 16;
567 arr[kc.x] += sharedView[0];
568 }
569 }
570
571 @CodeReflection
572 public static void privateAndLocal(@RO ComputeContext cc, @RW S32Array s32Array) {
573 cc.dispatchKernel(s32Array.length(),
574 kc -> squareKernelWithPrivateAndLocal(kc, s32Array)
575 );
576 }
577
578 @HatTest
579 public static void testPrivateAndLocal() {
580
581 var accelerator = new Accelerator(MethodHandles.lookup(), Backend.FIRST);//new JavaMultiThreadedBackend());
582 var arr = S32Array.create(accelerator, 32);
583 for (int i = 0; i < arr.length(); i++) {
584 arr.array(i, i);
585 }
586 accelerator.compute(
587 cc -> privateAndLocal(cc, arr) //QuotableComputeContextConsumer
588 ); // extends Quotable, Consumer<ComputeContext>
589 for (int i = 0; i < arr.length(); i++) {
590 HatAsserts.assertEquals(2 * i + 17, arr.array(i));
591 }
592 }
593 }