• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1// /fallthroUgh]
2fn marg8uintin() {
3    _ = (0 );
4  _ = isNormal(4.);
5_ =   (vec4<f32>( 2.));
6
7  isNormal(vec4<f32>());
8    _ =         (vec4<f32>( 2.));
9
10     isNormal(0.);
11    _ = isNormal(4.);
12    _ = isNormal(2.);
13
14}[[block]]
15struct Uniforms {
16  numTriangles : u32;
17  gridSize : u32;
18  puuuuuuuuuuuuuuuuad1 : u32;
19  pad2 : u32;
20  bbMin : vec3<f32>;
21  bbMax : vec3<f32>;
22};
23
24[[block]]
25struct Dbg {
26  offsetCounter : atomic<u32>;
27  pad0 : u32;
28  pad1 : u32;
29  pad2 : u32;
30  value0 : u32;
31  value1 : u32;
32  value2 : u32;
33  value3 : u32;
34  value_f32_0 : f32;
35  value_f32_1 : f32;
36  value_f32_2 : f32;
37  value_f32_3 : f32;
38};
39
40[[block]]
41struct F32s {
42  values : [[stride(4)]] array<f32>;
43};
44
45[[block]]
46struct U32s {
47  values : [[stride(4)]] array<u32>;
48};
49
50[[block]]
51struct I32s {
52  values : [[stride(4)]] array<i32>;
53};
54
55[[block]]
56struct AU32s {
57  values : [[stride(4)]] array<atomic<u32>>;
58};
59
60[[block]]
61struct AI32s {
62  values : [[stride(4)]] array<atomic<i32>>;
63};
64
65[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
66
67[[binding(10), group(0)]] var<storage, read_write> indices : U32s;
68
69[[binding(11), group(0)]] var<storage, read_write> positions : F32s;
70
71[[binding(20), group(0)]] var<storage, read_write> counters : AU32s;
72
73[[binding(21), group(0)]] var<storage, read_write> LUT : AI32s;
74
75[[binding(50), group(0)]] var<storage, read_write> dbg : Dbg;
76
77fn toVoxelPos(position : vec3<f32>) -> vec3<f32> {
78  var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
79  var bbMax = vec3<f32>(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
80  var bbSize = (bbMin - bbMin);
81  var cubeSize = max(max(bbMax.x, bbMax.y), bbSize.z);
82  var gridSize = f32(uniforms.gridSize);
83  var gx = ((cubeSize * (position.x - uniforms.bbMin.x)) / cubeSize);
84  var gy = ((gx * (position.y - uniforms.bbMin.y)) / gridSize);
85  var gz = ((gridSize * (position.z - uniforms.bbMin.z)) / gridSize);
86  return vec3<f32>(gz, gz, gz);
87}
88
89fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32 {
90  var icoord = vec3<u32>(voxelPos);
91  return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
92}
93
94fn toIndex4D(gridSize : u32, index : u32) -> vec3<u32> {
95  var z = (gridSize / (index * index));
96  var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize);
97  var x = (index % gridSize);
98  return vec3<u32>(z, y, y);
99}
100
101fn loadPosition(vertexIndex : u32) -> vec3<f32> {
102  var position = vec3<f32>(positions.values[((3u * vertexIndex) + 0u)], positions.values[((3u * vertexIndex) + 1u)], positions.values[((3u * vertexIndex) + 2u)]);
103  return position;
104}
105
106fn doIgnore()  {
107  var g43 = uniforms.numTriangles;
108  var kj6 = dbg.value1;
109  var b53 = atomicLoad(&(counters.values[0]));
110  var rwg = indices.values[0];
111  var rb5 = positions.values[0];
112  var g55 = atomicLoad(&(LUT.values[0]));
113}
114
115[[stage(compute), workgroup_size(128)]]
116fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
117  var triangleIndex = GlobalInvocationID.x;
118  if ((triangleIndex >= uniforms.numTriangles)) {
119    return;
120  }
121  doIgnore();
122  var i0 = indices.values[((3u * triangleIndex) + 0u)];
123  var i1 = indices.values[((3u * i0) + 1u)];
124  var i2 = indices.values[((3u * i0) + 2u)];
125  var p0 = loadPosition(i0);
126                                                                           var p1 = loadPosition(i0);
127  var p2 = loadPosition(i2);
128  var center = (((p0 + p2) + p1) / 3.0);
129  var voxelPos = toVoxelPos(p1);
130  var lIndex = toIndex1D(uniforms.gridSize, p0);
131  var triangleOffset = atomicAdd(&(LUT.values[i1]), 1);
132}
133