Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Investigating IGC generated kernel binary and dump info #419

Open
1e-to opened this issue Jun 2, 2021 · 3 comments
Open

Investigating IGC generated kernel binary and dump info #419

1e-to opened this issue Jun 2, 2021 · 3 comments
Labels
debug Related to #149

Comments

@1e-to
Copy link
Contributor

1e-to commented Jun 2, 2021 •
edited by PokhodenkoSA
Loading

When set to “1”, IGC will write number of dumps into /tmp/IntelIGC.

$ export IGC_ShaderDumpEnable=1

To read the DWARF of a kernel, we first need a copy of the IGC generated kernel binary.
To do that, run the Python script in a debugger, and set a breakpoint in the kernel:

$ gdb-oneapi -q python 
(gdb) break sum.py:13     # Assumes the kernel is in file sum.py, at line 13 
(gdb) run sum.py

Once the breakpoint hits, the kernel has been generated and offloaded.
At that point, the IGFX driver (i.e. our debugger driver) has copied the kernel into a file, and saved it at /tmp.
All files saved at /tmp/IntelIGC/pytho_xxxxxx/

Then, to read the DWARF in that kernel binary (elf), use tool llvm-dwarfdump.
The app outputs the contents of the DWARF.
So for example, to create a text file DWARF.dump with the contents of the DWARF data:

llvm-dwarfdump xxxx.elf > dwarf

!Nested function information is missing here!

OCL_asmefeefdc15e2ad3fe_simd8_dppyPy_dppy_py_devfn__5F__5F_main_5F__5F__2E_data_5F_parallel_5F_sum_24_1_2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_.elf: file format ELF64-unknown 

.debug_info contents: 

0x00000000: Compile Unit: length = 0x000002a1 version = 0x0004 abbr_offset = 0x0000 addr_size = 0x08 (next unit at 0x000002a5) 

0x0000000b: DW_TAG_compile_unit 

              DW_AT_producer ("spirv") 

              DW_AT_language (DW_LANG_Python) 

              DW_AT_name ("sum.py") 

              DW_AT_low_pc (0x0000000000000000) 

              DW_AT_high_pc (0x0000000000000810) 

              DW_AT_stmt_list (0x00000000) 

              DW_AT_unknown_2400 (0x0008) 

              DW_AT_comp_dir ("/localdisk/work/etotmeni/stepping") 

0x0000004d:   DW_TAG_subprogram 

                DW_AT_linkage_name ("dppy_py_devfn__5F__5F_main_5F__5F__2E_data_5F_parallel_5F_sum_24_1_2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29_") 

                DW_AT_name ("data_parallel_sum") 

                DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                DW_AT_decl_line (15) 

                DW_AT_unknown_2400 (0x0008) 

                DW_AT_external (true) 

                DW_AT_low_pc (0x0000000000000000) 

                DW_AT_high_pc (0x0000000000000810) 

0x00000138:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_gid0") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x22, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x00000160:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_gid1") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x25, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x00000188:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_gid2") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x28, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x000001b0:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_lid0") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x2b, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x000001d8:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_lid1") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x2e, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x00000200:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_lid2") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_WASM_location 0x32 +37, DW_OP_plus_uconst 0x31, <decoding error> ed 33 1a 0a 40 00 1e 0a 40 00 ef) 

0x00000228:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_grid0") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_reg26, DW_OP_bit_piece 0x40 0x0) 

0x00000244:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_grid1") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_regx 0x33, DW_OP_bit_piece 0x40 0x40) 

0x00000261:     DW_TAG_variable 

                  DW_AT_name ("__ocl_dbg_grid2") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (1) 

                  DW_AT_type (0x00000290 "long long") 

                  DW_AT_location (DW_OP_regx 0x34, DW_OP_bit_piece 0x40 0x0) 

0x0000027e:     DW_TAG_variable 

                  DW_AT_name ("i") 

                  DW_AT_decl_file ("/localdisk/work/etotmeni/stepping/sum.py") 

                  DW_AT_decl_line (17) 

                  DW_AT_type (0x0000029d "i64") 

                  DW_AT_const_value (0x0000000000000000) 

0x0000028f:     NULL 

0x00000290:   DW_TAG_base_type 

                DW_AT_name ("long long") 

                DW_AT_encoding (DW_ATE_signed) 

                DW_AT_byte_size (0x08) 

0x0000029d:   DW_TAG_base_type 

                DW_AT_name ("i64") 

                DW_AT_encoding (DW_ATE_unsigned) 

                DW_AT_byte_size (0x08) 

0x000002a4:   NULL
@1e-to
Copy link
Contributor Author

1e-to commented Jun 2, 2021 •
edited by PokhodenkoSA
Loading

Also in the .asm file at /tmp/IntelIGC/pytho_xxxxxx/ there is nothing interesting about nested functions and debugging information, only the code

//.kernel dppyPy_dppy_py_devfn__5F__5F_main_5F__5F__2E_data_5F_parallel_5F_sum_24_1_2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29_ 

//.platform SKL 

//.thread_config numGRF=128, numAcc=2 

//.options_string "-debug " 

//.full_options "-debug -emitLocation -enableStructurizer -enablePreemption -noStitchExternFunc -TotalGRFNum 128  4 -nopresched -nodpsendreorder -dontUseMultiThreadedLatencies -SBIDDepLoc -output -binary -dumpcommonisa -dumpvisa -noverifyCISA -generateDebugInfo -setstartbp -addKernelID " 

//.instCount 129 

//.RA type TRIVIAL_RA 

 
//.declare BuiltinR0 rf=r size=32 type=ud align=16 words (r0.0) 

//.declare  rf=r size=32 type=ud align=16 words (r11.0) 

//.declare BuiltinA0 rf=a size=4 type=ud align=1 words (a0.0) 

//.declare BuiltinA0Dot2 rf=a size=4 type=ud align=1 words (a0.2) 

//.declare hw_tid rf=r size=4 type=ud align=2 words (r5.0) 

//.declare %null rf=r size=4 type=ud align=2 words 

//.declare %local_id_x rf=r size=4 type=ud align=2 words (r3.3) 

//.declare %local_id_y rf=r size=4 type=ud align=2 words (r3.4) 

//.declare %local_size_x rf=r size=4 type=ud align=2 words (r2.7) 

//.declare %local_size_y rf=r size=4 type=ud align=2 words (r3.0) 

//.declare %group_id_x rf=r size=4 type=ud align=2 words (r0.1) 

//.declare %group_id_y rf=r size=4 type=ud align=2 words (r0.6) 

//.declare %group_id_z rf=r size=4 type=ud align=2 words (r0.7) 

//.declare %group_count_x rf=r size=4 type=ud align=2 words (r3.1) 

//.declare %group_count_y rf=r size=4 type=ud align=2 words (r3.2) 

//.declare %tsc rf=r size=20 type=ud align=2 words 

//.declare %arg rf=r size=0 type=ud align=16 words (r26.0) 

//.declare %retval rf=r size=0 type=ud align=16 words (r26.0) Output 

//.declare %sp rf=r size=8 type=uq align=4 words (r125.3) 

//.declare %fp rf=r size=8 type=uq align=4 words (r125.2) 

//.declare %sr0 rf=r size=16 type=ud align=2 words 

//.declare %cr0 rf=r size=12 type=ud align=2 words 

//.declare %ce0 rf=r size=4 type=ud align=2 words 

//.declare %dbg0 rf=r size=8 type=ud align=2 words 

//.declare V0033 rf=r size=32 type=d alias=+0 align=16 words (r11.0) 

//.declare V0034 rf=r size=8 type=uq align=4 words (r5.2) 

//.declare V0035 rf=r size=8 type=q align=4 words (r7.3) 

//.declare V0036 rf=r size=8 type=uq align=4 words (r6.1) 

//.declare V0037 rf=r size=8 type=q align=4 words (r8.3) 

//.declare V0038 rf=r size=8 type=uq align=4 words (r7.0) 

//.declare V0039 rf=r size=8 type=q align=4 words (r9.3) 

//.declare V0041 rf=r size=32 type=d alias=+0 align=16 words (r11.0) 

//.declare V0042 rf=r size=32 type=d align=16 words (r4.0) 

//.declare V0043 rf=r size=12 type=d align=2 words (r10.4) 

//.declare V0044 rf=r size=32 type=w align=16 words (r1.0) 

//.declare V0045 rf=r size=32 type=w align=16 words (r2.0) 

//.declare V0046 rf=r size=32 type=w align=16 words (r3.0) 

//.declare V0047 rf=r size=8 type=uq align=4 words (r10.1) Output 

//.declare V0057 rf=r size=16 type=w align=16 words (r8.0) 

//.declare V0058 rf=r size=4 type=d align=2 words (r5.1) 

//.declare V0060 rf=r size=4 type=ud alias=hw_tid+0 align=16 words (r5.0) 

