1 ## Notes about the PTX backend
2
3 Here is the Java source code for a kernel:
4
5 ```Java
6 public static void mandel(KernelContext kc, S32Array2D s32Array2D, S32Array pallette, float offsetx, float offsety, float scale) {
7 if (kc.x < kc.maxX) {
8 float width = s32Array2D.width();
9 float height = s32Array2D.height();
10 float x = ((kc.x % s32Array2D.width()) * scale - (scale / 2f * width)) / width + offsetx;
11 float y = ((kc.x / s32Array2D.width()) * scale - (scale / 2f * height)) / height + offsety;
12 float zx = x;
13 float zy = y;
14 float new_zx;
15 int colorIdx = 0;
16 while ((colorIdx < pallette.length()) && (((zx * zx) + (zy * zy)) < 4f)) {
17 new_zx = ((zx * zx) - (zy * zy)) + x;
18 zy = (2f * zx * zy) + y;
19 zx = new_zx;
20 colorIdx++;
21 }
22 int color = colorIdx < pallette.length() ? pallette.array(colorIdx) : 0;
23 s32Array2D.array(kc.x, color);
24 }
25 }
26
27 ```
28
29 And here is the corresponding Babylon code model:
30
31 ```
32 func @"mandel" @loc="39:5:file:/Users/grfrost/orahub/hat/examples/mandel/src/java/mandel/MandelCompute.java"
33 (%0 : hat.KernelContext, %1 : hat.buffer.S32Array2D, %2 : hat.buffer.S32Array, %3 : float, %4 : float, %5 : float)void -> {
34 %6 : Var<hat.KernelContext> = var %0 @"kc" ;
35 %7 : Var<hat.buffer.S32Array2D> = var %1 @"s32Array2D" ;
36 %8 : Var<hat.buffer.S32Array> = var %2 @"pallette" ;
37 %9 : Var<float> = var %3 @"offsetx" ;
38 %10 : Var<float> = var %4 @"offsety" ;
39 %11 : Var<float> = var %5 @"scale" ;
40 java.if
41 ()boolean -> {
42 %12 : hat.KernelContext = var.load %6 ;
43 %13 : int = field.load %12 @"hat.KernelContext::x()int" ;
44 %14 : hat.KernelContext = var.load %6 ;
45 %15 : int = field.load %14 @"hat.KernelContext::maxX()int" ;
46 %16 : boolean = lt %13 %15 ;
47 yield %16 ;
48 }
49 ()void -> {
50 %17 : hat.buffer.S32Array2D = var.load %7 ;
51 %18 : int = invoke %17 @"hat.buffer.S32Array2D::width()int" ;
52 %19 : float = conv %18 ;
53 %20 : Var<float> = var %19 @"width" ;
54 %21 : hat.buffer.S32Array2D = var.load %7 ;
55 %22 : int = invoke %21 @"hat.buffer.S32Array2D::height()int" ;
56 %23 : float = conv %22 ;
57 %24 : Var<float> = var %23 @"height" ;
58 %25 : hat.KernelContext = var.load %6 ;
59 %26 : int = field.load %25 @"hat.KernelContext::x()int" ;
60 %27 : hat.buffer.S32Array2D = var.load %7 ;
61 %28 : int = invoke %27 @"hat.buffer.S32Array2D::width()int" ;
62 %29 : int = mod %26 %28 ;
63 %30 : float = conv %29 ;
64 %31 : float = var.load %11 ;
65 %32 : float = mul %30 %31 ;
66 %33 : float = var.load %11 ;
67 %34 : float = constant @"2.0" ;
68 %35 : float = div %33 %34 ;
69 %36 : float = var.load %20 ;
70 %37 : float = mul %35 %36 ;
71 %38 : float = sub %32 %37 ;
72 %39 : float = var.load %20 ;
73 %40 : float = div %38 %39 ;
74 %41 : float = var.load %9 ;
75 %42 : float = add %40 %41 ;
76 %43 : Var<float> = var %42 @"x" ;
77 %44 : hat.KernelContext = var.load %6 ;
78 %45 : int = field.load %44 @"hat.KernelContext::x()int" ;
79 %46 : hat.buffer.S32Array2D = var.load %7 ;
80 %47 : int = invoke %46 @"hat.buffer.S32Array2D::width()int" ;
81 %48 : int = div %45 %47 ;
82 %49 : float = conv %48 ;
83 %50 : float = var.load %11 ;
84 %51 : float = mul %49 %50 ;
85 %52 : float = var.load %11 ;
86 %53 : float = constant @"2.0" ;
87 %54 : float = div %52 %53 ;
88 %55 : float = var.load %24 ;
89 %56 : float = mul %54 %55 ;
90 %57 : float = sub %51 %56 ;
91 %58 : float = var.load %24 ;
92 %59 : float = div %57 %58 ;
93 %60 : float = var.load %10 ;
94 %61 : float = add %59 %60 ;
95 %62 : Var<float> = var %61 @"y" ;
96 %63 : float = var.load %43 ;
97 %64 : Var<float> = var %63 @"zx" ;
98 %65 : float = var.load %62 ;
99 %66 : Var<float> = var %65 @"zy" ;
100 %67 : float = constant @"0.0" ;
101 %68 : Var<float> = var %67 @"new_zx" ;
102 %69 : int = constant @"0" ;
103 %70 : Var<int> = var %69 @"colorIdx" ;
104 java.while
105 ()boolean -> {
106 %71 : boolean = java.cand
107 ()boolean -> {
108 %72 : int = var.load %70 ;
109 %73 : hat.buffer.S32Array = var.load %8 ;
110 %74 : int = invoke %73 @"hat.buffer.S32Array::length()int" ;
111 %75 : boolean = lt %72 %74 ;
112 yield %75 ;
113 }
114 ()boolean -> {
115 %76 : float = var.load %64 ;
116 %77 : float = var.load %64 ;
117 %78 : float = mul %76 %77 ;
118 %79 : float = var.load %66 ;
119 %80 : float = var.load %66 ;
120 %81 : float = mul %79 %80 ;
121 %82 : float = add %78 %81 ;
122 %83 : float = constant @"4.0" ;
123 %84 : boolean = lt %82 %83 ;
124 yield %84 ;
125 };
126 yield %71 ;
127 }
128 ()void -> {
129 %85 : float = var.load %64 ;
130 %86 : float = var.load %64 ;
131 %87 : float = mul %85 %86 ;
132 %88 : float = var.load %66 ;
133 %89 : float = var.load %66 ;
134 %90 : float = mul %88 %89 ;
135 %91 : float = sub %87 %90 ;
136 %92 : float = var.load %43 ;
137 %93 : float = add %91 %92 ;
138 var.store %68 %93 ;
139 %94 : float = constant @"2.0" ;
140 %95 : float = var.load %64 ;
141 %96 : float = mul %94 %95 ;
142 %97 : float = var.load %66 ;
143 %98 : float = mul %96 %97 ;
144 %99 : float = var.load %62 ;
145 %100 : float = add %98 %99 ;
146 var.store %66 %100 ;
147 %101 : float = var.load %68 ;
148 var.store %64 %101 ;
149 %102 : int = var.load %70 ;
150 %103 : int = constant @"1" ;
151 %104 : int = add %102 %103 ;
152 var.store %70 %104 ;
153 java.continue ;
154 };
155 %105 : int = java.cexpression
156 ()boolean -> {
157 %106 : int = var.load %70 ;
158 %107 : hat.buffer.S32Array = var.load %8 ;
159 %108 : int = invoke %107 @"hat.buffer.S32Array::length()int" ;
160 %109 : boolean = lt %106 %108 ;
161 yield %109 ;
162 }
163 ()int -> {
164 %110 : hat.buffer.S32Array = var.load %8 ;
165 %111 : int = var.load %70 ;
166 %112 : long = conv %111 ;
167 %113 : int = invoke %110 %112 @"hat.buffer.S32Array::array(long)int" ;
168 yield %113 ;
169 }
170 ()int -> {
171 %114 : int = constant @"0" ;
172 yield %114 ;
173 };
174 %115 : Var<int> = var %105 @"color" ;
175 %116 : hat.buffer.S32Array2D = var.load %7 ;
176 %117 : hat.KernelContext = var.load %6 ;
177 %118 : int = field.load %117 @"hat.KernelContext::x()int" ;
178 %119 : long = conv %118 ;
179 %120 : int = var.load %115 ;
180 invoke %116 %119 %120 @"hat.buffer.S32Array2D::array(long, int)void" ;
181 yield ;
182 }
183 ()void -> {
184 yield;
185 };
186 return ;
187 };
188
189 ```
190 From the above we can generate C99 style CUDA code
191
192
193 Here is the Cuda C99 code generated
194
195 ```C
196 typedef struct KernelContext_s{
197 int x;
198 int maxX;
199 }KernelContext_t;
200
201 typedef struct S32Array2D_s{
202 int width;
203 int height;
204 int array[0];
205 }S32Array2D_t;
206
207 typedef struct S32Array_s{
208 int length;
209 int array[0];
210 }S32Array_t;
211
212
213 extern "C" __global__ void mandel(
214 S32Array2D_t* s32Array2D, S32Array_t* pallette, float offsetx, float offsety, float scale
215 ){
216 KernelContext_t kc;
217 kc.x=blockIdx.x*blockDim.x+threadIdx.x;
218 kc.maxX=gridDim.x*blockDim.x;
219 if(kc.x<kc.maxX){
220 float width = (float)s32Array2D->width;
221 float height = (float)s32Array2D->height;
222 float x = ((float)(kc.x%s32Array2D->width)*scale-scale/2.0*width)/width+offsetx;
223 float y = ((float)(kc.x/s32Array2D->width)*scale-scale/2.0*height)/height+offsety;
224 float zx = x;
225 float zy = y;
226 float new_zx = 0.0;
227 int colorIdx = 0;
228 while(colorIdx<pallette->length && zx*zx+zy*zy<4.0){
229 new_zx=zx*zx-zy*zy+x;
230 zy=2.0*zx*zy+y;
231 zx=new_zx;
232 colorIdx=colorIdx+1;
233 }
234 int color = colorIdx<pallette->length?pallette->array[(long)colorIdx]:0;
235 s32Array2D->array[(long)kc.x]=color;
236 }
237 return;
238 }
239 ```
240
241 Which we can convert to ptx using the nvcc compiler (on a NVDIA platform)
242
243 `nvcc -ptx mandel.cu -o mandel.ptx`
244
245 ```
246 //
247 // Generated by NVIDIA NVVM Compiler
248 //
249 // Compiler Build ID: CL-33191640
250 // Cuda compilation tools, release 12.2, V12.2.140
251 // Based on NVVM 7.0.1
252 //
253
254 .version 8.2
255 .target sm_52
256 .address_size 64
257
258 // .globl mandel
259
260 .visible .entry mandel(
261 .param .u64 mandel_param_0,
262 .param .u64 mandel_param_1,
263 .param .f32 mandel_param_2,
264 .param .f32 mandel_param_3,
265 .param .f32 mandel_param_4
266 )
267 {
268 .reg .pred %p<9>;
269 .reg .f32 %f<29>;
270 .reg .b32 %r<24>;
271 .reg .f64 %fd<22>;
272 .reg .b64 %rd<9>;
273
274
275 ld.param.u64 %rd3, [mandel_param_0];
276 ld.param.u64 %rd4, [mandel_param_1];
277 ld.param.f32 %f13, [mandel_param_2];
278 ld.param.f32 %f14, [mandel_param_3];
279 ld.param.f32 %f15, [mandel_param_4];
280 cvta.to.global.u64 %rd1, %rd4;
281 cvta.to.global.u64 %rd2, %rd3;
282 mov.u32 %r8, %ntid.x;
283 mov.u32 %r9, %ctaid.x;
284 mov.u32 %r10, %tid.x;
285 mad.lo.s32 %r1, %r9, %r8, %r10;
286 mov.u32 %r11, %nctaid.x;
287 mul.lo.s32 %r12, %r11, %r8;
288 setp.ge.s32 %p1, %r1, %r12;
289 @%p1 bra $L__BB0_7;
290
291 ld.global.u32 %r14, [%rd2];
292 cvt.rn.f32.s32 %f16, %r14;
293 ld.global.u32 %r15, [%rd2+4];
294 cvt.rn.f32.s32 %f17, %r15;
295 div.s32 %r16, %r1, %r14;
296 mul.lo.s32 %r17, %r16, %r14;
297 sub.s32 %r18, %r1, %r17;
298 cvt.rn.f32.s32 %f18, %r18;
299 mul.f32 %f19, %f18, %f15;
300 cvt.f64.f32 %fd2, %f19;
301 cvt.f64.f32 %fd3, %f15;
302 mul.f64 %fd4, %fd3, 0d3FE0000000000000;
303 cvt.f64.f32 %fd5, %f16;
304 mul.f64 %fd6, %fd4, %fd5;
305 sub.f64 %fd7, %fd2, %fd6;
306 div.rn.f64 %fd8, %fd7, %fd5;
307 cvt.f64.f32 %fd9, %f13;
308 add.f64 %fd10, %fd8, %fd9;
309 cvt.rn.f32.f64 %f1, %fd10;
310 cvt.rn.f32.s32 %f20, %r16;
311 mul.f32 %f21, %f20, %f15;
312 cvt.f64.f32 %fd11, %f21;
313 cvt.f64.f32 %fd12, %f17;
314 mul.f64 %fd13, %fd4, %fd12;
315 sub.f64 %fd14, %fd11, %fd13;
316 div.rn.f64 %fd15, %fd14, %fd12;
317 cvt.f64.f32 %fd16, %f14;
318 add.f64 %fd17, %fd15, %fd16;
319 cvt.rn.f32.f64 %f27, %fd17;
320 ld.global.u32 %r2, [%rd1];
321 setp.lt.s32 %p2, %r2, 1;
322 mov.u32 %r23, 0;
323 mul.f32 %f26, %f1, %f1;
324 mul.f32 %f25, %f27, %f27;
325 add.f32 %f22, %f26, %f25;
326 setp.geu.f32 %p3, %f22, 0f40800000;
327 or.pred %p4, %p2, %p3;
328 mov.u32 %r22, %r23;
329 @%p4 bra $L__BB0_4;
330
331 cvt.f64.f32 %fd1, %f27;
332 mov.f32 %f28, %f1;
333
334 $L__BB0_3:
335 sub.f32 %f23, %f26, %f25;
336 add.f32 %f9, %f23, %f1;
337 cvt.f64.f32 %fd18, %f28;
338 add.f64 %fd19, %fd18, %fd18;
339 cvt.f64.f32 %fd20, %f27;
340 fma.rn.f64 %fd21, %fd19, %fd20, %fd1;
341 cvt.rn.f32.f64 %f27, %fd21;
342 add.s32 %r22, %r22, 1;
343 setp.lt.s32 %p5, %r22, %r2;
344 mul.f32 %f26, %f9, %f9;
345 mul.f32 %f25, %f27, %f27;
346 add.f32 %f24, %f26, %f25;
347 setp.lt.f32 %p6, %f24, 0f40800000;
348 and.pred %p7, %p5, %p6;
349 mov.f32 %f28, %f9;
350 @%p7 bra $L__BB0_3;
351
352 $L__BB0_4:
353 setp.ge.s32 %p8, %r22, %r2;
354 @%p8 bra $L__BB0_6;
355
356 mul.wide.s32 %rd5, %r22, 4;
357 add.s64 %rd6, %rd1, %rd5;
358 ld.global.u32 %r23, [%rd6+4];
359
360 $L__BB0_6:
361 mul.wide.s32 %rd7, %r1, 4;
362 add.s64 %rd8, %rd2, %rd7;
363 st.global.u32 [%rd8+8], %r23;
364
365 $L__BB0_7:
366 ret;
367
368 }
369 ```
370 But we would like to create the ptx directly from the babylon model, without being on an NVidia platform. Probably the lowered model
371
372 We will probably use the lowered babylon model
373
374 ```
375 func @"mandel" @loc="39:5:file:/Users/grfrost/orahub/hat/examples/mandel/src/java/mandel/MandelCompute.java"
376 (%0 : hat.KernelContext, %1 : hat.buffer.S32Array2D, %2 : hat.buffer.S32Array, %3 : float, %4 : float, %5 : float)void -> {
377 %6 : Var<hat.KernelContext> = var %0 @"kc" ;
378 %7 : Var<hat.buffer.S32Array2D> = var %1 @"s32Array2D" ;
379 %8 : Var<hat.buffer.S32Array> = var %2 @"pallette" ;
380 %9 : Var<float> = var %3 @"offsetx" ;
381 %10 : Var<float> = var %4 @"offsety" ;
382 %11 : Var<float> = var %5 @"scale" ;
383 %12 : hat.KernelContext = var.load %6 ;
384 %13 : int = field.load %12 @"hat.KernelContext::x()int" ;
385 %14 : hat.KernelContext = var.load %6 ;
386 %15 : int = field.load %14 @"hat.KernelContext::maxX()int" ;
387 %16 : boolean = lt %13 %15 ;
388 cbranch %16 ^block_0 ^block_1;
389
390 ^block_0:
391 %17 : hat.buffer.S32Array2D = var.load %7 ;
392 %18 : int = invoke %17 @"hat.buffer.S32Array2D::width()int" ;
393 %19 : float = conv %18 ;
394 %20 : Var<float> = var %19 @"width" ;
395 %21 : hat.buffer.S32Array2D = var.load %7 ;
396 %22 : int = invoke %21 @"hat.buffer.S32Array2D::height()int" ;
397 %23 : float = conv %22 ;
398 %24 : Var<float> = var %23 @"height" ;
399 %25 : hat.KernelContext = var.load %6 ;
400 %26 : int = field.load %25 @"hat.KernelContext::x()int" ;
401 %27 : hat.buffer.S32Array2D = var.load %7 ;
402 %28 : int = invoke %27 @"hat.buffer.S32Array2D::width()int" ;
403 %29 : int = mod %26 %28 ;
404 %30 : float = conv %29 ;
405 %31 : float = var.load %11 ;
406 %32 : float = mul %30 %31 ;
407 %33 : float = var.load %11 ;
408 %34 : float = constant @"2.0" ;
409 %35 : float = div %33 %34 ;
410 %36 : float = var.load %20 ;
411 %37 : float = mul %35 %36 ;
412 %38 : float = sub %32 %37 ;
413 %39 : float = var.load %20 ;
414 %40 : float = div %38 %39 ;
415 %41 : float = var.load %9 ;
416 %42 : float = add %40 %41 ;
417 %43 : Var<float> = var %42 @"x" ;
418 %44 : hat.KernelContext = var.load %6 ;
419 %45 : int = field.load %44 @"hat.KernelContext::x()int" ;
420 %46 : hat.buffer.S32Array2D = var.load %7 ;
421 %47 : int = invoke %46 @"hat.buffer.S32Array2D::width()int" ;
422 %48 : int = div %45 %47 ;
423 %49 : float = conv %48 ;
424 %50 : float = var.load %11 ;
425 %51 : float = mul %49 %50 ;
426 %52 : float = var.load %11 ;
427 %53 : float = constant @"2.0" ;
428 %54 : float = div %52 %53 ;
429 %55 : float = var.load %24 ;
430 %56 : float = mul %54 %55 ;
431 %57 : float = sub %51 %56 ;
432 %58 : float = var.load %24 ;
433 %59 : float = div %57 %58 ;
434 %60 : float = var.load %10 ;
435 %61 : float = add %59 %60 ;
436 %62 : Var<float> = var %61 @"y" ;
437 %63 : float = var.load %43 ;
438 %64 : Var<float> = var %63 @"zx" ;
439 %65 : float = var.load %62 ;
440 %66 : Var<float> = var %65 @"zy" ;
441 %67 : float = constant @"0.0" ;
442 %68 : Var<float> = var %67 @"new_zx" ;
443 %69 : int = constant @"0" ;
444 %70 : Var<int> = var %69 @"colorIdx" ;
445 branch ^block_2;
446
447 ^block_2:
448 %71 : int = var.load %70 ;
449 %72 : hat.buffer.S32Array = var.load %8 ;
450 %73 : int = invoke %72 @"hat.buffer.S32Array::length()int" ;
451 %74 : boolean = lt %71 %73 ;
452 cbranch %74 ^block_3 ^block_4(%74);
453
454 ^block_3:
455 %75 : float = var.load %64 ;
456 %76 : float = var.load %64 ;
457 %77 : float = mul %75 %76 ;
458 %78 : float = var.load %66 ;
459 %79 : float = var.load %66 ;
460 %80 : float = mul %78 %79 ;
461 %81 : float = add %77 %80 ;
462 %82 : float = constant @"4.0" ;
463 %83 : boolean = lt %81 %82 ;
464 branch ^block_4(%83);
465
466 ^block_4(%84 : boolean):
467 cbranch %84 ^block_5 ^block_6;
468
469 ^block_5:
470 %85 : float = var.load %64 ;
471 %86 : float = var.load %64 ;
472 %87 : float = mul %85 %86 ;
473 %88 : float = var.load %66 ;
474 %89 : float = var.load %66 ;
475 %90 : float = mul %88 %89 ;
476 %91 : float = sub %87 %90 ;
477 %92 : float = var.load %43 ;
478 %93 : float = add %91 %92 ;
479 var.store %68 %93 ;
480 %94 : float = constant @"2.0" ;
481 %95 : float = var.load %64 ;
482 %96 : float = mul %94 %95 ;
483 %97 : float = var.load %66 ;
484 %98 : float = mul %96 %97 ;
485 %99 : float = var.load %62 ;
486 %100 : float = add %98 %99 ;
487 var.store %66 %100 ;
488 %101 : float = var.load %68 ;
489 var.store %64 %101 ;
490 %102 : int = var.load %70 ;
491 %103 : int = constant @"1" ;
492 %104 : int = add %102 %103 ;
493 var.store %70 %104 ;
494 branch ^block_2;
495
496 ^block_6:
497 %105 : int = var.load %70 ;
498 %106 : hat.buffer.S32Array = var.load %8 ;
499 %107 : int = invoke %106 @"hat.buffer.S32Array::length()int" ;
500 %108 : boolean = lt %105 %107 ;
501 cbranch %108 ^block_7 ^block_8;
502
503 ^block_7:
504 %109 : hat.buffer.S32Array = var.load %8 ;
505 %110 : int = var.load %70 ;
506 %111 : long = conv %110 ;
507 %112 : int = invoke %109 %111 @"hat.buffer.S32Array::array(long)int" ;
508 branch ^block_9(%112);
509
510 ^block_8:
511 %113 : int = constant @"0" ;
512 branch ^block_9(%113);
513
514 ^block_9(%114 : int):
515 %115 : Var<int> = var %114 @"color" ;
516 %116 : hat.buffer.S32Array2D = var.load %7 ;
517 %117 : hat.KernelContext = var.load %6 ;
518 %118 : int = field.load %117 @"hat.KernelContext::x()int" ;
519 %119 : long = conv %118 ;
520 %120 : int = var.load %115 ;
521 invoke %116 %119 %120 @"hat.buffer.S32Array2D::array(long, int)void" ;
522 branch ^block_10;
523
524 ^block_1:
525 branch ^block_10;
526
527 ^block_10:
528 return ;
529 };
530 ```
531