1
2 # Interface Mapping
3
4 ----
5
6 * [Contents](hat-00.md)
7 * House Keeping
8 * [Project Layout](hat-01-01-project-layout.md)
9 * [Building Babylon](hat-01-02-building-babylon.md)
10 * [Maven and CMake](hat-01-03-maven-cmake.md)
11 * Programming Model
12 * [Programming Model](hat-03-programming-model.md)
13 * Interface Mapping
14 * [Interface Mapping Overview](hat-04-01-interface-mapping.md)
15 * [Cascade Interface Mapping](hat-04-02-cascade-interface-mapping.md)
16 * Implementation Detail
17 * [Walkthrough Of Accelerator.compute()](hat-accelerator-compute.md)
18
19 ----
20 # Interface Mapping
21
22 ## or ... HAT from a Data POV
23
24 ### Or ... what is this `S32Array` thing and why can't I just pass `int[]` to my kernel
25
26 Again here is the canonical HAT 'hello world' kernel, weill use this to describe itgerface mapping
27
28 We implement this in HAT by collecting the kernel(s) and compute method(s) in a `Compute` class.
29
30 ```java
31 public class SquareCompute {
32 @CodeReflection
33 public static int square(int v) {
34 return v * v;
35 }
36
37 @CodeReflection
38 public static void squareKernel(KernelContext kc, S32Array s32Array) {
39 int value = s32Array.array(kc.x); // arr[cc.x]
40 s32Array.array(kc.x, square(value)); // arr[cc.x]=value*value
41 }
42
43 @CodeReflection
44 public static void square(ComputeContext cc, S32Array s32Array) {
45 cc.dispatchKernel(s32Array.length(),
46 kc -> squareKernel(kc, s32Array)
47 );
48 }
49 }
50 ```
51 Which 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.
52
53 ```java
54 // Create an accelerator bound to a particular backend
55
56 var accelerator = new Accelerator(
57 java.lang.invoke.MethodHandles.lookup(),
58 Backend.FIRST // Predicate<Backend>
59 );
60
61 // Ask the accelerator/backend to allocate an S32Array
62 var s32Array = S32Array.create(accelerator, 32);
63
64 // Fill it with data
65 for (int i = 0; i < s32Array.length(); i++) {
66 s32Array.array(i, i);
67 }
68
69 // Tell the accelerator to execute the square() compute entrypoint
70
71 accelerator.compute(
72 cc -> SquareCompute.square(cc, s32Array)
73 );
74
75 // Check the data
76 for (int i = 0; i < arr.length(); i++) {
77 System.out.println(i + " " + arr.array(i));
78 }
79 ```
80
81 HAT kernels only accept Java primitives and HAT buffers as parameters.
82
83 We don't directly support heap allocated data (such as int[])
84
85 From Java's point of view `S32Array` is a `hat.Buffer` and is defined as an interface.
86
87 ```java
88 public interface S32Array extends Buffer {
89 int length();
90 void length(int i);
91 int array(long idx);
92 void array(long idx, int i);
93 }
94 ```
95
96 From C99 style OpenCL/CUDA POV this will eventually be mapped to a typedef.
97
98 ```C++
99 typedef struct S32Array_s{
100 int length;
101 int array[]; //<-- ?
102 }S32Array_t;
103 ```
104
105 Our Java implementations should treat the interface as `data`, generally the only
106 methods that we include in a `hat.Buffer` should be
107
108 ```java
109 T name(); //getter for a field called name with type T, where T may be primitive or inner interface)
110 void name(T name); //setter for a field called name with type T, T must be primitive
111 T name(long idx); //get an array element [idx] where array is called name and T is either primitive or inner interface
112 void name(long idx, T name); //set an array element [idx] where array is called name and T is primitive
113 ```
114
115 Algorithms can assume that an interface is 'bound' to 'some' concrete data layout.
116
117 We could for example implement `S32Array` like this.
118
119 ```java
120 class JavaS32Array implements S32Array{
121 int[] arr;
122 int length(){ return arr.length;}
123 int array(long idx) {return arr[idx];}
124 void array(long idx, int value) {arr[idx] = value;}
125 void length(int len) ; // we'll come back to this ;)
126 }
127 ```
128
129 But for HAT to access native memory, allocated by the appropriate backend we need interfaces bound to MemorySegents/
130
131 HAT includes an API which allows us to take an interface which extends `hat.Buffer`, and 'bind' it to a Panama FFM MemorySegment.
132
133 This binding process automatically maps the accessors (for example `length()`, `array(long idx, int v)`) to low level Method and Var handel trickery underlying MemorySegments.
134
135 Conceptually we might imagine that HAT creates something like this
136
137 ```java
138 class PanamaS32Array implements S32Array{
139 MemorySegment segment;
140 final int SIZEOFINT = 4;
141 final long lenOffset = 0;
142 final long arrayOffset = lenOffset+SIZEOFINT;
143 int length(){ return segment.getInt(lenOffset);}
144 int array(long idx) {return segment.getInt(arrayOffset+idx*SIZEOFINT);}
145 void array(long idx, int value) {segment.setInt(arrayOffset+idx*SIZEOFINT,value);}
146 void length(int len) ; // we'll come back to this ;)
147 }
148 ```
149
150 Much like Java's `Proxy` class, the iface mapper creates an implementation of the interface 'on the fly', the new Classfile API is used to 'spin up' the new class and the accessors are are composed using Var/Method Handles and offsets derived from the size and order of fields.
151
152 Sadly an interface is not quite enough to establish exactly what is needed to complete the mapping. We need to tell the `iface mapper` the order and size of fields and possibly some padding information.
153
154 We do this by providing a 'layout description' using Panama's Layout api.
155
156 ```java
157 MemoryLayout s32ArrayLayout = MemoryLayout.structLayout(
158 JAVA_INT.withName("length"),
159 MemoryLayout.sequenceLayout(N, JAVA_INT.withName("length")).withName("array")
160 ).withName(S32Array.getSimpleName());
161 ```
162
163 Eventually we came to a common pattern for describing HAT buffers by adding a `create` method to our interface which hides the mapping detail
164
165 So the complete `S32Array` looks a like this. (....ish)
166
167 ```java
168 public interface S32Array extends Buffer {
169 int length();
170
171 void length(int i);
172
173 int array(long idx);
174
175 void array(long idx, int i);
176
177 S32Array create(Accelerator a, int len) {
178 MemoryLayout s32ArrayLayout = MemoryLayout.structLayout(
179 JAVA_INT.withName("length"),
180 MemoryLayout.sequenceLayout(len, JAVA_INT.withName("length")).withName("array")
181 ).withName(S32Array.getSimpleName());
182
183 S32Array s32Array = a.allocate(
184 SegmentMapper.of(MethodHandles.lookup(), S32Array.class, s32ArrayLayout, len)
185 );
186
187 return s32Array;
188 }
189 }
190 ```
191
192 So now hopefully this code makes more sense.
193
194 ```
195 var s32Array = S32Array.create(accelerator, 32);
196 ```
197
198 Whilst this code is much nicer than hand mapping each method to offsets. It is still quite verbose.
199
200 In the last few weeks we have been migrating to Schema builder which makes this code easier to express..
201
202 ```java
203 public interface S32Array extends Buffer {
204 int length();
205 void length(int i);
206 int array(long idx);
207 void array(long idx, int i);
208 Schema<S32Array> schema = Schema.of(S32Array.class, s->s
209 .arrayLen("length")
210 .array("array")
211 );
212 }
213 ```
214 The schema is embedded inside the interface and defines the order of fields. It also allows us to bind fields to each other (above we are telling the schema we have a `int length` field followed by an `int array[]` field and that the first defines the size of the second), we also can describe useful 'HAT' information for fields. Such as whether a field is 'atomic' ;)
215
216 Here is an example of a table of Results for the face detector.
217
218 ```java
219 public interface ResultTable extends Buffer{
220 interface Result extends Buffer.StructChild {
221 float x();
222 void x(float x);
223 float y();
224 void y(float y);
225 }
226 void count(int count);
227 int count();
228 int length();
229 Result result(long idx);
230
231 Schema<ResultTable> schema = Schema.of(ResultTable.class, s->s
232 .atomic("count")
233 .arrayLen("length")
234 .array("result", r->r
235 .field("x")
236 .field("y")
237 )
238 );
239 }
240 ```
241
242 Which in C99 OpenCL code will manifest as
243
244 ```C++
245 typedef Result_s{
246 int x,y
247 } Result_t;
248
249 typedef ResultTable_s{
250 int count;
251 int length;
252 Result_t result[0];
253 } Result_t;
254 ```
255
256 In our Java code this interface makes access to MemorySegments much cleaner
257
258 ```java
259 ResultTable resultTable = ResultTable.create(acc, 100);
260 for (int i=0; i<resultTable.length(); i++){
261 Result result = resultTable.result(i);
262 result.x(0);
263 result.y(0);
264 }
265 ```
266
267 The generated OpenCL/C99 code from Java kernel code is also quite clean
268
269 We might use a kernel to initialize the location of a bunch of Results
270
271 ```java
272 @CodeReflection public static void init(KernelContext kc, ResultTable resultTable) {
273 if (kc.x < kc.maxX){
274 Result result = resulTable.result(kc.x);
275 result.x(kc.x);
276 result.y(100);
277 }
278 }
279 ```
280
281 Whose Kernel code will look like this.
282
283 ```
284 typedef struct KernelContext_s{
285 int x;
286 int maxX;
287 }KernelContext_t;
288
289 typedef Result_s{
290 int x,y
291 } Result_t;
292
293 typedef ResultTable_s{
294 int count;
295 int length;
296 Result_t result[0];
297 } Result_t;
298
299 __kernel void init(
300 __global KernelContext_t *empty,
301 __global ResultTable_t* resultTable
302 ){
303 KernelContext_t kernelContext;
304 KernelContext_t *kc = &kernelContext;
305 kc->x=get_global_id(0);
306 kc->maxX = get_global_id(0);
307
308 if(kc->x<kc->maxX){
309 __global Result_t *result = &resultTable[kc->x];
310 result->x = kc->x;
311 }
312 return;
313 }
314 ```
315
316 A few notes from this generated code...
317
318 * `KernelContext` is itself just an iface mapped segment.
319 - But we don't pass `kc.x` o `kc.maxX` in the segment.
320 - Instead initialize using appropriate vendor calls
321
322 So for OpenCL all kernels start like this
323
324 ```
325 __kernel void init(__global KernelContext_t *empty , ....){
326 KernelContext_t kernelContext;
327 KernelContext_t *kc = &kernelContext;
328 kc->x=get_global_id(0);
329 kc->maxX = get_global_id(0);
330 ....
331 }
332 ```
333
334 Whereas CUDA ;)
335
336 ```
337 __kernel void init(__global KernelContext_t *empty , ....){
338 KernelContext_t kernelContext;
339 KernelContext_t *kc = &kernelContext;
340 kc->x=blockIdx.x*blockDim.x+threadIdx.x;
341 kc->maxX =gridDim.x*blockDim.x
342 ....
343 }
344 ```
345
346 This simplifies code gen. Generally the CUDA code and OpenCL code looks identical.
347
348 ----
349
350 The iface mapping code in hat is a modified form of the code hereWe have a copy of Per's segment mapping code from
351
352 https://github.com/minborg/panama-foreign/blob/segment-mapper/src/java.base/share/classes