//.declare V0061 rf=r size=4 type=ud alias=V0058+0 align=2 words (r5.1) 

//.declare V0062 rf=r size=4 type=d align=2 words (r5.2) Output 

//.declare V0063 rf=r size=4 type=d align=2 words (r5.3) 

//.declare V0064 rf=r size=32 type=d align=16 words (r12.0) 

//.declare V0065 rf=r size=16 type=uw alias=V0057+0 align=16 words (r8.0) 

//.declare V0066 rf=r size=8 type=q alias=V0047+0 align=4 words (r10.1) 

//.declare V0067 rf=r size=64 type=q align=16 words (r13.0) 

//.declare V0068 rf=r size=64 type=uq alias=V0067+0 align=16 words (r13.0) 

//.declare V0069 rf=r size=64 type=uq align=16 words (r15.0) 

//.declare P01 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0070 rf=r size=2 type=b align=1 words (r5.24) 

//.declare V0071 rf=r size=4 type=d align=2 words (r5.7) 

//.declare V0072 rf=r size=2 type=ub alias=V0070+0 align=1 words (r5.24) 

//.declare V0073 rf=r size=4 type=d align=2 words (r6.0) 

//.declare V0074 rf=r size=4 type=d align=2 words (r6.1) 

//.declare V0075 rf=r size=4 type=uw alias=V0073+0 align=2 words (r6.0) 

//.declare V0076 rf=r size=2 type=uw align=1 words (r5.13) 

//.declare A00 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0077 rf=r size=4 type=d align=2 words (r6.4) 

//.declare V0078 rf=r size=32 type=d align=16 words (r17.0) 

//.declare V0079 rf=r size=32 type=uw alias=V0044+0 align=16 words (r1.0) 

//.declare V0080 rf=r size=64 type=q align=16 words (r18.0) 

//.declare V0081 rf=r size=32 type=ud alias=V0078+0 align=16 words (r17.0) 

//.declare P02 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0082 rf=r size=2 type=b align=1 words (r6.20) 

//.declare V0083 rf=r size=4 type=d align=2 words (r6.6) 

//.declare V0084 rf=r size=2 type=ub alias=V0082+0 align=1 words (r6.20) 

//.declare V0085 rf=r size=4 type=d align=2 words (r6.7) 

//.declare V0086 rf=r size=4 type=d align=2 words (r7.2) 

//.declare V0087 rf=r size=4 type=uw alias=V0085+0 align=2 words (r6.14) 

//.declare V0088 rf=r size=2 type=uw align=1 words (r6.11) 

//.declare A01 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0089 rf=r size=4 type=d align=2 words (r7.3) 

//.declare V0090 rf=r size=32 type=d align=16 words (r20.0) 

//.declare V0091 rf=r size=32 type=uw alias=V0045+0 align=16 words (r2.0) 

//.declare V0092 rf=r size=64 type=q align=16 words (r21.0) 

//.declare V0093 rf=r size=32 type=ud alias=V0090+0 align=16 words (r20.0) 

//.declare P03 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0094 rf=r size=2 type=b align=1 words (r7.16) 

//.declare V0095 rf=r size=4 type=d align=2 words (r7.5) 

//.declare V0096 rf=r size=2 type=ub alias=V0094+0 align=1 words (r7.16) 

//.declare V0097 rf=r size=4 type=d align=2 words (r8.4) 

//.declare V0098 rf=r size=4 type=d align=2 words (r8.5) 

//.declare V0099 rf=r size=4 type=uw alias=V0097+0 align=2 words (r8.8) 

//.declare V0100 rf=r size=2 type=uw align=1 words (r7.9) 

//.declare A02 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0101 rf=r size=4 type=d align=2 words (r9.0) 

//.declare V0102 rf=r size=32 type=d align=16 words (r23.0) 

//.declare V0103 rf=r size=32 type=uw alias=V0046+0 align=16 words (r3.0) 

//.declare V0104 rf=r size=64 type=q align=16 words (r24.0) 

//.declare V0105 rf=r size=32 type=ud alias=V0102+0 align=16 words (r23.0) 

//.declare V0106 rf=r size=32 type=d align=16 words (r26.0) 

//.declare V0107 rf=r size=64 type=q align=16 words (r27.0) 

//.declare V0108 rf=r size=32 type=ud alias=V0106+0 align=16 words (r26.0) 

//.declare V0109 rf=r size=32 type=d align=16 words (r29.0) 

//.declare V0110 rf=r size=64 type=q align=16 words (r30.0) 

//.declare V0111 rf=r size=32 type=ud alias=V0109+0 align=16 words (r29.0) 

//.declare V0112 rf=r size=32 type=d align=16 words (r32.0) 

//.declare V0113 rf=r size=64 type=q align=16 words (r33.0) 

//.declare V0114 rf=r size=32 type=ud alias=V0112+0 align=16 words (r32.0) 

//.declare P04 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0115 rf=r size=2 type=b align=1 words (r9.4) 

//.declare V0116 rf=r size=4 type=d align=2 words (r9.2) 

//.declare V0117 rf=r size=2 type=ub alias=V0115+0 align=1 words (r9.4) 

//.declare V0118 rf=r size=4 type=d align=2 words (r9.3) 

//.declare V0119 rf=r size=4 type=d align=2 words (r9.4) 

//.declare V0120 rf=r size=4 type=uw alias=V0118+0 align=2 words (r9.6) 

//.declare V0121 rf=r size=2 type=uw align=1 words (r9.3) 

//.declare A03 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0122 rf=r size=8 type=q align=4 words (r10.0) 

//.declare V0123 rf=r size=4 type=ud alias=V0119+0 align=2 words (r9.4) 

//.declare P05 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0124 rf=r size=2 type=b align=1 words (r9.20) 

//.declare V0125 rf=r size=4 type=d align=2 words (r10.7) 

//.declare V0126 rf=r size=2 type=ub alias=V0124+0 align=1 words (r9.20) 

//.declare V0127 rf=r size=4 type=d align=2 words (r35.0) 

//.declare V0128 rf=r size=4 type=d align=2 words (r35.1) 

//.declare V0129 rf=r size=4 type=uw alias=V0127+0 align=2 words (r35.0) 

//.declare V0130 rf=r size=2 type=uw align=1 words (r9.11) 

//.declare A04 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0131 rf=r size=8 type=q align=4 words (r35.1) 

//.declare V0132 rf=r size=4 type=ud alias=V0128+0 align=2 words (r35.1) 

//.declare P06 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0133 rf=r size=2 type=b align=1 words (r35.16) 

//.declare V0134 rf=r size=4 type=d align=2 words (r35.5) 

//.declare V0135 rf=r size=2 type=ub alias=V0133+0 align=1 words (r35.16) 

//.declare V0136 rf=r size=4 type=d align=2 words (r35.6) 

//.declare V0137 rf=r size=4 type=d align=2 words (r35.7) 

//.declare V0138 rf=r size=4 type=uw alias=V0136+0 align=2 words (r35.12) 

//.declare V0139 rf=r size=2 type=uw align=1 words (r35.9) 

//.declare A05 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0140 rf=r size=8 type=q align=4 words (r36.0) 

//.declare V0141 rf=r size=4 type=ud alias=V0137+0 align=2 words (r35.7) 

//.declare P07 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0142 rf=r size=2 type=b align=1 words (r36.8) 

//.declare V0143 rf=r size=4 type=d align=2 words (r36.3) 

//.declare V0144 rf=r size=2 type=ub alias=V0142+0 align=1 words (r36.8) 

//.declare V0145 rf=r size=4 type=d align=2 words (r36.4) 

//.declare V0146 rf=r size=4 type=d align=2 words (r36.5) 

//.declare V0147 rf=r size=4 type=uw alias=V0145+0 align=2 words (r36.8) 

//.declare V0148 rf=r size=2 type=uw align=1 words (r36.5) 

//.declare A06 rf=a size=2 type=uw align=1 words (a0.0) 

//.declare V0149 rf=r size=8 type=q align=4 words (r36.3) 

//.declare V0150 rf=r size=12 type=ud alias=V0043+0 align=2 words (r10.4) 

//.declare V0151 rf=r size=4 type=ud alias=V0146+0 align=2 words (r36.5) 

//.declare V0152 rf=r size=32 type=d align=16 words (r37.0) 

//.declare V0153 rf=r size=64 type=q align=16 words (r38.0) 

//.declare V0154 rf=r size=32 type=ud alias=V0152+0 align=16 words (r37.0) 

//.declare V0155 rf=r size=8 type=q align=4 words (r40.0) 

//.declare V0156 rf=r size=32 type=ud alias=V0042+0 align=16 words (r4.0) 

//.declare V0157 rf=r size=64 type=q align=16 words (r41.0) 

