mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-02-20 02:42:43 +00:00
[msl-out] zero init variables in function address space
This commit is contained in:
teoxoy
Teodor Tanasoaia
parent
239bbbbed4
commit
b3d5e6d807
@ -3269,15 +3269,20 @@ impl<W: Write> Writer<W> {
|
||||
};
|
||||
let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)];
|
||||
write!(self.out, "{}{} {}", back::INDENT, ty_name, local_name)?;
|
||||
if let Some(value) = local.init {
|
||||
let coco = ConstantContext {
|
||||
handle: value,
|
||||
arena: &module.constants,
|
||||
names: &self.names,
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, " = {}", coco)?;
|
||||
}
|
||||
match local.init {
|
||||
Some(value) => {
|
||||
let coco = ConstantContext {
|
||||
handle: value,
|
||||
arena: &module.constants,
|
||||
names: &self.names,
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, " = {}", coco)?;
|
||||
}
|
||||
None => {
|
||||
write!(self.out, " = {{}}")?;
|
||||
}
|
||||
};
|
||||
writeln!(self.out, ";")?;
|
||||
}
|
||||
|
||||
@ -3801,15 +3806,20 @@ impl<W: Write> Writer<W> {
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
|
||||
if let Some(value) = local.init {
|
||||
let coco = ConstantContext {
|
||||
handle: value,
|
||||
arena: &module.constants,
|
||||
names: &self.names,
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, " = {}", coco)?;
|
||||
}
|
||||
match local.init {
|
||||
Some(value) => {
|
||||
let coco = ConstantContext {
|
||||
handle: value,
|
||||
arena: &module.constants,
|
||||
names: &self.names,
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, " = {}", coco)?;
|
||||
}
|
||||
None => {
|
||||
write!(self.out, " = {{}}")?;
|
||||
}
|
||||
};
|
||||
writeln!(self.out, ";")?;
|
||||
}
|
||||
|
||||
|
||||
@ -45,7 +45,7 @@ void test_matrix_within_struct_accesses(
|
||||
constant Baz& baz
|
||||
) {
|
||||
int idx = 9;
|
||||
Baz t;
|
||||
Baz t = {};
|
||||
int _e4 = idx;
|
||||
idx = _e4 - 1;
|
||||
metal::float3x2 unnamed = baz.m;
|
||||
@ -103,7 +103,7 @@ vertex foo_vertOutput foo_vert(
|
||||
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
|
||||
) {
|
||||
float foo = 0.0;
|
||||
type_16 c;
|
||||
type_16 c = {};
|
||||
float baz_1 = foo;
|
||||
foo = 1.0;
|
||||
test_matrix_within_struct_accesses(baz);
|
||||
@ -139,7 +139,7 @@ kernel void atomics(
|
||||
device Bar& bar [[buffer(0)]]
|
||||
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
|
||||
) {
|
||||
int tmp;
|
||||
int tmp = {};
|
||||
int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed);
|
||||
int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed);
|
||||
tmp = _e7;
|
||||
|
||||
@ -37,9 +37,9 @@ fragment main_Output main_(
|
||||
) {
|
||||
const FragmentIn fragment_in = { varyings.index };
|
||||
int i1_ = 0;
|
||||
metal::int2 i2_;
|
||||
metal::int2 i2_ = {};
|
||||
float v1_ = 0.0;
|
||||
metal::float4 v4_;
|
||||
metal::float4 v4_ = {};
|
||||
uint uniform_index = uni.index;
|
||||
uint non_uniform_index = fragment_in.index;
|
||||
i2_ = metal::int2(0);
|
||||
|
||||
@ -8,15 +8,15 @@ using metal::uint;
|
||||
kernel void main_(
|
||||
) {
|
||||
int i = 0;
|
||||
metal::int2 i2_;
|
||||
metal::int3 i3_;
|
||||
metal::int4 i4_;
|
||||
metal::int2 i2_ = {};
|
||||
metal::int3 i3_ = {};
|
||||
metal::int4 i4_ = {};
|
||||
uint u = 0u;
|
||||
metal::uint2 u2_;
|
||||
metal::uint3 u3_;
|
||||
metal::uint4 u4_;
|
||||
metal::float2 f2_;
|
||||
metal::float4 f4_;
|
||||
metal::uint2 u2_ = {};
|
||||
metal::uint3 u3_ = {};
|
||||
metal::uint4 u4_ = {};
|
||||
metal::float2 f2_ = {};
|
||||
metal::float4 f4_ = {};
|
||||
i2_ = metal::int2(0);
|
||||
i3_ = metal::int3(0);
|
||||
i4_ = metal::int4(0);
|
||||
|
||||
@ -37,15 +37,15 @@ kernel void main_(
|
||||
, device Particles& particlesDst [[buffer(2)]]
|
||||
, constant _mslBufferSizes& _buffer_sizes [[buffer(3)]]
|
||||
) {
|
||||
metal::float2 vPos;
|
||||
metal::float2 vVel;
|
||||
metal::float2 cMass;
|
||||
metal::float2 cVel;
|
||||
metal::float2 colVel;
|
||||
metal::float2 vPos = {};
|
||||
metal::float2 vVel = {};
|
||||
metal::float2 cMass = {};
|
||||
metal::float2 cVel = {};
|
||||
metal::float2 colVel = {};
|
||||
int cMassCount = 0;
|
||||
int cVelCount = 0;
|
||||
metal::float2 pos;
|
||||
metal::float2 vel;
|
||||
metal::float2 pos = {};
|
||||
metal::float2 vel = {};
|
||||
uint i = 0u;
|
||||
uint index = global_invocation_id.x;
|
||||
if (index >= NUM_PARTICLES) {
|
||||
|
||||
@ -16,7 +16,7 @@ struct PrimeIndices {
|
||||
uint collatz_iterations(
|
||||
uint n_base
|
||||
) {
|
||||
uint n;
|
||||
uint n = {};
|
||||
uint i = 0u;
|
||||
n = n_base;
|
||||
while(true) {
|
||||
|
||||
@ -49,7 +49,7 @@ struct main_Input {
|
||||
kernel void main_(
|
||||
metal::uint3 global_id [[thread_position_in_grid]]
|
||||
) {
|
||||
int pos;
|
||||
int pos = {};
|
||||
metal::threadgroup_barrier(metal::mem_flags::mem_device);
|
||||
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
||||
switch(1) {
|
||||
|
||||
@ -27,7 +27,7 @@ struct vert_mainOutput {
|
||||
};
|
||||
vertex vert_mainOutput vert_main(
|
||||
) {
|
||||
FragmentInput out;
|
||||
FragmentInput out = {};
|
||||
out.position = metal::float4(2.0, 4.0, 5.0, 6.0);
|
||||
out._flat = 8u;
|
||||
out._linear = 27.0;
|
||||
|
||||
@ -58,7 +58,7 @@ metal::float3 bool_cast(
|
||||
|
||||
float constructors(
|
||||
) {
|
||||
Foo foo;
|
||||
Foo foo = {};
|
||||
foo = Foo {metal::float4(1.0), 1};
|
||||
metal::float2x2 mat2comp = metal::float2x2(metal::float2(1.0, 0.0), metal::float2(0.0, 1.0));
|
||||
metal::float4x4 mat4comp = metal::float4x4(metal::float4(1.0, 0.0, 0.0, 0.0), metal::float4(0.0, 1.0, 0.0, 0.0), metal::float4(0.0, 0.0, 1.0, 0.0), metal::float4(0.0, 0.0, 0.0, 1.0));
|
||||
|
||||
@ -42,7 +42,7 @@ metal::float4 mock_function(
|
||||
threadgroup type_5 const& in_workgroup,
|
||||
thread type_6 const& in_private
|
||||
) {
|
||||
type_9 in_function;
|
||||
type_9 in_function = {};
|
||||
for(int _i=0; _i<2; ++_i) in_function.inner[_i] = type_9 {metal::float4(0.7070000171661377, 0.0, 0.0, 1.0), metal::float4(0.0, 0.7070000171661377, 0.0, 1.0)}.inner[_i];
|
||||
metal::float4 _e22 = in_storage.a.inner[i];
|
||||
metal::float4 _e25 = in_uniform.a.inner[i];
|
||||
|
||||
@ -65,7 +65,7 @@ vertex vs_mainOutput vs_main(
|
||||
) {
|
||||
const auto position = varyings.position;
|
||||
const auto normal = varyings.normal;
|
||||
VertexOutput out;
|
||||
VertexOutput out = {};
|
||||
metal::float4x4 w = u_entity.world;
|
||||
metal::float4x4 _e7 = u_entity.world;
|
||||
metal::float4 world_pos = _e7 * static_cast<metal::float4>(position);
|
||||
|
||||
@ -23,8 +23,8 @@ vertex vs_mainOutput vs_main(
|
||||
uint vertex_index [[vertex_id]]
|
||||
, constant Data& r_data [[buffer(0)]]
|
||||
) {
|
||||
int tmp1_;
|
||||
int tmp2_;
|
||||
int tmp1_ = {};
|
||||
int tmp2_ = {};
|
||||
tmp1_ = static_cast<int>(vertex_index) / 2;
|
||||
tmp2_ = static_cast<int>(vertex_index) & 1;
|
||||
int _e10 = tmp1_;
|
||||
|
||||
Loading…
Reference in New Issue
Block a user