1
2 # HAT's Programming Model
3 ----
4
5 * [Contents](hat-00.md)
6 * House Keeping
7 * [Project Layout](hat-01-01-project-layout.md)
8 * [Building Babylon](hat-01-02-building-babylon.md)
9 * [Maven and CMake](hat-01-03-maven-cmake.md)
10 * Programming Model
11 * [Programming Model](hat-03-programming-model.md)
12 * Interface Mapping
13 * [Interface Mapping Overview](hat-04-01-interface-mapping.md)
14 * [Cascade Interface Mapping](hat-04-02-cascade-interface-mapping.md)
15 * Implementation Detail
16 * [Walkthrough Of Accelerator.compute()](hat-accelerator-compute.md)
17
18 ---
19
20 # HAT's Programming model
21
22 Let's consider a trivial opencl kernel which squares each element in an int buffer
23
24 ```java
25 int square(int value){
26 return value*value;
27 }
28
29 __kernel void squareKernel( __global int* s32Array){
30 int value = s32Array[get_global_id(0)];
31 s32Array[get_global_id(0)]=square(value);
32 return;
33 }
34
35 ```
36
37 We implement this in HAT by collecting the kernel(s) and compute method(s) in a `Compute` class.
38
39 ```java
40 public class SquareCompute {
41 @CodeReflection
42 public static int square(int v) {
43 return v * v;
44 }
45
46 @CodeReflection
47 public static void squareKernel(KernelContext kc, S32Array s32Array) {
48 int value = s32Array.array(kc.x); // arr[cc.x]
49 s32Array.array(kc.x, square(value)); // arr[cc.x]=value*value
50 }
51
52 @CodeReflection
53 public static void square(ComputeContext cc, S32Array s32Array) {
54 cc.dispatchKernel(s32Array.length(),
55 kc -> squareKernel(kc, s32Array)
56 );
57 }
58 }
59 ```
60 And we dispatch by creating the appropriate data buffer and then asking an `Accelerator` (bound to a typical vendor backend) to execute the compute method.. which in turn coordinates the dispatch of the various kernels.
61
62 ```java
63 // Create an accelerator bound to a particular backend
64
65 var accelerator = new Accelerator(
66 java.lang.invoke.MethodHandles.lookup(),
67 Backend.FIRST // Predicate<Backend>
68 );
69
70 // Ask the accelerator/backend to allocate an S32Array
71 var s32Array = S32Array.create(accelerator, 32);
72
73 // Fill it with data
74 for (int i = 0; i < s32Array.length(); i++) {
75 s32Array.array(i, i);
76 }
77
78 // Tell the accelerator to execute the square() compute entrypoint
79
80 accelerator.compute(
81 cc -> SquareCompute.square(cc, s32Array)
82 );
83
84 // Check the data
85 for (int i = 0; i < arr.length(); i++) {
86 System.out.println(i + " " + arr.array(i));
87 }
88 ```
89
90 ## Programming model notes
91
92 The most important concept here is that we separate `normal java` code,
93 from `compute` code from `kernel` code
94
95 We must not assume that Compute or Kernel code are ever executed by the JVM
96
97 ### Kernel Code (kernel entrypoints and kernel reachable methods)
98 Kernel's and any kernel reachable methods will naturally be restricted to subset of Java.
99
100 * No exceptions (no exceptions! :) )
101 * No heap access (no `new`)
102 * No access to static or instance fields from this or any other classes )
103 * Except `final static primitives` (which generally get constant pooled)
104 * Except fields of `KernelContext` (thread identity `.x`, `.maxX`, `.groups`... )
105 - We may even decide to access these via methods (`.x()`);
106 * The only methods that can be called are either :-
107 * Kernel reachable methods
108 - Technically you can call a kernel entrypoint, but must pass your KernelContext
109 * `ifaceMappedSegment` accessor/mutators (see later)
110 * Calls on `KernelContext` (backend kernel features)
111 - `KernelContext.barrier()`
112 - `kernelContext.I32.hypot(x,y)`
113 #### Kernel Entrypoints
114 * Declared `@CodeReflection static public void`
115 * Later we may allow reductions to return data...
116 * Parameters
117 * 0 is always a `KernelContext` (KernelContext2D, KernelContext3D logically follow)
118 * 1..n are restricted to uniform primitive values and Panama FFM `ifaceMappedSegments`
119
120 #### Kernel Reachable Methods
121 * Declared `@CodeReflection static public`
122 * All Parameters are restricted to uniform primitive values and Panama FFM `ifaceMappedSegments`
123
124 ### Compute Code (Compute entry points and compute reachable methods)
125 Code within the `compute entrypoint` and `compute reachable
126 methods` have much fewer Java restrictions than kernels but generally...
127
128 * Exceptions are discouraged
129 * Java Synchronization is discouraged
130 * Don't assume any allocation of local `ifaceMappedSegmants` are allocated
131 * Java accesses/mutations to `ifaceMappedSegment` will likely impact performance
132 * Code should ideally just contain simple control flow and kernel dispatches.
133 * Data movements (to and from backend) will automatically be derived from control flow and `ifaceMappedSegment` accesses
134 - We hope to never have to add `cc.moveToDevice(hatBuffer)`
135 * All methods reachable from a `compute entrypoint` are either :-
136 * Compute Reachable Methods
137 - Technically methods can be compute reachable and kernel reachable.
138 * `ifaceMappedSegment` accessor/mutators (see later)
139 * Calls on the `ComputeContext` to generate ranges, or dispatch kernels.
140
141 #### Compute Entry Points
142 * Declared `@CodeReflection static public void`
143 * Parameter 0 is `ComputeContext`
144
145
146 #### Compute Reachable Methods
147 * Declared `@CodeReflection static public `