//.declare P08 rf=f size=2 type=uw align=1 words (f0.0) 

//.declare V0158 rf=r size=64 type=q align=16 words (r43.0) 

//.declare V0159 rf=r size=8 type=q alias=V0034+0 align=4 words (r5.2) 

//.declare V0160 rf=r size=64 type=uq alias=V0158+0 align=16 words (r43.0) 

//.declare V0161 rf=r size=32 type=f align=16 words (r45.0) 

//.declare V0162 rf=r size=64 type=q align=16 words (r46.0) 

//.declare V0163 rf=r size=8 type=q alias=V0036+0 align=4 words (r6.1) 

//.declare V0164 rf=r size=64 type=uq alias=V0162+0 align=16 words (r46.0) 

//.declare V0165 rf=r size=32 type=f align=16 words (r48.0) 

//.declare V0166 rf=r size=64 type=q align=16 words (r49.0) 

//.declare V0167 rf=r size=8 type=q alias=V0038+0 align=4 words (r7.0) 

//.declare V0168 rf=r size=64 type=uq alias=V0166+0 align=16 words (r49.0) 

//.declare V0169 rf=r size=64 type=uq align=16 words (r51.0) 

//.declare V0170 rf=r size=8 type=uq align=4 words (r5.0) 

//.declare V0171 rf=r size=8 type=uq align=4 words (r5.1) 

//.declare V0172 rf=r size=8 type=uq align=4 words (r5.3) 

//.declare V0173 rf=r size=8 type=uq align=4 words (r6.0) 

//.declare V0174 rf=r size=8 type=uq align=4 words (r6.2) 

//.declare V0175 rf=r size=8 type=uq align=4 words (r6.3) 

//.declare V0176 rf=r size=8 type=q align=4 words (r7.1) 

//.declare V0177 rf=r size=8 type=q align=4 words (r7.2) 

//.declare V0178 rf=r size=8 type=q align=4 words (r8.0) 

//.declare V0179 rf=r size=8 type=q align=4 words (r8.1) 

//.declare V0180 rf=r size=8 type=q align=4 words (r8.2) 

//.declare V0181 rf=r size=8 type=q align=4 words (r9.0) 

//.declare V0182 rf=r size=8 type=q align=4 words (r9.1) 

//.declare V0183 rf=r size=8 type=q align=4 words (r9.2) 

//.declare V0184 rf=r size=8 type=q align=4 words (r10.0) 

//.declare  rf=r size=32 type=ud align=16 words (r127.0) 

//.declare  rf=r size=2 type=w align=1 words (r40.4) 

//.declare  rf=r size=2 type=w align=1 words (r40.5) 

//.declare  rf=r size=2 type=w align=1 words (r40.6) 

//.declare  rf=r size=2 type=w align=1 words (r40.7) 

//.declare  rf=r size=2 type=w align=1 words (r40.8) 

//.declare  rf=r size=2 type=w align=1 words (r40.9) 

//.declare  rf=r size=2 type=w align=1 words (r40.10) 

//.declare  rf=r size=2 type=w align=1 words (r40.11) 

//.declare  rf=r size=2 type=w align=1 words (r40.12) 

//.declare  rf=r size=2 type=w align=1 words (r40.13) 

//.declare  rf=r size=2 type=w align=1 words (r40.14) 

//.declare  rf=r size=2 type=w align=1 words (r53.0) 

//.declare  rf=r size=2 type=w align=1 words (r53.1) 

//.declare  rf=r size=2 type=w align=1 words (r53.2) 


// .inputs 

// +----------+----------+--------+----------+------------+ 

// | id       | type     |  bytes | at       | class      | 

// +----------+----------+--------+----------+------------+ 

// | V0044    | :w x 16  |     32 | r1       | general    | 

// | V0045    | :w x 16  |     32 | r2       | general    | 

// | V0046    | :w x 16  |     32 | r3       | general    | 

// | V0042    | :d x 8   |     32 | r4       | general    | 

// | V0170    | :uq      |      8 | r5       | general    | 

// | V0171    | :uq      |      8 | r5+8     | general    | 

// | V0034    | :uq      |      8 | r5+16    | general    | 

// | V0172    | :uq      |      8 | r5+24    | general    | 

// | V0173    | :uq      |      8 | r6       | general    | 

// | V0036    | :uq      |      8 | r6+8     | general    | 

// | V0174    | :uq      |      8 | r6+16    | general    | 

// | V0175    | :uq      |      8 | r6+24    | general    | 

// | V0038    | :uq      |      8 | r7       | general    | 

// | V0176    | :q       |      8 | r7+8     | general    | 

// | V0177    | :q       |      8 | r7+16    | general    | 

// | V0035    | :q       |      8 | r7+24    | general    | 

// | V0178    | :q       |      8 | r8       | general    | 

// | V0179    | :q       |      8 | r8+8     | general    | 

// | V0180    | :q       |      8 | r8+16    | general    | 

// | V0037    | :q       |      8 | r8+24    | general    | 

// | V0181    | :q       |      8 | r9       | general    | 

// | V0182    | :q       |      8 | r9+8     | general    | 

// | V0183    | :q       |      8 | r9+16    | general    | 

// | V0039    | :q       |      8 | r9+24    | general    | 

// | V0184    | :q       |      8 | r10      | general    | 

// | V0047    | :uq      |      8 | r10+8    | general    | 

// | V0043    | :d x 3   |     12 | r10+16   | general    | 

// +----------+----------+--------+----------+------------+ 


//_main: 

L0: 

        mov (1|M0)               null<1>:ud    0x76948161:ud                              {Breakpoint} // [00000] 

(W)     mov (8|M0)               r11.0<1>:ud   r0.0<1;1,0>:ud                   // [00010] 

(W)     and (1|M0)               r5.0<1>:ud    r0.5<0;1,0>:ud    0x1FF:ud              // &1:[00020] 

(W)     or (1|M0)                cr0.0<1>:ud   cr0.0<0;1,0>:ud   0x4C0:uw              {Switch} // $1:&2:[00030] 

(W)     mov (8|M0)               r8.0<1>:w     0x76543210:v                               // $2:&3:[00040] 

(W)     mov (1|M0)               r5.1<1>:ud    r5.0<0;1,0>:ud                   // $3:&4:[00050] 

(W)     mul (1|M0)               r5.2<1>:d     r5.1<0;1,0>:d     64:w               // $4:&5:[00060] 

(W)     add (1|M0)               r5.3<1>:d     r5.2<0;1,0>:d     0:w               // $5:&6:[00070] 

        mov (8|M0)               r12.0<1>:d    r8.0<8;8,1>:uw                   // $6:&7:[00080] 

        mul (8|M0)               r12.0<1>:d    r12.0<8;8,1>:d    8:w               // $7:&8:[00090] 

        add (8|M0)               r12.0<1>:d    r5.3<0;1,0>:d     r12.0<8;8,1>:d   // $8:&9:[000a0] 

        mov (8|M0)               r13.0<1>:q    r12.0<8;8,1>:d                   // $9:&10:[000b0] 

        add (8|M0)               r13.0<1>:q    r10.1<0;1,0>:q    r13.0<4;4,1>:q   // $10:&11:[000c0] 

        mov (8|M0)               r15.0<1>:uq   0x0:uw                              // $11:&12:[000d0] 

        sends (8|M0)             null:ud  r13     r15     0x8C            0x040682FF           // wr:2+2, rd:0; hdc.dc1; a64 qword scattering write // $12:&13:[000e0] 

(W)     mov (1|M0)               r40.4<1>:w    0:w                               // $13:&14:[000f0] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.4<0;1,0>:w    0:w               // $13:&15:[00100] 

(W)     mov (1|M0)               r40.5<1>:w    1:w                               // $14:&16:[00110] 

(W&f0.0) sel (1|M0)              r5.24<2>:b    r40.5<0;1,0>:w    5:w               // $14:&17:[00120] 

(W)     mov (1|M0)               r5.7<1>:d     r5.24<0;1,0>:ub                  // $15:&18:[00130] 

(W)     add (1|M0)               r6.0<1>:d     r5.7<0;1,0>:d     0:w               // $16:&19:[00140] 

(W)     mul (1|M0)               r5.13<1>:uw   r6.0<0;1,0>:uw    0x4:uw              // $17:&20:[00150] 

(W)     add (1|M0)               a0.0<1>:uw    r5.13<0;1,0>:uw   0x160:uw              // $18:&21:[00160] 

(W)     mov (1|M0)               r6.1<1>:d     r[a0.0]<0;1,0>:d                 // $19:&22:[00170] 

(W)     mul (1|M0)               r6.4<1>:d     r10.4<0;1,0>:d    r6.1<0;1,0>:d    // $20:&23:[00180] 

        mov (8|M0)               r17.0<1>:d    r1.0<8;8,1>:uw                   // $21:&24:[00190] 

        add (8|M0)               r17.0<1>:d    r17.0<8;8,1>:d    r6.4<0;1,0>:d    // $22:&25:[001a0] 

        add (8|M0)               r17.0<1>:d    r17.0<8;8,1>:d    r4.0<0;1,0>:d    // $23:&26:[001b0] 

        mov (8|M0)               r18.0<1>:q    r17.0<8;8,1>:ud                  // $24:&27:[001c0] 

(W)     mov (1|M0)               r40.6<1>:w    1:w                               // $25:&28:[001d0] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.6<0;1,0>:w    0:w               // $25:&29:[001e0] 

(W)     mov (1|M0)               r40.7<1>:w    1:w                               // $26:&30:[001f0] 

(W&f0.0) sel (1|M0)              r6.20<2>:b    r40.7<0;1,0>:w    5:w               // $26:&31:[00200] 

(W)     mov (1|M0)               r6.6<1>:d     r6.20<0;1,0>:ub                  // $27:&32:[00210] 

(W)     add (1|M0)               r6.7<1>:d     r6.6<0;1,0>:d     1:w               // $28:&33:[00220] 

(W)     mul (1|M0)               r6.11<1>:uw   r6.14<0;1,0>:uw   0x4:uw              // $29:&34:[00230] 

(W)     add (1|M0)               a0.0<1>:uw    r6.11<0;1,0>:uw   0x160:uw              // $30:&35:[00240] 

(W)     mov (1|M0)               r7.2<1>:d     r[a0.0]<0;1,0>:d                 // $31:&36:[00250] 

(W)     mul (1|M0)               r7.3<1>:d     r10.5<0;1,0>:d    r7.2<0;1,0>:d    // $32:&37:[00260] 

        mov (8|M0)               r20.0<1>:d    r2.0<8;8,1>:uw                   // $33:&38:[00270] 

        add (8|M0)               r20.0<1>:d    r20.0<8;8,1>:d    r7.3<0;1,0>:d    // $34:&39:[00280] 

        add (8|M0)               r20.0<1>:d    r20.0<8;8,1>:d    r4.1<0;1,0>:d    // $35:&40:[00290] 

        mov (8|M0)               r21.0<1>:q    r20.0<8;8,1>:ud                  // $36:&41:[002a0] 

(W)     mov (1|M0)               r40.8<1>:w    2:w                               // $37:&42:[002b0] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.8<0;1,0>:w    0:w               // $37:&43:[002c0] 

(W)     mov (1|M0)               r40.9<1>:w    1:w                               // $38:&44:[002d0] 

(W&f0.0) sel (1|M0)              r7.16<2>:b    r40.9<0;1,0>:w    5:w               // $38:&45:[002e0] 

(W)     mov (1|M0)               r7.5<1>:d     r7.16<0;1,0>:ub                  // $39:&46:[002f0] 

(W)     add (1|M0)               r8.4<1>:d     r7.5<0;1,0>:d     2:w               // $40:&47:[00300] 

(W)     mul (1|M0)               r7.9<1>:uw    r8.8<0;1,0>:uw    0x4:uw              // $41:&48:[00310] 

(W)     add (1|M0)               a0.0<1>:uw    r7.9<0;1,0>:uw    0x160:uw              // $42:&49:[00320] 

(W)     mov (1|M0)               r8.5<1>:d     r[a0.0]<0;1,0>:d                 // $43:&50:[00330] 

(W)     mul (1|M0)               r9.0<1>:d     r10.6<0;1,0>:d    r8.5<0;1,0>:d    // $44:&51:[00340] 

        mov (8|M0)               r23.0<1>:d    r3.0<8;8,1>:uw                   // $45:&52:[00350] 

        add (8|M0)               r23.0<1>:d    r23.0<8;8,1>:d    r9.0<0;1,0>:d    // $46:&53:[00360] 

        add (8|M0)               r23.0<1>:d    r23.0<8;8,1>:d    r4.2<0;1,0>:d    // $47:&54:[00370] 

        mov (8|M0)               r24.0<1>:q    r23.0<8;8,1>:ud                  // $48:&55:[00380] 

// File: /localdisk/work/etotmeni/stepping/sum.py 

 
// Line 15: @dppy.kernel 


        mov (8|M0)               r26.0<1>:d    r1.0<8;8,1>:uw                   // #15:$51:&56:[00390] 

        mov (8|M0)               r27.0<1>:q    r26.0<8;8,1>:ud                  // #15:$52:&57:[003a0] 

        mov (8|M0)               r29.0<1>:d    r2.0<8;8,1>:uw                   // #15:$53:&58:[003b0] 

        mov (8|M0)               r30.0<1>:q    r29.0<8;8,1>:ud                  // #15:$54:&59:[003c0] 

        mov (8|M0)               r32.0<1>:d    r3.0<8;8,1>:uw                   // #15:$55:&60:[003d0] 

        mov (8|M0)               r33.0<1>:q    r32.0<8;8,1>:ud                  // #15:$56:&61:[003e0] 

(W)     mov (1|M0)               r40.10<1>:w   0:w                               // #15:$57:&62:[003f0] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.10<0;1,0>:w   0:w               // #15:$57:&63:[00400] 

(W)     mov (1|M0)               r40.11<1>:w   1:w                               // #15:$58:&64:[00410] 

(W&f0.0) sel (1|M0)              r9.4<2>:b     r40.11<0;1,0>:w   5:w               // #15:$58:&65:[00420] 

(W)     mov (1|M0)               r9.2<1>:d     r9.4<0;1,0>:ub                   // #15:$59:&66:[00430] 

(W)     add (1|M0)               r9.3<1>:d     r9.2<0;1,0>:d     0:w               // #15:$60:&67:[00440] 

(W)     mul (1|M0)               r9.3<1>:uw    r9.6<0;1,0>:uw    0x4:uw              // #15:$61:&68:[00450] 

(W)     add (1|M0)               a0.0<1>:uw    r9.3<0;1,0>:uw    0x160:uw              // #15:$62:&69:[00460] 

(W)     mov (1|M0)               r9.4<1>:d     r[a0.0]<0;1,0>:d                 // #15:$63:&70:[00470] 

(W)     mov (1|M0)               r10.0<1>:q    r9.4<0;1,0>:ud                   // #15:$64:&71:[00480] 

(W)     mov (1|M0)               r40.12<1>:w   1:w                               // #15:$65:&72:[00490] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.12<0;1,0>:w   0:w               // #15:$65:&73:[004a0] 

(W)     mov (1|M0)               r40.13<1>:w   1:w                               // #15:$66:&74:[004b0] 

(W&f0.0) sel (1|M0)              r9.20<2>:b    r40.13<0;1,0>:w   5:w               // #15:$66:&75:[004c0] 

(W)     mov (1|M0)               r10.7<1>:d    r9.20<0;1,0>:ub                  // #15:$67:&76:[004d0] 

(W)     add (1|M0)               r35.0<1>:d    r10.7<0;1,0>:d    1:w               // #15:$68:&77:[004e0] 

(W)     mul (1|M0)               r9.11<1>:uw   r35.0<0;1,0>:uw   0x4:uw              // #15:$69:&78:[004f0] 

(W)     add (1|M0)               a0.0<1>:uw    r9.11<0;1,0>:uw   0x160:uw              // #15:$70:&79:[00500] 

(W)     mov (1|M0)               r35.1<1>:d    r[a0.0]<0;1,0>:d                 // #15:$71:&80:[00510] 

(W)     mov (1|M0)               r35.1<1>:q    r35.1<0;1,0>:ud                  // #15:$72:&81:[00520] 

(W)     mov (1|M0)               r40.14<1>:w   2:w                               // #15:$73:&82:[00530] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r40.14<0;1,0>:w   0:w               // #15:$73:&83:[00540] 

(W)     mov (1|M0)               r53.0<1>:w    1:w                               // #15:$74:&84:[00550] 

(W&f0.0) sel (1|M0)              r35.16<2>:b   r53.0<0;1,0>:w    5:w               // #15:$74:&85:[00560] 

(W)     mov (1|M0)               r35.5<1>:d    r35.16<0;1,0>:ub                 // #15:$75:&86:[00570] 

(W)     add (1|M0)               r35.6<1>:d    r35.5<0;1,0>:d    2:w               // #15:$76:&87:[00580] 

(W)     mul (1|M0)               r35.9<1>:uw   r35.12<0;1,0>:uw  0x4:uw              // #15:$77:&88:[00590] 

(W)     add (1|M0)               a0.0<1>:uw    r35.9<0;1,0>:uw   0x160:uw              // #15:$78:&89:[005a0] 

(W)     mov (1|M0)               r35.7<1>:d    r[a0.0]<0;1,0>:d                 // #15:$79:&90:[005b0] 

