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