encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-08-04 03:05:42 +00:00
Code Issues Packages Projects Releases Wiki Activity
dawn-cmake/test/tint/bug/tint/943.spvasm.expected.msl
David Neto a3f2bf6c60 spirv-reader: phis as a particular case of hoisting to a var 
We already compute the "first" and "last" basic block that
uses a value, so we could know when to hoist a value into
a var declaration. You have to do this sometimes to make
sure all uses are in scope of the declaration.

Until now we tracked Phis with an entirely different mechanism.
But there are cases which broke down. That's what happens
in crbug.com/tint/1649.

Additionally, GraphicsFuzz cases generarte similar weirdness.

Also, be more careful about ensuring that the assignments
generated to feed phis behave as if they occur in parallel.
Within a single batch of such assignments, generate and
use intermediate let-declarations for phis that that batch
will overwrite.

Also, unwrap-references when rectifying the signedness of
binary operators.

Skip tests that fail due to crbug.comt/tint/98:
  test/tint/bug/tint/749.spvasm.*

Fixed: tint:1649
Change-Id: I7314c351b74a10bfa9a18011f3d80a520568011c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/101220
Auto-Submit: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
2022-09-16 20:18:39 +00:00

517 lines
20 KiB
Plaintext
Raw Blame History

warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
note: 'workgroupBarrier' must only be called from uniform control flow
note: reading from module-scope private variable 'dimInner_1' may result in a non-uniform value
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct Uniforms {
/* 0x0000 */ float tint_symbol;
/* 0x0004 */ tint_array<int8_t, 12> tint_pad;
/* 0x0010 */ packed_int3 aShape;
/* 0x001c */ tint_array<int8_t, 4> tint_pad_1;
/* 0x0020 */ packed_int3 bShape;
/* 0x002c */ tint_array<int8_t, 4> tint_pad_2;
/* 0x0030 */ packed_int3 outShape;
/* 0x003c */ tint_array<int8_t, 4> tint_pad_3;
/* 0x0040 */ int2 outShapeStrides;
};
struct ssbOut {
/* 0x0000 */ tint_array<float, 1> result;
};
struct ssbA {
/* 0x0000 */ tint_array<float, 1> A;
};
struct ssbB {
/* 0x0000 */ tint_array<float, 1> B;
};
bool coordsInBounds_vi2_vi2_(thread int2* const coord, thread int2* const shape) {
bool x_87 = false;
bool x_88 = false;
int2 const x_76 = *(coord);
bool const x_81 = all((x_76 >= int2(0)));
x_88 = x_81;
if (x_81) {
int2 const x_84 = *(coord);
int2 const x_85 = *(shape);
x_87 = all((x_84 < x_85));
x_88 = x_87;
}
return x_88;
}
float mm_readA_i1_i1_(thread int* const row, thread int* const col, const constant Uniforms* const tint_symbol_2, thread int* const tint_symbol_3, thread int* const tint_symbol_4, thread int* const tint_symbol_5, const device ssbA* const tint_symbol_6) {
int batchASize = 0;
int2 param_10 = 0;
int2 param_11 = 0;
float x_430 = 0.0f;
int const x_417 = (*(tint_symbol_2)).aShape[1];
int const x_419 = (*(tint_symbol_2)).aShape[2];
batchASize = as_type<int>((as_type<uint>(x_417) * as_type<uint>(x_419)));
int const x_421 = *(row);
int const x_422 = *(col);
int const x_424 = *(tint_symbol_3);
int const x_425 = *(tint_symbol_4);
param_10 = int2(x_421, x_422);
param_11 = int2(x_424, x_425);
bool const x_429 = coordsInBounds_vi2_vi2_(&(param_10), &(param_11));
if (x_429) {
int const x_438 = *(tint_symbol_5);
int const x_439 = batchASize;
int const x_441 = *(row);
int const x_442 = *(tint_symbol_4);
int const x_445 = *(col);
float const x_448 = (*(tint_symbol_6)).A[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_438) * as_type<uint>(x_439)))) + as_type<uint>(as_type<int>((as_type<uint>(x_441) * as_type<uint>(x_442))))))) + as_type<uint>(x_445)))];
x_430 = x_448;
} else {
x_430 = 0.0f;
}
float const x_450 = x_430;
return x_450;
}
float mm_readB_i1_i1_(thread int* const row_1, thread int* const col_1, const constant Uniforms* const tint_symbol_7, thread int* const tint_symbol_8, thread int* const tint_symbol_9, thread int* const tint_symbol_10, const device ssbB* const tint_symbol_11) {
int batchBSize = 0;
int2 param_12 = 0;
int2 param_13 = 0;
float x_468 = 0.0f;
int const x_455 = (*(tint_symbol_7)).bShape[1];
int const x_457 = (*(tint_symbol_7)).bShape[2];
batchBSize = as_type<int>((as_type<uint>(x_455) * as_type<uint>(x_457)));
int const x_459 = *(row_1);
int const x_460 = *(col_1);
int const x_462 = *(tint_symbol_8);
int const x_463 = *(tint_symbol_9);
param_12 = int2(x_459, x_460);
param_13 = int2(x_462, x_463);
bool const x_467 = coordsInBounds_vi2_vi2_(&(param_12), &(param_13));
if (x_467) {
int const x_475 = *(tint_symbol_10);
int const x_476 = batchBSize;
int const x_478 = *(row_1);
int const x_479 = *(tint_symbol_9);
int const x_482 = *(col_1);
float const x_485 = (*(tint_symbol_11)).B[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_475) * as_type<uint>(x_476)))) + as_type<uint>(as_type<int>((as_type<uint>(x_478) * as_type<uint>(x_479))))))) + as_type<uint>(x_482)))];
x_468 = x_485;
} else {
x_468 = 0.0f;
}
float const x_487 = x_468;
return x_487;
}
int getOutputFlatIndex_vi3_(thread int3* const coords, const constant Uniforms* const tint_symbol_12) {
int3 const x_99 = *(coords);
int const x_105 = (*(tint_symbol_12)).outShapeStrides[0];
int const x_107 = (*(tint_symbol_12)).outShapeStrides[1];
return int(dot(float3(x_99), float3(int3(x_105, x_107, 1))));
}
void setOutput_i1_f1_(thread int* const flatIndex, thread float* const value, device ssbOut* const tint_symbol_13) {
int const x_95 = *(flatIndex);
float const x_96 = *(value);
(*(tint_symbol_13)).result[x_95] = x_96;
return;
}
void setOutput_i1_i1_i1_f1_(thread int* const d0, thread int* const d1, thread int* const d2, thread float* const value_1, const constant Uniforms* const tint_symbol_14, device ssbOut* const tint_symbol_15) {
int flatIndex_1 = 0;
int3 param = 0;
int param_1 = 0;
float param_2 = 0.0f;
int const x_115 = *(d0);
int const x_116 = *(d1);
int const x_117 = *(d2);
param = int3(x_115, x_116, x_117);
int const x_120 = getOutputFlatIndex_vi3_(&(param), tint_symbol_14);
flatIndex_1 = x_120;
int const x_122 = flatIndex_1;
param_1 = x_122;
float const x_124 = *(value_1);
param_2 = x_124;
setOutput_i1_f1_(&(param_1), &(param_2), tint_symbol_15);
return;
}
void mm_write_i1_i1_f1_(thread int* const row_2, thread int* const col_2, thread float* const value_2, thread int* const tint_symbol_16, const constant Uniforms* const tint_symbol_17, device ssbOut* const tint_symbol_18) {
int3 outCoord = 0;
int param_14 = 0;
int param_15 = 0;
int param_16 = 0;
float param_17 = 0.0f;
int const x_491 = *(tint_symbol_16);
int const x_492 = *(row_2);
int const x_493 = *(col_2);
outCoord = int3(x_491, x_492, x_493);
int const x_496 = *(tint_symbol_16);
param_14 = x_496;
int const x_498 = *(row_2);
param_15 = x_498;
int const x_500 = *(col_2);
param_16 = x_500;
float const x_502 = *(value_2);
param_17 = x_502;
setOutput_i1_i1_i1_f1_(&(param_14), &(param_15), &(param_16), &(param_17), tint_symbol_17, tint_symbol_18);
return;
}
void mm_matMul_i1_i1_i1_(thread int* const dimAOuter, thread int* const dimInner, thread int* const dimBOuter, thread uint3* const tint_symbol_19, thread uint3* const tint_symbol_20, const constant Uniforms* const tint_symbol_21, thread int* const tint_symbol_22, thread int* const tint_symbol_23, thread int* const tint_symbol_24, const device ssbA* const tint_symbol_25, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_26, thread int* const tint_symbol_27, const device ssbB* const tint_symbol_28, threadgroup tint_array<tint_array<float, 1>, 64>* const tint_symbol_29, device ssbOut* const tint_symbol_30) {
int tileRow = 0;
int tileCol = 0;
int globalRow = 0;
int globalCol = 0;
int numTiles = 0;
int innerRow = 0;
int innerCol = 0;
tint_array<tint_array<float, 1>, 1> acc = {};
int tileColA = 0;
int tileRowB = 0;
int t = 0;
int innerRow_1 = 0;
int innerCol_1 = 0;
int inputRow = 0;
int inputCol = 0;
int param_3 = 0;
int param_4 = 0;
int innerRow_2 = 0;
int innerCol_2 = 0;
int inputRow_1 = 0;
int inputCol_1 = 0;
int param_5 = 0;
int param_6 = 0;
int k = 0;
int inner = 0;
tint_array<float, 1> BCached = {};
int innerRow_3 = 0;
float ACached = 0.0f;
int innerCol_3 = 0;
int innerRow_4 = 0;
int innerCol_4 = 0;
int param_7 = 0;
int param_8 = 0;
float param_9 = 0.0f;
uint const x_132 = (*(tint_symbol_19))[1];
tileRow = as_type<int>((as_type<uint>(as_type<int>(x_132)) * as_type<uint>(1)));
uint const x_137 = (*(tint_symbol_19))[0];
tileCol = as_type<int>((as_type<uint>(as_type<int>(x_137)) * as_type<uint>(1)));
uint const x_143 = (*(tint_symbol_20))[1];
globalRow = as_type<int>((as_type<uint>(as_type<int>(x_143)) * as_type<uint>(1)));
uint const x_148 = (*(tint_symbol_20))[0];
globalCol = as_type<int>((as_type<uint>(as_type<int>(x_148)) * as_type<uint>(1)));
int const x_152 = *(dimInner);
numTiles = as_type<int>((as_type<uint>((as_type<int>((as_type<uint>(x_152) - as_type<uint>(1))) / 64)) + as_type<uint>(1)));
innerRow = 0;
while (true) {
int const x_163 = innerRow;
if ((x_163 < 1)) {
} else {
break;
}
innerCol = 0;
while (true) {
int const x_171 = innerCol;
if ((x_171 < 1)) {
} else {
break;
}
int const x_177 = innerRow;
int const x_178 = innerCol;
acc[x_177][x_178] = 0.0f;
{
int const x_181 = innerCol;
innerCol = as_type<int>((as_type<uint>(x_181) + as_type<uint>(1)));
}
}
{
int const x_183 = innerRow;
innerRow = as_type<int>((as_type<uint>(x_183) + as_type<uint>(1)));
}
}
uint const x_187 = (*(tint_symbol_19))[0];
tileColA = as_type<int>((as_type<uint>(as_type<int>(x_187)) * as_type<uint>(64)));
uint const x_192 = (*(tint_symbol_19))[1];
tileRowB = as_type<int>((as_type<uint>(as_type<int>(x_192)) * as_type<uint>(1)));
t = 0;
while (true) {
int const x_201 = t;
int const x_202 = numTiles;
if ((x_201 < x_202)) {
} else {
break;
}
innerRow_1 = 0;
while (true) {
int const x_210 = innerRow_1;
if ((x_210 < 1)) {
} else {
break;
}
innerCol_1 = 0;
while (true) {
int const x_218 = innerCol_1;
if ((x_218 < 64)) {
} else {
break;
}
int const x_221 = tileRow;
int const x_222 = innerRow_1;
inputRow = as_type<int>((as_type<uint>(x_221) + as_type<uint>(x_222)));
int const x_225 = tileColA;
int const x_226 = innerCol_1;
inputCol = as_type<int>((as_type<uint>(x_225) + as_type<uint>(x_226)));
int const x_233 = inputRow;
int const x_234 = inputCol;
int const x_235 = globalRow;
int const x_236 = innerRow_1;
int const x_238 = t;
int const x_240 = inputCol;
param_3 = as_type<int>((as_type<uint>(x_235) + as_type<uint>(x_236)));
param_4 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_238) * as_type<uint>(64)))) + as_type<uint>(x_240)));
float const x_244 = mm_readA_i1_i1_(&(param_3), &(param_4), tint_symbol_21, tint_symbol_22, tint_symbol_23, tint_symbol_24, tint_symbol_25);
(*(tint_symbol_26))[x_233][x_234] = x_244;
{
int const x_247 = innerCol_1;
innerCol_1 = as_type<int>((as_type<uint>(x_247) + as_type<uint>(1)));
}
}
{
int const x_249 = innerRow_1;
innerRow_1 = as_type<int>((as_type<uint>(x_249) + as_type<uint>(1)));
}
}
innerRow_2 = 0;
while (true) {
int const x_257 = innerRow_2;
if ((x_257 < 1)) {
} else {
break;
}
innerCol_2 = 0;
while (true) {
int const x_265 = innerCol_2;
if ((x_265 < 1)) {
} else {
break;
}
int const x_268 = tileRowB;
int const x_269 = innerRow_2;
inputRow_1 = as_type<int>((as_type<uint>(x_268) + as_type<uint>(x_269)));
int const x_272 = tileCol;
int const x_273 = innerCol_2;
inputCol_1 = as_type<int>((as_type<uint>(x_272) + as_type<uint>(x_273)));
int const x_278 = inputRow_1;
int const x_279 = inputCol_1;
int const x_280 = t;
int const x_282 = inputRow_1;
int const x_284 = globalCol;
int const x_285 = innerCol_2;
param_5 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_280) * as_type<uint>(64)))) + as_type<uint>(x_282)));
param_6 = as_type<int>((as_type<uint>(x_284) + as_type<uint>(x_285)));
float const x_289 = mm_readB_i1_i1_(&(param_5), &(param_6), tint_symbol_21, tint_symbol_23, tint_symbol_27, tint_symbol_24, tint_symbol_28);
(*(tint_symbol_29))[x_278][x_279] = x_289;
{
int const x_291 = innerCol_2;
innerCol_2 = as_type<int>((as_type<uint>(x_291) + as_type<uint>(1)));
}
}
{
int const x_293 = innerRow_2;
innerRow_2 = as_type<int>((as_type<uint>(x_293) + as_type<uint>(1)));
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
k = 0;
while (true) {
int const x_302 = k;
if ((x_302 < 64)) {
} else {
break;
}
inner = 0;
while (true) {
int const x_310 = inner;
if ((x_310 < 1)) {
} else {
break;
}
int const x_314 = inner;
int const x_315 = k;
int const x_316 = tileCol;
int const x_317 = inner;
float const x_320 = (*(tint_symbol_29))[x_315][as_type<int>((as_type<uint>(x_316) + as_type<uint>(x_317)))];
BCached[x_314] = x_320;
{
int const x_322 = inner;
inner = as_type<int>((as_type<uint>(x_322) + as_type<uint>(1)));
}
}
innerRow_3 = 0;
while (true) {
int const x_330 = innerRow_3;
if ((x_330 < 1)) {
} else {
break;
}
int const x_333 = tileRow;
int const x_334 = innerRow_3;
int const x_336 = k;
float const x_338 = (*(tint_symbol_26))[as_type<int>((as_type<uint>(x_333) + as_type<uint>(x_334)))][x_336];
ACached = x_338;
innerCol_3 = 0;
while (true) {
int const x_345 = innerCol_3;
if ((x_345 < 1)) {
} else {
break;
}
int const x_347 = innerRow_3;
int const x_348 = innerCol_3;
float const x_349 = ACached;
int const x_350 = innerCol_3;
float const x_352 = BCached[x_350];
float const x_355 = acc[x_347][x_348];
acc[x_347][x_348] = (x_355 + (x_349 * x_352));
{
int const x_358 = innerCol_3;
innerCol_3 = as_type<int>((as_type<uint>(x_358) + as_type<uint>(1)));
}
}
{
int const x_360 = innerRow_3;
innerRow_3 = as_type<int>((as_type<uint>(x_360) + as_type<uint>(1)));
}
}
{
int const x_362 = k;
k = as_type<int>((as_type<uint>(x_362) + as_type<uint>(1)));
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
{
int const x_364 = t;
t = as_type<int>((as_type<uint>(x_364) + as_type<uint>(1)));
}
}
innerRow_4 = 0;
while (true) {
int const x_372 = innerRow_4;
if ((x_372 < 1)) {
} else {
break;
}
innerCol_4 = 0;
while (true) {
bool x_393 = false;
bool x_394 = false;
int const x_380 = innerCol_4;
if ((x_380 < 1)) {
} else {
break;
}
int const x_382 = globalCol;
int const x_383 = innerCol_4;
int const x_385 = *(dimBOuter);
bool const x_386 = (as_type<int>((as_type<uint>(x_382) + as_type<uint>(x_383))) < x_385);
x_394 = x_386;
if (x_386) {
int const x_389 = globalRow;
int const x_390 = innerRow_4;
int const x_392 = *(dimAOuter);
x_393 = (as_type<int>((as_type<uint>(x_389) + as_type<uint>(x_390))) < x_392);
x_394 = x_393;
}
if (x_394) {
int const x_397 = globalRow;
int const x_398 = innerRow_4;
int const x_400 = globalCol;
int const x_401 = innerCol_4;
int const x_403 = innerRow_4;
int const x_404 = innerCol_4;
param_7 = as_type<int>((as_type<uint>(x_397) + as_type<uint>(x_398)));
param_8 = as_type<int>((as_type<uint>(x_400) + as_type<uint>(x_401)));
float const x_409 = acc[x_403][x_404];
param_9 = x_409;
mm_write_i1_i1_f1_(&(param_7), &(param_8), &(param_9), tint_symbol_24, tint_symbol_21, tint_symbol_30);
}
{
int const x_411 = innerCol_4;
innerCol_4 = as_type<int>((as_type<uint>(x_411) + as_type<uint>(1)));
}
}
{
int const x_413 = innerRow_4;
innerRow_4 = as_type<int>((as_type<uint>(x_413) + as_type<uint>(1)));
}
}
return;
}
void main_1(const constant Uniforms* const tint_symbol_31, thread int* const tint_symbol_32, thread int* const tint_symbol_33, thread int* const tint_symbol_34, thread uint3* const tint_symbol_35, thread int* const tint_symbol_36, thread uint3* const tint_symbol_37, const device ssbA* const tint_symbol_38, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_39, const device ssbB* const tint_symbol_40, threadgroup tint_array<tint_array<float, 1>, 64>* const tint_symbol_41, device ssbOut* const tint_symbol_42) {
int param_18 = 0;
int param_19 = 0;
int param_20 = 0;
int const x_67 = (*(tint_symbol_31)).aShape[1];
*(tint_symbol_32) = x_67;
int const x_71 = (*(tint_symbol_31)).aShape[2];
*(tint_symbol_33) = x_71;
int const x_75 = (*(tint_symbol_31)).bShape[2];
*(tint_symbol_34) = x_75;
uint const x_505 = (*(tint_symbol_35))[2];
*(tint_symbol_36) = as_type<int>(x_505);
int const x_508 = *(tint_symbol_32);
param_18 = x_508;
int const x_510 = *(tint_symbol_33);
param_19 = x_510;
int const x_512 = *(tint_symbol_34);
param_20 = x_512;
mm_matMul_i1_i1_i1_(&(param_18), &(param_19), &(param_20), tint_symbol_37, tint_symbol_35, tint_symbol_31, tint_symbol_32, tint_symbol_33, tint_symbol_36, tint_symbol_38, tint_symbol_39, tint_symbol_34, tint_symbol_40, tint_symbol_41, tint_symbol_42);
return;
}
void tint_symbol_1_inner(uint3 gl_LocalInvocationID_param, uint3 gl_GlobalInvocationID_param, uint local_invocation_index, threadgroup tint_array<tint_array<float, 1>, 64>* const tint_symbol_43, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_44, thread uint3* const tint_symbol_45, thread uint3* const tint_symbol_46, const constant Uniforms* const tint_symbol_47, thread int* const tint_symbol_48, thread int* const tint_symbol_49, thread int* const tint_symbol_50, thread int* const tint_symbol_51, const device ssbA* const tint_symbol_52, const device ssbB* const tint_symbol_53, device ssbOut* const tint_symbol_54) {
{
uint const i_1 = local_invocation_index;
uint const i_2 = (local_invocation_index % 1u);
(*(tint_symbol_43))[i_1][i_2] = 0.0f;
}
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 64u)) {
uint const i = (idx / 64u);
uint const i_1 = (idx % 64u);
(*(tint_symbol_44))[i][i_1] = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_45) = gl_LocalInvocationID_param;
*(tint_symbol_46) = gl_GlobalInvocationID_param;
main_1(tint_symbol_47, tint_symbol_48, tint_symbol_49, tint_symbol_50, tint_symbol_46, tint_symbol_51, tint_symbol_45, tint_symbol_52, tint_symbol_44, tint_symbol_53, tint_symbol_43, tint_symbol_54);
}
kernel void tint_symbol_1(const constant Uniforms* tint_symbol_59 [[buffer(0)]], const device ssbA* tint_symbol_64 [[buffer(2)]], const device ssbB* tint_symbol_65 [[buffer(3)]], device ssbOut* tint_symbol_66 [[buffer(1)]], uint3 gl_LocalInvocationID_param [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array<tint_array<float, 1>, 64> tint_symbol_55;
threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_56;
thread uint3 tint_symbol_57 = 0u;
thread uint3 tint_symbol_58 = 0u;
thread int tint_symbol_60 = 0;
thread int tint_symbol_61 = 0;
thread int tint_symbol_62 = 0;
thread int tint_symbol_63 = 0;
tint_symbol_1_inner(gl_LocalInvocationID_param, gl_GlobalInvocationID_param, local_invocation_index, &(tint_symbol_55), &(tint_symbol_56), &(tint_symbol_57), &(tint_symbol_58), tint_symbol_59, &(tint_symbol_60), &(tint_symbol_61), &(tint_symbol_62), &(tint_symbol_63), tint_symbol_64, tint_symbol_65, tint_symbol_66);
return;
}
Reference in New Issue View Git Blame Copy Permalink