Skip to content

Instantly share code, notes, and snippets.

@slp

slp/llama-vulkan-msl.txt

Created February 11, 2024 08:29
Show Gist options
  • Save slp/42bf29394f1a4c4c785cace7406e7b00 to your computer and use it in GitHub Desktop.
Save slp/42bf29394f1a4c4c785cace7406e7b00 to your computer and use it in GitHub Desktop.
Download ZIP
llama.cpp Vulkan MATMUT shader translation to MSL
Raw
llama-vulkan-msl.txt
[mvk-info] Converting SPIR-V:
; SPIR-V
; Version: 1.5
; Generator: Google Shaderc over Glslang; 11
; Bound: 869
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability StorageBuffer16BitAccess
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %4 "main" %11 %18 %61 %76 %235 %250 %384 %399 %710 %783
OpExecutionMode %4 LocalSize 1 1 1
OpDecorate %11 BuiltIn GlobalInvocationId
OpMemberDecorate %16 0 Offset 0
OpMemberDecorate %16 1 Offset 4
OpMemberDecorate %16 2 Offset 8
OpMemberDecorate %16 3 Offset 12
OpMemberDecorate %16 4 Offset 16
OpMemberDecorate %16 5 Offset 20
OpMemberDecorate %16 6 Offset 24
OpMemberDecorate %16 7 Offset 28
OpMemberDecorate %16 8 Offset 32
OpMemberDecorate %16 9 Offset 36
OpMemberDecorate %16 10 Offset 40
OpMemberDecorate %16 11 Offset 44
OpMemberDecorate %16 12 Offset 48
OpMemberDecorate %16 13 Offset 52
OpDecorate %16 Block
OpDecorate %55 SpecId 1
OpDecorate %61 BuiltIn WorkgroupId
OpDecorate %76 BuiltIn LocalInvocationId
OpDecorate %79 SpecId 9
OpDecorate %83 SpecId 4
OpDecorate %96 SpecId 6
OpDecorate %98 SpecId 7
OpDecorate %108 SpecId 3
OpDecorate %157 SpecId 2
OpDecorate %175 SpecId 5
OpDecorate %178 SpecId 8
OpDecorate %247 ArrayStride 16
OpMemberDecorate %248 0 ColMajor
OpMemberDecorate %248 0 NonWritable
OpMemberDecorate %248 0 Offset 0
OpMemberDecorate %248 0 MatrixStride 8
OpDecorate %248 Block
OpDecorate %250 DescriptorSet 0
OpDecorate %250 Binding 0
OpDecorate %353 SpecId 0
OpDecorate %354 BuiltIn WorkgroupSize
OpDecorate %396 ArrayStride 32
OpMemberDecorate %397 0 ColMajor
OpMemberDecorate %397 0 NonWritable
OpMemberDecorate %397 0 Offset 0
OpMemberDecorate %397 0 MatrixStride 16
OpDecorate %397 Block
OpDecorate %399 DescriptorSet 0
OpDecorate %399 Binding 1
OpDecorate %710 BuiltIn NumWorkgroups
OpDecorate %780 ArrayStride 4
OpMemberDecorate %781 0 NonReadable
OpMemberDecorate %781 0 Offset 0
OpDecorate %781 Block
OpDecorate %783 DescriptorSet 0
OpDecorate %783 Binding 2
%2 = OpTypeVoid
%3 = OpTypeFunction %2
%6 = OpTypeInt 32 0
%9 = OpTypeVector %6 3
%10 = OpTypePointer Input %9
%11 = OpVariable %10 Input
%12 = OpConstant %6 2
%13 = OpTypePointer Input %6
%16 = OpTypeStruct %6 %6 %6 %6 %6 %6 %6 %6 %6 %6 %6 %6 %6 %6
%17 = OpTypePointer PushConstant %16
%18 = OpVariable %17 PushConstant
%19 = OpTypeInt 32 1
%20 = OpConstant %19 8
%21 = OpTypePointer PushConstant %6
%33 = OpConstant %19 10
%39 = OpConstant %19 9
%45 = OpConstant %19 7
%52 = OpConstant %19 0
%55 = OpSpecConstant %6 64
%57 = OpConstant %6 1
%61 = OpVariable %10 Input
%62 = OpConstant %6 0
%76 = OpVariable %10 Input
%79 = OpSpecConstant %6 32
%83 = OpSpecConstant %6 32
%84 = OpSpecConstantOp %6 UDiv %55 %83
%88 = OpSpecConstantOp %6 UDiv %55 %83
%96 = OpSpecConstant %6 2
%97 = OpSpecConstantOp %6 UDiv %83 %96
%98 = OpSpecConstant %6 4
%99 = OpSpecConstantOp %6 UDiv %97 %98
%103 = OpSpecConstantOp %6 UDiv %97 %98
%108 = OpSpecConstant %6 16
%109 = OpConstant %6 8
%110 = OpSpecConstantOp %6 UDiv %108 %109
%115 = OpSpecConstantOp %6 UDiv %108 %109
%119 = OpConstant %19 6
%124 = OpConstant %19 2
%135 = OpConstant %19 11
%141 = OpConstant %19 3
%152 = OpConstant %19 12
%157 = OpSpecConstant %6 64
%159 = OpConstant %19 4
%174 = OpSpecConstantOp %6 IMul %96 %98
%175 = OpSpecConstant %6 32
%176 = OpSpecConstantOp %6 IMul %83 %175
%177 = OpSpecConstantOp %6 IMul %79 %98
%178 = OpSpecConstant %6 2
%179 = OpSpecConstantOp %6 IMul %177 %178
%180 = OpSpecConstantOp %6 IMul %179 %96
%181 = OpSpecConstantOp %6 UDiv %176 %180
%182 = OpSpecConstantOp %6 IMul %174 %181
%183 = OpSpecConstantOp %6 IMul %182 %178
%184 = OpTypeBool
%186 = OpTypeFloat 32
%187 = OpSpecConstantOp %6 IMul %96 %98
%188 = OpSpecConstantOp %6 IMul %187 %181
%189 = OpSpecConstantOp %6 IMul %188 %178
%190 = OpTypeArray %186 %189
%191 = OpTypePointer Function %190
%194 = OpConstant %186 0
%195 = OpTypePointer Function %186
%198 = OpConstant %19 1
%230 = OpTypeFloat 16
%231 = OpSpecConstantOp %6 IAdd %108 %57
%232 = OpSpecConstantOp %6 IMul %55 %231
%233 = OpTypeArray %230 %232
%234 = OpTypePointer Workgroup %233
%235 = OpVariable %234 Workgroup
%239 = OpSpecConstantOp %6 IAdd %108 %57
%245 = OpTypeVector %230 4
%246 = OpTypeMatrix %245 2
%247 = OpTypeRuntimeArray %246
%248 = OpTypeStruct %247
%249 = OpTypePointer StorageBuffer %248
%250 = OpVariable %249 StorageBuffer
%252 = OpTypePointer StorageBuffer %230
%255 = OpTypePointer Workgroup %230
%260 = OpSpecConstantOp %6 IAdd %108 %57
%273 = OpSpecConstantOp %6 IAdd %108 %57
%286 = OpSpecConstantOp %6 IAdd %108 %57
%291 = OpConstant %6 3
%300 = OpSpecConstantOp %6 IAdd %108 %57
%305 = OpConstant %6 4
%314 = OpSpecConstantOp %6 IAdd %108 %57
%319 = OpConstant %6 5
%328 = OpSpecConstantOp %6 IAdd %108 %57
%333 = OpConstant %6 6
%342 = OpSpecConstantOp %6 IAdd %108 %57
%347 = OpConstant %6 7
%353 = OpSpecConstant %6 1
%354 = OpSpecConstantComposite %9 %353 %57 %57
%355 = OpSpecConstantOp %6 CompositeExtract %354 0
%356 = OpSpecConstantOp %6 IMul %355 %109
%357 = OpSpecConstantOp %6 UDiv %356 %108
%380 = OpSpecConstantOp %6 IAdd %108 %57
%381 = OpSpecConstantOp %6 IMul %157 %380
%382 = OpTypeArray %230 %381
%383 = OpTypePointer Workgroup %382
%384 = OpVariable %383 Workgroup
%388 = OpSpecConstantOp %6 IAdd %108 %57
%394 = OpTypeVector %186 4
%395 = OpTypeMatrix %394 2
%396 = OpTypeRuntimeArray %395
%397 = OpTypeStruct %396
%398 = OpTypePointer StorageBuffer %397
%399 = OpVariable %398 StorageBuffer
%401 = OpTypePointer StorageBuffer %186
%409 = OpSpecConstantOp %6 IAdd %108 %57
%423 = OpSpecConstantOp %6 IAdd %108 %57
%437 = OpSpecConstantOp %6 IAdd %108 %57
%451 = OpSpecConstantOp %6 IAdd %108 %57
%465 = OpSpecConstantOp %6 IAdd %108 %57
%479 = OpSpecConstantOp %6 IAdd %108 %57
%493 = OpSpecConstantOp %6 IAdd %108 %57
%506 = OpConstant %6 264
%507 = OpSpecConstantOp %6 UDiv %108 %109
%510 = OpSpecConstantOp %6 UDiv %108 %109
%537 = OpSpecConstantOp %6 IMul %96 %98
%538 = OpTypeArray %230 %537
%539 = OpTypePointer Function %538
%555 = OpSpecConstantOp %6 IAdd %108 %57
%561 = OpTypePointer Function %230
%583 = OpSpecConstantOp %6 IMul %181 %178
%584 = OpTypeArray %230 %583
%585 = OpTypePointer Function %584
%594 = OpSpecConstantOp %6 UDiv %175 %181
%602 = OpSpecConstantOp %6 IAdd %108 %57
%649 = OpSpecConstantOp %6 IMul %96 %98
%702 = OpConstant %19 13
%710 = OpVariable %10 Input
%780 = OpTypeRuntimeArray %186
%781 = OpTypeStruct %780
%782 = OpTypePointer StorageBuffer %781
%783 = OpVariable %782 StorageBuffer
%788 = OpConstant %19 5
%801 = OpSpecConstantOp %6 IMul %96 %98
%4 = OpFunction %2 None %3
%5 = OpLabel
%192 = OpVariable %191 Function
%540 = OpVariable %539 Function
%586 = OpVariable %585 Function
%14 = OpAccessChain %13 %11 %12
%15 = OpLoad %6 %14
%22 = OpAccessChain %21 %18 %20
%23 = OpLoad %6 %22
%24 = OpUDiv %6 %15 %23
%30 = OpUMod %6 %15 %23
%34 = OpAccessChain %21 %18 %33
%35 = OpLoad %6 %34
%36 = OpUDiv %6 %24 %35
%40 = OpAccessChain %21 %18 %39
%41 = OpLoad %6 %40
%42 = OpUDiv %6 %30 %41
%46 = OpAccessChain %21 %18 %45
%47 = OpLoad %6 %46
%48 = OpIMul %6 %36 %47
%50 = OpIAdd %6 %48 %42
%53 = OpAccessChain %21 %18 %52
%54 = OpLoad %6 %53
%56 = OpIAdd %6 %54 %55
%58 = OpISub %6 %56 %57
%59 = OpUDiv %6 %58 %55
%63 = OpAccessChain %13 %61 %62
%64 = OpLoad %6 %63
%66 = OpUMod %6 %64 %59
%71 = OpUDiv %6 %64 %59
%73 = OpAccessChain %13 %61 %57
%74 = OpLoad %6 %73
%77 = OpAccessChain %13 %76 %62
%78 = OpLoad %6 %77
%80 = OpUDiv %6 %78 %79
%85 = OpUMod %6 %80 %84
%89 = OpUDiv %6 %80 %88
%93 = OpUMod %6 %78 %79
%100 = OpUMod %6 %93 %99
%104 = OpUDiv %6 %93 %103
%111 = OpUMod %6 %78 %110
%116 = OpUDiv %6 %78 %115
%120 = OpAccessChain %21 %18 %119
%121 = OpLoad %6 %120
%122 = OpIMul %6 %71 %121
%125 = OpAccessChain %21 %18 %124
%126 = OpLoad %6 %125
%128 = OpIAdd %6 %71 %57
%131 = OpIMul %6 %128 %121
%132 = OpExtInst %6 %1 UMin %126 %131
%136 = OpAccessChain %21 %18 %135
%137 = OpLoad %6 %136
%138 = OpIMul %6 %50 %137
%140 = OpIMul %6 %66 %55
%142 = OpAccessChain %21 %18 %141
%143 = OpLoad %6 %142
%144 = OpIMul %6 %140 %143
%145 = OpIAdd %6 %138 %144
%147 = OpIAdd %6 %145 %122
%148 = OpUDiv %6 %147 %109
%153 = OpAccessChain %21 %18 %152
%154 = OpLoad %6 %153
%155 = OpIMul %6 %15 %154
%158 = OpIMul %6 %74 %157
%160 = OpAccessChain %21 %18 %159
%161 = OpLoad %6 %160
%162 = OpIMul %6 %158 %161
%163 = OpIAdd %6 %155 %162
%165 = OpIAdd %6 %163 %122
%166 = OpUDiv %6 %165 %109
OpBranch %168
%168 = OpLabel
%819 = OpPhi %6 %62 %5 %199 %169
%185 = OpULessThan %184 %819 %183
OpLoopMerge %170 %169 Unroll
OpBranchConditional %185 %169 %170
%169 = OpLabel
%196 = OpAccessChain %195 %192 %819
OpStore %196 %194
%199 = OpIAdd %6 %819 %198
OpBranch %168
%170 = OpLabel
OpBranch %202
%202 = OpLabel
%844 = OpPhi %6 %166 %170 %512 %205
%840 = OpPhi %6 %148 %170 %509 %205
%820 = OpPhi %6 %122 %170 %686 %205
%209 = OpULessThan %184 %820 %132
OpLoopMerge %204 %205 Unroll
OpBranchConditional %209 %203 %204
%203 = OpLabel
OpBranch %211
%211 = OpLabel
%836 = OpPhi %6 %62 %203 %359 %212
%217 = OpULessThan %184 %836 %55
OpLoopMerge %213 %212 Unroll
OpBranchConditional %217 %212 %213
%212 = OpLabel
%222 = OpIAdd %6 %116 %836
%225 = OpIMul %6 %222 %143
%226 = OpUDiv %6 %225 %109
%227 = OpIAdd %6 %840 %226
%229 = OpIAdd %6 %227 %111
%240 = OpIMul %6 %222 %239
%242 = OpIMul %6 %111 %109
%243 = OpIAdd %6 %240 %242
%253 = OpAccessChain %252 %250 %52 %229 %52 %62
%254 = OpLoad %230 %253
%256 = OpAccessChain %255 %235 %243
OpStore %256 %254
%261 = OpIMul %6 %222 %260
%264 = OpIAdd %6 %261 %242
%265 = OpIAdd %6 %264 %57
%267 = OpAccessChain %252 %250 %52 %229 %52 %57
%268 = OpLoad %230 %267
%269 = OpAccessChain %255 %235 %265
OpStore %269 %268
%274 = OpIMul %6 %222 %273
%277 = OpIAdd %6 %274 %242
%278 = OpIAdd %6 %277 %12
%280 = OpAccessChain %252 %250 %52 %229 %52 %12
%281 = OpLoad %230 %280
%282 = OpAccessChain %255 %235 %278
OpStore %282 %281
%287 = OpIMul %6 %222 %286
%290 = OpIAdd %6 %287 %242
%292 = OpIAdd %6 %290 %291
%294 = OpAccessChain %252 %250 %52 %229 %52 %291
%295 = OpLoad %230 %294
%296 = OpAccessChain %255 %235 %292
OpStore %296 %295
%301 = OpIMul %6 %222 %300
%304 = OpIAdd %6 %301 %242
%306 = OpIAdd %6 %304 %305
%308 = OpAccessChain %252 %250 %52 %229 %198 %62
%309 = OpLoad %230 %308
%310 = OpAccessChain %255 %235 %306
OpStore %310 %309
%315 = OpIMul %6 %222 %314
%318 = OpIAdd %6 %315 %242
%320 = OpIAdd %6 %318 %319
%322 = OpAccessChain %252 %250 %52 %229 %198 %57
%323 = OpLoad %230 %322
%324 = OpAccessChain %255 %235 %320
OpStore %324 %323
%329 = OpIMul %6 %222 %328
%332 = OpIAdd %6 %329 %242
%334 = OpIAdd %6 %332 %333
%336 = OpAccessChain %252 %250 %52 %229 %198 %12
%337 = OpLoad %230 %336
%338 = OpAccessChain %255 %235 %334
OpStore %338 %337
%343 = OpIMul %6 %222 %342
%346 = OpIAdd %6 %343 %242
%348 = OpIAdd %6 %346 %347
%350 = OpAccessChain %252 %250 %52 %229 %198 %291
%351 = OpLoad %230 %350
%352 = OpAccessChain %255 %235 %348
OpStore %352 %351
%359 = OpIAdd %6 %836 %357
OpBranch %211
%213 = OpLabel
OpBranch %361
%361 = OpLabel
%837 = OpPhi %6 %62 %213 %505 %362
%367 = OpULessThan %184 %837 %157
OpLoopMerge %363 %362 Unroll
OpBranchConditional %367 %362 %363
%362 = OpLabel
%372 = OpIAdd %6 %116 %837
%375 = OpIMul %6 %372 %161
%376 = OpUDiv %6 %375 %109
%377 = OpIAdd %6 %844 %376
%379 = OpIAdd %6 %377 %111
%389 = OpIMul %6 %372 %388
%391 = OpIMul %6 %111 %109
%392 = OpIAdd %6 %389 %391
%402 = OpAccessChain %401 %399 %52 %379 %52 %62
%403 = OpLoad %186 %402
%404 = OpFConvert %230 %403
%405 = OpAccessChain %255 %384 %392
OpStore %405 %404
%410 = OpIMul %6 %372 %409
%413 = OpIAdd %6 %410 %391
%414 = OpIAdd %6 %413 %57
%416 = OpAccessChain %401 %399 %52 %379 %52 %57
%417 = OpLoad %186 %416
%418 = OpFConvert %230 %417
%419 = OpAccessChain %255 %384 %414
OpStore %419 %418
%424 = OpIMul %6 %372 %423
%427 = OpIAdd %6 %424 %391
%428 = OpIAdd %6 %427 %12
%430 = OpAccessChain %401 %399 %52 %379 %52 %12
%431 = OpLoad %186 %430
%432 = OpFConvert %230 %431
%433 = OpAccessChain %255 %384 %428
OpStore %433 %432
%438 = OpIMul %6 %372 %437
%441 = OpIAdd %6 %438 %391
%442 = OpIAdd %6 %441 %291
%444 = OpAccessChain %401 %399 %52 %379 %52 %291
%445 = OpLoad %186 %444
%446 = OpFConvert %230 %445
%447 = OpAccessChain %255 %384 %442
OpStore %447 %446
%452 = OpIMul %6 %372 %451
%455 = OpIAdd %6 %452 %391
%456 = OpIAdd %6 %455 %305
%458 = OpAccessChain %401 %399 %52 %379 %198 %62
%459 = OpLoad %186 %458
%460 = OpFConvert %230 %459
%461 = OpAccessChain %255 %384 %456
OpStore %461 %460
%466 = OpIMul %6 %372 %465
%469 = OpIAdd %6 %466 %391
%470 = OpIAdd %6 %469 %319
%472 = OpAccessChain %401 %399 %52 %379 %198 %57
%473 = OpLoad %186 %472
%474 = OpFConvert %230 %473
%475 = OpAccessChain %255 %384 %470
OpStore %475 %474
%480 = OpIMul %6 %372 %479
%483 = OpIAdd %6 %480 %391
%484 = OpIAdd %6 %483 %333
%486 = OpAccessChain %401 %399 %52 %379 %198 %12
%487 = OpLoad %186 %486
%488 = OpFConvert %230 %487
%489 = OpAccessChain %255 %384 %484
OpStore %489 %488
%494 = OpIMul %6 %372 %493
%497 = OpIAdd %6 %494 %391
%498 = OpIAdd %6 %497 %347
%500 = OpAccessChain %401 %399 %52 %379 %198 %291
%501 = OpLoad %186 %500
%502 = OpFConvert %230 %501
%503 = OpAccessChain %255 %384 %498
OpStore %503 %502
%505 = OpIAdd %6 %837 %357
OpBranch %361
%363 = OpLabel
OpControlBarrier %12 %12 %506
%509 = OpIAdd %6 %840 %507
%512 = OpIAdd %6 %844 %510
OpBranch %514
%514 = OpLabel
%846 = OpPhi %6 %62 %363 %684 %517
%520 = OpULessThan %184 %846 %108
OpLoopMerge %516 %517 None
OpBranchConditional %520 %515 %516
%515 = OpLabel
OpBranch %522
%522 = OpLabel
%850 = OpPhi %6 %62 %515 %566 %525
%528 = OpULessThan %184 %850 %96
OpLoopMerge %524 %525 Unroll
OpBranchConditional %528 %523 %524
%523 = OpLabel
OpBranch %530
%530 = OpLabel
%868 = OpPhi %6 %62 %523 %564 %531
%536 = OpULessThan %184 %868 %98
OpLoopMerge %532 %531 Unroll
OpBranchConditional %536 %531 %532
%531 = OpLabel
%542 = OpIMul %6 %850 %98
%544 = OpIAdd %6 %542 %868
%546 = OpIMul %6 %85 %83
%548 = OpIMul %6 %850 %97
%549 = OpIAdd %6 %546 %548
%551 = OpIMul %6 %100 %98
%552 = OpIAdd %6 %549 %551
%554 = OpIAdd %6 %552 %868
%556 = OpIMul %6 %554 %555
%558 = OpIAdd %6 %556 %846
%559 = OpAccessChain %255 %235 %558
%560 = OpLoad %230 %559
%562 = OpAccessChain %561 %540 %544
OpStore %562 %560
%564 = OpIAdd %6 %868 %198
OpBranch %530
%532 = OpLabel
OpBranch %525
%525 = OpLabel
%566 = OpIAdd %6 %850 %198
OpBranch %522
%524 = OpLabel
OpBranch %568
%568 = OpLabel
%851 = OpPhi %6 %62 %524 %612 %571
%574 = OpULessThan %184 %851 %181
OpLoopMerge %570 %571 Unroll
OpBranchConditional %574 %569 %570
%569 = OpLabel
OpBranch %576
%576 = OpLabel
%865 = OpPhi %6 %62 %569 %610 %577
%582 = OpULessThan %184 %865 %178
OpLoopMerge %578 %577 Unroll
OpBranchConditional %582 %577 %578
%577 = OpLabel
%588 = OpIMul %6 %851 %178
%590 = OpIAdd %6 %588 %865
%592 = OpIMul %6 %89 %175
%595 = OpIMul %6 %851 %594
%596 = OpIAdd %6 %592 %595
%598 = OpIMul %6 %104 %178
%599 = OpIAdd %6 %596 %598
%601 = OpIAdd %6 %599 %865
%603 = OpIMul %6 %601 %602
%605 = OpIAdd %6 %603 %846
%606 = OpAccessChain %255 %384 %605
%607 = OpLoad %230 %606
%608 = OpAccessChain %561 %586 %590
OpStore %608 %607
%610 = OpIAdd %6 %865 %198
OpBranch %576
%578 = OpLabel
OpBranch %571
%571 = OpLabel
%612 = OpIAdd %6 %851 %198
OpBranch %568
%570 = OpLabel
OpBranch %614
%614 = OpLabel
%852 = OpPhi %6 %62 %570 %682 %617
%620 = OpULessThan %184 %852 %181
OpLoopMerge %616 %617 Unroll
OpBranchConditional %620 %615 %616
%615 = OpLabel
OpBranch %622
%622 = OpLabel
%856 = OpPhi %6 %62 %615 %680 %625
%628 = OpULessThan %184 %856 %96
OpLoopMerge %624 %625 Unroll
OpBranchConditional %628 %623 %624
%623 = OpLabel
OpBranch %630
%630 = OpLabel
%858 = OpPhi %6 %62 %623 %678 %633
%636 = OpULessThan %184 %858 %178
OpLoopMerge %632 %633 Unroll
OpBranchConditional %636 %631 %632
%631 = OpLabel
OpBranch %638
%638 = OpLabel
%860 = OpPhi %6 %62 %631 %676 %639
%644 = OpULessThan %184 %860 %98
OpLoopMerge %640 %639 Unroll
OpBranchConditional %644 %639 %640
%639 = OpLabel
%646 = OpIMul %6 %852 %178
%648 = OpIAdd %6 %646 %858
%650 = OpIMul %6 %648 %649
%652 = OpIMul %6 %856 %98
%653 = OpIAdd %6 %650 %652
%655 = OpIAdd %6 %653 %860
%659 = OpIAdd %6 %652 %860
%660 = OpAccessChain %561 %540 %659
%661 = OpLoad %230 %660
%662 = OpFConvert %186 %661
%667 = OpAccessChain %561 %586 %648
%668 = OpLoad %230 %667
%669 = OpFConvert %186 %668
%671 = OpAccessChain %195 %192 %655
%672 = OpLoad %186 %671
%673 = OpExtInst %186 %1 Fma %662 %669 %672
OpStore %671 %673
%676 = OpIAdd %6 %860 %198
OpBranch %638
%640 = OpLabel
OpBranch %633
%633 = OpLabel
%678 = OpIAdd %6 %858 %198
OpBranch %630
%632 = OpLabel
OpBranch %625
%625 = OpLabel
%680 = OpIAdd %6 %856 %198
OpBranch %622
%624 = OpLabel
OpBranch %617
%617 = OpLabel
%682 = OpIAdd %6 %852 %198
OpBranch %614
%616 = OpLabel
OpBranch %517
%517 = OpLabel
%684 = OpIAdd %6 %846 %198
OpBranch %514
%516 = OpLabel
OpControlBarrier %12 %12 %506
OpBranch %205
%205 = OpLabel
%686 = OpIAdd %6 %820 %108
OpBranch %202
%204 = OpLabel
%691 = OpIMul %6 %85 %83
%692 = OpIAdd %6 %140 %691
%697 = OpIMul %6 %89 %175
%698 = OpIAdd %6 %158 %697
%703 = OpAccessChain %21 %18 %702
%704 = OpLoad %6 %703
%705 = OpIMul %6 %15 %704
%709 = OpIMul %6 %71 %704
%711 = OpAccessChain %13 %710 %12
%712 = OpLoad %6 %711
%713 = OpIMul %6 %709 %712
%714 = OpIAdd %6 %705 %713
OpBranch %716
%716 = OpLabel
%821 = OpPhi %6 %62 %204 %818 %719
%722 = OpULessThan %184 %821 %181
OpLoopMerge %718 %719 Unroll
OpBranchConditional %722 %717 %718
%717 = OpLabel
OpBranch %724
%724 = OpLabel
%822 = OpPhi %6 %62 %717 %816 %727
%730 = OpULessThan %184 %822 %96
OpLoopMerge %726 %727 Unroll
OpBranchConditional %730 %725 %726
%725 = OpLabel
%734 = OpIMul %6 %822 %97
%735 = OpIAdd %6 %692 %734
%737 = OpIMul %6 %100 %98
%738 = OpIAdd %6 %735 %737
%742 = OpIMul %6 %821 %594
%743 = OpIAdd %6 %698 %742
%745 = OpIMul %6 %104 %178
%746 = OpIAdd %6 %743 %745
OpBranch %748
%748 = OpLabel
%824 = OpPhi %6 %62 %725 %814 %751
%754 = OpULessThan %184 %824 %178
OpLoopMerge %750 %751 Unroll
OpBranchConditional %754 %749 %750
%749 = OpLabel
OpBranch %756
%756 = OpLabel
%826 = OpPhi %6 %62 %749 %812 %759
%762 = OpULessThan %184 %826 %98
OpLoopMerge %758 %759 Unroll
OpBranchConditional %762 %757 %758
%757 = OpLabel
%765 = OpIAdd %6 %738 %826
%768 = OpULessThan %184 %765 %54
OpSelectionMerge %770 None
OpBranchConditional %768 %769 %770
%769 = OpLabel
%773 = OpIAdd %6 %746 %824
%774 = OpAccessChain %21 %18 %198
%775 = OpLoad %6 %774
%776 = OpULessThan %184 %773 %775
OpBranch %770
%770 = OpLabel
%777 = OpPhi %184 %768 %757 %776 %769
OpSelectionMerge %779 None
OpBranchConditional %777 %778 %779
%778 = OpLabel
%787 = OpIAdd %6 %746 %824
%789 = OpAccessChain %21 %18 %788
%790 = OpLoad %6 %789
%791 = OpIMul %6 %787 %790
%792 = OpIAdd %6 %714 %791
%794 = OpIAdd %6 %792 %738
%796 = OpIAdd %6 %794 %826
%798 = OpIMul %6 %821 %178
%800 = OpIAdd %6 %798 %824
%802 = OpIMul %6 %800 %801
%804 = OpIMul %6 %822 %98
%805 = OpIAdd %6 %802 %804
%807 = OpIAdd %6 %805 %826
%808 = OpAccessChain %195 %192 %807
%809 = OpLoad %186 %808
%810 = OpAccessChain %401 %783 %52 %796
OpStore %810 %809
OpBranch %779
%779 = OpLabel
OpBranch %759
%759 = OpLabel
%812 = OpIAdd %6 %826 %198
OpBranch %756
%758 = OpLabel
OpBranch %751
%751 = OpLabel
%814 = OpIAdd %6 %824 %198
OpBranch %748
%750 = OpLabel
OpBranch %727
%727 = OpLabel
%816 = OpIAdd %6 %822 %198
OpBranch %724
%726 = OpLabel
OpBranch %719
%719 = OpLabel
%818 = OpIAdd %6 %821 %198
OpBranch %716
%718 = OpLabel
OpReturn
OpFunctionEnd
End SPIR-V
Converted MSL:
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wmissing-braces"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
template<typename T, size_t Num>
struct spvUnsafeArray
{
T elements[Num ? Num : 1];
thread T& operator [] (size_t pos) thread
{
return elements[pos];
}
constexpr const thread T& operator [] (size_t pos) const thread
{
return elements[pos];
}
device T& operator [] (size_t pos) device
{
return elements[pos];
}
constexpr const device T& operator [] (size_t pos) const device
{
return elements[pos];
}
constexpr const constant T& operator [] (size_t pos) const constant
{
return elements[pos];
}
threadgroup T& operator [] (size_t pos) threadgroup
{
return elements[pos];
}
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
{
return elements[pos];
}
};
struct _16
{
uint _m0;
uint _m1;
uint _m2;
uint _m3;
uint _m4;
uint _m5;
uint _m6;
uint _m7;
uint _m8;
uint _m9;
uint _m10;
uint _m11;
uint _m12;
uint _m13;
};
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 64u
#endif
constant uint _55 = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_9
#define SPIRV_CROSS_CONSTANT_ID_9 32u
#endif
constant uint _79 = SPIRV_CROSS_CONSTANT_ID_9;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 32u
#endif
constant uint _83 = SPIRV_CROSS_CONSTANT_ID_4;
constant uint _84 = (_55 / _83);
constant uint _88 = (_55 / _83);
#ifndef SPIRV_CROSS_CONSTANT_ID_6
#define SPIRV_CROSS_CONSTANT_ID_6 2u
#endif
constant uint _96 = SPIRV_CROSS_CONSTANT_ID_6;
constant uint _97 = (_83 / _96);
#ifndef SPIRV_CROSS_CONSTANT_ID_7
#define SPIRV_CROSS_CONSTANT_ID_7 4u
#endif
constant uint _98 = SPIRV_CROSS_CONSTANT_ID_7;
constant uint _99 = (_97 / _98);
constant uint _103 = (_97 / _98);
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 16u
#endif
constant uint _108 = SPIRV_CROSS_CONSTANT_ID_3;
constant uint _110 = (_108 / 8u);
constant uint _115 = (_108 / 8u);
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 64u
#endif
constant uint _157 = SPIRV_CROSS_CONSTANT_ID_2;
constant uint _174 = (_96 * _98);
#ifndef SPIRV_CROSS_CONSTANT_ID_5
#define SPIRV_CROSS_CONSTANT_ID_5 32u
#endif
constant uint _175 = SPIRV_CROSS_CONSTANT_ID_5;
constant uint _176 = (_83 * _175);
constant uint _177 = (_79 * _98);
#ifndef SPIRV_CROSS_CONSTANT_ID_8
#define SPIRV_CROSS_CONSTANT_ID_8 2u
#endif
constant uint _178 = SPIRV_CROSS_CONSTANT_ID_8;
constant uint _179 = (_177 * _178);
constant uint _180 = (_179 * _96);
constant uint _181 = (_176 / _180);
constant uint _182 = (_174 * _181);
constant uint _183 = (_182 * _178);
constant uint _187 = (_96 * _98);
constant uint _188 = (_187 * _181);
constant uint _189 = (_188 * _178);
constant uint _231 = (_108 + 1u);
constant uint _232 = (_55 * _231);
constant uint _239 = (_108 + 1u);
struct _248
{
half2x4 _m0[1];
};
constant uint _260 = (_108 + 1u);
constant uint _273 = (_108 + 1u);
constant uint _286 = (_108 + 1u);
constant uint _300 = (_108 + 1u);
constant uint _314 = (_108 + 1u);
constant uint _328 = (_108 + 1u);
constant uint _342 = (_108 + 1u);
constant uint _353_tmp [[function_constant(0)]];
constant uint _353 = is_function_constant_defined(_353_tmp) ? _353_tmp : 1u;
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_353, 1u, 1u);
constant uint _355 = gl_WorkGroupSize.x;
constant uint _356 = (_355 * 8u);
constant uint _357 = (_356 / _108);
constant uint _380 = (_108 + 1u);
constant uint _381 = (_157 * _380);
constant uint _388 = (_108 + 1u);
struct _397
{
float2x4 _m0[1];
};
constant uint _409 = (_108 + 1u);
constant uint _423 = (_108 + 1u);
constant uint _437 = (_108 + 1u);
constant uint _451 = (_108 + 1u);
constant uint _465 = (_108 + 1u);
constant uint _479 = (_108 + 1u);
constant uint _493 = (_108 + 1u);
constant uint _507 = (_108 / 8u);
constant uint _510 = (_108 / 8u);
constant uint _537 = (_96 * _98);
constant uint _555 = (_108 + 1u);
constant uint _583 = (_181 * _178);
constant uint _594 = (_175 / _181);
constant uint _602 = (_108 + 1u);
constant uint _649 = (_96 * _98);
struct _781
{
float _m0[1];
};
constant uint _801 = (_96 * _98);
kernel void main0(constant _16& _18 [[buffer(0)]], const device _248& _250 [[buffer(1)]], const device _397& _399 [[buffer(2)]], device _781& _783 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
{
threadgroup spvUnsafeArray<half, _232> _235;
threadgroup spvUnsafeArray<half, _381> _384;
uint _59 = ((_18._m0 + _55) - 1u) / _55;
uint _71 = gl_WorkGroupID.x / _59;
uint _80 = gl_LocalInvocationID.x / _79;
uint _85 = _80 % _84;
uint _89 = _80 / _88;
uint _93 = gl_LocalInvocationID.x % _79;
uint _100 = _93 % _99;
uint _104 = _93 / _103;
uint _111 = gl_LocalInvocationID.x % _110;
uint _116 = gl_LocalInvocationID.x / _115;
uint _122 = _71 * _18._m6;
uint _132 = min(_18._m2, ((_71 + 1u) * _18._m6));
uint _140 = (gl_WorkGroupID.x % _59) * _55;
uint _158 = gl_WorkGroupID.y * _157;
spvUnsafeArray<float, _189> _192;
for (uint _819 = 0u; _819 < _183; )
{
_192[_819] = 0.0;
_819++;
continue;
}
uint _840;
uint _844;
_844 = (((gl_GlobalInvocationID.z * _18._m12) + (_158 * _18._m4)) + _122) / 8u;
_840 = (((((((gl_GlobalInvocationID.z / _18._m8) / _18._m10) * _18._m7) + ((gl_GlobalInvocationID.z % _18._m8) / _18._m9)) * _18._m11) + (_140 * _18._m3)) + _122) / 8u;
uint _509;
uint _512;
spvUnsafeArray<half, _537> _540;
spvUnsafeArray<half, _583> _586;
for (uint _820 = _122; _820 < _132; _844 = _512, _840 = _509, _820 += _108)
{
for (uint _836 = 0u; _836 < _55; )
{
uint _222 = _116 + _836;
uint _229 = (_840 + ((_222 * _18._m3) / 8u)) + _111;
uint _242 = _111 * 8u;
_235[(_222 * _239) + _242] = ((device half*)&_250._m0[_229][0])[0u];
_235[((_222 * _260) + _242) + 1u] = ((device half*)&_250._m0[_229][0])[1u];
_235[((_222 * _273) + _242) + 2u] = ((device half*)&_250._m0[_229][0])[2u];
_235[((_222 * _286) + _242) + 3u] = ((device half*)&_250._m0[_229][0])[3u];
_235[((_222 * _300) + _242) + 4u] = ((device half*)&_250._m0[_229][1])[0u];
_235[((_222 * _314) + _242) + 5u] = ((device half*)&_250._m0[_229][1])[1u];
_235[((_222 * _328) + _242) + 6u] = ((device half*)&_250._m0[_229][1])[2u];
_235[((_222 * _342) + _242) + 7u] = ((device half*)&_250._m0[_229][1])[3u];
_836 += _357;
continue;
}
for (uint _837 = 0u; _837 < _157; )
{
uint _372 = _116 + _837;
uint _379 = (_844 + ((_372 * _18._m4) / 8u)) + _111;
uint _391 = _111 * 8u;
_384[(_372 * _388) + _391] = half(((device float*)&_399._m0[_379][0])[0u]);
_384[((_372 * _409) + _391) + 1u] = half(((device float*)&_399._m0[_379][0])[1u]);
_384[((_372 * _423) + _391) + 2u] = half(((device float*)&_399._m0[_379][0])[2u]);
_384[((_372 * _437) + _391) + 3u] = half(((device float*)&_399._m0[_379][0])[3u]);
_384[((_372 * _451) + _391) + 4u] = half(((device float*)&_399._m0[_379][1])[0u]);
_384[((_372 * _465) + _391) + 5u] = half(((device float*)&_399._m0[_379][1])[1u]);
_384[((_372 * _479) + _391) + 6u] = half(((device float*)&_399._m0[_379][1])[2u]);
_384[((_372 * _493) + _391) + 7u] = half(((device float*)&_399._m0[_379][1])[3u]);
_837 += _357;
continue;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
_509 = _840 + _507;
_512 = _844 + _510;
uint _846;
_846 = 0u;
for (; _846 < _108; _846++)
{
uint _850;
_850 = 0u;
for (; _850 < _96; _850++)
{
for (uint _868 = 0u; _868 < _98; )
{
_540[(_850 * _98) + _868] = _235[(((((_85 * _83) + (_850 * _97)) + (_100 * _98)) + _868) * _555) + _846];
_868++;
continue;
}
}
uint _851;
_851 = 0u;
for (; _851 < _181; _851++)
{
for (uint _865 = 0u; _865 < _178; )
{
_586[(_851 * _178) + _865] = _384[(((((_89 * _175) + (_851 * _594)) + (_104 * _178)) + _865) * _602) + _846];
_865++;
continue;
}
}
uint _852;
_852 = 0u;
for (; _852 < _181; _852++)
{
uint _856;
_856 = 0u;
for (; _856 < _96; _856++)
{
uint _858;
_858 = 0u;
for (; _858 < _178; _858++)
{
for (uint _860 = 0u; _860 < _98; )
{
uint _648 = (_852 * _178) + _858;
uint _652 = _856 * _98;
uint _655 = ((_648 * _649) + _652) + _860;
_192[_655] = fma(float(_540[_652 + _860]), float(_586[_648]), _192[_655]);
_860++;
continue;
}
}
}
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
uint _692 = _140 + (_85 * _83);
uint _698 = _158 + (_89 * _175);
uint _714 = (gl_GlobalInvocationID.z * _18._m13) + ((_71 * _18._m13) * gl_NumWorkGroups.z);
for (uint _821 = 0u; _821 < _181; _821++)
{
for (uint _822 = 0u; _822 < _96; _822++)
{
uint _738 = (_692 + (_822 * _97)) + (_100 * _98);
uint _746 = (_698 + (_821 * _594)) + (_104 * _178);
for (uint _824 = 0u; _824 < _178; _824++)
{
for (uint _826 = 0u; _826 < _98; _826++)
{
bool _768 = (_738 + _826) < _18._m0;
bool _777;
if (_768)
{
_777 = (_746 + _824) < _18._m1;
}
else
{
_777 = _768;
}
if (_777)
{
_783._m0[((_714 + ((_746 + _824) * _18._m5)) + _738) + _826] = _192[((((_821 * _178) + _824) * _801) + (_822 * _98)) + _826];
}
}
}
}
}
}
End MSL
Estimated original GLSL:
#version 450
#if defined(GL_AMD_gpu_shader_half_float)
#extension GL_AMD_gpu_shader_half_float : require
#elif defined(GL_EXT_shader_explicit_arithmetic_types_float16)
#extension GL_EXT_shader_explicit_arithmetic_types_float16 : require
#else
#error No extension available for FP16.
#endif
#extension GL_EXT_shader_16bit_storage : require
#if defined(GL_EXT_control_flow_attributes)
#extension GL_EXT_control_flow_attributes : require
#define SPIRV_CROSS_FLATTEN [[flatten]]
#define SPIRV_CROSS_BRANCH [[dont_flatten]]
#define SPIRV_CROSS_UNROLL [[unroll]]
#define SPIRV_CROSS_LOOP [[dont_unroll]]
#else
#define SPIRV_CROSS_FLATTEN
#define SPIRV_CROSS_BRANCH
#define SPIRV_CROSS_UNROLL
#define SPIRV_CROSS_LOOP
#endif
layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
layout(constant_id = 1) const uint _55 = 64u;
layout(constant_id = 9) const uint _79 = 32u;
layout(constant_id = 4) const uint _83 = 32u;
const uint _84 = (_55 / _83);
const uint _88 = (_55 / _83);
layout(constant_id = 6) const uint _96 = 2u;
const uint _97 = (_83 / _96);
layout(constant_id = 7) const uint _98 = 4u;
const uint _99 = (_97 / _98);
const uint _103 = (_97 / _98);
layout(constant_id = 3) const uint _108 = 16u;
const uint _110 = (_108 / 8u);
const uint _115 = (_108 / 8u);
layout(constant_id = 2) const uint _157 = 64u;
const uint _174 = (_96 * _98);
layout(constant_id = 5) const uint _175 = 32u;
const uint _176 = (_83 * _175);
const uint _177 = (_79 * _98);
layout(constant_id = 8) const uint _178 = 2u;
const uint _179 = (_177 * _178);
const uint _180 = (_179 * _96);
const uint _181 = (_176 / _180);
const uint _182 = (_174 * _181);
const uint _183 = (_182 * _178);
const uint _187 = (_96 * _98);
const uint _188 = (_187 * _181);
const uint _189 = (_188 * _178);
const uint _231 = (_108 + 1u);
const uint _232 = (_55 * _231);
const uint _239 = (_108 + 1u);
const uint _260 = (_108 + 1u);
const uint _273 = (_108 + 1u);
const uint _286 = (_108 + 1u);
const uint _300 = (_108 + 1u);
const uint _314 = (_108 + 1u);
const uint _328 = (_108 + 1u);
const uint _342 = (_108 + 1u);
const uint _355 = gl_WorkGroupSize.x;
const uint _356 = (_355 * 8u);
const uint _357 = (_356 / _108);
const uint _380 = (_108 + 1u);
const uint _381 = (_157 * _380);
const uint _388 = (_108 + 1u);
const uint _409 = (_108 + 1u);
const uint _423 = (_108 + 1u);
const uint _437 = (_108 + 1u);
const uint _451 = (_108 + 1u);
const uint _465 = (_108 + 1u);
const uint _479 = (_108 + 1u);
const uint _493 = (_108 + 1u);
const uint _507 = (_108 / 8u);
const uint _510 = (_108 / 8u);
const uint _537 = (_96 * _98);
const uint _555 = (_108 + 1u);
const uint _583 = (_181 * _178);
const uint _594 = (_175 / _181);
const uint _602 = (_108 + 1u);
const uint _649 = (_96 * _98);
const uint _801 = (_96 * _98);
layout(set = 0, binding = 0, std430) readonly buffer _248_250
{
f16mat2x4 _m0[];
} _250;
layout(set = 0, binding = 1, std430) readonly buffer _397_399
{
mat2x4 _m0[];
} _399;
layout(set = 0, binding = 2, std430) writeonly buffer _781_783
{
float _m0[];
} _783;
layout(push_constant, std430) uniform _16_18
{
uint _m0;
uint _m1;
uint _m2;
uint _m3;
uint _m4;
uint _m5;
uint _m6;
uint _m7;
uint _m8;
uint _m9;
uint _m10;
uint _m11;
uint _m12;
uint _m13;
} _18;
shared float16_t _235[_232];
shared float16_t _384[_381];
void main()
{
uint _59 = ((_18._m0 + _55) - 1u) / _55;
uint _71 = gl_WorkGroupID.x / _59;
uint _80 = gl_LocalInvocationID.x / _79;
uint _85 = _80 % _84;
uint _89 = _80 / _88;
uint _93 = gl_LocalInvocationID.x % _79;
uint _100 = _93 % _99;
uint _104 = _93 / _103;
uint _111 = gl_LocalInvocationID.x % _110;
uint _116 = gl_LocalInvocationID.x / _115;
uint _122 = _71 * _18._m6;
uint _132 = min(_18._m2, ((_71 + 1u) * _18._m6));
uint _140 = (gl_WorkGroupID.x % _59) * _55;
uint _158 = gl_WorkGroupID.y * _157;
float _192[_189];
SPIRV_CROSS_UNROLL
for (uint _819 = 0u; _819 < _183; )
{
_192[_819] = 0.0;
_819++;
continue;
}
uint _840;
uint _844;
_844 = (((gl_GlobalInvocationID.z * _18._m12) + (_158 * _18._m4)) + _122) / 8u;
_840 = (((((((gl_GlobalInvocationID.z / _18._m8) / _18._m10) * _18._m7) + ((gl_GlobalInvocationID.z % _18._m8) / _18._m9)) * _18._m11) + (_140 * _18._m3)) + _122) / 8u;
uint _509;
uint _512;
float16_t _540[_537];
float16_t _586[_583];
SPIRV_CROSS_UNROLL
for (uint _820 = _122; _820 < _132; _844 = _512, _840 = _509, _820 += _108)
{
SPIRV_CROSS_UNROLL
for (uint _836 = 0u; _836 < _55; )
{
uint _222 = _116 + _836;
uint _229 = (_840 + ((_222 * _18._m3) / 8u)) + _111;
uint _242 = _111 * 8u;
_235[(_222 * _239) + _242] = _250._m0[_229][0].x;
_235[((_222 * _260) + _242) + 1u] = _250._m0[_229][0].y;
_235[((_222 * _273) + _242) + 2u] = _250._m0[_229][0].z;
_235[((_222 * _286) + _242) + 3u] = _250._m0[_229][0].w;
_235[((_222 * _300) + _242) + 4u] = _250._m0[_229][1].x;
_235[((_222 * _314) + _242) + 5u] = _250._m0[_229][1].y;
_235[((_222 * _328) + _242) + 6u] = _250._m0[_229][1].z;
_235[((_222 * _342) + _242) + 7u] = _250._m0[_229][1].w;
_836 += _357;
continue;
}
SPIRV_CROSS_UNROLL
for (uint _837 = 0u; _837 < _157; )
{
uint _372 = _116 + _837;
uint _379 = (_844 + ((_372 * _18._m4) / 8u)) + _111;
uint _391 = _111 * 8u;
_384[(_372 * _388) + _391] = float16_t(_399._m0[_379][0].x);
_384[((_372 * _409) + _391) + 1u] = float16_t(_399._m0[_379][0].y);
_384[((_372 * _423) + _391) + 2u] = float16_t(_399._m0[_379][0].z);
_384[((_372 * _437) + _391) + 3u] = float16_t(_399._m0[_379][0].w);
_384[((_372 * _451) + _391) + 4u] = float16_t(_399._m0[_379][1].x);
_384[((_372 * _465) + _391) + 5u] = float16_t(_399._m0[_379][1].y);
_384[((_372 * _479) + _391) + 6u] = float16_t(_399._m0[_379][1].z);
_384[((_372 * _493) + _391) + 7u] = float16_t(_399._m0[_379][1].w);
_837 += _357;
continue;
}
barrier();
_509 = _840 + _507;
_512 = _844 + _510;
uint _846;
_846 = 0u;
for (; _846 < _108; _846++)
{
uint _850;
_850 = 0u;
SPIRV_CROSS_UNROLL
for (; _850 < _96; _850++)
{
SPIRV_CROSS_UNROLL
for (uint _868 = 0u; _868 < _98; )
{
_540[(_850 * _98) + _868] = _235[(((((_85 * _83) + (_850 * _97)) + (_100 * _98)) + _868) * _555) + _846];
_868++;
continue;
}
}
uint _851;
_851 = 0u;
SPIRV_CROSS_UNROLL
for (; _851 < _181; _851++)
{
SPIRV_CROSS_UNROLL
for (uint _865 = 0u; _865 < _178; )
{
_586[(_851 * _178) + _865] = _384[(((((_89 * _175) + (_851 * _594)) + (_104 * _178)) + _865) * _602) + _846];
_865++;
continue;
}
}
uint _852;
_852 = 0u;
SPIRV_CROSS_UNROLL
for (; _852 < _181; _852++)
{
uint _856;
_856 = 0u;
SPIRV_CROSS_UNROLL
for (; _856 < _96; _856++)
{
uint _858;
_858 = 0u;
SPIRV_CROSS_UNROLL
for (; _858 < _178; _858++)
{
SPIRV_CROSS_UNROLL
for (uint _860 = 0u; _860 < _98; )
{
uint _648 = (_852 * _178) + _858;
uint _652 = _856 * _98;
uint _655 = ((_648 * _649) + _652) + _860;
_192[_655] = fma(float(_540[_652 + _860]), float(_586[_648]), _192[_655]);
_860++;
continue;
}
}
}
}
}
barrier();
}
uint _692 = _140 + (_85 * _83);
uint _698 = _158 + (_89 * _175);
uint _714 = (gl_GlobalInvocationID.z * _18._m13) + ((_71 * _18._m13) * gl_NumWorkGroups.z);
SPIRV_CROSS_UNROLL
for (uint _821 = 0u; _821 < _181; _821++)
{
SPIRV_CROSS_UNROLL
for (uint _822 = 0u; _822 < _96; _822++)
{
uint _738 = (_692 + (_822 * _97)) + (_100 * _98);
uint _746 = (_698 + (_821 * _594)) + (_104 * _178);
SPIRV_CROSS_UNROLL
for (uint _824 = 0u; _824 < _178; _824++)
{
SPIRV_CROSS_UNROLL
for (uint _826 = 0u; _826 < _98; _826++)
{
bool _768 = (_738 + _826) < _18._m0;
bool _777;
if (_768)
{
_777 = (_746 + _824) < _18._m1;
}
else
{
_777 = _768;
}
if (_777)
{
_783._m0[((_714 + ((_746 + _824) * _18._m5)) + _738) + _826] = _192[((((_821 * _178) + _824) * _801) + (_822 * _98)) + _826];
}
}
}
}
}
}
End GLSL
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment