144 lines
5.4 KiB
Plaintext
144 lines
5.4 KiB
Plaintext
bug/chromium/1273230.wgsl:4:7 warning: use of deprecated intrinsic
|
|
_ = isNormal(4.);
|
|
^^^^^^^^
|
|
|
|
bug/chromium/1273230.wgsl:7:3 warning: use of deprecated intrinsic
|
|
isNormal(vec4<f32>());
|
|
^^^^^^^^
|
|
|
|
bug/chromium/1273230.wgsl:10:6 warning: use of deprecated intrinsic
|
|
isNormal(0.);
|
|
^^^^^^^^
|
|
|
|
bug/chromium/1273230.wgsl:11:9 warning: use of deprecated intrinsic
|
|
_ = isNormal(4.);
|
|
^^^^^^^^
|
|
|
|
bug/chromium/1273230.wgsl:12:9 warning: use of deprecated intrinsic
|
|
_ = isNormal(2.);
|
|
^^^^^^^^
|
|
|
|
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
|
|
template<typename T, int N, int M>
|
|
inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
|
|
return lhs * vec<T, N>(rhs);
|
|
}
|
|
|
|
template<typename T, int N, int M>
|
|
inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
|
return vec<T, M>(lhs) * rhs;
|
|
}
|
|
|
|
struct Uniforms {
|
|
/* 0x0000 */ uint numTriangles;
|
|
/* 0x0004 */ uint gridSize;
|
|
/* 0x0008 */ uint puuuuuuuuuuuuuuuuad1;
|
|
/* 0x000c */ uint pad2;
|
|
/* 0x0010 */ packed_float3 bbMin;
|
|
/* 0x001c */ int8_t tint_pad[4];
|
|
/* 0x0020 */ packed_float3 bbMax;
|
|
/* 0x002c */ int8_t tint_pad_1[4];
|
|
};
|
|
struct Dbg {
|
|
/* 0x0000 */ atomic_uint offsetCounter;
|
|
/* 0x0004 */ uint pad0;
|
|
/* 0x0008 */ uint pad1;
|
|
/* 0x000c */ uint pad2;
|
|
/* 0x0010 */ uint value0;
|
|
/* 0x0014 */ uint value1;
|
|
/* 0x0018 */ uint value2;
|
|
/* 0x001c */ uint value3;
|
|
/* 0x0020 */ float value_f32_0;
|
|
/* 0x0024 */ float value_f32_1;
|
|
/* 0x0028 */ float value_f32_2;
|
|
/* 0x002c */ float value_f32_3;
|
|
};
|
|
struct F32s {
|
|
/* 0x0000 */ float values[1];
|
|
};
|
|
struct U32s {
|
|
/* 0x0000 */ uint values[1];
|
|
};
|
|
struct I32s {
|
|
int values[1];
|
|
};
|
|
struct AU32s {
|
|
/* 0x0000 */ atomic_uint values[1];
|
|
};
|
|
struct AI32s {
|
|
/* 0x0000 */ atomic_int values[1];
|
|
};
|
|
|
|
void marg8uintin() {
|
|
isnormal(4.0f);
|
|
isnormal(float4());
|
|
isnormal(0.0f);
|
|
isnormal(4.0f);
|
|
isnormal(2.0f);
|
|
}
|
|
|
|
float3 toVoxelPos(float3 position, const constant Uniforms* const tint_symbol) {
|
|
float3 bbMin = float3((*(tint_symbol)).bbMin[0], (*(tint_symbol)).bbMin[1], (*(tint_symbol)).bbMin[2]);
|
|
float3 bbMax = float3((*(tint_symbol)).bbMax[0], (*(tint_symbol)).bbMax[1], (*(tint_symbol)).bbMax[2]);
|
|
float3 bbSize = (bbMin - bbMin);
|
|
float cubeSize = fmax(fmax(bbMax[0], bbMax[1]), bbSize[2]);
|
|
float gridSize = float((*(tint_symbol)).gridSize);
|
|
float gx = ((cubeSize * (position[0] - (*(tint_symbol)).bbMin[0])) / cubeSize);
|
|
float gy = ((gx * (position[1] - (*(tint_symbol)).bbMin[1])) / gridSize);
|
|
float gz = ((gridSize * (position[2] - (*(tint_symbol)).bbMin[2])) / gridSize);
|
|
return float3(gz, gz, gz);
|
|
}
|
|
|
|
uint toIndex1D(uint gridSize, float3 voxelPos) {
|
|
uint3 icoord = uint3(voxelPos);
|
|
return ((icoord[0] + (gridSize * icoord[1])) + ((gridSize * gridSize) * icoord[2]));
|
|
}
|
|
|
|
uint3 toIndex4D(uint gridSize, uint index) {
|
|
uint z_1 = (gridSize / (index * index));
|
|
uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / gridSize);
|
|
uint x_1 = (index % gridSize);
|
|
return uint3(z_1, y_1, y_1);
|
|
}
|
|
|
|
float3 loadPosition(uint vertexIndex, device F32s* const tint_symbol_1) {
|
|
float3 position = float3((*(tint_symbol_1)).values[((3u * vertexIndex) + 0u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 1u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 2u)]);
|
|
return position;
|
|
}
|
|
|
|
void doIgnore(const constant Uniforms* const tint_symbol_2, device Dbg* const tint_symbol_3, device AU32s* const tint_symbol_4, device U32s* const tint_symbol_5, device F32s* const tint_symbol_6, device AI32s* const tint_symbol_7) {
|
|
uint g43 = (*(tint_symbol_2)).numTriangles;
|
|
uint kj6 = (*(tint_symbol_3)).value1;
|
|
uint b53 = atomic_load_explicit(&((*(tint_symbol_4)).values[0]), memory_order_relaxed);
|
|
uint rwg = (*(tint_symbol_5)).values[0];
|
|
float rb5 = (*(tint_symbol_6)).values[0];
|
|
int g55 = atomic_load_explicit(&((*(tint_symbol_7)).values[0]), memory_order_relaxed);
|
|
}
|
|
|
|
void main_count_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_8, device Dbg* const tint_symbol_9, device AU32s* const tint_symbol_10, device U32s* const tint_symbol_11, device F32s* const tint_symbol_12, device AI32s* const tint_symbol_13) {
|
|
uint triangleIndex = GlobalInvocationID[0];
|
|
if ((triangleIndex >= (*(tint_symbol_8)).numTriangles)) {
|
|
return;
|
|
}
|
|
doIgnore(tint_symbol_8, tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12, tint_symbol_13);
|
|
uint i0 = (*(tint_symbol_11)).values[((3u * triangleIndex) + 0u)];
|
|
uint i1 = (*(tint_symbol_11)).values[((3u * i0) + 1u)];
|
|
uint i2 = (*(tint_symbol_11)).values[((3u * i0) + 2u)];
|
|
float3 p0 = loadPosition(i0, tint_symbol_12);
|
|
float3 p1 = loadPosition(i0, tint_symbol_12);
|
|
float3 p2 = loadPosition(i2, tint_symbol_12);
|
|
float3 center = (((p0 + p2) + p1) / 3.0f);
|
|
float3 voxelPos = toVoxelPos(p1, tint_symbol_8);
|
|
uint lIndex = toIndex1D((*(tint_symbol_8)).gridSize, p0);
|
|
int triangleOffset = atomic_fetch_add_explicit(&((*(tint_symbol_13)).values[i1]), 1, memory_order_relaxed);
|
|
}
|
|
|
|
kernel void main_count(const constant Uniforms* tint_symbol_14 [[buffer(0)]], device Dbg* tint_symbol_15 [[buffer(1)]], device AU32s* tint_symbol_16 [[buffer(2)]], device U32s* tint_symbol_17 [[buffer(3)]], device F32s* tint_symbol_18 [[buffer(4)]], device AI32s* tint_symbol_19 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
|
|
main_count_inner(GlobalInvocationID, tint_symbol_14, tint_symbol_15, tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19);
|
|
return;
|
|
}
|
|
|