(W)     mov (1|M0)               r36.0<1>:q    r35.7<0;1,0>:ud                  // #15:$80:&91:[005c0] 

// Line 17: i = dppy.get_global_id(0) 

 
(W)     mov (1|M0)               r53.1<1>:w    0:w                               // #17:$83:&92:[005d0] 

(W)     cmp (1|M0)    (eq)f0.0   null<1>:w     r53.1<0;1,0>:w    0:w               // #17:$83:&93:[005e0] 

(W)     mov (1|M0)               r53.2<1>:w    1:w                               // #17:$84:&94:[005f0] 

(W&f0.0) sel (1|M0)              r36.8<2>:b    r53.2<0;1,0>:w    5:w               // #17:$84:&95:[00600] 

(W)     mov (1|M0)               r36.3<1>:d    r36.8<0;1,0>:ub                  // #17:$85:&96:[00610] 

(W)     add (1|M0)               r36.4<1>:d    r36.3<0;1,0>:d    0:w               // #17:$86:&97:[00620] 

(W)     mul (1|M0)               r36.5<1>:uw   r36.8<0;1,0>:uw   0x4:uw              // #17:$87:&98:[00630] 

(W)     add (1|M0)               a0.0<1>:uw    r36.5<0;1,0>:uw   0x160:uw              // #17:$88:&99:[00640] 

(W)     mov (1|M0)               r36.5<1>:d    r[a0.0]<0;1,0>:d                 // #17:$89:&100:[00650] 

(W)     mul (1|M0)               r36.3<1>:q    r10.4<0;1,0>:ud   r36.5<0;1,0>:ud  // #17:$90:&101:[00660] 

        mov (8|M0)               r37.0<1>:d    r1.0<8;8,1>:uw                   // #17:$91:&102:[00670] 

        mov (8|M0)               r38.0<1>:q    r37.0<8;8,1>:ud                  // #17:$92:&103:[00680] 

        add (8|M0)               r38.0<1>:q    r38.0<4;4,1>:q    r36.3<0;1,0>:q   // #17:$93:&104:[00690] 

(W)     mov (1|M0)               r40.0<1>:q    r4.0<0;1,0>:ud                   // #17:$94:&105:[006a0] 

        add (8|M0)               r41.0<1>:q    r38.0<4;4,1>:q    r40.0<0;1,0>:q   // #17:$95:&106:[006b0] 

// Line 18: c[i] = kernel_sum(a[i], b[i]) 

 

 

        cmp (8|M0)    (lt)f0.0   null<1>:q     r41.0<4;4,1>:q    0:w               // #18:$98:&107:[006c0] 

(f0.0)  sel (8|M0)               r43.0<1>:q    r7.3<0;1,0>:q     0:w               // #18:$99:&108:[006d0] 

        add (8|M0)               r43.0<1>:q    r43.0<4;4,1>:q    r41.0<4;4,1>:q   // #18:$100:&109:[006e0] 

        shl (8|M0)               r43.0<1>:q    r43.0<4;4,1>:q    2:w               // #18:$101:&110:[006f0] 

        add (8|M0)               r43.0<1>:q    r5.2<0;1,0>:q     r43.0<4;4,1>:q   // #18:$102:&111:[00700] 

        send (8|M0)              r45:f    r43:uq  0xC            0x041401FF           // wr:2+0, rd:1; hdc.dc1; a64 dword gathering read // #18:$103:&112:[00710] 

(f0.0)  sel (8|M0)               r46.0<1>:q    r8.3<0;1,0>:q     0:w               // #18:$104:&113:[00720] 

        add (8|M0)               r46.0<1>:q    r46.0<4;4,1>:q    r41.0<4;4,1>:q   // #18:$105:&114:[00730] 

        shl (8|M0)               r46.0<1>:q    r46.0<4;4,1>:q    2:w               // #18:$106:&115:[00740] 

        add (8|M0)               r46.0<1>:q    r6.1<0;1,0>:q     r46.0<4;4,1>:q   // #18:$107:&116:[00750] 

        send (8|M0)              r48:f    r46:uq  0xC            0x041401FF           // wr:2+0, rd:1; hdc.dc1; a64 dword gathering read // #18:$108:&117:[00760] 

        add (8|M0)               r45.0<1>:f    r45.0<8;8,1>:f    r48.0<8;8,1>:f   // #18:$109:&118:[00770] 

(f0.0)  sel (8|M0)               r49.0<1>:q    r9.3<0;1,0>:q     0:w               // #18:$110:&119:[00780] 

        add (8|M0)               r49.0<1>:q    r49.0<4;4,1>:q    r41.0<4;4,1>:q   // #18:$111:&120:[00790] 

        shl (8|M0)               r49.0<1>:q    r49.0<4;4,1>:q    2:w               // #18:$112:&121:[007a0] 

        add (8|M0)               r49.0<1>:q    r7.0<0;1,0>:q     r49.0<4;4,1>:q   // #18:$113:&122:[007b0] 

        sends (8|M0)             null:ud  r49     r45     0x4C            0x040681FF           // wr:2+1, rd:0; hdc.dc1; a64 dword scattering write // #18:$114:&123:[007c0] 

        mov (8|M0)               r51.0<1>:uq   0x0:uw                              // #18:$117:&124:[007d0] 

        sends (8|M0)             null:ud  r13     r51     0x8C            0x040682FF           // wr:2+2, rd:0; hdc.dc1; a64 qword scattering write // #18:$118:&125:[007e0] 

(W)     mov (8|M0)               r127.0<1>:ud  r11.0<8;8,1>:ud                  // #18:$119:&126:[007f0] 

(W)     send (8|M0)              null     r127    0x27            0x02000010           {EOT} // wr:1+0, rd:0; spawner; end of thread // #18:$119:&127:[00800]// Bank Conflict Statistics:  

// -- GOOD: 0 
// --  BAD: 0 
// --   OK: 0

@1e-to
Copy link
Contributor Author

1e-to commented Jun 2, 2021 •
edited by PokhodenkoSA
Loading

Also I checked the dbg file (type data), but there is nothing there either

hexdump -C xxxx.dbg > dump.txt

Output

00000000  10 d0 ad de 01 00 ca 00  64 70 70 79 50 79 5f 64  |........dppyPy_d| 

00000010  70 70 79 5f 70 79 5f 64  65 76 66 6e 5f 5f 35 46  |ppy_py_devfn__5F| 

00000020  5f 5f 35 46 5f 6d 61 69  6e 5f 35 46 5f 5f 35 46  |__5F_main_5F__5F| 

00000030  5f 5f 32 45 5f 64 61 74  61 5f 35 46 5f 70 61 72  |__2E_data_5F_par| 

00000040  61 6c 6c 65 6c 5f 35 46  5f 73 75 6d 5f 32 34 5f  |allel_5F_sum_24_| 

00000050  31 5f 32 45 5f 61 72 72  61 79 5f 32 38 5f 66 6c  |1_2E_array_28_fl| 

00000060  6f 61 74 33 32 5f 32 43  5f 5f 32 30 5f 31 64 5f  |oat32_2C__20_1d_| 

00000070  32 43 5f 5f 32 30 5f 43  5f 32 39 5f 5f 32 45 5f  |2C__20_C_29__2E_| 

00000080  61 72 72 61 79 5f 32 38  5f 66 6c 6f 61 74 33 32  |array_28_float32| 

00000090  5f 32 43 5f 5f 32 30 5f  31 64 5f 32 43 5f 5f 32  |_2C__20_1d_2C__2| 

000000a0  30 5f 43 5f 32 39 5f 5f  32 45 5f 61 72 72 61 79  |0_C_29__2E_array| 

000000b0  5f 32 38 5f 66 6c 6f 61  74 33 32 5f 32 43 5f 5f  |_28_float32_2C__| 

000000c0  32 30 5f 31 64 5f 32 43  5f 5f 32 30 5f 43 5f 32  |20_1d_2C__20_C_2| 

000000d0  39 5f 00 00 00 00 00 00  00 00 70 00 00 00 01 00  |9_........p.....| 

000000e0  00 00 30 00 00 00 02 00  00 00 40 00 00 00 03 00  |..0.......@.....| 

000000f0  00 00 50 00 00 00 04 00  00 00 60 00 00 00 05 00  |..P.......`.....| 

00000100  00 00 70 00 00 00 06 00  00 00 80 00 00 00 07 00  |..p.............| 

00000110  00 00 90 00 00 00 08 00  00 00 a0 00 00 00 09 00  |................| 

00000120  00 00 b0 00 00 00 0a 00  00 00 c0 00 00 00 0b 00  |................| 

00000130  00 00 d0 00 00 00 0c 00  00 00 e0 00 00 00 0d 00  |................| 

00000140  00 00 f0 00 00 00 0e 00  00 00 10 01 00 00 0f 00  |................| 

00000150  00 00 30 01 00 00 10 00  00 00 40 01 00 00 11 00  |..0.......@.....| 

00000160  00 00 50 01 00 00 12 00  00 00 60 01 00 00 13 00  |..P.......`.....| 

