[hlsl] transform: Zero init arrays with a loop
If the array size is greater than a threshold. This is a work around for FXC stalling when initializing large arrays with a single zero-init assignment. Bug: tint:936 Fixed: tint:943 Fixed: tint:942 Change-Id: Ie93c8f373874b8d6d020d041fa48b38fb1352f71 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56775 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
b0455217fa
commit
b4ff73e250
|
|
@ -66,6 +66,10 @@ Output Hlsl::Run(const Program* in, const DataMap&) {
|
||||||
manager.Add<ExternalTextureTransform>();
|
manager.Add<ExternalTextureTransform>();
|
||||||
manager.Add<PromoteInitializersToConstVar>();
|
manager.Add<PromoteInitializersToConstVar>();
|
||||||
manager.Add<PadArrayElements>();
|
manager.Add<PadArrayElements>();
|
||||||
|
|
||||||
|
ZeroInitWorkgroupMemory::Config zero_init_cfg;
|
||||||
|
zero_init_cfg.init_arrays_with_loop_size_threshold = 32; // 8 scalars
|
||||||
|
data.Add<ZeroInitWorkgroupMemory::Config>(zero_init_cfg);
|
||||||
data.Add<CanonicalizeEntryPointIO::Config>(
|
data.Add<CanonicalizeEntryPointIO::Config>(
|
||||||
CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
|
CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
|
||||||
auto out = manager.Run(in, data);
|
auto out = manager.Run(in, data);
|
||||||
|
|
|
||||||
|
|
@ -24,6 +24,7 @@
|
||||||
#include "src/utils/get_or_create.h"
|
#include "src/utils/get_or_create.h"
|
||||||
|
|
||||||
TINT_INSTANTIATE_TYPEINFO(tint::transform::ZeroInitWorkgroupMemory);
|
TINT_INSTANTIATE_TYPEINFO(tint::transform::ZeroInitWorkgroupMemory);
|
||||||
|
TINT_INSTANTIATE_TYPEINFO(tint::transform::ZeroInitWorkgroupMemory::Config);
|
||||||
|
|
||||||
namespace tint {
|
namespace tint {
|
||||||
namespace transform {
|
namespace transform {
|
||||||
|
|
@ -32,14 +33,16 @@ namespace transform {
|
||||||
struct ZeroInitWorkgroupMemory::State {
|
struct ZeroInitWorkgroupMemory::State {
|
||||||
/// The clone context
|
/// The clone context
|
||||||
CloneContext& ctx;
|
CloneContext& ctx;
|
||||||
/// The built statements
|
/// The config
|
||||||
ast::StatementList& stmts;
|
Config cfg;
|
||||||
|
|
||||||
/// Zero() generates the statements required to zero initialize the workgroup
|
/// Zero() generates the statements required to zero initialize the workgroup
|
||||||
/// storage expression of type `ty`.
|
/// storage expression of type `ty`.
|
||||||
/// @param ty the expression type
|
/// @param ty the expression type
|
||||||
|
/// @param stmts the built statements
|
||||||
/// @param get_expr a function that builds the AST nodes for the expression
|
/// @param get_expr a function that builds the AST nodes for the expression
|
||||||
void Zero(const sem::Type* ty,
|
void Zero(const sem::Type* ty,
|
||||||
|
ast::StatementList& stmts,
|
||||||
const std::function<ast::Expression*()>& get_expr) {
|
const std::function<ast::Expression*()>& get_expr) {
|
||||||
if (CanZero(ty)) {
|
if (CanZero(ty)) {
|
||||||
auto* var = get_expr();
|
auto* var = get_expr();
|
||||||
|
|
@ -61,21 +64,32 @@ struct ZeroInitWorkgroupMemory::State {
|
||||||
if (auto* str = ty->As<sem::Struct>()) {
|
if (auto* str = ty->As<sem::Struct>()) {
|
||||||
for (auto* member : str->Members()) {
|
for (auto* member : str->Members()) {
|
||||||
auto name = ctx.Clone(member->Declaration()->symbol());
|
auto name = ctx.Clone(member->Declaration()->symbol());
|
||||||
Zero(member->Type(),
|
Zero(member->Type(), stmts,
|
||||||
[&] { return ctx.dst->MemberAccessor(get_expr(), name); });
|
[&] { return ctx.dst->MemberAccessor(get_expr(), name); });
|
||||||
}
|
}
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|
||||||
if (auto* arr = ty->As<sem::Array>()) {
|
if (auto* arr = ty->As<sem::Array>()) {
|
||||||
// TODO(bclayton): If array sizes become pipeline-overridable then this
|
if (ShouldEmitForLoop(arr)) {
|
||||||
// will need to emit code for a loop.
|
auto i = ctx.dst->Symbols().New("i");
|
||||||
// See https://github.com/gpuweb/gpuweb/pull/1792
|
auto* i_decl = ctx.dst->Decl(ctx.dst->Var(i, ctx.dst->ty.i32()));
|
||||||
for (size_t i = 0; i < arr->Count(); i++) {
|
auto* cond = ctx.dst->create<ast::BinaryExpression>(
|
||||||
Zero(arr->ElemType(), [&] {
|
ast::BinaryOp::kLessThan, ctx.dst->Expr(i),
|
||||||
return ctx.dst->IndexAccessor(get_expr(),
|
ctx.dst->Expr(static_cast<int>(arr->Count())));
|
||||||
static_cast<ProgramBuilder::u32>(i));
|
auto* inc = ctx.dst->Assign(i, ctx.dst->Add(i, 1));
|
||||||
});
|
ast::StatementList for_stmts;
|
||||||
|
Zero(arr->ElemType(), for_stmts,
|
||||||
|
[&] { return ctx.dst->IndexAccessor(get_expr(), i); });
|
||||||
|
auto* body = ctx.dst->Block(for_stmts);
|
||||||
|
stmts.emplace_back(ctx.dst->For(i_decl, cond, inc, body));
|
||||||
|
} else {
|
||||||
|
for (size_t i = 0; i < arr->Count(); i++) {
|
||||||
|
Zero(arr->ElemType(), stmts, [&] {
|
||||||
|
return ctx.dst->IndexAccessor(get_expr(),
|
||||||
|
static_cast<ProgramBuilder::u32>(i));
|
||||||
|
});
|
||||||
|
}
|
||||||
}
|
}
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|
@ -89,7 +103,7 @@ struct ZeroInitWorkgroupMemory::State {
|
||||||
/// CanZero() returns false, then the type needs to be initialized by
|
/// CanZero() returns false, then the type needs to be initialized by
|
||||||
/// decomposing the initialization into multiple sub-initializations.
|
/// decomposing the initialization into multiple sub-initializations.
|
||||||
/// @param ty the type to inspect
|
/// @param ty the type to inspect
|
||||||
static bool CanZero(const sem::Type* ty) {
|
bool CanZero(const sem::Type* ty) {
|
||||||
if (ty->Is<sem::Atomic>()) {
|
if (ty->Is<sem::Atomic>()) {
|
||||||
return false;
|
return false;
|
||||||
}
|
}
|
||||||
|
|
@ -101,21 +115,39 @@ struct ZeroInitWorkgroupMemory::State {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
if (auto* arr = ty->As<sem::Array>()) {
|
if (auto* arr = ty->As<sem::Array>()) {
|
||||||
if (!CanZero(arr->ElemType())) {
|
if (ShouldEmitForLoop(arr) || !CanZero(arr->ElemType())) {
|
||||||
return false;
|
return false;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
return true;
|
return true;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/// @returns true if the array should be emitted as a for-loop instead of
|
||||||
|
/// using zero-initializer statements.
|
||||||
|
/// @param array the array
|
||||||
|
bool ShouldEmitForLoop(const sem::Array* array) {
|
||||||
|
// TODO(bclayton): If array sizes become pipeline-overridable then this
|
||||||
|
// we need to return true for these arrays.
|
||||||
|
// See https://github.com/gpuweb/gpuweb/pull/1792
|
||||||
|
return (cfg.init_arrays_with_loop_size_threshold != 0) &&
|
||||||
|
(array->SizeInBytes() >= cfg.init_arrays_with_loop_size_threshold);
|
||||||
|
}
|
||||||
};
|
};
|
||||||
|
|
||||||
ZeroInitWorkgroupMemory::ZeroInitWorkgroupMemory() = default;
|
ZeroInitWorkgroupMemory::ZeroInitWorkgroupMemory() = default;
|
||||||
|
|
||||||
ZeroInitWorkgroupMemory::~ZeroInitWorkgroupMemory() = default;
|
ZeroInitWorkgroupMemory::~ZeroInitWorkgroupMemory() = default;
|
||||||
|
|
||||||
void ZeroInitWorkgroupMemory::Run(CloneContext& ctx, const DataMap&, DataMap&) {
|
void ZeroInitWorkgroupMemory::Run(CloneContext& ctx,
|
||||||
|
const DataMap& inputs,
|
||||||
|
DataMap&) {
|
||||||
auto& sem = ctx.src->Sem();
|
auto& sem = ctx.src->Sem();
|
||||||
|
|
||||||
|
Config cfg;
|
||||||
|
if (auto* c = inputs.Get<Config>()) {
|
||||||
|
cfg = *c;
|
||||||
|
}
|
||||||
|
|
||||||
for (auto* ast_func : ctx.src->AST().Functions()) {
|
for (auto* ast_func : ctx.src->AST().Functions()) {
|
||||||
if (!ast_func->IsEntryPoint()) {
|
if (!ast_func->IsEntryPoint()) {
|
||||||
continue;
|
continue;
|
||||||
|
|
@ -129,7 +161,7 @@ void ZeroInitWorkgroupMemory::Run(CloneContext& ctx, const DataMap&, DataMap&) {
|
||||||
if (var->StorageClass() != ast::StorageClass::kWorkgroup) {
|
if (var->StorageClass() != ast::StorageClass::kWorkgroup) {
|
||||||
continue;
|
continue;
|
||||||
}
|
}
|
||||||
State{ctx, stmts}.Zero(var->Type()->UnwrapRef(), [&] {
|
State{ctx, cfg}.Zero(var->Type()->UnwrapRef(), stmts, [&] {
|
||||||
auto var_name = ctx.Clone(var->Declaration()->symbol());
|
auto var_name = ctx.Clone(var->Declaration()->symbol());
|
||||||
return ctx.dst->Expr(var_name);
|
return ctx.dst->Expr(var_name);
|
||||||
});
|
});
|
||||||
|
|
@ -193,5 +225,11 @@ void ZeroInitWorkgroupMemory::Run(CloneContext& ctx, const DataMap&, DataMap&) {
|
||||||
ctx.Clone();
|
ctx.Clone();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
ZeroInitWorkgroupMemory::Config::Config() = default;
|
||||||
|
ZeroInitWorkgroupMemory::Config::Config(const Config&) = default;
|
||||||
|
ZeroInitWorkgroupMemory::Config::~Config() = default;
|
||||||
|
ZeroInitWorkgroupMemory::Config& ZeroInitWorkgroupMemory::Config::operator=(
|
||||||
|
const Config&) = default;
|
||||||
|
|
||||||
} // namespace transform
|
} // namespace transform
|
||||||
} // namespace tint
|
} // namespace tint
|
||||||
|
|
|
||||||
|
|
@ -26,6 +26,27 @@ namespace transform {
|
||||||
class ZeroInitWorkgroupMemory
|
class ZeroInitWorkgroupMemory
|
||||||
: public Castable<ZeroInitWorkgroupMemory, Transform> {
|
: public Castable<ZeroInitWorkgroupMemory, Transform> {
|
||||||
public:
|
public:
|
||||||
|
/// Configuration options for the transform
|
||||||
|
struct Config : public Castable<Config, Data> {
|
||||||
|
/// Constructor
|
||||||
|
Config();
|
||||||
|
|
||||||
|
/// Copy constructor
|
||||||
|
Config(const Config&);
|
||||||
|
|
||||||
|
/// Destructor
|
||||||
|
~Config() override;
|
||||||
|
|
||||||
|
/// Assignment operator
|
||||||
|
/// @returns this Config
|
||||||
|
Config& operator=(const Config&);
|
||||||
|
|
||||||
|
/// If greater than 0, then arrays of at least this size in bytes will be
|
||||||
|
/// zero initialized using a for loop. If 0, then the array is assigned a
|
||||||
|
/// zero initialized array with a single statement.
|
||||||
|
uint32_t init_arrays_with_loop_size_threshold = 0;
|
||||||
|
};
|
||||||
|
|
||||||
/// Constructor
|
/// Constructor
|
||||||
ZeroInitWorkgroupMemory();
|
ZeroInitWorkgroupMemory();
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -558,6 +558,56 @@ fn f([[builtin(local_invocation_index)]] local_invocation_index : u32) {
|
||||||
EXPECT_EQ(expect, str(got));
|
EXPECT_EQ(expect, str(got));
|
||||||
}
|
}
|
||||||
|
|
||||||
|
TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArray_InitWithLoop) {
|
||||||
|
auto* src = R"(
|
||||||
|
struct S {
|
||||||
|
a : array<i32, 3>; // size: 12, less than the loop threshold
|
||||||
|
b : array<i32, 4>; // size: 16, equal to the loop threshold
|
||||||
|
c : array<i32, 5>; // size: 20, greater than the loop threshold
|
||||||
|
};
|
||||||
|
|
||||||
|
var<workgroup> w : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn f() {
|
||||||
|
ignore(w); // Initialization should be inserted above this statement
|
||||||
|
}
|
||||||
|
)";
|
||||||
|
auto* expect = R"(
|
||||||
|
struct S {
|
||||||
|
a : array<i32, 3>;
|
||||||
|
b : array<i32, 4>;
|
||||||
|
c : array<i32, 5>;
|
||||||
|
};
|
||||||
|
|
||||||
|
var<workgroup> w : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn f([[builtin(local_invocation_index)]] local_invocation_index : u32) {
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
w.a = array<i32, 3>();
|
||||||
|
for(var i : i32; (i < 4); i = (i + 1)) {
|
||||||
|
w.b[i] = i32();
|
||||||
|
}
|
||||||
|
for(var i_1 : i32; (i_1 < 5); i_1 = (i_1 + 1)) {
|
||||||
|
w.c[i_1] = i32();
|
||||||
|
}
|
||||||
|
}
|
||||||
|
workgroupBarrier();
|
||||||
|
ignore(w);
|
||||||
|
}
|
||||||
|
)";
|
||||||
|
|
||||||
|
ZeroInitWorkgroupMemory::Config cfg;
|
||||||
|
cfg.init_arrays_with_loop_size_threshold = 16;
|
||||||
|
|
||||||
|
DataMap data;
|
||||||
|
data.Add<ZeroInitWorkgroupMemory::Config>(cfg);
|
||||||
|
auto got = Run<ZeroInitWorkgroupMemory>(src, data);
|
||||||
|
|
||||||
|
EXPECT_EQ(expect, str(got));
|
||||||
|
}
|
||||||
|
|
||||||
} // namespace
|
} // namespace
|
||||||
} // namespace transform
|
} // namespace transform
|
||||||
} // namespace tint
|
} // namespace tint
|
||||||
|
|
|
||||||
|
|
@ -69,10 +69,16 @@ void main(tint_symbol_1 tint_symbol) {
|
||||||
const uint3 global_id = tint_symbol.global_id;
|
const uint3 global_id = tint_symbol.global_id;
|
||||||
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
if ((local_invocation_index == 0u)) {
|
if ((local_invocation_index == 0u)) {
|
||||||
const float tint_symbol_5[64][64] = (float[64][64])0;
|
for(int i = 0; (i < 64); i = (i + 1)) {
|
||||||
mm_Asub = tint_symbol_5;
|
for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
|
||||||
const float tint_symbol_6[64][64] = (float[64][64])0;
|
mm_Asub[i][i_1] = 0.0f;
|
||||||
mm_Bsub = tint_symbol_6;
|
}
|
||||||
|
}
|
||||||
|
for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
|
||||||
|
for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
|
||||||
|
mm_Bsub[i_2][i_3] = 0.0f;
|
||||||
|
}
|
||||||
|
}
|
||||||
}
|
}
|
||||||
GroupMemoryBarrierWithGroupSync();
|
GroupMemoryBarrierWithGroupSync();
|
||||||
const uint tileRow = (local_id.y * RowPerThread);
|
const uint tileRow = (local_id.y * RowPerThread);
|
||||||
|
|
|
||||||
|
|
@ -0,0 +1,80 @@
|
||||||
|
[[block]] struct Params {
|
||||||
|
filterDim : u32;
|
||||||
|
blockDim : u32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[group(0), binding(0)]] var samp : sampler;
|
||||||
|
[[group(0), binding(1)]] var<uniform> params : Params;
|
||||||
|
[[group(1), binding(1)]] var inputTex : texture_2d<f32>;
|
||||||
|
[[group(1), binding(2)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
|
||||||
|
|
||||||
|
[[block]] struct Flip {
|
||||||
|
value : u32;
|
||||||
|
};
|
||||||
|
[[group(1), binding(3)]] var<uniform> flip : Flip;
|
||||||
|
|
||||||
|
// This shader blurs the input texture in one direction, depending on whether
|
||||||
|
// |flip.value| is 0 or 1.
|
||||||
|
// It does so by running (256 / 4) threads per workgroup to load 256
|
||||||
|
// texels into 4 rows of shared memory. Each thread loads a
|
||||||
|
// 4 x 4 block of texels to take advantage of the texture sampling
|
||||||
|
// hardware.
|
||||||
|
// Then, each thread computes the blur result by averaging the adjacent texel values
|
||||||
|
// in shared memory.
|
||||||
|
// Because we're operating on a subset of the texture, we cannot compute all of the
|
||||||
|
// results since not all of the neighbors are available in shared memory.
|
||||||
|
// Specifically, with 256 x 256 tiles, we can only compute and write out
|
||||||
|
// square blocks of size 256 - (filterSize - 1). We compute the number of blocks
|
||||||
|
// needed in Javascript and dispatch that amount.
|
||||||
|
|
||||||
|
var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(64, 1, 1)]]
|
||||||
|
fn main(
|
||||||
|
[[builtin(workgroup_id)]] WorkGroupID : vec3<u32>,
|
||||||
|
[[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>
|
||||||
|
) {
|
||||||
|
let filterOffset : u32 = (params.filterDim - 1u) / 2u;
|
||||||
|
let dims : vec2<i32> = textureDimensions(inputTex, 0);
|
||||||
|
|
||||||
|
let baseIndex = vec2<i32>(
|
||||||
|
WorkGroupID.xy * vec2<u32>(params.blockDim, 4u) +
|
||||||
|
LocalInvocationID.xy * vec2<u32>(4u, 1u)
|
||||||
|
) - vec2<i32>(i32(filterOffset), 0);
|
||||||
|
|
||||||
|
for (var r : u32 = 0u; r < 4u; r = r + 1u) {
|
||||||
|
for (var c : u32 = 0u; c < 4u; c = c + 1u) {
|
||||||
|
var loadIndex = baseIndex + vec2<i32>(i32(c), i32(r));
|
||||||
|
if (flip.value != 0u) {
|
||||||
|
loadIndex = loadIndex.yx;
|
||||||
|
}
|
||||||
|
|
||||||
|
tile[r][4u * LocalInvocationID.x + c] =
|
||||||
|
textureSampleLevel(inputTex, samp,
|
||||||
|
(vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims), 0.0).rgb;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
workgroupBarrier();
|
||||||
|
|
||||||
|
for (var r : u32 = 0u; r < 4u; r = r + 1u) {
|
||||||
|
for (var c : u32 = 0u; c < 4u; c = c + 1u) {
|
||||||
|
var writeIndex = baseIndex + vec2<i32>(i32(c), i32(r));
|
||||||
|
if (flip.value != 0u) {
|
||||||
|
writeIndex = writeIndex.yx;
|
||||||
|
}
|
||||||
|
|
||||||
|
let center : u32 = 4u * LocalInvocationID.x + c;
|
||||||
|
if (center >= filterOffset &&
|
||||||
|
center < 256u - filterOffset &&
|
||||||
|
all(writeIndex < dims)) {
|
||||||
|
var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
|
||||||
|
for (var f : u32 = 0u; f < params.filterDim; f = f + 1u) {
|
||||||
|
var i : u32 = center + f - filterOffset;
|
||||||
|
acc = acc + (1.0 / f32(params.filterDim)) * tile[r][i];
|
||||||
|
}
|
||||||
|
textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,96 @@
|
||||||
|
SamplerState samp : register(s0, space0);
|
||||||
|
cbuffer cbuffer_params : register(b1, space0) {
|
||||||
|
uint4 params[1];
|
||||||
|
};
|
||||||
|
Texture2D<float4> inputTex : register(t1, space1);
|
||||||
|
RWTexture2D<float4> outputTex : register(u2, space1);
|
||||||
|
|
||||||
|
cbuffer cbuffer_flip : register(b3, space1) {
|
||||||
|
uint4 flip[1];
|
||||||
|
};
|
||||||
|
groupshared float3 tile[4][256];
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint3 LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
uint3 WorkGroupID : SV_GroupID;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(64, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint3 WorkGroupID = tint_symbol.WorkGroupID;
|
||||||
|
const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
|
||||||
|
for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
|
||||||
|
tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
const uint scalar_offset = (0u) / 4;
|
||||||
|
const uint filterOffset = ((params[scalar_offset / 4][scalar_offset % 4] - 1u) / 2u);
|
||||||
|
int3 tint_tmp;
|
||||||
|
inputTex.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
|
||||||
|
const int2 dims = tint_tmp.xy;
|
||||||
|
const uint scalar_offset_1 = (4u) / 4;
|
||||||
|
const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
|
||||||
|
{
|
||||||
|
uint r = 0u;
|
||||||
|
for(; !(!((r < 4u))); r = (r + 1u)) {
|
||||||
|
{
|
||||||
|
uint c = 0u;
|
||||||
|
for(; !(!((c < 4u))); c = (c + 1u)) {
|
||||||
|
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
|
||||||
|
const uint scalar_offset_2 = (0u) / 4;
|
||||||
|
if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
|
||||||
|
loadIndex = loadIndex.yx;
|
||||||
|
}
|
||||||
|
tile[r][((4u * LocalInvocationID.x) + c)] = inputTex.SampleLevel(samp, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), 0.0f).rgb;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
{
|
||||||
|
uint r = 0u;
|
||||||
|
for(; !(!((r < 4u))); r = (r + 1u)) {
|
||||||
|
{
|
||||||
|
uint c = 0u;
|
||||||
|
for(; !(!((c < 4u))); c = (c + 1u)) {
|
||||||
|
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
|
||||||
|
const uint scalar_offset_3 = (0u) / 4;
|
||||||
|
if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
|
||||||
|
writeIndex = writeIndex.yx;
|
||||||
|
}
|
||||||
|
const uint center = ((4u * LocalInvocationID.x) + c);
|
||||||
|
bool tint_tmp_2 = (center >= filterOffset);
|
||||||
|
if (tint_tmp_2) {
|
||||||
|
tint_tmp_2 = (center < (256u - filterOffset));
|
||||||
|
}
|
||||||
|
bool tint_tmp_1 = (tint_tmp_2);
|
||||||
|
if (tint_tmp_1) {
|
||||||
|
tint_tmp_1 = all((writeIndex < dims));
|
||||||
|
}
|
||||||
|
if ((tint_tmp_1)) {
|
||||||
|
float3 acc = float3(0.0f, 0.0f, 0.0f);
|
||||||
|
{
|
||||||
|
uint f = 0u;
|
||||||
|
while (true) {
|
||||||
|
const uint scalar_offset_4 = (0u) / 4;
|
||||||
|
if (!(!(!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))))) { break; }
|
||||||
|
uint i = ((center + f) - filterOffset);
|
||||||
|
const uint scalar_offset_5 = (0u) / 4;
|
||||||
|
acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
|
||||||
|
f = (f + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
outputTex[writeIndex] = float4(acc, 1.0f);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,102 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct Params {
|
||||||
|
/* 0x0000 */ uint filterDim;
|
||||||
|
/* 0x0004 */ uint blockDim;
|
||||||
|
};
|
||||||
|
struct Flip {
|
||||||
|
/* 0x0000 */ uint value;
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper_1 {
|
||||||
|
float3 arr[256];
|
||||||
|
};
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
tint_array_wrapper_1 arr[4];
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[texture(1)]], sampler tint_symbol_5 [[sampler(0)]], texture2d<float, access::write> tint_symbol_6 [[texture(2)]], uint3 WorkGroupID [[threadgroup_position_in_grid]], uint3 LocalInvocationID [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Params& params [[buffer(1)]], constant Flip& flip [[buffer(3)]]) {
|
||||||
|
threadgroup tint_array_wrapper tint_symbol_3;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
tint_array_wrapper const tint_symbol_2 = {.arr={}};
|
||||||
|
tint_symbol_3 = tint_symbol_2;
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
uint const filterOffset = ((params.filterDim - 1u) / 2u);
|
||||||
|
int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
|
||||||
|
int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
|
||||||
|
{
|
||||||
|
uint r = 0u;
|
||||||
|
while (true) {
|
||||||
|
if (!((r < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
{
|
||||||
|
uint c = 0u;
|
||||||
|
while (true) {
|
||||||
|
if (!((c < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
|
||||||
|
if ((flip.value != 0u)) {
|
||||||
|
loadIndex = loadIndex.yx;
|
||||||
|
}
|
||||||
|
tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
|
||||||
|
{
|
||||||
|
c = (c + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
r = (r + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
{
|
||||||
|
uint r = 0u;
|
||||||
|
while (true) {
|
||||||
|
if (!((r < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
{
|
||||||
|
uint c = 0u;
|
||||||
|
while (true) {
|
||||||
|
if (!((c < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
|
||||||
|
if ((flip.value != 0u)) {
|
||||||
|
writeIndex = writeIndex.yx;
|
||||||
|
}
|
||||||
|
uint const center = ((4u * LocalInvocationID.x) + c);
|
||||||
|
if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
|
||||||
|
float3 acc = float3(0.0f, 0.0f, 0.0f);
|
||||||
|
{
|
||||||
|
uint f = 0u;
|
||||||
|
while (true) {
|
||||||
|
if (!((f < params.filterDim))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
uint i = ((center + f) - filterOffset);
|
||||||
|
acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
|
||||||
|
{
|
||||||
|
f = (f + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
|
||||||
|
}
|
||||||
|
{
|
||||||
|
c = (c + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
{
|
||||||
|
r = (r + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,374 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 239
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpCapability ImageQuery
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol_2 %tint_symbol %tint_symbol_1
|
||||||
|
OpExecutionMode %main LocalSize 64 1 1
|
||||||
|
OpName %samp "samp"
|
||||||
|
OpName %Params "Params"
|
||||||
|
OpMemberName %Params 0 "filterDim"
|
||||||
|
OpMemberName %Params 1 "blockDim"
|
||||||
|
OpName %params "params"
|
||||||
|
OpName %inputTex "inputTex"
|
||||||
|
OpName %outputTex "outputTex"
|
||||||
|
OpName %Flip "Flip"
|
||||||
|
OpMemberName %Flip 0 "value"
|
||||||
|
OpName %flip "flip"
|
||||||
|
OpName %tile "tile"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %tint_symbol_1 "tint_symbol_1"
|
||||||
|
OpName %tint_symbol_2 "tint_symbol_2"
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %r "r"
|
||||||
|
OpName %c "c"
|
||||||
|
OpName %loadIndex "loadIndex"
|
||||||
|
OpName %r_0 "r"
|
||||||
|
OpName %c_0 "c"
|
||||||
|
OpName %writeIndex "writeIndex"
|
||||||
|
OpName %acc "acc"
|
||||||
|
OpName %f "f"
|
||||||
|
OpName %i "i"
|
||||||
|
OpDecorate %samp DescriptorSet 0
|
||||||
|
OpDecorate %samp Binding 0
|
||||||
|
OpDecorate %Params Block
|
||||||
|
OpMemberDecorate %Params 0 Offset 0
|
||||||
|
OpMemberDecorate %Params 1 Offset 4
|
||||||
|
OpDecorate %params NonWritable
|
||||||
|
OpDecorate %params DescriptorSet 0
|
||||||
|
OpDecorate %params Binding 1
|
||||||
|
OpDecorate %inputTex DescriptorSet 1
|
||||||
|
OpDecorate %inputTex Binding 1
|
||||||
|
OpDecorate %outputTex NonReadable
|
||||||
|
OpDecorate %outputTex DescriptorSet 1
|
||||||
|
OpDecorate %outputTex Binding 2
|
||||||
|
OpDecorate %Flip Block
|
||||||
|
OpMemberDecorate %Flip 0 Offset 0
|
||||||
|
OpDecorate %flip NonWritable
|
||||||
|
OpDecorate %flip DescriptorSet 1
|
||||||
|
OpDecorate %flip Binding 3
|
||||||
|
OpDecorate %_arr_v3float_uint_256 ArrayStride 16
|
||||||
|
OpDecorate %_arr__arr_v3float_uint_256_uint_4 ArrayStride 4096
|
||||||
|
OpDecorate %tint_symbol BuiltIn WorkgroupId
|
||||||
|
OpDecorate %tint_symbol_1 BuiltIn LocalInvocationId
|
||||||
|
OpDecorate %tint_symbol_2 BuiltIn LocalInvocationIndex
|
||||||
|
%3 = OpTypeSampler
|
||||||
|
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
|
||||||
|
%samp = OpVariable %_ptr_UniformConstant_3 UniformConstant
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%Params = OpTypeStruct %uint %uint
|
||||||
|
%_ptr_Uniform_Params = OpTypePointer Uniform %Params
|
||||||
|
%params = OpVariable %_ptr_Uniform_Params Uniform
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%10 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||||
|
%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
|
||||||
|
%inputTex = OpVariable %_ptr_UniformConstant_10 UniformConstant
|
||||||
|
%14 = OpTypeImage %float 2D 0 0 0 2 Rgba8
|
||||||
|
%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
|
||||||
|
%outputTex = OpVariable %_ptr_UniformConstant_14 UniformConstant
|
||||||
|
%Flip = OpTypeStruct %uint
|
||||||
|
%_ptr_Uniform_Flip = OpTypePointer Uniform %Flip
|
||||||
|
%flip = OpVariable %_ptr_Uniform_Flip Uniform
|
||||||
|
%v3float = OpTypeVector %float 3
|
||||||
|
%uint_256 = OpConstant %uint 256
|
||||||
|
%_arr_v3float_uint_256 = OpTypeArray %v3float %uint_256
|
||||||
|
%uint_4 = OpConstant %uint 4
|
||||||
|
%_arr__arr_v3float_uint_256_uint_4 = OpTypeArray %_arr_v3float_uint_256 %uint_4
|
||||||
|
%_ptr_Workgroup__arr__arr_v3float_uint_256_uint_4 = OpTypePointer Workgroup %_arr__arr_v3float_uint_256_uint_4
|
||||||
|
%tile = OpVariable %_ptr_Workgroup__arr__arr_v3float_uint_256_uint_4 Workgroup
|
||||||
|
%v3uint = OpTypeVector %uint 3
|
||||||
|
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||||
|
%tint_symbol_1 = OpVariable %_ptr_Input_v3uint Input
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol_2 = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%31 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%41 = OpConstantNull %_arr__arr_v3float_uint_256_uint_4
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||||
|
%uint_1 = OpConstant %uint 1
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%v2int = OpTypeVector %int 2
|
||||||
|
%int_0 = OpConstant %int 0
|
||||||
|
%v2uint = OpTypeVector %uint 2
|
||||||
|
%66 = OpConstantComposite %v2uint %uint_4 %uint_1
|
||||||
|
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||||
|
%74 = OpConstantNull %uint
|
||||||
|
%_ptr_Function_v2int = OpTypePointer Function %v2int
|
||||||
|
%102 = OpConstantNull %v2int
|
||||||
|
%_ptr_Workgroup_v3float = OpTypePointer Workgroup %v3float
|
||||||
|
%v4float = OpTypeVector %float 4
|
||||||
|
%122 = OpTypeSampledImage %10
|
||||||
|
%v2float = OpTypeVector %float 2
|
||||||
|
%float_0_25 = OpConstant %float 0.25
|
||||||
|
%128 = OpConstantComposite %v2float %float_0_25 %float_0_25
|
||||||
|
%float_0 = OpConstant %float 0
|
||||||
|
%v2bool = OpTypeVector %bool 2
|
||||||
|
%193 = OpConstantComposite %v3float %float_0 %float_0 %float_0
|
||||||
|
%_ptr_Function_v3float = OpTypePointer Function %v3float
|
||||||
|
%196 = OpConstantNull %v3float
|
||||||
|
%float_1 = OpConstant %float 1
|
||||||
|
%main = OpFunction %void None %31
|
||||||
|
%34 = OpLabel
|
||||||
|
%r = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%c = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%loadIndex = OpVariable %_ptr_Function_v2int Function %102
|
||||||
|
%r_0 = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%c_0 = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%writeIndex = OpVariable %_ptr_Function_v2int Function %102
|
||||||
|
%acc = OpVariable %_ptr_Function_v3float Function %196
|
||||||
|
%f = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%i = OpVariable %_ptr_Function_uint Function %74
|
||||||
|
%35 = OpLoad %uint %tint_symbol_2
|
||||||
|
%37 = OpIEqual %bool %35 %uint_0
|
||||||
|
OpSelectionMerge %39 None
|
||||||
|
OpBranchConditional %37 %40 %39
|
||||||
|
%40 = OpLabel
|
||||||
|
OpStore %tile %41
|
||||||
|
OpBranch %39
|
||||||
|
%39 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%46 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
|
||||||
|
%47 = OpLoad %uint %46
|
||||||
|
%49 = OpISub %uint %47 %uint_1
|
||||||
|
%50 = OpUDiv %uint %49 %uint_2
|
||||||
|
%54 = OpLoad %10 %inputTex
|
||||||
|
%51 = OpImageQuerySizeLod %v2int %54 %int_0
|
||||||
|
%58 = OpLoad %v3uint %tint_symbol
|
||||||
|
%59 = OpVectorShuffle %v2uint %58 %58 0 1
|
||||||
|
%60 = OpAccessChain %_ptr_Uniform_uint %params %uint_1
|
||||||
|
%61 = OpLoad %uint %60
|
||||||
|
%62 = OpCompositeConstruct %v2uint %61 %uint_4
|
||||||
|
%63 = OpIMul %v2uint %59 %62
|
||||||
|
%64 = OpLoad %v3uint %tint_symbol_1
|
||||||
|
%65 = OpVectorShuffle %v2uint %64 %64 0 1
|
||||||
|
%67 = OpIMul %v2uint %65 %66
|
||||||
|
%68 = OpIAdd %v2uint %63 %67
|
||||||
|
%56 = OpBitcast %v2int %68
|
||||||
|
%69 = OpBitcast %int %50
|
||||||
|
%70 = OpCompositeConstruct %v2int %69 %int_0
|
||||||
|
%71 = OpISub %v2int %56 %70
|
||||||
|
OpStore %r %uint_0
|
||||||
|
OpBranch %75
|
||||||
|
%75 = OpLabel
|
||||||
|
OpLoopMerge %76 %77 None
|
||||||
|
OpBranch %78
|
||||||
|
%78 = OpLabel
|
||||||
|
%80 = OpLoad %uint %r
|
||||||
|
%81 = OpULessThan %bool %80 %uint_4
|
||||||
|
%79 = OpLogicalNot %bool %81
|
||||||
|
OpSelectionMerge %82 None
|
||||||
|
OpBranchConditional %79 %83 %82
|
||||||
|
%83 = OpLabel
|
||||||
|
OpBranch %76
|
||||||
|
%82 = OpLabel
|
||||||
|
OpStore %c %uint_0
|
||||||
|
OpBranch %85
|
||||||
|
%85 = OpLabel
|
||||||
|
OpLoopMerge %86 %87 None
|
||||||
|
OpBranch %88
|
||||||
|
%88 = OpLabel
|
||||||
|
%90 = OpLoad %uint %c
|
||||||
|
%91 = OpULessThan %bool %90 %uint_4
|
||||||
|
%89 = OpLogicalNot %bool %91
|
||||||
|
OpSelectionMerge %92 None
|
||||||
|
OpBranchConditional %89 %93 %92
|
||||||
|
%93 = OpLabel
|
||||||
|
OpBranch %86
|
||||||
|
%92 = OpLabel
|
||||||
|
%95 = OpLoad %uint %c
|
||||||
|
%94 = OpBitcast %int %95
|
||||||
|
%97 = OpLoad %uint %r
|
||||||
|
%96 = OpBitcast %int %97
|
||||||
|
%98 = OpCompositeConstruct %v2int %94 %96
|
||||||
|
%99 = OpIAdd %v2int %71 %98
|
||||||
|
OpStore %loadIndex %99
|
||||||
|
%103 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
|
||||||
|
%104 = OpLoad %uint %103
|
||||||
|
%105 = OpINotEqual %bool %104 %uint_0
|
||||||
|
OpSelectionMerge %106 None
|
||||||
|
OpBranchConditional %105 %107 %106
|
||||||
|
%107 = OpLabel
|
||||||
|
%108 = OpLoad %v2int %loadIndex
|
||||||
|
%109 = OpVectorShuffle %v2int %108 %108 1 0
|
||||||
|
OpStore %loadIndex %109
|
||||||
|
OpBranch %106
|
||||||
|
%106 = OpLabel
|
||||||
|
%110 = OpLoad %uint %r
|
||||||
|
%111 = OpAccessChain %_ptr_Input_uint %tint_symbol_1 %uint_0
|
||||||
|
%112 = OpLoad %uint %111
|
||||||
|
%113 = OpIMul %uint %uint_4 %112
|
||||||
|
%114 = OpLoad %uint %c
|
||||||
|
%115 = OpIAdd %uint %113 %114
|
||||||
|
%117 = OpAccessChain %_ptr_Workgroup_v3float %tile %110 %115
|
||||||
|
%120 = OpLoad %3 %samp
|
||||||
|
%121 = OpLoad %10 %inputTex
|
||||||
|
%123 = OpSampledImage %122 %121 %120
|
||||||
|
%126 = OpLoad %v2int %loadIndex
|
||||||
|
%124 = OpConvertSToF %v2float %126
|
||||||
|
%129 = OpFAdd %v2float %124 %128
|
||||||
|
%130 = OpConvertSToF %v2float %51
|
||||||
|
%131 = OpFDiv %v2float %129 %130
|
||||||
|
%118 = OpImageSampleExplicitLod %v4float %123 %131 Lod %float_0
|
||||||
|
%133 = OpVectorShuffle %v3float %118 %118 0 1 2
|
||||||
|
OpStore %117 %133
|
||||||
|
OpBranch %87
|
||||||
|
%87 = OpLabel
|
||||||
|
%134 = OpLoad %uint %c
|
||||||
|
%135 = OpIAdd %uint %134 %uint_1
|
||||||
|
OpStore %c %135
|
||||||
|
OpBranch %85
|
||||||
|
%86 = OpLabel
|
||||||
|
OpBranch %77
|
||||||
|
%77 = OpLabel
|
||||||
|
%136 = OpLoad %uint %r
|
||||||
|
%137 = OpIAdd %uint %136 %uint_1
|
||||||
|
OpStore %r %137
|
||||||
|
OpBranch %75
|
||||||
|
%76 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
OpStore %r_0 %uint_0
|
||||||
|
OpBranch %140
|
||||||
|
%140 = OpLabel
|
||||||
|
OpLoopMerge %141 %142 None
|
||||||
|
OpBranch %143
|
||||||
|
%143 = OpLabel
|
||||||
|
%145 = OpLoad %uint %r_0
|
||||||
|
%146 = OpULessThan %bool %145 %uint_4
|
||||||
|
%144 = OpLogicalNot %bool %146
|
||||||
|
OpSelectionMerge %147 None
|
||||||
|
OpBranchConditional %144 %148 %147
|
||||||
|
%148 = OpLabel
|
||||||
|
OpBranch %141
|
||||||
|
%147 = OpLabel
|
||||||
|
OpStore %c_0 %uint_0
|
||||||
|
OpBranch %150
|
||||||
|
%150 = OpLabel
|
||||||
|
OpLoopMerge %151 %152 None
|
||||||
|
OpBranch %153
|
||||||
|
%153 = OpLabel
|
||||||
|
%155 = OpLoad %uint %c_0
|
||||||
|
%156 = OpULessThan %bool %155 %uint_4
|
||||||
|
%154 = OpLogicalNot %bool %156
|
||||||
|
OpSelectionMerge %157 None
|
||||||
|
OpBranchConditional %154 %158 %157
|
||||||
|
%158 = OpLabel
|
||||||
|
OpBranch %151
|
||||||
|
%157 = OpLabel
|
||||||
|
%160 = OpLoad %uint %c_0
|
||||||
|
%159 = OpBitcast %int %160
|
||||||
|
%162 = OpLoad %uint %r_0
|
||||||
|
%161 = OpBitcast %int %162
|
||||||
|
%163 = OpCompositeConstruct %v2int %159 %161
|
||||||
|
%164 = OpIAdd %v2int %71 %163
|
||||||
|
OpStore %writeIndex %164
|
||||||
|
%166 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
|
||||||
|
%167 = OpLoad %uint %166
|
||||||
|
%168 = OpINotEqual %bool %167 %uint_0
|
||||||
|
OpSelectionMerge %169 None
|
||||||
|
OpBranchConditional %168 %170 %169
|
||||||
|
%170 = OpLabel
|
||||||
|
%171 = OpLoad %v2int %writeIndex
|
||||||
|
%172 = OpVectorShuffle %v2int %171 %171 1 0
|
||||||
|
OpStore %writeIndex %172
|
||||||
|
OpBranch %169
|
||||||
|
%169 = OpLabel
|
||||||
|
%173 = OpAccessChain %_ptr_Input_uint %tint_symbol_1 %uint_0
|
||||||
|
%174 = OpLoad %uint %173
|
||||||
|
%175 = OpIMul %uint %uint_4 %174
|
||||||
|
%176 = OpLoad %uint %c_0
|
||||||
|
%177 = OpIAdd %uint %175 %176
|
||||||
|
%178 = OpUGreaterThanEqual %bool %177 %50
|
||||||
|
OpSelectionMerge %179 None
|
||||||
|
OpBranchConditional %178 %180 %179
|
||||||
|
%180 = OpLabel
|
||||||
|
%181 = OpISub %uint %uint_256 %50
|
||||||
|
%182 = OpULessThan %bool %177 %181
|
||||||
|
OpBranch %179
|
||||||
|
%179 = OpLabel
|
||||||
|
%183 = OpPhi %bool %178 %169 %182 %180
|
||||||
|
OpSelectionMerge %184 None
|
||||||
|
OpBranchConditional %183 %185 %184
|
||||||
|
%185 = OpLabel
|
||||||
|
%187 = OpLoad %v2int %writeIndex
|
||||||
|
%188 = OpSLessThan %v2bool %187 %51
|
||||||
|
%186 = OpAll %bool %188
|
||||||
|
OpBranch %184
|
||||||
|
%184 = OpLabel
|
||||||
|
%190 = OpPhi %bool %183 %179 %186 %185
|
||||||
|
OpSelectionMerge %191 None
|
||||||
|
OpBranchConditional %190 %192 %191
|
||||||
|
%192 = OpLabel
|
||||||
|
OpStore %acc %193
|
||||||
|
OpStore %f %uint_0
|
||||||
|
OpBranch %198
|
||||||
|
%198 = OpLabel
|
||||||
|
OpLoopMerge %199 %200 None
|
||||||
|
OpBranch %201
|
||||||
|
%201 = OpLabel
|
||||||
|
%203 = OpLoad %uint %f
|
||||||
|
%204 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
|
||||||
|
%205 = OpLoad %uint %204
|
||||||
|
%206 = OpULessThan %bool %203 %205
|
||||||
|
%202 = OpLogicalNot %bool %206
|
||||||
|
OpSelectionMerge %207 None
|
||||||
|
OpBranchConditional %202 %208 %207
|
||||||
|
%208 = OpLabel
|
||||||
|
OpBranch %199
|
||||||
|
%207 = OpLabel
|
||||||
|
%209 = OpLoad %uint %f
|
||||||
|
%210 = OpIAdd %uint %177 %209
|
||||||
|
%211 = OpISub %uint %210 %50
|
||||||
|
OpStore %i %211
|
||||||
|
%213 = OpLoad %v3float %acc
|
||||||
|
%216 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
|
||||||
|
%217 = OpLoad %uint %216
|
||||||
|
%215 = OpConvertUToF %float %217
|
||||||
|
%218 = OpFDiv %float %float_1 %215
|
||||||
|
%219 = OpLoad %uint %r_0
|
||||||
|
%220 = OpLoad %uint %i
|
||||||
|
%221 = OpAccessChain %_ptr_Workgroup_v3float %tile %219 %220
|
||||||
|
%222 = OpLoad %v3float %221
|
||||||
|
%223 = OpVectorTimesScalar %v3float %222 %218
|
||||||
|
%224 = OpFAdd %v3float %213 %223
|
||||||
|
OpStore %acc %224
|
||||||
|
OpBranch %200
|
||||||
|
%200 = OpLabel
|
||||||
|
%225 = OpLoad %uint %f
|
||||||
|
%226 = OpIAdd %uint %225 %uint_1
|
||||||
|
OpStore %f %226
|
||||||
|
OpBranch %198
|
||||||
|
%199 = OpLabel
|
||||||
|
%228 = OpLoad %14 %outputTex
|
||||||
|
%229 = OpLoad %v2int %writeIndex
|
||||||
|
%230 = OpLoad %v3float %acc
|
||||||
|
%231 = OpCompositeExtract %float %230 0
|
||||||
|
%232 = OpCompositeExtract %float %230 1
|
||||||
|
%233 = OpCompositeExtract %float %230 2
|
||||||
|
%234 = OpCompositeConstruct %v4float %231 %232 %233 %float_1
|
||||||
|
OpImageWrite %228 %229 %234
|
||||||
|
OpBranch %191
|
||||||
|
%191 = OpLabel
|
||||||
|
OpBranch %152
|
||||||
|
%152 = OpLabel
|
||||||
|
%235 = OpLoad %uint %c_0
|
||||||
|
%236 = OpIAdd %uint %235 %uint_1
|
||||||
|
OpStore %c_0 %236
|
||||||
|
OpBranch %150
|
||||||
|
%151 = OpLabel
|
||||||
|
OpBranch %142
|
||||||
|
%142 = OpLabel
|
||||||
|
%237 = OpLoad %uint %r_0
|
||||||
|
%238 = OpIAdd %uint %237 %uint_1
|
||||||
|
OpStore %r_0 %238
|
||||||
|
OpBranch %140
|
||||||
|
%141 = OpLabel
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,106 @@
|
||||||
|
[[block]]
|
||||||
|
struct Params {
|
||||||
|
filterDim : u32;
|
||||||
|
blockDim : u32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[group(0), binding(0)]] var samp : sampler;
|
||||||
|
|
||||||
|
[[group(0), binding(1)]] var<uniform> params : Params;
|
||||||
|
|
||||||
|
[[group(1), binding(1)]] var inputTex : texture_2d<f32>;
|
||||||
|
|
||||||
|
[[group(1), binding(2)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
|
||||||
|
|
||||||
|
[[block]]
|
||||||
|
struct Flip {
|
||||||
|
value : u32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[group(1), binding(3)]] var<uniform> flip : Flip;
|
||||||
|
|
||||||
|
var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(64, 1, 1)]]
|
||||||
|
fn main([[builtin(workgroup_id)]] WorkGroupID : vec3<u32>, [[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
|
||||||
|
let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
|
||||||
|
let dims : vec2<i32> = textureDimensions(inputTex, 0);
|
||||||
|
let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
|
||||||
|
{
|
||||||
|
var r : u32 = 0u;
|
||||||
|
loop {
|
||||||
|
if (!((r < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
{
|
||||||
|
var c : u32 = 0u;
|
||||||
|
loop {
|
||||||
|
if (!((c < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
|
||||||
|
if ((flip.value != 0u)) {
|
||||||
|
loadIndex = loadIndex.yx;
|
||||||
|
}
|
||||||
|
tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
c = (c + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
r = (r + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
workgroupBarrier();
|
||||||
|
{
|
||||||
|
var r : u32 = 0u;
|
||||||
|
loop {
|
||||||
|
if (!((r < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
{
|
||||||
|
var c : u32 = 0u;
|
||||||
|
loop {
|
||||||
|
if (!((c < 4u))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
|
||||||
|
if ((flip.value != 0u)) {
|
||||||
|
writeIndex = writeIndex.yx;
|
||||||
|
}
|
||||||
|
let center : u32 = ((4u * LocalInvocationID.x) + c);
|
||||||
|
if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
|
||||||
|
var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
|
||||||
|
{
|
||||||
|
var f : u32 = 0u;
|
||||||
|
loop {
|
||||||
|
if (!((f < params.filterDim))) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
var i : u32 = ((center + f) - filterOffset);
|
||||||
|
acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
f = (f + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
|
||||||
|
}
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
c = (c + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
r = (r + 1u);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
@ -336,10 +336,15 @@ void main(tint_symbol_1 tint_symbol) {
|
||||||
const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
|
const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
|
||||||
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
if ((local_invocation_index == 0u)) {
|
if ((local_invocation_index == 0u)) {
|
||||||
const float tint_symbol_6[64][64] = (float[64][64])0;
|
for(int i = 0; (i < 64); i = (i + 1)) {
|
||||||
mm_Asub = tint_symbol_6;
|
for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
|
||||||
const float tint_symbol_7[64][1] = (float[64][1])0;
|
mm_Asub[i][i_1] = 0.0f;
|
||||||
mm_Bsub = tint_symbol_7;
|
}
|
||||||
|
}
|
||||||
|
for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
|
||||||
|
const float tint_symbol_6[1] = (float[1])0;
|
||||||
|
mm_Bsub[i_2] = tint_symbol_6;
|
||||||
|
}
|
||||||
}
|
}
|
||||||
GroupMemoryBarrierWithGroupSync();
|
GroupMemoryBarrierWithGroupSync();
|
||||||
gl_LocalInvocationID = gl_LocalInvocationID_param;
|
gl_LocalInvocationID = gl_LocalInvocationID_param;
|
||||||
|
|
|
||||||
Loading…
Reference in New Issue