[writer/hlsl,msl] Correctly generate inf and nan literals
Also add missing msl macros to the renamer. Bug: tint:951 Change-Id: I543e6eae885c979596ca63f9d6c7378dd5425e8a Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56941 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@chromium.org> Auto-Submit: Ben Clayton <bclayton@google.com> Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
5d8eb4a758
commit
692fc20797
|
|
@ -32,7 +32,8 @@ namespace transform {
|
||||||
namespace {
|
namespace {
|
||||||
|
|
||||||
// This list is used for a binary search and must be kept in sorted order.
|
// This list is used for a binary search and must be kept in sorted order.
|
||||||
const char* kReservedKeywordsHLSL[] = {"AddressU",
|
const char* kReservedKeywordsHLSL[] = {
|
||||||
|
"AddressU",
|
||||||
"AddressV",
|
"AddressV",
|
||||||
"AddressW",
|
"AddressW",
|
||||||
"AllMemoryBarrier",
|
"AllMemoryBarrier",
|
||||||
|
|
@ -594,10 +595,44 @@ const char* kReservedKeywordsHLSL[] = {"AddressU",
|
||||||
"virtual",
|
"virtual",
|
||||||
"void",
|
"void",
|
||||||
"volatile",
|
"volatile",
|
||||||
"while"};
|
"while",
|
||||||
|
};
|
||||||
|
|
||||||
// This list is used for a binary search and must be kept in sorted order.
|
// This list is used for a binary search and must be kept in sorted order.
|
||||||
const char* kReservedKeywordsMSL[] = {"access",
|
const char* kReservedKeywordsMSL[] = {
|
||||||
|
"HUGE_VALF",
|
||||||
|
"HUGE_VALH",
|
||||||
|
"INFINITY",
|
||||||
|
"MAXFLOAT",
|
||||||
|
"MAXHALF",
|
||||||
|
"M_1_PI_F",
|
||||||
|
"M_1_PI_H",
|
||||||
|
"M_2_PI_F",
|
||||||
|
"M_2_PI_H",
|
||||||
|
"M_2_SQRTPI_F",
|
||||||
|
"M_2_SQRTPI_H",
|
||||||
|
"M_E_F",
|
||||||
|
"M_E_H",
|
||||||
|
"M_LN10_F",
|
||||||
|
"M_LN10_H",
|
||||||
|
"M_LN2_F",
|
||||||
|
"M_LN2_H",
|
||||||
|
"M_LOG10E_F",
|
||||||
|
"M_LOG10E_H",
|
||||||
|
"M_LOG2E_F",
|
||||||
|
"M_LOG2E_H",
|
||||||
|
"M_PI_2_F",
|
||||||
|
"M_PI_2_H",
|
||||||
|
"M_PI_4_F",
|
||||||
|
"M_PI_4_H",
|
||||||
|
"M_PI_F",
|
||||||
|
"M_PI_H",
|
||||||
|
"M_SQRT1_2_F",
|
||||||
|
"M_SQRT1_2_H",
|
||||||
|
"M_SQRT2_F",
|
||||||
|
"M_SQRT2_H",
|
||||||
|
"NAN",
|
||||||
|
"access",
|
||||||
"alignas",
|
"alignas",
|
||||||
"alignof",
|
"alignof",
|
||||||
"and",
|
"and",
|
||||||
|
|
@ -687,6 +722,7 @@ const char* kReservedKeywordsMSL[] = {"access",
|
||||||
"half4x4",
|
"half4x4",
|
||||||
"if",
|
"if",
|
||||||
"imageblock",
|
"imageblock",
|
||||||
|
"infinity",
|
||||||
"inline",
|
"inline",
|
||||||
"int",
|
"int",
|
||||||
"int16_t",
|
"int16_t",
|
||||||
|
|
@ -832,7 +868,8 @@ const char* kReservedKeywordsMSL[] = {"access",
|
||||||
"wchar_t",
|
"wchar_t",
|
||||||
"while",
|
"while",
|
||||||
"xor",
|
"xor",
|
||||||
"xor_eq"};
|
"xor_eq",
|
||||||
|
};
|
||||||
|
|
||||||
} // namespace
|
} // namespace
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -847,7 +847,8 @@ INSTANTIATE_TEST_SUITE_P(RenamerTestHlsl,
|
||||||
"volatile"));
|
"volatile"));
|
||||||
// "while" // WGSL reserved keyword
|
// "while" // WGSL reserved keyword
|
||||||
|
|
||||||
INSTANTIATE_TEST_SUITE_P(RenamerTestMsl,
|
INSTANTIATE_TEST_SUITE_P(
|
||||||
|
RenamerTestMsl,
|
||||||
RenamerTestMsl,
|
RenamerTestMsl,
|
||||||
testing::Values(
|
testing::Values(
|
||||||
// c++14 spec
|
// c++14 spec
|
||||||
|
|
@ -1088,7 +1089,44 @@ INSTANTIATE_TEST_SUITE_P(RenamerTestMsl,
|
||||||
"ushort3",
|
"ushort3",
|
||||||
"ushort4",
|
"ushort4",
|
||||||
// "vec", // WGSL reserved keyword
|
// "vec", // WGSL reserved keyword
|
||||||
"vertex"));
|
"vertex",
|
||||||
|
|
||||||
|
// https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
|
||||||
|
// Table 6.5. Constants for single-precision floating-point math
|
||||||
|
// functions
|
||||||
|
"MAXFLOAT",
|
||||||
|
"HUGE_VALF",
|
||||||
|
"INFINITY",
|
||||||
|
"infinity",
|
||||||
|
"NAN",
|
||||||
|
"M_E_F",
|
||||||
|
"M_LOG2E_F",
|
||||||
|
"M_LOG10E_F",
|
||||||
|
"M_LN2_F",
|
||||||
|
"M_LN10_F",
|
||||||
|
"M_PI_F",
|
||||||
|
"M_PI_2_F",
|
||||||
|
"M_PI_4_F",
|
||||||
|
"M_1_PI_F",
|
||||||
|
"M_2_PI_F",
|
||||||
|
"M_2_SQRTPI_F",
|
||||||
|
"M_SQRT2_F",
|
||||||
|
"M_SQRT1_2_F",
|
||||||
|
"MAXHALF",
|
||||||
|
"HUGE_VALH",
|
||||||
|
"M_E_H",
|
||||||
|
"M_LOG2E_H",
|
||||||
|
"M_LOG10E_H",
|
||||||
|
"M_LN2_H",
|
||||||
|
"M_LN10_H",
|
||||||
|
"M_PI_H",
|
||||||
|
"M_PI_2_H",
|
||||||
|
"M_PI_4_H",
|
||||||
|
"M_1_PI_H",
|
||||||
|
"M_2_PI_H",
|
||||||
|
"M_2_SQRTPI_H",
|
||||||
|
"M_SQRT2_H",
|
||||||
|
"M_SQRT1_2_H"));
|
||||||
|
|
||||||
} // namespace
|
} // namespace
|
||||||
} // namespace transform
|
} // namespace transform
|
||||||
|
|
|
||||||
|
|
@ -15,6 +15,7 @@
|
||||||
#include "src/writer/hlsl/generator_impl.h"
|
#include "src/writer/hlsl/generator_impl.h"
|
||||||
|
|
||||||
#include <algorithm>
|
#include <algorithm>
|
||||||
|
#include <cmath>
|
||||||
#include <iomanip>
|
#include <iomanip>
|
||||||
#include <set>
|
#include <set>
|
||||||
#include <utility>
|
#include <utility>
|
||||||
|
|
@ -2474,7 +2475,14 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, ast::Literal* lit) {
|
||||||
if (auto* l = lit->As<ast::BoolLiteral>()) {
|
if (auto* l = lit->As<ast::BoolLiteral>()) {
|
||||||
out << (l->IsTrue() ? "true" : "false");
|
out << (l->IsTrue() ? "true" : "false");
|
||||||
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
|
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
|
||||||
|
if (std::isinf(fl->value())) {
|
||||||
|
out << (fl->value() >= 0 ? "asfloat(0x7f800000u)"
|
||||||
|
: "asfloat(0xff800000u)");
|
||||||
|
} else if (std::isnan(fl->value())) {
|
||||||
|
out << "asfloat(0x7fc00000u)";
|
||||||
|
} else {
|
||||||
out << FloatToString(fl->value()) << "f";
|
out << FloatToString(fl->value()) << "f";
|
||||||
|
}
|
||||||
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
|
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
|
||||||
out << sl->value();
|
out << sl->value();
|
||||||
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
|
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
|
||||||
|
|
|
||||||
|
|
@ -15,6 +15,7 @@
|
||||||
#include "src/writer/msl/generator_impl.h"
|
#include "src/writer/msl/generator_impl.h"
|
||||||
|
|
||||||
#include <algorithm>
|
#include <algorithm>
|
||||||
|
#include <cmath>
|
||||||
#include <iomanip>
|
#include <iomanip>
|
||||||
#include <utility>
|
#include <utility>
|
||||||
#include <vector>
|
#include <vector>
|
||||||
|
|
@ -1086,7 +1087,13 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, ast::Literal* lit) {
|
||||||
if (auto* l = lit->As<ast::BoolLiteral>()) {
|
if (auto* l = lit->As<ast::BoolLiteral>()) {
|
||||||
out << (l->IsTrue() ? "true" : "false");
|
out << (l->IsTrue() ? "true" : "false");
|
||||||
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
|
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
|
||||||
|
if (std::isinf(fl->value())) {
|
||||||
|
out << (fl->value() >= 0 ? "INFINITY" : "-INFINITY");
|
||||||
|
} else if (std::isnan(fl->value())) {
|
||||||
|
out << "NAN";
|
||||||
|
} else {
|
||||||
out << FloatToString(fl->value()) << "f";
|
out << FloatToString(fl->value()) << "f";
|
||||||
|
}
|
||||||
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
|
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
|
||||||
out << sl->value();
|
out << sl->value();
|
||||||
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
|
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
|
||||||
|
|
|
||||||
|
|
@ -1,32 +1,506 @@
|
||||||
SKIP: FAILED
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct Uniforms {
|
||||||
|
/* 0x0000 */ float tint_symbol;
|
||||||
|
/* 0x0004 */ int8_t tint_pad[12];
|
||||||
|
/* 0x0010 */ packed_int3 aShape;
|
||||||
|
/* 0x001c */ int8_t tint_pad_1[4];
|
||||||
|
/* 0x0020 */ packed_int3 bShape;
|
||||||
|
/* 0x002c */ int8_t tint_pad_2[4];
|
||||||
|
/* 0x0030 */ packed_int3 outShape;
|
||||||
|
/* 0x003c */ int8_t tint_pad_3[4];
|
||||||
|
/* 0x0040 */ packed_int2 outShapeStrides;
|
||||||
|
};
|
||||||
|
struct ssbOut {
|
||||||
|
/* 0x0000 */ float result[1];
|
||||||
|
};
|
||||||
|
struct ssbA {
|
||||||
|
/* 0x0000 */ float A[1];
|
||||||
|
};
|
||||||
|
struct ssbB {
|
||||||
|
/* 0x0000 */ float B[1];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper_1 {
|
||||||
|
float arr[64];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
tint_array_wrapper_1 arr[64];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper_3 {
|
||||||
|
float arr[1];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper_2 {
|
||||||
|
tint_array_wrapper_3 arr[64];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper_4 {
|
||||||
|
tint_array_wrapper_3 arr[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
bool coordsInBounds_vi2_vi2_(thread int2* const coord, thread int2* const shape) {
|
||||||
|
bool x_87 = false;
|
||||||
|
bool x_88_phi = false;
|
||||||
|
int2 const x_76 = *(coord);
|
||||||
|
bool const x_81 = all((x_76 >= int2(0, 0)));
|
||||||
|
x_88_phi = x_81;
|
||||||
|
if (x_81) {
|
||||||
|
int2 const x_84 = *(coord);
|
||||||
|
int2 const x_85 = *(shape);
|
||||||
|
x_87 = all((x_84 < x_85));
|
||||||
|
x_88_phi = x_87;
|
||||||
|
}
|
||||||
|
bool const x_88 = x_88_phi;
|
||||||
|
return x_88;
|
||||||
|
}
|
||||||
|
|
||||||
Validation Failure:
|
float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread int* const row, thread int* const col, thread int* const tint_symbol_5, thread int* const tint_symbol_6, thread int* const tint_symbol_7) {
|
||||||
|
int batchASize = 0;
|
||||||
|
int2 param_10 = 0;
|
||||||
|
int2 param_11 = 0;
|
||||||
|
float x_430 = 0.0f;
|
||||||
|
int const x_417 = x_48.aShape.y;
|
||||||
|
int const x_419 = x_48.aShape.z;
|
||||||
|
batchASize = (x_417 * x_419);
|
||||||
|
int const x_421 = *(row);
|
||||||
|
int const x_422 = *(col);
|
||||||
|
int const x_424 = *(tint_symbol_5);
|
||||||
|
int const x_425 = *(tint_symbol_6);
|
||||||
|
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_7);
|
||||||
|
int const x_439 = batchASize;
|
||||||
|
int const x_441 = *(row);
|
||||||
|
int const x_442 = *(tint_symbol_6);
|
||||||
|
int const x_445 = *(col);
|
||||||
|
float const x_448 = x_165.A[(((x_438 * x_439) + (x_441 * x_442)) + x_445)];
|
||||||
|
x_430 = x_448;
|
||||||
|
} else {
|
||||||
|
x_430 = 0.0f;
|
||||||
|
}
|
||||||
|
float const x_450 = x_430;
|
||||||
|
return x_450;
|
||||||
|
}
|
||||||
|
|
||||||
Compilation failed:
|
float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread int* const row_1, thread int* const col_1, thread int* const tint_symbol_8, thread int* const tint_symbol_9, thread int* const tint_symbol_10) {
|
||||||
|
int batchBSize = 0;
|
||||||
|
int2 param_12 = 0;
|
||||||
|
int2 param_13 = 0;
|
||||||
|
float x_468 = 0.0f;
|
||||||
|
int const x_455 = x_48.bShape.y;
|
||||||
|
int const x_457 = x_48.bShape.z;
|
||||||
|
batchBSize = (x_455 * 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 = x_185.B[(((x_475 * x_476) + (x_478 * x_479)) + x_482)];
|
||||||
|
x_468 = x_485;
|
||||||
|
} else {
|
||||||
|
x_468 = 0.0f;
|
||||||
|
}
|
||||||
|
float const x_487 = x_468;
|
||||||
|
return x_487;
|
||||||
|
}
|
||||||
|
|
||||||
program_source:5:22: error: expected parameter declarator
|
int getOutputFlatIndex_vi3_(constant Uniforms& x_48, thread int3* const coords) {
|
||||||
/* 0x0000 */ float NAN;
|
int3 const x_99 = *(coords);
|
||||||
^
|
int const x_105 = x_48.outShapeStrides.x;
|
||||||
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:28: note: expanded from macro 'NAN'
|
int const x_107 = x_48.outShapeStrides.y;
|
||||||
#define NAN __builtin_nanf("")
|
return int(dot(float3(x_99), float3(int3(x_105, x_107, 1))));
|
||||||
^
|
}
|
||||||
program_source:5:22: error: expected ')'
|
|
||||||
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:28: note: expanded from macro 'NAN'
|
void setOutput_i1_f1_(device ssbOut& x_54, thread int* const flatIndex, thread float* const value) {
|
||||||
#define NAN __builtin_nanf("")
|
int const x_95 = *(flatIndex);
|
||||||
^
|
float const x_96 = *(value);
|
||||||
program_source:5:22: note: to match this '('
|
x_54.result[x_95] = x_96;
|
||||||
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:27: note: expanded from macro 'NAN'
|
return;
|
||||||
#define NAN __builtin_nanf("")
|
}
|
||||||
^
|
|
||||||
program_source:494:31: warning: equality comparison with extraneous parentheses
|
void setOutput_i1_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const d0, thread int* const d1, thread int* const d2, thread float* const value_1) {
|
||||||
|
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_(x_48, &(param));
|
||||||
|
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_(x_54, &(param_1), &(param_2));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
void mm_write_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const row_2, thread int* const col_2, thread float* const value_2, thread int* const tint_symbol_11) {
|
||||||
|
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_11);
|
||||||
|
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_11);
|
||||||
|
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_(x_48, x_54, &(param_14), &(param_15), &(param_16), &(param_17));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const dimAOuter, thread int* const dimInner, thread int* const dimBOuter, thread uint3* const tint_symbol_12, thread uint3* const tint_symbol_13, thread int* const tint_symbol_14, thread int* const tint_symbol_15, thread int* const tint_symbol_16, threadgroup tint_array_wrapper* const tint_symbol_17, thread int* const tint_symbol_18, threadgroup tint_array_wrapper_2* const tint_symbol_19) {
|
||||||
|
int tileRow = 0;
|
||||||
|
int tileCol = 0;
|
||||||
|
int globalRow = 0;
|
||||||
|
int globalCol = 0;
|
||||||
|
int numTiles = 0;
|
||||||
|
int innerRow = 0;
|
||||||
|
int innerCol = 0;
|
||||||
|
tint_array_wrapper_4 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_wrapper_3 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_12)).y;
|
||||||
|
tileRow = (as_type<int>(x_132) * 1);
|
||||||
|
uint const x_137 = (*(tint_symbol_12)).x;
|
||||||
|
tileCol = (as_type<int>(x_137) * 1);
|
||||||
|
uint const x_143 = (*(tint_symbol_13)).y;
|
||||||
|
globalRow = (as_type<int>(x_143) * 1);
|
||||||
|
uint const x_148 = (*(tint_symbol_13)).x;
|
||||||
|
globalCol = (as_type<int>(x_148) * 1);
|
||||||
|
int const x_152 = *(dimInner);
|
||||||
|
numTiles = (((x_152 - 1) / 64) + 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.arr[x_177].arr[x_178] = 0.0f;
|
||||||
|
{
|
||||||
|
int const x_181 = innerCol;
|
||||||
|
innerCol = (x_181 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_183 = innerRow;
|
||||||
|
innerRow = (x_183 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
uint const x_187 = (*(tint_symbol_12)).x;
|
||||||
|
tileColA = (as_type<int>(x_187) * 64);
|
||||||
|
uint const x_192 = (*(tint_symbol_12)).y;
|
||||||
|
tileRowB = (as_type<int>(x_192) * 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 = (x_221 + x_222);
|
||||||
|
int const x_225 = tileColA;
|
||||||
|
int const x_226 = innerCol_1;
|
||||||
|
inputCol = (x_225 + 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 = (x_235 + x_236);
|
||||||
|
param_4 = ((x_238 * 64) + x_240);
|
||||||
|
float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_14, tint_symbol_15, tint_symbol_16);
|
||||||
|
(*(tint_symbol_17)).arr[x_233].arr[x_234] = x_244;
|
||||||
|
{
|
||||||
|
int const x_247 = innerCol_1;
|
||||||
|
innerCol_1 = (x_247 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_249 = innerRow_1;
|
||||||
|
innerRow_1 = (x_249 + 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 = (x_268 + x_269);
|
||||||
|
int const x_272 = tileCol;
|
||||||
|
int const x_273 = innerCol_2;
|
||||||
|
inputCol_1 = (x_272 + 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 = ((x_280 * 64) + x_282);
|
||||||
|
param_6 = (x_284 + x_285);
|
||||||
|
float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_15, tint_symbol_18, tint_symbol_16);
|
||||||
|
(*(tint_symbol_19)).arr[x_278].arr[x_279] = x_289;
|
||||||
|
{
|
||||||
|
int const x_291 = innerCol_2;
|
||||||
|
innerCol_2 = (x_291 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_293 = innerRow_2;
|
||||||
|
innerRow_2 = (x_293 + 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_19)).arr[x_315].arr[(x_316 + x_317)];
|
||||||
|
BCached.arr[x_314] = x_320;
|
||||||
|
{
|
||||||
|
int const x_322 = inner;
|
||||||
|
inner = (x_322 + 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_17)).arr[(x_333 + x_334)].arr[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.arr[x_350];
|
||||||
|
float const x_355 = acc.arr[x_347].arr[x_348];
|
||||||
|
acc.arr[x_347].arr[x_348] = (x_355 + (x_349 * x_352));
|
||||||
|
{
|
||||||
|
int const x_358 = innerCol_3;
|
||||||
|
innerCol_3 = (x_358 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_360 = innerRow_3;
|
||||||
|
innerRow_3 = (x_360 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_362 = k;
|
||||||
|
k = (x_362 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
{
|
||||||
|
int const x_364 = t;
|
||||||
|
t = (x_364 + 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_phi = 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 = ((x_382 + x_383) < x_385);
|
||||||
|
x_394_phi = x_386;
|
||||||
|
if (x_386) {
|
||||||
|
int const x_389 = globalRow;
|
||||||
|
int const x_390 = innerRow_4;
|
||||||
|
int const x_392 = *(dimAOuter);
|
||||||
|
x_393 = ((x_389 + x_390) < x_392);
|
||||||
|
x_394_phi = x_393;
|
||||||
|
}
|
||||||
|
bool const x_394 = x_394_phi;
|
||||||
|
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 = (x_397 + x_398);
|
||||||
|
param_8 = (x_400 + x_401);
|
||||||
|
float const x_409 = acc.arr[x_403].arr[x_404];
|
||||||
|
param_9 = x_409;
|
||||||
|
mm_write_i1_i1_f1_(x_48, x_54, &(param_7), &(param_8), &(param_9), tint_symbol_16);
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_411 = innerCol_4;
|
||||||
|
innerCol_4 = (x_411 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
int const x_413 = innerRow_4;
|
||||||
|
innerRow_4 = (x_413 + 1);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
void main_1(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const tint_symbol_20, thread int* const tint_symbol_21, thread int* const tint_symbol_22, thread uint3* const tint_symbol_23, thread int* const tint_symbol_24, thread uint3* const tint_symbol_25, threadgroup tint_array_wrapper* const tint_symbol_26, threadgroup tint_array_wrapper_2* const tint_symbol_27) {
|
||||||
|
int param_18 = 0;
|
||||||
|
int param_19 = 0;
|
||||||
|
int param_20 = 0;
|
||||||
|
int const x_67 = x_48.aShape.y;
|
||||||
|
*(tint_symbol_20) = x_67;
|
||||||
|
int const x_71 = x_48.aShape.z;
|
||||||
|
*(tint_symbol_21) = x_71;
|
||||||
|
int const x_75 = x_48.bShape.z;
|
||||||
|
*(tint_symbol_22) = x_75;
|
||||||
|
uint const x_505 = (*(tint_symbol_23)).z;
|
||||||
|
*(tint_symbol_24) = as_type<int>(x_505);
|
||||||
|
int const x_508 = *(tint_symbol_20);
|
||||||
|
param_18 = x_508;
|
||||||
|
int const x_510 = *(tint_symbol_21);
|
||||||
|
param_19 = x_510;
|
||||||
|
int const x_512 = *(tint_symbol_22);
|
||||||
|
param_20 = x_512;
|
||||||
|
mm_matMul_i1_i1_i1_(x_48, x_165, x_185, x_54, &(param_18), &(param_19), &(param_20), tint_symbol_25, tint_symbol_23, tint_symbol_20, tint_symbol_21, tint_symbol_24, tint_symbol_26, tint_symbol_22, tint_symbol_27);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void tint_symbol_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]], constant Uniforms& x_48 [[buffer(3)]], const device ssbA& x_165 [[buffer(1)]], const device ssbB& x_185 [[buffer(2)]], device ssbOut& x_54 [[buffer(0)]]) {
|
||||||
|
threadgroup tint_array_wrapper tint_symbol_28;
|
||||||
|
threadgroup tint_array_wrapper_2 tint_symbol_29;
|
||||||
|
thread uint3 tint_symbol_30 = 0u;
|
||||||
|
thread uint3 tint_symbol_31 = 0u;
|
||||||
|
thread int tint_symbol_32 = 0;
|
||||||
|
thread int tint_symbol_33 = 0;
|
||||||
|
thread int tint_symbol_34 = 0;
|
||||||
|
thread int tint_symbol_35 = 0;
|
||||||
if ((local_invocation_index == 0u)) {
|
if ((local_invocation_index == 0u)) {
|
||||||
~~~~~~~~~~~~~~~~~~~~~~~^~~~~
|
tint_array_wrapper const tint_symbol_3 = {.arr={}};
|
||||||
program_source:494:31: note: remove extraneous parentheses around the comparison to silence this warning
|
tint_symbol_28 = tint_symbol_3;
|
||||||
if ((local_invocation_index == 0u)) {
|
tint_array_wrapper_2 const tint_symbol_4 = {.arr={}};
|
||||||
~ ^ ~
|
tint_symbol_29 = tint_symbol_4;
|
||||||
program_source:494:31: note: use '=' to turn this equality comparison into an assignment
|
}
|
||||||
if ((local_invocation_index == 0u)) {
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
^~
|
tint_symbol_30 = gl_LocalInvocationID_param;
|
||||||
=
|
tint_symbol_31 = gl_GlobalInvocationID_param;
|
||||||
|
main_1(x_48, x_165, x_185, x_54, &(tint_symbol_32), &(tint_symbol_33), &(tint_symbol_34), &(tint_symbol_31), &(tint_symbol_35), &(tint_symbol_30), &(tint_symbol_28), &(tint_symbol_29));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -0,0 +1,158 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.0
|
||||||
|
; Generator: Khronos Glslang Reference Front End; 10
|
||||||
|
; Bound: 86
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
%1 = OpExtInstImport "GLSL.std.450"
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
|
||||||
|
OpExecutionMode %main LocalSize 128 1 1
|
||||||
|
OpSource GLSL 450
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %setOutput_i1_f1_ "setOutput(i1;f1;"
|
||||||
|
OpName %flatIndex "flatIndex"
|
||||||
|
OpName %value "value"
|
||||||
|
OpName %getAAtOutCoords_ "getAAtOutCoords("
|
||||||
|
OpName %unaryOperation_f1_ "unaryOperation(f1;"
|
||||||
|
OpName %a "a"
|
||||||
|
OpName %ssbOut "ssbOut"
|
||||||
|
OpMemberName %ssbOut 0 "result"
|
||||||
|
OpName %_ ""
|
||||||
|
OpName %ssbA "ssbA"
|
||||||
|
OpMemberName %ssbA 0 "A"
|
||||||
|
OpName %__0 ""
|
||||||
|
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
|
||||||
|
OpName %index "index"
|
||||||
|
OpName %Uniforms "Uniforms"
|
||||||
|
OpMemberName %Uniforms 0 "NAN"
|
||||||
|
OpMemberName %Uniforms 1 "aShape"
|
||||||
|
OpMemberName %Uniforms 2 "outShape"
|
||||||
|
OpMemberName %Uniforms 3 "outShapeStrides"
|
||||||
|
OpMemberName %Uniforms 4 "size"
|
||||||
|
OpName %__1 ""
|
||||||
|
OpName %a_0 "a"
|
||||||
|
OpName %param "param"
|
||||||
|
OpName %param_0 "param"
|
||||||
|
OpName %param_1 "param"
|
||||||
|
OpDecorate %_runtimearr_float ArrayStride 4
|
||||||
|
OpMemberDecorate %ssbOut 0 NonReadable
|
||||||
|
OpMemberDecorate %ssbOut 0 Offset 0
|
||||||
|
OpDecorate %ssbOut BufferBlock
|
||||||
|
OpDecorate %_ DescriptorSet 0
|
||||||
|
OpDecorate %_ Binding 0
|
||||||
|
OpDecorate %_runtimearr_float_0 ArrayStride 4
|
||||||
|
OpMemberDecorate %ssbA 0 NonWritable
|
||||||
|
OpMemberDecorate %ssbA 0 Offset 0
|
||||||
|
OpDecorate %ssbA BufferBlock
|
||||||
|
OpDecorate %__0 DescriptorSet 0
|
||||||
|
OpDecorate %__0 Binding 1
|
||||||
|
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||||
|
OpMemberDecorate %Uniforms 0 Offset 0
|
||||||
|
OpMemberDecorate %Uniforms 1 Offset 4
|
||||||
|
OpMemberDecorate %Uniforms 2 Offset 8
|
||||||
|
OpMemberDecorate %Uniforms 3 Offset 12
|
||||||
|
OpMemberDecorate %Uniforms 4 Offset 16
|
||||||
|
OpDecorate %Uniforms Block
|
||||||
|
OpDecorate %__1 DescriptorSet 0
|
||||||
|
OpDecorate %__1 Binding 2
|
||||||
|
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%3 = OpTypeFunction %void
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%_ptr_Function_int = OpTypePointer Function %int
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%_ptr_Function_float = OpTypePointer Function %float
|
||||||
|
%10 = OpTypeFunction %void %_ptr_Function_int %_ptr_Function_float
|
||||||
|
%15 = OpTypeFunction %float
|
||||||
|
%18 = OpTypeFunction %float %_ptr_Function_float
|
||||||
|
%_runtimearr_float = OpTypeRuntimeArray %float
|
||||||
|
%ssbOut = OpTypeStruct %_runtimearr_float
|
||||||
|
%_ptr_Uniform_ssbOut = OpTypePointer Uniform %ssbOut
|
||||||
|
%_ = OpVariable %_ptr_Uniform_ssbOut Uniform
|
||||||
|
%int_0 = OpConstant %int 0
|
||||||
|
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||||
|
%_runtimearr_float_0 = OpTypeRuntimeArray %float
|
||||||
|
%ssbA = OpTypeStruct %_runtimearr_float_0
|
||||||
|
%_ptr_Uniform_ssbA = OpTypePointer Uniform %ssbA
|
||||||
|
%__0 = OpVariable %_ptr_Uniform_ssbA Uniform
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%v3uint = OpTypeVector %uint 3
|
||||||
|
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||||
|
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%float_0 = OpConstant %float 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%float_0x1p_128 = OpConstant %float 0x1p+128
|
||||||
|
%Uniforms = OpTypeStruct %float %int %int %int %int
|
||||||
|
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||||
|
%__1 = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||||
|
%int_4 = OpConstant %int 4
|
||||||
|
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||||
|
%uint_128 = OpConstant %uint 128
|
||||||
|
%uint_1 = OpConstant %uint 1
|
||||||
|
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_128 %uint_1 %uint_1
|
||||||
|
%main = OpFunction %void None %3
|
||||||
|
%5 = OpLabel
|
||||||
|
%index = OpVariable %_ptr_Function_int Function
|
||||||
|
%a_0 = OpVariable %_ptr_Function_float Function
|
||||||
|
%param = OpVariable %_ptr_Function_float Function
|
||||||
|
%param_0 = OpVariable %_ptr_Function_int Function
|
||||||
|
%param_1 = OpVariable %_ptr_Function_float Function
|
||||||
|
%60 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||||
|
%61 = OpLoad %uint %60
|
||||||
|
%62 = OpBitcast %int %61
|
||||||
|
OpStore %index %62
|
||||||
|
%63 = OpLoad %int %index
|
||||||
|
%69 = OpAccessChain %_ptr_Uniform_int %__1 %int_4
|
||||||
|
%70 = OpLoad %int %69
|
||||||
|
%71 = OpSLessThan %bool %63 %70
|
||||||
|
OpSelectionMerge %73 None
|
||||||
|
OpBranchConditional %71 %72 %73
|
||||||
|
%72 = OpLabel
|
||||||
|
%75 = OpFunctionCall %float %getAAtOutCoords_
|
||||||
|
OpStore %a_0 %75
|
||||||
|
%77 = OpLoad %float %a_0
|
||||||
|
OpStore %param %77
|
||||||
|
%78 = OpFunctionCall %float %unaryOperation_f1_ %param
|
||||||
|
%80 = OpLoad %int %index
|
||||||
|
OpStore %param_0 %80
|
||||||
|
OpStore %param_1 %78
|
||||||
|
%82 = OpFunctionCall %void %setOutput_i1_f1_ %param_0 %param_1
|
||||||
|
OpBranch %73
|
||||||
|
%73 = OpLabel
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%setOutput_i1_f1_ = OpFunction %void None %10
|
||||||
|
%flatIndex = OpFunctionParameter %_ptr_Function_int
|
||||||
|
%value = OpFunctionParameter %_ptr_Function_float
|
||||||
|
%14 = OpLabel
|
||||||
|
%27 = OpLoad %int %flatIndex
|
||||||
|
%28 = OpLoad %float %value
|
||||||
|
%30 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %27
|
||||||
|
OpStore %30 %28
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%getAAtOutCoords_ = OpFunction %float None %15
|
||||||
|
%17 = OpLabel
|
||||||
|
%41 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||||
|
%42 = OpLoad %uint %41
|
||||||
|
%43 = OpAccessChain %_ptr_Uniform_float %__0 %int_0 %42
|
||||||
|
%44 = OpLoad %float %43
|
||||||
|
OpReturnValue %44
|
||||||
|
OpFunctionEnd
|
||||||
|
%unaryOperation_f1_ = OpFunction %float None %18
|
||||||
|
%a = OpFunctionParameter %_ptr_Function_float
|
||||||
|
%21 = OpLabel
|
||||||
|
%47 = OpLoad %float %a
|
||||||
|
%50 = OpFOrdLessThan %bool %47 %float_0
|
||||||
|
OpSelectionMerge %52 None
|
||||||
|
OpBranchConditional %50 %51 %52
|
||||||
|
%51 = OpLabel
|
||||||
|
OpReturnValue %float_0x1p_128
|
||||||
|
%52 = OpLabel
|
||||||
|
%55 = OpLoad %float %a
|
||||||
|
%56 = OpExtInst %float %1 Log %55
|
||||||
|
OpReturnValue %56
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,63 @@
|
||||||
|
RWByteAddressBuffer x_16 : register(u0, space0);
|
||||||
|
ByteAddressBuffer x_20 : register(t1, space0);
|
||||||
|
static uint3 gl_GlobalInvocationID = uint3(0u, 0u, 0u);
|
||||||
|
cbuffer cbuffer_x_24 : register(b2, space0) {
|
||||||
|
uint4 x_24[2];
|
||||||
|
};
|
||||||
|
|
||||||
|
float getAAtOutCoords_() {
|
||||||
|
const uint x_42 = gl_GlobalInvocationID.x;
|
||||||
|
const float x_44 = asfloat(x_20.Load((4u * x_42)));
|
||||||
|
return x_44;
|
||||||
|
}
|
||||||
|
|
||||||
|
float unaryOperation_f1_(inout float a) {
|
||||||
|
const float x_47 = a;
|
||||||
|
if ((x_47 < 0.0f)) {
|
||||||
|
return asfloat(0x7f800000u);
|
||||||
|
}
|
||||||
|
const float x_55 = a;
|
||||||
|
return log(x_55);
|
||||||
|
}
|
||||||
|
|
||||||
|
void setOutput_i1_f1_(inout int flatIndex, inout float value) {
|
||||||
|
const int x_27 = flatIndex;
|
||||||
|
const float x_28 = value;
|
||||||
|
x_16.Store((4u * uint(x_27)), asuint(x_28));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
void main_1() {
|
||||||
|
int index = 0;
|
||||||
|
float a_1 = 0.0f;
|
||||||
|
float param = 0.0f;
|
||||||
|
int param_1 = 0;
|
||||||
|
float param_2 = 0.0f;
|
||||||
|
const uint x_61 = gl_GlobalInvocationID.x;
|
||||||
|
index = asint(x_61);
|
||||||
|
const int x_63 = index;
|
||||||
|
const uint scalar_offset = (16u) / 4;
|
||||||
|
const int x_70 = asint(x_24[scalar_offset / 4][scalar_offset % 4]);
|
||||||
|
if ((x_63 < x_70)) {
|
||||||
|
const float x_75 = getAAtOutCoords_();
|
||||||
|
a_1 = x_75;
|
||||||
|
param = a_1;
|
||||||
|
const float x_78 = unaryOperation_f1_(param);
|
||||||
|
param_1 = index;
|
||||||
|
param_2 = x_78;
|
||||||
|
setOutput_i1_f1_(param_1, param_2);
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(128, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
|
||||||
|
gl_GlobalInvocationID = gl_GlobalInvocationID_param;
|
||||||
|
main_1();
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,70 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct ssbOut {
|
||||||
|
/* 0x0000 */ float result[1];
|
||||||
|
};
|
||||||
|
struct ssbA {
|
||||||
|
/* 0x0000 */ float A[1];
|
||||||
|
};
|
||||||
|
struct Uniforms {
|
||||||
|
/* 0x0000 */ float tint_symbol;
|
||||||
|
/* 0x0004 */ int aShape;
|
||||||
|
/* 0x0008 */ int outShape;
|
||||||
|
/* 0x000c */ int outShapeStrides;
|
||||||
|
/* 0x0010 */ int size;
|
||||||
|
};
|
||||||
|
|
||||||
|
float getAAtOutCoords_(const device ssbA& x_20, thread uint3* const tint_symbol_3) {
|
||||||
|
uint const x_42 = (*(tint_symbol_3)).x;
|
||||||
|
float const x_44 = x_20.A[x_42];
|
||||||
|
return x_44;
|
||||||
|
}
|
||||||
|
|
||||||
|
float unaryOperation_f1_(thread float* const a) {
|
||||||
|
float const x_47 = *(a);
|
||||||
|
if ((x_47 < 0.0f)) {
|
||||||
|
return INFINITY;
|
||||||
|
}
|
||||||
|
float const x_55 = *(a);
|
||||||
|
return log(x_55);
|
||||||
|
}
|
||||||
|
|
||||||
|
void setOutput_i1_f1_(device ssbOut& x_16, thread int* const flatIndex, thread float* const value) {
|
||||||
|
int const x_27 = *(flatIndex);
|
||||||
|
float const x_28 = *(value);
|
||||||
|
x_16.result[x_27] = x_28;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
void main_1(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_16, thread uint3* const tint_symbol_4) {
|
||||||
|
int index = 0;
|
||||||
|
float a_1 = 0.0f;
|
||||||
|
float param = 0.0f;
|
||||||
|
int param_1 = 0;
|
||||||
|
float param_2 = 0.0f;
|
||||||
|
uint const x_61 = (*(tint_symbol_4)).x;
|
||||||
|
index = as_type<int>(x_61);
|
||||||
|
int const x_63 = index;
|
||||||
|
int const x_70 = x_24.size;
|
||||||
|
if ((x_63 < x_70)) {
|
||||||
|
float const x_75 = getAAtOutCoords_(x_20, tint_symbol_4);
|
||||||
|
a_1 = x_75;
|
||||||
|
float const x_77 = a_1;
|
||||||
|
param = x_77;
|
||||||
|
float const x_78 = unaryOperation_f1_(&(param));
|
||||||
|
int const x_80 = index;
|
||||||
|
param_1 = x_80;
|
||||||
|
param_2 = x_78;
|
||||||
|
setOutput_i1_f1_(x_16, &(param_1), &(param_2));
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void tint_symbol_1(uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], constant Uniforms& x_24 [[buffer(2)]], const device ssbA& x_20 [[buffer(1)]], device ssbOut& x_16 [[buffer(0)]]) {
|
||||||
|
thread uint3 tint_symbol_5 = 0u;
|
||||||
|
tint_symbol_5 = gl_GlobalInvocationID_param;
|
||||||
|
main_1(x_24, x_20, x_16, &(tint_symbol_5));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,164 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 92
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
%46 = OpExtInstImport "GLSL.std.450"
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 128 1 1
|
||||||
|
OpName %ssbOut "ssbOut"
|
||||||
|
OpMemberName %ssbOut 0 "result"
|
||||||
|
OpName %x_16 "x_16"
|
||||||
|
OpName %ssbA "ssbA"
|
||||||
|
OpMemberName %ssbA 0 "A"
|
||||||
|
OpName %x_20 "x_20"
|
||||||
|
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
|
||||||
|
OpName %Uniforms "Uniforms"
|
||||||
|
OpMemberName %Uniforms 0 "NAN"
|
||||||
|
OpMemberName %Uniforms 1 "aShape"
|
||||||
|
OpMemberName %Uniforms 2 "outShape"
|
||||||
|
OpMemberName %Uniforms 3 "outShapeStrides"
|
||||||
|
OpMemberName %Uniforms 4 "size"
|
||||||
|
OpName %x_24 "x_24"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %getAAtOutCoords_ "getAAtOutCoords_"
|
||||||
|
OpName %unaryOperation_f1_ "unaryOperation_f1_"
|
||||||
|
OpName %a "a"
|
||||||
|
OpName %setOutput_i1_f1_ "setOutput_i1_f1_"
|
||||||
|
OpName %flatIndex "flatIndex"
|
||||||
|
OpName %value "value"
|
||||||
|
OpName %main_1 "main_1"
|
||||||
|
OpName %index "index"
|
||||||
|
OpName %a_1 "a_1"
|
||||||
|
OpName %param "param"
|
||||||
|
OpName %param_1 "param_1"
|
||||||
|
OpName %param_2 "param_2"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %ssbOut Block
|
||||||
|
OpMemberDecorate %ssbOut 0 Offset 0
|
||||||
|
OpDecorate %_runtimearr_float ArrayStride 4
|
||||||
|
OpDecorate %x_16 DescriptorSet 0
|
||||||
|
OpDecorate %x_16 Binding 0
|
||||||
|
OpDecorate %ssbA Block
|
||||||
|
OpMemberDecorate %ssbA 0 Offset 0
|
||||||
|
OpDecorate %x_20 NonWritable
|
||||||
|
OpDecorate %x_20 DescriptorSet 0
|
||||||
|
OpDecorate %x_20 Binding 1
|
||||||
|
OpDecorate %Uniforms Block
|
||||||
|
OpMemberDecorate %Uniforms 0 Offset 0
|
||||||
|
OpMemberDecorate %Uniforms 1 Offset 4
|
||||||
|
OpMemberDecorate %Uniforms 2 Offset 8
|
||||||
|
OpMemberDecorate %Uniforms 3 Offset 12
|
||||||
|
OpMemberDecorate %Uniforms 4 Offset 16
|
||||||
|
OpDecorate %x_24 NonWritable
|
||||||
|
OpDecorate %x_24 DescriptorSet 0
|
||||||
|
OpDecorate %x_24 Binding 2
|
||||||
|
OpDecorate %tint_symbol BuiltIn GlobalInvocationId
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%_runtimearr_float = OpTypeRuntimeArray %float
|
||||||
|
%ssbOut = OpTypeStruct %_runtimearr_float
|
||||||
|
%_ptr_StorageBuffer_ssbOut = OpTypePointer StorageBuffer %ssbOut
|
||||||
|
%x_16 = OpVariable %_ptr_StorageBuffer_ssbOut StorageBuffer
|
||||||
|
%ssbA = OpTypeStruct %_runtimearr_float
|
||||||
|
%_ptr_StorageBuffer_ssbA = OpTypePointer StorageBuffer %ssbA
|
||||||
|
%x_20 = OpVariable %_ptr_StorageBuffer_ssbA StorageBuffer
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%v3uint = OpTypeVector %uint 3
|
||||||
|
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||||
|
%13 = OpConstantNull %v3uint
|
||||||
|
%gl_GlobalInvocationID = OpVariable %_ptr_Private_v3uint Private %13
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%Uniforms = OpTypeStruct %float %int %int %int %int
|
||||||
|
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||||
|
%x_24 = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||||
|
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||||
|
%20 = OpTypeFunction %float
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%_ptr_Private_uint = OpTypePointer Private %uint
|
||||||
|
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
|
||||||
|
%_ptr_Function_float = OpTypePointer Function %float
|
||||||
|
%30 = OpTypeFunction %float %_ptr_Function_float
|
||||||
|
%float_0 = OpConstant %float 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%float_0x1p_128 = OpConstant %float 0x1p+128
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%_ptr_Function_int = OpTypePointer Function %int
|
||||||
|
%47 = OpTypeFunction %void %_ptr_Function_int %_ptr_Function_float
|
||||||
|
%59 = OpTypeFunction %void
|
||||||
|
%63 = OpConstantNull %int
|
||||||
|
%65 = OpConstantNull %float
|
||||||
|
%uint_4 = OpConstant %uint 4
|
||||||
|
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||||
|
%getAAtOutCoords_ = OpFunction %float None %20
|
||||||
|
%22 = OpLabel
|
||||||
|
%25 = OpAccessChain %_ptr_Private_uint %gl_GlobalInvocationID %uint_0
|
||||||
|
%26 = OpLoad %uint %25
|
||||||
|
%28 = OpAccessChain %_ptr_StorageBuffer_float %x_20 %uint_0 %26
|
||||||
|
%29 = OpLoad %float %28
|
||||||
|
OpReturnValue %29
|
||||||
|
OpFunctionEnd
|
||||||
|
%unaryOperation_f1_ = OpFunction %float None %30
|
||||||
|
%a = OpFunctionParameter %_ptr_Function_float
|
||||||
|
%34 = OpLabel
|
||||||
|
%36 = OpLoad %float %a
|
||||||
|
%38 = OpFOrdLessThan %bool %36 %float_0
|
||||||
|
OpSelectionMerge %40 None
|
||||||
|
OpBranchConditional %38 %41 %40
|
||||||
|
%41 = OpLabel
|
||||||
|
OpReturnValue %float_0x1p_128
|
||||||
|
%40 = OpLabel
|
||||||
|
%44 = OpLoad %float %a
|
||||||
|
%45 = OpExtInst %float %46 Log %44
|
||||||
|
OpReturnValue %45
|
||||||
|
OpFunctionEnd
|
||||||
|
%setOutput_i1_f1_ = OpFunction %void None %47
|
||||||
|
%flatIndex = OpFunctionParameter %_ptr_Function_int
|
||||||
|
%value = OpFunctionParameter %_ptr_Function_float
|
||||||
|
%53 = OpLabel
|
||||||
|
%55 = OpLoad %int %flatIndex
|
||||||
|
%57 = OpLoad %float %value
|
||||||
|
%58 = OpAccessChain %_ptr_StorageBuffer_float %x_16 %uint_0 %55
|
||||||
|
OpStore %58 %57
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%main_1 = OpFunction %void None %59
|
||||||
|
%61 = OpLabel
|
||||||
|
%index = OpVariable %_ptr_Function_int Function %63
|
||||||
|
%a_1 = OpVariable %_ptr_Function_float Function %65
|
||||||
|
%param = OpVariable %_ptr_Function_float Function %65
|
||||||
|
%param_1 = OpVariable %_ptr_Function_int Function %63
|
||||||
|
%param_2 = OpVariable %_ptr_Function_float Function %65
|
||||||
|
%69 = OpAccessChain %_ptr_Private_uint %gl_GlobalInvocationID %uint_0
|
||||||
|
%70 = OpLoad %uint %69
|
||||||
|
%71 = OpBitcast %int %70
|
||||||
|
OpStore %index %71
|
||||||
|
%72 = OpLoad %int %index
|
||||||
|
%75 = OpAccessChain %_ptr_Uniform_int %x_24 %uint_4
|
||||||
|
%76 = OpLoad %int %75
|
||||||
|
%77 = OpSLessThan %bool %72 %76
|
||||||
|
OpSelectionMerge %78 None
|
||||||
|
OpBranchConditional %77 %79 %78
|
||||||
|
%79 = OpLabel
|
||||||
|
%80 = OpFunctionCall %float %getAAtOutCoords_
|
||||||
|
OpStore %a_1 %80
|
||||||
|
%81 = OpLoad %float %a_1
|
||||||
|
OpStore %param %81
|
||||||
|
%82 = OpFunctionCall %float %unaryOperation_f1_ %param
|
||||||
|
%84 = OpLoad %int %index
|
||||||
|
OpStore %param_1 %84
|
||||||
|
OpStore %param_2 %82
|
||||||
|
%85 = OpFunctionCall %void %setOutput_i1_f1_ %param_1 %param_2
|
||||||
|
OpBranch %78
|
||||||
|
%78 = OpLabel
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%main = OpFunction %void None %59
|
||||||
|
%89 = OpLabel
|
||||||
|
%90 = OpLoad %v3uint %tint_symbol
|
||||||
|
OpStore %gl_GlobalInvocationID %90
|
||||||
|
%91 = OpFunctionCall %void %main_1
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,82 @@
|
||||||
|
type RTArr = [[stride(4)]] array<f32>;
|
||||||
|
|
||||||
|
type RTArr_1 = [[stride(4)]] array<f32>;
|
||||||
|
|
||||||
|
[[block]]
|
||||||
|
struct ssbOut {
|
||||||
|
result : RTArr_1;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[block]]
|
||||||
|
struct ssbA {
|
||||||
|
A : RTArr_1;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[block]]
|
||||||
|
struct Uniforms {
|
||||||
|
NAN : f32;
|
||||||
|
aShape : i32;
|
||||||
|
outShape : i32;
|
||||||
|
outShapeStrides : i32;
|
||||||
|
size : i32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[group(0), binding(0)]] var<storage, read_write> x_16 : ssbOut;
|
||||||
|
|
||||||
|
[[group(0), binding(1)]] var<storage, read> x_20 : ssbA;
|
||||||
|
|
||||||
|
var<private> gl_GlobalInvocationID : vec3<u32>;
|
||||||
|
|
||||||
|
[[group(0), binding(2)]] var<uniform> x_24 : Uniforms;
|
||||||
|
|
||||||
|
fn getAAtOutCoords_() -> f32 {
|
||||||
|
let x_42 : u32 = gl_GlobalInvocationID.x;
|
||||||
|
let x_44 : f32 = x_20.A[x_42];
|
||||||
|
return x_44;
|
||||||
|
}
|
||||||
|
|
||||||
|
fn unaryOperation_f1_(a : ptr<function, f32>) -> f32 {
|
||||||
|
let x_47 : f32 = *(a);
|
||||||
|
if ((x_47 < 0.0)) {
|
||||||
|
return 0x1p+128;
|
||||||
|
}
|
||||||
|
let x_55 : f32 = *(a);
|
||||||
|
return log(x_55);
|
||||||
|
}
|
||||||
|
|
||||||
|
fn setOutput_i1_f1_(flatIndex : ptr<function, i32>, value : ptr<function, f32>) {
|
||||||
|
let x_27 : i32 = *(flatIndex);
|
||||||
|
let x_28 : f32 = *(value);
|
||||||
|
x_16.result[x_27] = x_28;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
fn main_1() {
|
||||||
|
var index : i32;
|
||||||
|
var a_1 : f32;
|
||||||
|
var param : f32;
|
||||||
|
var param_1 : i32;
|
||||||
|
var param_2 : f32;
|
||||||
|
let x_61 : u32 = gl_GlobalInvocationID.x;
|
||||||
|
index = bitcast<i32>(x_61);
|
||||||
|
let x_63 : i32 = index;
|
||||||
|
let x_70 : i32 = x_24.size;
|
||||||
|
if ((x_63 < x_70)) {
|
||||||
|
let x_75 : f32 = getAAtOutCoords_();
|
||||||
|
a_1 = x_75;
|
||||||
|
let x_77 : f32 = a_1;
|
||||||
|
param = x_77;
|
||||||
|
let x_78 : f32 = unaryOperation_f1_(&(param));
|
||||||
|
let x_80 : i32 = index;
|
||||||
|
param_1 = x_80;
|
||||||
|
param_2 = x_78;
|
||||||
|
setOutput_i1_f1_(&(param_1), &(param_2));
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(128, 1, 1)]]
|
||||||
|
fn main([[builtin(global_invocation_id)]] gl_GlobalInvocationID_param : vec3<u32>) {
|
||||||
|
gl_GlobalInvocationID = gl_GlobalInvocationID_param;
|
||||||
|
main_1();
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,21 @@
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpSource HLSL 600
|
||||||
|
OpName %out_var_SV_TARGET "out.var.SV_TARGET"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %out_var_SV_TARGET Location 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%float_0x1p_128 = OpConstant %float -0x1p+128
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%10 = OpLabel
|
||||||
|
%12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
|
||||||
|
OpStore %out_var_SV_TARGET %12
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,20 @@
|
||||||
|
static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
|
||||||
|
void main_1() {
|
||||||
|
out_var_SV_TARGET = float4(asfloat(0xff800000u), asfloat(0xff800000u), asfloat(0xff800000u), asfloat(0xff800000u));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol {
|
||||||
|
float4 out_var_SV_TARGET_1 : SV_Target0;
|
||||||
|
};
|
||||||
|
|
||||||
|
tint_symbol main() {
|
||||||
|
main_1();
|
||||||
|
const main_out tint_symbol_1 = {out_var_SV_TARGET};
|
||||||
|
const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_2;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,23 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
float4 out_var_SV_TARGET_1 [[color(0)]];
|
||||||
|
};
|
||||||
|
|
||||||
|
void main_1(thread float4* const tint_symbol_4) {
|
||||||
|
*(tint_symbol_4) = float4(-INFINITY, -INFINITY, -INFINITY, -INFINITY);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
fragment tint_symbol_1 tint_symbol() {
|
||||||
|
thread float4 tint_symbol_5 = 0.0f;
|
||||||
|
main_1(&(tint_symbol_5));
|
||||||
|
main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
|
||||||
|
tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_3;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,52 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 26
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %tint_symbol_1
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpName %out_var_SV_TARGET "out_var_SV_TARGET"
|
||||||
|
OpName %tint_symbol_1 "tint_symbol_1"
|
||||||
|
OpName %main_1 "main_1"
|
||||||
|
OpName %main_out "main_out"
|
||||||
|
OpMemberName %main_out 0 "out_var_SV_TARGET_1"
|
||||||
|
OpName %tint_symbol_2 "tint_symbol_2"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol_1 Location 0
|
||||||
|
OpMemberDecorate %main_out 0 Offset 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Private_v4float = OpTypePointer Private %v4float
|
||||||
|
%5 = OpConstantNull %v4float
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%8 = OpTypeFunction %void
|
||||||
|
%float_n0x1p_128 = OpConstant %float -0x1p+128
|
||||||
|
%13 = OpConstantComposite %v4float %float_n0x1p_128 %float_n0x1p_128 %float_n0x1p_128 %float_n0x1p_128
|
||||||
|
%main_out = OpTypeStruct %v4float
|
||||||
|
%14 = OpTypeFunction %void %main_out
|
||||||
|
%main_1 = OpFunction %void None %8
|
||||||
|
%11 = OpLabel
|
||||||
|
OpStore %out_var_SV_TARGET %13
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%tint_symbol_2 = OpFunction %void None %14
|
||||||
|
%tint_symbol = OpFunctionParameter %main_out
|
||||||
|
%18 = OpLabel
|
||||||
|
%19 = OpCompositeExtract %v4float %tint_symbol 0
|
||||||
|
OpStore %tint_symbol_1 %19
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%main = OpFunction %void None %8
|
||||||
|
%21 = OpLabel
|
||||||
|
%22 = OpFunctionCall %void %main_1
|
||||||
|
%24 = OpLoad %v4float %out_var_SV_TARGET
|
||||||
|
%25 = OpCompositeConstruct %main_out %24
|
||||||
|
%23 = OpFunctionCall %void %tint_symbol_2 %25
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,17 @@
|
||||||
|
var<private> out_var_SV_TARGET : vec4<f32>;
|
||||||
|
|
||||||
|
fn main_1() {
|
||||||
|
out_var_SV_TARGET = vec4<f32>(-0x1p+128, -0x1p+128, -0x1p+128, -0x1p+128);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
[[location(0)]]
|
||||||
|
out_var_SV_TARGET_1 : vec4<f32>;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[stage(fragment)]]
|
||||||
|
fn main() -> main_out {
|
||||||
|
main_1();
|
||||||
|
return main_out(out_var_SV_TARGET);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,21 @@
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpSource HLSL 600
|
||||||
|
OpName %out_var_SV_TARGET "out.var.SV_TARGET"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %out_var_SV_TARGET Location 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%float_0x1p_128 = OpConstant %float 0x1p+128
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%10 = OpLabel
|
||||||
|
%12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
|
||||||
|
OpStore %out_var_SV_TARGET %12
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,20 @@
|
||||||
|
static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
|
||||||
|
void main_1() {
|
||||||
|
out_var_SV_TARGET = float4(asfloat(0x7f800000u), asfloat(0x7f800000u), asfloat(0x7f800000u), asfloat(0x7f800000u));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol {
|
||||||
|
float4 out_var_SV_TARGET_1 : SV_Target0;
|
||||||
|
};
|
||||||
|
|
||||||
|
tint_symbol main() {
|
||||||
|
main_1();
|
||||||
|
const main_out tint_symbol_1 = {out_var_SV_TARGET};
|
||||||
|
const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_2;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,23 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
float4 out_var_SV_TARGET_1 [[color(0)]];
|
||||||
|
};
|
||||||
|
|
||||||
|
void main_1(thread float4* const tint_symbol_4) {
|
||||||
|
*(tint_symbol_4) = float4(INFINITY, INFINITY, INFINITY, INFINITY);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
fragment tint_symbol_1 tint_symbol() {
|
||||||
|
thread float4 tint_symbol_5 = 0.0f;
|
||||||
|
main_1(&(tint_symbol_5));
|
||||||
|
main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
|
||||||
|
tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_3;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,52 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 26
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %tint_symbol_1
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpName %out_var_SV_TARGET "out_var_SV_TARGET"
|
||||||
|
OpName %tint_symbol_1 "tint_symbol_1"
|
||||||
|
OpName %main_1 "main_1"
|
||||||
|
OpName %main_out "main_out"
|
||||||
|
OpMemberName %main_out 0 "out_var_SV_TARGET_1"
|
||||||
|
OpName %tint_symbol_2 "tint_symbol_2"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol_1 Location 0
|
||||||
|
OpMemberDecorate %main_out 0 Offset 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Private_v4float = OpTypePointer Private %v4float
|
||||||
|
%5 = OpConstantNull %v4float
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%8 = OpTypeFunction %void
|
||||||
|
%float_0x1p_128 = OpConstant %float 0x1p+128
|
||||||
|
%13 = OpConstantComposite %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
|
||||||
|
%main_out = OpTypeStruct %v4float
|
||||||
|
%14 = OpTypeFunction %void %main_out
|
||||||
|
%main_1 = OpFunction %void None %8
|
||||||
|
%11 = OpLabel
|
||||||
|
OpStore %out_var_SV_TARGET %13
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%tint_symbol_2 = OpFunction %void None %14
|
||||||
|
%tint_symbol = OpFunctionParameter %main_out
|
||||||
|
%18 = OpLabel
|
||||||
|
%19 = OpCompositeExtract %v4float %tint_symbol 0
|
||||||
|
OpStore %tint_symbol_1 %19
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%main = OpFunction %void None %8
|
||||||
|
%21 = OpLabel
|
||||||
|
%22 = OpFunctionCall %void %main_1
|
||||||
|
%24 = OpLoad %v4float %out_var_SV_TARGET
|
||||||
|
%25 = OpCompositeConstruct %main_out %24
|
||||||
|
%23 = OpFunctionCall %void %tint_symbol_2 %25
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,17 @@
|
||||||
|
var<private> out_var_SV_TARGET : vec4<f32>;
|
||||||
|
|
||||||
|
fn main_1() {
|
||||||
|
out_var_SV_TARGET = vec4<f32>(0x1p+128, 0x1p+128, 0x1p+128, 0x1p+128);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
[[location(0)]]
|
||||||
|
out_var_SV_TARGET_1 : vec4<f32>;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[stage(fragment)]]
|
||||||
|
fn main() -> main_out {
|
||||||
|
main_1();
|
||||||
|
return main_out(out_var_SV_TARGET);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,21 @@
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpSource HLSL 600
|
||||||
|
OpName %out_var_SV_TARGET "out.var.SV_TARGET"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %out_var_SV_TARGET Location 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%float_0x1p_128 = OpConstant %float 0x1.1p+128
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%10 = OpLabel
|
||||||
|
%12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
|
||||||
|
OpStore %out_var_SV_TARGET %12
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,20 @@
|
||||||
|
static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
|
||||||
|
void main_1() {
|
||||||
|
out_var_SV_TARGET = float4(asfloat(0x7fc00000u), asfloat(0x7fc00000u), asfloat(0x7fc00000u), asfloat(0x7fc00000u));
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol {
|
||||||
|
float4 out_var_SV_TARGET_1 : SV_Target0;
|
||||||
|
};
|
||||||
|
|
||||||
|
tint_symbol main() {
|
||||||
|
main_1();
|
||||||
|
const main_out tint_symbol_1 = {out_var_SV_TARGET};
|
||||||
|
const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_2;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,23 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct main_out {
|
||||||
|
float4 out_var_SV_TARGET_1;
|
||||||
|
};
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
float4 out_var_SV_TARGET_1 [[color(0)]];
|
||||||
|
};
|
||||||
|
|
||||||
|
void main_1(thread float4* const tint_symbol_4) {
|
||||||
|
*(tint_symbol_4) = float4(NAN, NAN, NAN, NAN);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
fragment tint_symbol_1 tint_symbol() {
|
||||||
|
thread float4 tint_symbol_5 = 0.0f;
|
||||||
|
main_1(&(tint_symbol_5));
|
||||||
|
main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
|
||||||
|
tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
|
||||||
|
return tint_symbol_3;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,52 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 26
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint Fragment %main "main" %tint_symbol_1
|
||||||
|
OpExecutionMode %main OriginUpperLeft
|
||||||
|
OpName %out_var_SV_TARGET "out_var_SV_TARGET"
|
||||||
|
OpName %tint_symbol_1 "tint_symbol_1"
|
||||||
|
OpName %main_1 "main_1"
|
||||||
|
OpName %main_out "main_out"
|
||||||
|
OpMemberName %main_out 0 "out_var_SV_TARGET_1"
|
||||||
|
OpName %tint_symbol_2 "tint_symbol_2"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol_1 Location 0
|
||||||
|
OpMemberDecorate %main_out 0 Offset 0
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%_ptr_Private_v4float = OpTypePointer Private %v4float
|
||||||
|
%5 = OpConstantNull %v4float
|
||||||
|
%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
|
||||||
|
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||||
|
%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%8 = OpTypeFunction %void
|
||||||
|
%float_0x1_1p_128 = OpConstant %float 0x1.1p+128
|
||||||
|
%13 = OpConstantComposite %v4float %float_0x1_1p_128 %float_0x1_1p_128 %float_0x1_1p_128 %float_0x1_1p_128
|
||||||
|
%main_out = OpTypeStruct %v4float
|
||||||
|
%14 = OpTypeFunction %void %main_out
|
||||||
|
%main_1 = OpFunction %void None %8
|
||||||
|
%11 = OpLabel
|
||||||
|
OpStore %out_var_SV_TARGET %13
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%tint_symbol_2 = OpFunction %void None %14
|
||||||
|
%tint_symbol = OpFunctionParameter %main_out
|
||||||
|
%18 = OpLabel
|
||||||
|
%19 = OpCompositeExtract %v4float %tint_symbol 0
|
||||||
|
OpStore %tint_symbol_1 %19
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
%main = OpFunction %void None %8
|
||||||
|
%21 = OpLabel
|
||||||
|
%22 = OpFunctionCall %void %main_1
|
||||||
|
%24 = OpLoad %v4float %out_var_SV_TARGET
|
||||||
|
%25 = OpCompositeConstruct %main_out %24
|
||||||
|
%23 = OpFunctionCall %void %tint_symbol_2 %25
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,17 @@
|
||||||
|
var<private> out_var_SV_TARGET : vec4<f32>;
|
||||||
|
|
||||||
|
fn main_1() {
|
||||||
|
out_var_SV_TARGET = vec4<f32>(0x1.1p+128, 0x1.1p+128, 0x1.1p+128, 0x1.1p+128);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct main_out {
|
||||||
|
[[location(0)]]
|
||||||
|
out_var_SV_TARGET_1 : vec4<f32>;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[stage(fragment)]]
|
||||||
|
fn main() -> main_out {
|
||||||
|
main_1();
|
||||||
|
return main_out(out_var_SV_TARGET);
|
||||||
|
}
|
||||||
Loading…
Reference in New Issue