00000170  00 00 70 01 00 00 14 00  00 00 80 01 00 00 15 00  |..p.............| 

00000180  00 00 90 01 00 00 16 00  00 00 a0 01 00 00 17 00  |................| 

00000190  00 00 b0 01 00 00 18 00  00 00 c0 01 00 00 19 00  |................| 

000001a0  00 00 d0 01 00 00 1a 00  00 00 f0 01 00 00 1b 00  |................| 

000001b0  00 00 10 02 00 00 1c 00  00 00 20 02 00 00 1d 00  |.......... .....| 

000001c0  00 00 30 02 00 00 1e 00  00 00 40 02 00 00 1f 00  |..0.......@.....| 

000001d0  00 00 50 02 00 00 20 00  00 00 60 02 00 00 21 00  |..P... ...`...!.| 

000001e0  00 00 70 02 00 00 22 00  00 00 80 02 00 00 23 00  |..p...".......#.| 

000001f0  00 00 90 02 00 00 24 00  00 00 a0 02 00 00 25 00  |......$.......%.| 

00000200  00 00 b0 02 00 00 26 00  00 00 d0 02 00 00 27 00  |......&.......'.| 

00000210  00 00 f0 02 00 00 28 00  00 00 00 03 00 00 29 00  |......(.......).| 

00000220  00 00 10 03 00 00 2a 00  00 00 20 03 00 00 2b 00  |......*... ...+.| 

00000230  00 00 30 03 00 00 2c 00  00 00 40 03 00 00 2d 00  |..0...,...@...-.| 

00000240  00 00 50 03 00 00 2e 00  00 00 60 03 00 00 2f 00  |..P.......`.../.| 

00000250  00 00 70 03 00 00 30 00  00 00 80 03 00 00 33 00  |..p...0.......3.| 

00000260  00 00 90 03 00 00 34 00  00 00 a0 03 00 00 35 00  |......4.......5.| 

00000270  00 00 b0 03 00 00 36 00  00 00 c0 03 00 00 37 00  |......6.......7.| 

00000280  00 00 d0 03 00 00 38 00  00 00 e0 03 00 00 39 00  |......8.......9.| 

00000290  00 00 f0 03 00 00 3a 00  00 00 10 04 00 00 3b 00  |......:.......;.| 

000002a0  00 00 30 04 00 00 3c 00  00 00 40 04 00 00 3d 00  |..0...<...@...=.| 

000002b0  00 00 50 04 00 00 3e 00  00 00 60 04 00 00 3f 00  |..P...>...`...?.| 

000002c0  00 00 70 04 00 00 40 00  00 00 80 04 00 00 41 00  |..p...@.......A.| 

000002d0  00 00 90 04 00 00 42 00  00 00 b0 04 00 00 43 00  |......B.......C.| 

000002e0  00 00 d0 04 00 00 44 00  00 00 e0 04 00 00 45 00  |......D.......E.| 

000002f0  00 00 f0 04 00 00 46 00  00 00 00 05 00 00 47 00  |......F.......G.| 

00000300  00 00 10 05 00 00 48 00  00 00 20 05 00 00 49 00  |......H... ...I.| 

00000310  00 00 30 05 00 00 4a 00  00 00 50 05 00 00 4b 00  |..0...J...P...K.| 

00000320  00 00 70 05 00 00 4c 00  00 00 80 05 00 00 4d 00  |..p...L.......M.| 

00000330  00 00 90 05 00 00 4e 00  00 00 a0 05 00 00 4f 00  |......N.......O.| 

00000340  00 00 b0 05 00 00 50 00  00 00 c0 05 00 00 53 00  |......P.......S.| 

00000350  00 00 d0 05 00 00 54 00  00 00 f0 05 00 00 55 00  |......T.......U.| 

00000360  00 00 10 06 00 00 56 00  00 00 20 06 00 00 57 00  |......V... ...W.| 

00000370  00 00 30 06 00 00 58 00  00 00 40 06 00 00 59 00  |..0...X...@...Y.| 

00000380  00 00 50 06 00 00 5a 00  00 00 60 06 00 00 5b 00  |..P...Z...`...[.| 

00000390  00 00 70 06 00 00 5c 00  00 00 80 06 00 00 5d 00  |..p...\.......].| 

000003a0  00 00 90 06 00 00 5e 00  00 00 a0 06 00 00 5f 00  |......^......._.| 

000003b0  00 00 b0 06 00 00 62 00  00 00 c0 06 00 00 63 00  |......b.......c.| 

000003c0  00 00 d0 06 00 00 64 00  00 00 e0 06 00 00 65 00  |......d.......e.| 

000003d0  00 00 f0 06 00 00 66 00  00 00 00 07 00 00 67 00  |......f.......g.| 

000003e0  00 00 10 07 00 00 68 00  00 00 20 07 00 00 69 00  |......h... ...i.| 

000003f0  00 00 30 07 00 00 6a 00  00 00 40 07 00 00 6b 00  |..0...j...@...k.| 

00000400  00 00 50 07 00 00 6c 00  00 00 60 07 00 00 6d 00  |..P...l...`...m.| 

00000410  00 00 70 07 00 00 6e 00  00 00 80 07 00 00 6f 00  |..p...n.......o.| 

00000420  00 00 90 07 00 00 70 00  00 00 a0 07 00 00 71 00  |......p.......q.| 

00000430  00 00 b0 07 00 00 72 00  00 00 c0 07 00 00 75 00  |......r.......u.| 

00000440  00 00 d0 07 00 00 76 00  00 00 e0 07 00 00 77 00  |......v.......w.| 

00000450  00 00 f0 07 00 00 78 00  00 00 10 08 00 00 9c 00  |......x.........| 

00000460  00 00 03 00 56 33 33 01  00 00 00 77 00 02 02 0b  |....V33....w....| 

00000470  00 00 00 03 00 56 33 34  01 00 00 00 77 00 02 02  |.....V34....w...| 

00000480  05 00 10 00 03 00 56 33  35 01 00 00 00 77 00 02  |......V35....w..| 

00000490  02 07 00 18 00 03 00 56  33 36 01 00 00 00 77 00  |.......V36....w.| 

000004a0  02 02 06 00 08 00 03 00  56 33 37 01 00 00 00 77  |........V37....w| 

000004b0  00 02 02 08 00 18 00 03  00 56 33 38 01 00 00 00  |.........V38....| 

000004c0  77 00 02 02 07 00 00 00  03 00 56 33 39 01 00 00  |w.........V39...| 

000004d0  00 77 00 02 02 09 00 18  00 03 00 56 34 31 01 00  |.w.........V41..| 

000004e0  00 00 77 00 02 02 0b 00  00 00 03 00 56 34 32 01  |..w.........V42.| 

000004f0  00 00 00 77 00 02 02 04  00 00 00 03 00 56 34 33  |...w.........V43| 

00000500  01 00 00 00 77 00 02 02  0a 00 10 00 03 00 56 34  |....w.........V4| 

00000510  34 01 00 00 00 77 00 02  02 01 00 00 00 03 00 56  |4....w.........V| 

00000520  34 35 01 00 00 00 77 00  02 02 02 00 00 00 03 00  |45....w.........| 

00000530  56 34 36 01 00 00 00 77  00 02 02 03 00 00 00 03  |V46....w........| 

00000540  00 56 34 37 01 00 00 00  77 00 02 02 0a 00 08 00  |.V47....w.......| 

00000550  03 00 56 35 37 01 00 00  00 77 00 02 02 08 00 00  |..V57....w......| 

00000560  00 03 00 56 35 38 01 00  00 00 77 00 02 02 05 00  |...V58....w.....| 

00000570  04 00 03 00 56 36 30 01  00 00 00 77 00 02 02 05  |....V60....w....| 

00000580  00 00 00 03 00 56 36 31  01 00 00 00 77 00 02 02  |.....V61....w...| 

00000590  05 00 04 00 03 00 56 36  32 01 00 00 00 77 00 02  |......V62....w..| 

000005a0  02 05 00 08 00 03 00 56  36 33 01 00 00 00 77 00  |.......V63....w.| 

000005b0  02 02 05 00 0c 00 03 00  56 36 34 01 00 00 00 77  |........V64....w| 

000005c0  00 02 02 0c 00 00 00 03  00 56 36 35 01 00 00 00  |.........V65....| 

000005d0  77 00 02 02 08 00 00 00  03 00 56 36 36 01 00 00  |w.........V66...| 

000005e0  00 77 00 02 02 0a 00 08  00 03 00 56 36 37 01 00  |.w.........V67..| 

000005f0  00 00 77 00 02 02 0d 00  00 00 03 00 56 36 38 01  |..w.........V68.| 

00000600  00 00 00 77 00 02 02 0d  00 00 00 03 00 56 36 39  |...w.........V69| 

00000610  01 00 00 00 77 00 02 02  0f 00 00 00 02 00 50 31  |....w.........P1| 

00000620  01 00 0e 00 0e 00 01 01  00 00 00 00 03 00 56 37  |..............V7| 

00000630  30 01 00 00 00 77 00 02  02 05 00 18 00 03 00 56  |0....w.........V| 

00000640  37 31 01 00 00 00 77 00  02 02 05 00 1c 00 03 00  |71....w.........| 

00000650  56 37 32 01 00 00 00 77  00 02 02 05 00 18 00 03  |V72....w........| 

00000660  00 56 37 33 01 00 00 00  77 00 02 02 06 00 00 00  |.V73....w.......| 

00000670  03 00 56 37 34 01 00 00  00 77 00 02 02 06 00 04  |..V74....w......| 

00000680  00 03 00 56 37 35 01 00  00 00 77 00 02 02 06 00  |...V75....w.....| 

00000690  00 00 03 00 56 37 36 01  00 00 00 77 00 02 02 05  |....V76....w....| 

000006a0  00 1a 00 02 00 41 30 01  00 13 00 13 00 00 00 00  |.....A0.........| 

000006b0  00 00 00 03 00 56 37 37  01 00 00 00 77 00 02 02  |.....V77....w...| 

000006c0  06 00 10 00 03 00 56 37  38 01 00 00 00 77 00 02  |......V78....w..| 

000006d0  02 11 00 00 00 03 00 56  37 39 01 00 00 00 77 00  |.......V79....w.| 

000006e0  02 02 01 00 00 00 03 00  56 38 30 01 00 00 00 77  |........V80....w| 

000006f0  00 02 02 12 00 00 00 03  00 56 38 31 01 00 00 00  |.........V81....| 

00000700  77 00 02 02 11 00 00 00  02 00 50 32 01 00 1a 00  |w.........P2....| 

00000710  1a 00 01 01 00 00 00 00  03 00 56 38 32 01 00 00  |..........V82...| 

00000720  00 77 00 02 02 06 00 14  00 03 00 56 38 33 01 00  |.w.........V83..| 

00000730  00 00 77 00 02 02 06 00  18 00 03 00 56 38 34 01  |..w.........V84.| 

00000740  00 00 00 77 00 02 02 06  00 14 00 03 00 56 38 35  |...w.........V85| 

00000750  01 00 00 00 77 00 02 02  06 00 1c 00 03 00 56 38  |....w.........V8| 

00000760  36 01 00 00 00 77 00 02  02 07 00 08 00 03 00 56  |6....w.........V| 

00000770  38 37 01 00 00 00 77 00  02 02 06 00 1c 00 03 00  |87....w.........| 

00000780  56 38 38 01 00 00 00 77  00 02 02 06 00 16 00 02  |V88....w........| 

00000790  00 41 31 01 00 1f 00 1f  00 00 00 00 00 00 00 03  |.A1.............| 

000007a0  00 56 38 39 01 00 00 00  77 00 02 02 07 00 0c 00  |.V89....w.......| 

000007b0  03 00 56 39 30 01 00 00  00 77 00 02 02 14 00 00  |..V90....w......| 

000007c0  00 03 00 56 39 31 01 00  00 00 77 00 02 02 02 00  |...V91....w.....| 

000007d0  00 00 03 00 56 39 32 01  00 00 00 77 00 02 02 15  |....V92....w....| 

000007e0  00 00 00 03 00 56 39 33  01 00 00 00 77 00 02 02  |.....V93....w...| 

000007f0  14 00 00 00 02 00 50 33  01 00 26 00 26 00 01 01  |......P3..&.&...| 

00000800  00 00 00 00 03 00 56 39  34 01 00 00 00 77 00 02  |......V94....w..| 

00000810  02 07 00 10 00 03 00 56  39 35 01 00 00 00 77 00  |.......V95....w.| 

00000820  02 02 07 00 14 00 03 00  56 39 36 01 00 00 00 77  |........V96....w| 

00000830  00 02 02 07 00 10 00 03  00 56 39 37 01 00 00 00  |.........V97....| 

00000840  77 00 02 02 08 00 10 00  03 00 56 39 38 01 00 00  |w.........V98...| 

00000850  00 77 00 02 02 08 00 14  00 03 00 56 39 39 01 00  |.w.........V99..| 

00000860  00 00 77 00 02 02 08 00  10 00 04 00 56 31 30 30  |..w.........V100| 

00000870  01 00 00 00 77 00 02 02  07 00 12 00 02 00 41 32  |....w.........A2| 

00000880  01 00 2b 00 2b 00 00 00  00 00 00 00 04 00 56 31  |..+.+.........V1| 

00000890  30 31 01 00 00 00 77 00  02 02 09 00 00 00 04 00  |01....w.........| 

000008a0  56 31 30 32 01 00 00 00  77 00 02 02 17 00 00 00  |V102....w.......| 

000008b0  04 00 56 31 30 33 01 00  00 00 77 00 02 02 03 00  |..V103....w.....| 

000008c0  00 00 04 00 56 31 30 34  01 00 00 00 77 00 02 02  |....V104....w...| 

000008d0  18 00 00 00 04 00 56 31  30 35 01 00 00 00 77 00  |......V105....w.| 

000008e0  02 02 17 00 00 00 04 00  56 31 30 36 01 00 00 00  |........V106....| 

000008f0  77 00 02 02 1a 00 00 00  04 00 56 31 30 37 01 00  |w.........V107..| 

00000900  00 00 77 00 02 02 1b 00  00 00 04 00 56 31 30 38  |..w.........V108| 

00000910  01 00 00 00 77 00 02 02  1a 00 00 00 04 00 56 31  |....w.........V1| 

00000920  30 39 01 00 00 00 77 00  02 02 1d 00 00 00 04 00  |09....w.........| 

00000930  56 31 31 30 01 00 00 00  77 00 02 02 1e 00 00 00  |V110....w.......| 

00000940  04 00 56 31 31 31 01 00  00 00 77 00 02 02 1d 00  |..V111....w.....| 

00000950  00 00 04 00 56 31 31 32  01 00 00 00 77 00 02 02  |....V112....w...| 

00000960  20 00 00 00 04 00 56 31  31 33 01 00 00 00 77 00  | .....V113....w.| 

00000970  02 02 21 00 00 00 04 00  56 31 31 34 01 00 00 00  |..!.....V114....| 

00000980  77 00 02 02 20 00 00 00  02 00 50 34 01 00 3a 00  |w... .....P4..:.| 

00000990  3a 00 01 01 00 00 00 00  04 00 56 31 31 35 01 00  |:.........V115..| 

000009a0  00 00 77 00 02 02 09 00  04 00 04 00 56 31 31 36  |..w.........V116| 

000009b0  01 00 00 00 77 00 02 02  09 00 08 00 04 00 56 31  |....w.........V1| 

000009c0  31 37 01 00 00 00 77 00  02 02 09 00 04 00 04 00  |17....w.........| 

000009d0  56 31 31 38 01 00 00 00  77 00 02 02 09 00 0c 00  |V118....w.......| 

000009e0  04 00 56 31 31 39 01 00  00 00 77 00 02 02 09 00  |..V119....w.....| 

000009f0  10 00 04 00 56 31 32 30  01 00 00 00 77 00 02 02  |....V120....w...| 

00000a00  09 00 0c 00 04 00 56 31  32 31 01 00 00 00 77 00  |......V121....w.| 

00000a10  02 02 09 00 06 00 02 00  41 33 01 00 3f 00 3f 00  |........A3..?.?.| 

00000a20  00 00 00 00 00 00 04 00  56 31 32 32 01 00 00 00  |........V122....| 

00000a30  77 00 02 02 0a 00 00 00  04 00 56 31 32 33 01 00  |w.........V123..| 

00000a40  00 00 77 00 02 02 09 00  10 00 02 00 50 35 01 00  |..w.........P5..| 

00000a50  42 00 42 00 01 01 00 00  00 00 04 00 56 31 32 34  |B.B.........V124| 

00000a60  01 00 00 00 77 00 02 02  09 00 14 00 04 00 56 31  |....w.........V1| 

00000a70  32 35 01 00 00 00 77 00  02 02 0a 00 1c 00 04 00  |25....w.........| 

00000a80  56 31 32 36 01 00 00 00  77 00 02 02 09 00 14 00  |V126....w.......| 

00000a90  04 00 56 31 32 37 01 00  00 00 77 00 02 02 23 00  |..V127....w...#.| 

00000aa0  00 00 04 00 56 31 32 38  01 00 00 00 77 00 02 02  |....V128....w...| 

00000ab0  23 00 04 00 04 00 56 31  32 39 01 00 00 00 77 00  |#.....V129....w.| 

00000ac0  02 02 23 00 00 00 04 00  56 31 33 30 01 00 00 00  |..#.....V130....| 

00000ad0  77 00 02 02 09 00 16 00  02 00 41 34 01 00 47 00  |w.........A4..G.| 

00000ae0  47 00 00 00 00 00 00 00  04 00 56 31 33 31 01 00  |G.........V131..| 

00000af0  00 00 77 00 02 02 23 00  08 00 04 00 56 31 33 32  |..w...#.....V132| 

00000b00  01 00 00 00 77 00 02 02  23 00 04 00 02 00 50 36  |....w...#.....P6| 

00000b10  01 00 4a 00 4a 00 01 01  00 00 00 00 04 00 56 31  |..J.J.........V1| 

00000b20  33 33 01 00 00 00 77 00  02 02 23 00 10 00 04 00  |33....w...#.....| 

00000b30  56 31 33 34 01 00 00 00  77 00 02 02 23 00 14 00  |V134....w...#...| 

00000b40  04 00 56 31 33 35 01 00  00 00 77 00 02 02 23 00  |..V135....w...#.| 

00000b50  10 00 04 00 56 31 33 36  01 00 00 00 77 00 02 02  |....V136....w...| 

00000b60  23 00 18 00 04 00 56 31  33 37 01 00 00 00 77 00  |#.....V137....w.| 

00000b70  02 02 23 00 1c 00 04 00  56 31 33 38 01 00 00 00  |..#.....V138....| 

00000b80  77 00 02 02 23 00 18 00  04 00 56 31 33 39 01 00  |w...#.....V139..| 

00000b90  00 00 77 00 02 02 23 00  12 00 02 00 41 35 01 00  |..w...#.....A5..| 

00000ba0  4f 00 4f 00 00 00 00 00  00 00 04 00 56 31 34 30  |O.O.........V140| 

00000bb0  01 00 00 00 77 00 02 02  24 00 00 00 04 00 56 31  |....w...$.....V1| 

00000bc0  34 31 01 00 00 00 77 00  02 02 23 00 1c 00 02 00  |41....w...#.....| 

00000bd0  50 37 01 00 54 00 54 00  01 01 00 00 00 00 04 00  |P7..T.T.........| 

00000be0  56 31 34 32 01 00 00 00  77 00 02 02 24 00 08 00  |V142....w...$...| 

00000bf0  04 00 56 31 34 33 01 00  00 00 77 00 02 02 24 00  |..V143....w...$.| 

00000c00  0c 00 04 00 56 31 34 34  01 00 00 00 77 00 02 02  |....V144....w...| 

00000c10  24 00 08 00 04 00 56 31  34 35 01 00 00 00 77 00  |$.....V145....w.| 

00000c20  02 02 24 00 10 00 04 00  56 31 34 36 01 00 00 00  |..$.....V146....| 

00000c30  77 00 02 02 24 00 14 00  04 00 56 31 34 37 01 00  |w...$.....V147..| 

00000c40  00 00 77 00 02 02 24 00  10 00 04 00 56 31 34 38  |..w...$.....V148| 

00000c50  01 00 00 00 77 00 02 02  24 00 0a 00 02 00 41 36  |....w...$.....A6| 

00000c60  01 00 59 00 59 00 00 00  00 00 00 00 04 00 56 31  |..Y.Y.........V1| 

00000c70  34 39 01 00 00 00 77 00  02 02 24 00 18 00 04 00  |49....w...$.....| 

00000c80  56 31 35 30 01 00 00 00  77 00 02 02 0a 00 10 00  |V150....w.......| 

00000c90  04 00 56 31 35 31 01 00  00 00 77 00 02 02 24 00  |..V151....w...$.| 

00000ca0  14 00 04 00 56 31 35 32  01 00 00 00 77 00 02 02  |....V152....w...| 

00000cb0  25 00 00 00 04 00 56 31  35 33 01 00 00 00 77 00  |%.....V153....w.| 

00000cc0  02 02 26 00 00 00 04 00  56 31 35 34 01 00 00 00  |..&.....V154....| 

00000cd0  77 00 02 02 25 00 00 00  04 00 56 31 35 35 01 00  |w...%.....V155..| 

00000ce0  00 00 77 00 02 02 28 00  00 00 04 00 56 31 35 36  |..w...(.....V156| 

00000cf0  01 00 00 00 77 00 02 02  04 00 00 00 04 00 56 31  |....w.........V1| 

00000d00  35 37 01 00 00 00 77 00  02 02 29 00 00 00 02 00  |57....w...).....| 

00000d10  50 38 01 00 63 00 6e 00  01 01 00 00 00 00 04 00  |P8..c.n.........| 

00000d20  56 31 35 38 01 00 00 00  77 00 02 02 2b 00 00 00  |V158....w...+...| 

00000d30  04 00 56 31 35 39 01 00  00 00 77 00 02 02 05 00  |..V159....w.....| 

00000d40  10 00 04 00 56 31 36 30  01 00 00 00 77 00 02 02  |....V160....w...| 

00000d50  2b 00 00 00 04 00 56 31  36 31 01 00 00 00 77 00  |+.....V161....w.| 

00000d60  02 02 2d 00 00 00 04 00  56 31 36 32 01 00 00 00  |..-.....V162....| 

00000d70  77 00 02 02 2e 00 00 00  04 00 56 31 36 33 01 00  |w.........V163..| 

00000d80  00 00 77 00 02 02 06 00  08 00 04 00 56 31 36 34  |..w.........V164| 

00000d90  01 00 00 00 77 00 02 02  2e 00 00 00 04 00 56 31  |....w.........V1| 

00000da0  36 35 01 00 00 00 77 00  02 02 30 00 00 00 04 00  |65....w...0.....| 

00000db0  56 31 36 36 01 00 00 00  77 00 02 02 31 00 00 00  |V166....w...1...| 

00000dc0  04 00 56 31 36 37 01 00  00 00 77 00 02 02 07 00  |..V167....w.....| 

00000dd0  00 00 04 00 56 31 36 38  01 00 00 00 77 00 02 02  |....V168....w...| 

00000de0  31 00 00 00 04 00 56 31  36 39 01 00 00 00 77 00  |1.....V169....w.| 

00000df0  02 02 33 00 00 00 04 00  56 31 37 30 01 00 00 00  |..3.....V170....| 

00000e00  77 00 02 02 05 00 00 00  04 00 56 31 37 31 01 00  |w.........V171..| 

00000e10  00 00 77 00 02 02 05 00  08 00 04 00 56 31 37 32  |..w.........V172| 

00000e20  01 00 00 00 77 00 02 02  05 00 18 00 04 00 56 31  |....w.........V1| 

00000e30  37 33 01 00 00 00 77 00  02 02 06 00 00 00 04 00  |73....w.........| 

00000e40  56 31 37 34 01 00 00 00  77 00 02 02 06 00 10 00  |V174....w.......| 

00000e50  04 00 56 31 37 35 01 00  00 00 77 00 02 02 06 00  |..V175....w.....| 

00000e60  18 00 04 00 56 31 37 36  01 00 00 00 77 00 02 02  |....V176....w...| 

00000e70  07 00 08 00 04 00 56 31  37 37 01 00 00 00 77 00  |......V177....w.| 

00000e80  02 02 07 00 10 00 04 00  56 31 37 38 01 00 00 00  |........V178....| 

00000e90  77 00 02 02 08 00 00 00  04 00 56 31 37 39 01 00  |w.........V179..| 

00000ea0  00 00 77 00 02 02 08 00  08 00 04 00 56 31 38 30  |..w.........V180| 

00000eb0  01 00 00 00 77 00 02 02  08 00 10 00 04 00 56 31  |....w.........V1| 

00000ec0  38 31 01 00 00 00 77 00  02 02 09 00 00 00 04 00  |81....w.........| 

00000ed0  56 31 38 32 01 00 00 00  77 00 02 02 09 00 08 00  |V182....w.......| 

00000ee0  04 00 56 31 38 33 01 00  00 00 77 00 02 02 09 00  |..V183....w.....| 

00000ef0  10 00 04 00 56 31 38 34  01 00 00 00 77 00 02 02  |....V184....w...| 

00000f00  0a 00 00 00 00 00 00 00  00 00 00 00 00 00 00     |...............| 

00000f0f

@1e-to 1e-to mentioned this issue Jun 2, 2021
Closed
2 tasks
@diptorupd
Copy link
Contributor

@mingjie-intel the information in the ticket are useful if we want to go one level lower than SPIR-V and look at the IGC compiled binaries.

@diptorupd diptorupd added the debug Related to #149 label Dec 19, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
debug Related to #149
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@diptorupd @1e-to @mingjie-intel