msl: Overload matrix-vector arithmetic operators
These operators are not defined in the metal namespace when the vector operands are packed. Fixed: tint:1121 Change-Id: I2e8f4302e08117ca41bac6c05fb24a70d1215740 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62480 Kokoro: Kokoro <noreply+kokoro@google.com> Auto-Submit: James Price <jrprice@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
This commit is contained in:
James Price
parent
46978033a7
commit
85d2e448de
|
|
@ -2310,6 +2310,26 @@ bool GeneratorImpl::EmitPackedType(std::ostream& out,
|
|||
if (!EmitType(out, vec, "")) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (vec->is_float_vector() && !matrix_packed_vector_overloads_) {
|
||||
// Overload operators for matrix-vector arithmetic where the vector
|
||||
// operand is packed, as these overloads to not exist in the metal
|
||||
// namespace.
|
||||
TextBuffer b;
|
||||
TINT_DEFER(helpers_.Append(b));
|
||||
line(&b) << R"(template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
)";
|
||||
matrix_packed_vector_overloads_ = true;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -355,6 +355,9 @@ class GeneratorImpl : public TextGenerator {
|
|||
/// True if an invariant attribute has been generated.
|
||||
bool has_invariant_ = false;
|
||||
|
||||
/// True if matrix-packed_vector operator overloads have been generated.
|
||||
bool matrix_packed_vector_overloads_ = false;
|
||||
|
||||
std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
|
||||
std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ int4 arr[4];
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ int4 arr[4];
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct Inner {
|
||||
/* 0x0000 */ int x;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct Inner {
|
||||
/* 0x0000 */ int x;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ int4 arr[4];
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct Inner {
|
||||
/* 0x0000 */ int x;
|
||||
};
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto 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;
|
||||
|
|
|
|||
|
|
@ -0,0 +1,127 @@
|
|||
// Take from here:
|
||||
// https://github.com/shrekshao/webgpu-deferred-renderer/blob/4f8bf0910793100aa8d60dbd1319bddb5357b1fa/renderer/LightCulling.js
|
||||
// With these token replacements:
|
||||
// $NUM_TILE_LIGHT_SLOT = 64
|
||||
// $NUM_TILES = 4
|
||||
// $TILE_COUNT_Y = 2
|
||||
// $TILE_COUNT_X = 2
|
||||
// $TILE_SIZE = 16
|
||||
|
||||
struct LightData {
|
||||
position : vec4<f32>;
|
||||
color : vec3<f32>;
|
||||
radius : f32;
|
||||
};
|
||||
[[block]] struct LightsBuffer {
|
||||
lights: array<LightData>;
|
||||
};
|
||||
[[group(0), binding(0)]] var<storage, read_write> lightsBuffer: LightsBuffer;
|
||||
struct TileLightIdData {
|
||||
count: atomic<u32>;
|
||||
lightId: array<u32, 64>;
|
||||
};
|
||||
[[block]] struct Tiles {
|
||||
data: array<TileLightIdData, 4>;
|
||||
};
|
||||
[[group(1), binding(0)]] var<storage, read_write> tileLightId: Tiles;
|
||||
|
||||
[[block]] struct Config {
|
||||
numLights : u32;
|
||||
numTiles : u32;
|
||||
tileCountX : u32;
|
||||
tileCountY : u32;
|
||||
numTileLightSlot : u32;
|
||||
tileSize : u32;
|
||||
};
|
||||
[[group(2), binding(0)]] var<uniform> config: Config;
|
||||
[[block]] struct Uniforms {
|
||||
min : vec4<f32>;
|
||||
max : vec4<f32>;
|
||||
// camera
|
||||
viewMatrix : mat4x4<f32>;
|
||||
projectionMatrix : mat4x4<f32>;
|
||||
// Tile info
|
||||
fullScreenSize : vec4<f32>; // width, height
|
||||
};
|
||||
[[group(3), binding(0)]] var<uniform> uniforms: Uniforms;
|
||||
[[stage(compute), workgroup_size(64, 1, 1)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
var index = GlobalInvocationID.x;
|
||||
if (index >= config.numLights) {
|
||||
return;
|
||||
}
|
||||
// Light position updating
|
||||
lightsBuffer.lights[index].position.y = lightsBuffer.lights[index].position.y - 0.1 + 0.001 * (f32(index) - 64.0 * floor(f32(index) / 64.0));
|
||||
|
||||
if (lightsBuffer.lights[index].position.y < uniforms.min.y) {
|
||||
lightsBuffer.lights[index].position.y = uniforms.max.y;
|
||||
}
|
||||
// Light culling
|
||||
// Implementation here is Tiled without per tile min-max depth
|
||||
// You could also implement cluster culling
|
||||
// Feel free to add more compute passes if necessary
|
||||
// some math reference: http://www.txutxi.com/?p=444
|
||||
var M: mat4x4<f32> = uniforms.projectionMatrix;
|
||||
var viewNear: f32 = - M[3][2] / ( -1.0 + M[2][2]);
|
||||
var viewFar: f32 = - M[3][2] / (1.0 + M[2][2]);
|
||||
var lightPos = lightsBuffer.lights[index].position;
|
||||
lightPos = uniforms.viewMatrix * lightPos;
|
||||
lightPos = lightPos / lightPos.w;
|
||||
var lightRadius: f32 = lightsBuffer.lights[index].radius;
|
||||
var boxMin: vec4<f32> = lightPos - vec4<f32>(vec3<f32>(lightRadius), 0.0);
|
||||
var boxMax: vec4<f32> = lightPos + vec4<f32>(vec3<f32>(lightRadius), 0.0);
|
||||
var frustumPlanes: array<vec4<f32>, 6>;
|
||||
frustumPlanes[4] = vec4<f32>(0.0, 0.0, -1.0, viewNear); // near
|
||||
frustumPlanes[5] = vec4<f32>(0.0, 0.0, 1.0, -viewFar); // far
|
||||
let TILE_SIZE: i32 = 16;
|
||||
let TILE_COUNT_X: i32 = 2;
|
||||
let TILE_COUNT_Y: i32 = 2;
|
||||
for (var y : i32 = 0; y < TILE_COUNT_Y; y = y + 1) {
|
||||
for (var x : i32 = 0; x < TILE_COUNT_X; x = x + 1) {
|
||||
var tilePixel0Idx : vec2<i32> = vec2<i32>(x * TILE_SIZE, y * TILE_SIZE);
|
||||
// tile position in NDC space
|
||||
var floorCoord: vec2<f32> = 2.0 * vec2<f32>(tilePixel0Idx) / uniforms.fullScreenSize.xy - vec2<f32>(1.0); // -1, 1
|
||||
var ceilCoord: vec2<f32> = 2.0 * vec2<f32>(tilePixel0Idx + vec2<i32>(TILE_SIZE)) / uniforms.fullScreenSize.xy - vec2<f32>(1.0); // -1, 1
|
||||
var viewFloorCoord: vec2<f32> = vec2<f32>( (- viewNear * floorCoord.x - M[2][0] * viewNear) / M[0][0] , (- viewNear * floorCoord.y - M[2][1] * viewNear) / M[1][1] );
|
||||
var viewCeilCoord: vec2<f32> = vec2<f32>( (- viewNear * ceilCoord.x - M[2][0] * viewNear) / M[0][0] , (- viewNear * ceilCoord.y - M[2][1] * viewNear) / M[1][1] );
|
||||
frustumPlanes[0] = vec4<f32>(1.0, 0.0, - viewFloorCoord.x / viewNear, 0.0); // left
|
||||
frustumPlanes[1] = vec4<f32>(-1.0, 0.0, viewCeilCoord.x / viewNear, 0.0); // right
|
||||
frustumPlanes[2] = vec4<f32>(0.0, 1.0, - viewFloorCoord.y / viewNear, 0.0); // bottom
|
||||
frustumPlanes[3] = vec4<f32>(0.0, -1.0, viewCeilCoord.y / viewNear, 0.0); // top
|
||||
var dp: f32 = 0.0; // dot product
|
||||
for (var i: u32 = 0u; i < 6u; i = i + 1u)
|
||||
{
|
||||
var p: vec4<f32>;
|
||||
if (frustumPlanes[i].x > 0.0) {
|
||||
p.x = boxMax.x;
|
||||
} else {
|
||||
p.x = boxMin.x;
|
||||
}
|
||||
if (frustumPlanes[i].y > 0.0) {
|
||||
p.y = boxMax.y;
|
||||
} else {
|
||||
p.y = boxMin.y;
|
||||
}
|
||||
if (frustumPlanes[i].z > 0.0) {
|
||||
p.z = boxMax.z;
|
||||
} else {
|
||||
p.z = boxMin.z;
|
||||
}
|
||||
p.w = 1.0;
|
||||
dp = dp + min(0.0, dot(p, frustumPlanes[i]));
|
||||
}
|
||||
if (dp >= 0.0) {
|
||||
// light is overlapping with the tile
|
||||
var tileId: u32 = u32(x + y * TILE_COUNT_X);
|
||||
if (tileId < 0u || tileId >= config.numTiles) {
|
||||
continue;
|
||||
}
|
||||
var offset: u32 = atomicAdd(&(tileLightId.data[tileId].count), 1u);
|
||||
if (offset >= config.numTileLightSlot) {
|
||||
continue;
|
||||
}
|
||||
tileLightId.data[tileId].lightId[offset] = GlobalInvocationID.x;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,115 @@
|
|||
uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
|
||||
uint original_value = 0;
|
||||
buffer.InterlockedAdd(offset, value, original_value);
|
||||
return original_value;
|
||||
}
|
||||
|
||||
RWByteAddressBuffer lightsBuffer : register(u0, space0);
|
||||
|
||||
RWByteAddressBuffer tileLightId : register(u0, space1);
|
||||
|
||||
cbuffer cbuffer_config : register(b0, space2) {
|
||||
uint4 config[2];
|
||||
};
|
||||
|
||||
cbuffer cbuffer_uniforms : register(b0, space3) {
|
||||
uint4 uniforms[11];
|
||||
};
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint3 GlobalInvocationID : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
float4x4 tint_symbol_6(uint4 buffer[11], uint offset) {
|
||||
const uint scalar_offset = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_1 = ((offset + 16u)) / 4;
|
||||
const uint scalar_offset_2 = ((offset + 32u)) / 4;
|
||||
const uint scalar_offset_3 = ((offset + 48u)) / 4;
|
||||
return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
|
||||
}
|
||||
|
||||
void main_inner(uint3 GlobalInvocationID) {
|
||||
uint index = GlobalInvocationID.x;
|
||||
if ((index >= config[0].x)) {
|
||||
return;
|
||||
}
|
||||
lightsBuffer.Store(((32u * index) + 4u), asuint(((asfloat(lightsBuffer.Load(((32u * index) + 4u))) - 0.100000001f) + (0.001f * (float(index) - (64.0f * floor((float(index) / 64.0f))))))));
|
||||
if ((asfloat(lightsBuffer.Load(((32u * index) + 4u))) < asfloat(uniforms[0].y))) {
|
||||
lightsBuffer.Store(((32u * index) + 4u), asuint(asfloat(uniforms[1].y)));
|
||||
}
|
||||
float4x4 M = tint_symbol_6(uniforms, 96u);
|
||||
float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
|
||||
float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
|
||||
float4 lightPos = asfloat(lightsBuffer.Load4((32u * index)));
|
||||
lightPos = mul(lightPos, tint_symbol_6(uniforms, 32u));
|
||||
lightPos = (lightPos / lightPos.w);
|
||||
float lightRadius = asfloat(lightsBuffer.Load(((32u * index) + 28u)));
|
||||
float4 boxMin = (lightPos - float4(float3((lightRadius).xxx), 0.0f));
|
||||
float4 boxMax = (lightPos + float4(float3((lightRadius).xxx), 0.0f));
|
||||
float4 frustumPlanes[6] = (float4[6])0;
|
||||
frustumPlanes[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
|
||||
frustumPlanes[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
|
||||
const int TILE_SIZE = 16;
|
||||
const int TILE_COUNT_X = 2;
|
||||
{
|
||||
for(int y_1 = 0; (y_1 < 2); y_1 = (y_1 + 1)) {
|
||||
{
|
||||
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = (x_1 + 1)) {
|
||||
int2 tilePixel0Idx = int2((x_1 * TILE_SIZE), (y_1 * TILE_SIZE));
|
||||
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / asfloat(uniforms[10]).xy) - float2((1.0f).xx));
|
||||
float2 ceilCoord = (((2.0f * float2((tilePixel0Idx + int2((TILE_SIZE).xx)))) / asfloat(uniforms[10]).xy) - float2((1.0f).xx));
|
||||
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord.x) / viewNear), 0.0f);
|
||||
frustumPlanes[1] = float4(-1.0f, 0.0f, (viewCeilCoord.x / viewNear), 0.0f);
|
||||
frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord.y) / viewNear), 0.0f);
|
||||
frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord.y / viewNear), 0.0f);
|
||||
float dp = 0.0f;
|
||||
{
|
||||
for(uint i = 0u; (i < 6u); i = (i + 1u)) {
|
||||
float4 p = float4(0.0f, 0.0f, 0.0f, 0.0f);
|
||||
if ((frustumPlanes[i].x > 0.0f)) {
|
||||
p.x = boxMax.x;
|
||||
} else {
|
||||
p.x = boxMin.x;
|
||||
}
|
||||
if ((frustumPlanes[i].y > 0.0f)) {
|
||||
p.y = boxMax.y;
|
||||
} else {
|
||||
p.y = boxMin.y;
|
||||
}
|
||||
if ((frustumPlanes[i].z > 0.0f)) {
|
||||
p.z = boxMax.z;
|
||||
} else {
|
||||
p.z = boxMin.z;
|
||||
}
|
||||
p.w = 1.0f;
|
||||
dp = (dp + min(0.0f, dot(p, frustumPlanes[i])));
|
||||
}
|
||||
}
|
||||
if ((dp >= 0.0f)) {
|
||||
uint tileId = uint((x_1 + (y_1 * TILE_COUNT_X)));
|
||||
bool tint_tmp = (tileId < 0u);
|
||||
if (!tint_tmp) {
|
||||
tint_tmp = (tileId >= config[0].y);
|
||||
}
|
||||
if ((tint_tmp)) {
|
||||
continue;
|
||||
}
|
||||
uint offset = atomicAdd_1(tileLightId, (260u * tileId), 1u);
|
||||
if ((offset >= config[1].x)) {
|
||||
continue;
|
||||
}
|
||||
tileLightId.Store((((260u * tileId) + 4u) + (4u * offset)), asuint(GlobalInvocationID.x));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
[numthreads(64, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
main_inner(tint_symbol.GlobalInvocationID);
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,130 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct LightData {
|
||||
/* 0x0000 */ float4 position;
|
||||
/* 0x0010 */ packed_float3 color;
|
||||
/* 0x001c */ float radius;
|
||||
};
|
||||
struct LightsBuffer {
|
||||
/* 0x0000 */ LightData lights[1];
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ uint arr[64];
|
||||
};
|
||||
struct TileLightIdData {
|
||||
/* 0x0000 */ atomic_uint count;
|
||||
/* 0x0004 */ tint_array_wrapper lightId;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
/* 0x0000 */ TileLightIdData arr[4];
|
||||
};
|
||||
struct Tiles {
|
||||
/* 0x0000 */ tint_array_wrapper_1 data;
|
||||
};
|
||||
struct Config {
|
||||
/* 0x0000 */ uint numLights;
|
||||
/* 0x0004 */ uint numTiles;
|
||||
/* 0x0008 */ uint tileCountX;
|
||||
/* 0x000c */ uint tileCountY;
|
||||
/* 0x0010 */ uint numTileLightSlot;
|
||||
/* 0x0014 */ uint tileSize;
|
||||
};
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ float4 min;
|
||||
/* 0x0010 */ float4 max;
|
||||
/* 0x0020 */ float4x4 viewMatrix;
|
||||
/* 0x0060 */ float4x4 projectionMatrix;
|
||||
/* 0x00a0 */ float4 fullScreenSize;
|
||||
};
|
||||
struct tint_array_wrapper_2 {
|
||||
float4 arr[6];
|
||||
};
|
||||
|
||||
void tint_symbol_inner(constant Config& config, constant Uniforms& uniforms, device LightsBuffer& lightsBuffer, device Tiles& tileLightId, uint3 GlobalInvocationID) {
|
||||
uint index = GlobalInvocationID.x;
|
||||
if ((index >= config.numLights)) {
|
||||
return;
|
||||
}
|
||||
lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001f) + (0.001f * (float(index) - (64.0f * floor((float(index) / 64.0f))))));
|
||||
if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) {
|
||||
lightsBuffer.lights[index].position.y = uniforms.max.y;
|
||||
}
|
||||
float4x4 M = uniforms.projectionMatrix;
|
||||
float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
|
||||
float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
|
||||
float4 lightPos = lightsBuffer.lights[index].position;
|
||||
lightPos = (uniforms.viewMatrix * lightPos);
|
||||
lightPos = (lightPos / lightPos.w);
|
||||
float lightRadius = lightsBuffer.lights[index].radius;
|
||||
float4 boxMin = (lightPos - float4(float3(lightRadius), 0.0f));
|
||||
float4 boxMax = (lightPos + float4(float3(lightRadius), 0.0f));
|
||||
tint_array_wrapper_2 frustumPlanes = {};
|
||||
frustumPlanes.arr[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
|
||||
frustumPlanes.arr[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
|
||||
int const TILE_SIZE = 16;
|
||||
int const TILE_COUNT_X = 2;
|
||||
int const TILE_COUNT_Y = 2;
|
||||
for(int y_1 = 0; (y_1 < TILE_COUNT_Y); y_1 = as_type<int>((as_type<uint>(y_1) + as_type<uint>(1)))) {
|
||||
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
|
||||
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
|
||||
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - float2(1.0f));
|
||||
float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / uniforms.fullScreenSize.xy) - float2(1.0f));
|
||||
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
frustumPlanes.arr[0] = float4(1.0f, 0.0f, (-(viewFloorCoord.x) / viewNear), 0.0f);
|
||||
frustumPlanes.arr[1] = float4(-1.0f, 0.0f, (viewCeilCoord.x / viewNear), 0.0f);
|
||||
frustumPlanes.arr[2] = float4(0.0f, 1.0f, (-(viewFloorCoord.y) / viewNear), 0.0f);
|
||||
frustumPlanes.arr[3] = float4(0.0f, -1.0f, (viewCeilCoord.y / viewNear), 0.0f);
|
||||
float dp = 0.0f;
|
||||
for(uint i = 0u; (i < 6u); i = (i + 1u)) {
|
||||
float4 p = 0.0f;
|
||||
if ((frustumPlanes.arr[i].x > 0.0f)) {
|
||||
p.x = boxMax.x;
|
||||
} else {
|
||||
p.x = boxMin.x;
|
||||
}
|
||||
if ((frustumPlanes.arr[i].y > 0.0f)) {
|
||||
p.y = boxMax.y;
|
||||
} else {
|
||||
p.y = boxMin.y;
|
||||
}
|
||||
if ((frustumPlanes.arr[i].z > 0.0f)) {
|
||||
p.z = boxMax.z;
|
||||
} else {
|
||||
p.z = boxMin.z;
|
||||
}
|
||||
p.w = 1.0f;
|
||||
dp = (dp + fmin(0.0f, dot(p, frustumPlanes.arr[i])));
|
||||
}
|
||||
if ((dp >= 0.0f)) {
|
||||
uint tileId = uint(as_type<int>((as_type<uint>(x_1) + as_type<uint>(as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_COUNT_X)))))));
|
||||
if (((tileId < 0u) || (tileId >= config.numTiles))) {
|
||||
continue;
|
||||
}
|
||||
uint offset = atomic_fetch_add_explicit(&(tileLightId.data.arr[tileId].count), 1u, memory_order_relaxed);
|
||||
if ((offset >= config.numTileLightSlot)) {
|
||||
continue;
|
||||
}
|
||||
tileLightId.data.arr[tileId].lightId.arr[offset] = GlobalInvocationID.x;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Config& config [[buffer(0)]], constant Uniforms& uniforms [[buffer(1)]], device LightsBuffer& lightsBuffer [[buffer(2)]], device Tiles& tileLightId [[buffer(3)]]) {
|
||||
tint_symbol_inner(config, uniforms, lightsBuffer, tileLightId, GlobalInvocationID);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,617 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 417
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%60 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %GlobalInvocationID_1
|
||||
OpExecutionMode %main LocalSize 64 1 1
|
||||
OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
|
||||
OpName %LightsBuffer "LightsBuffer"
|
||||
OpMemberName %LightsBuffer 0 "lights"
|
||||
OpName %LightData "LightData"
|
||||
OpMemberName %LightData 0 "position"
|
||||
OpMemberName %LightData 1 "color"
|
||||
OpMemberName %LightData 2 "radius"
|
||||
OpName %lightsBuffer "lightsBuffer"
|
||||
OpName %Tiles "Tiles"
|
||||
OpMemberName %Tiles 0 "data"
|
||||
OpName %TileLightIdData "TileLightIdData"
|
||||
OpMemberName %TileLightIdData 0 "count"
|
||||
OpMemberName %TileLightIdData 1 "lightId"
|
||||
OpName %tileLightId "tileLightId"
|
||||
OpName %Config "Config"
|
||||
OpMemberName %Config 0 "numLights"
|
||||
OpMemberName %Config 1 "numTiles"
|
||||
OpMemberName %Config 2 "tileCountX"
|
||||
OpMemberName %Config 3 "tileCountY"
|
||||
OpMemberName %Config 4 "numTileLightSlot"
|
||||
OpMemberName %Config 5 "tileSize"
|
||||
OpName %config "config"
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "min"
|
||||
OpMemberName %Uniforms 1 "max"
|
||||
OpMemberName %Uniforms 2 "viewMatrix"
|
||||
OpMemberName %Uniforms 3 "projectionMatrix"
|
||||
OpMemberName %Uniforms 4 "fullScreenSize"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main_inner "main_inner"
|
||||
OpName %GlobalInvocationID "GlobalInvocationID"
|
||||
OpName %index "index"
|
||||
OpName %M "M"
|
||||
OpName %viewNear "viewNear"
|
||||
OpName %viewFar "viewFar"
|
||||
OpName %lightPos "lightPos"
|
||||
OpName %lightRadius "lightRadius"
|
||||
OpName %boxMin "boxMin"
|
||||
OpName %boxMax "boxMax"
|
||||
OpName %frustumPlanes "frustumPlanes"
|
||||
OpName %y "y"
|
||||
OpName %x "x"
|
||||
OpName %tilePixel0Idx "tilePixel0Idx"
|
||||
OpName %floorCoord "floorCoord"
|
||||
OpName %ceilCoord "ceilCoord"
|
||||
OpName %viewFloorCoord "viewFloorCoord"
|
||||
OpName %viewCeilCoord "viewCeilCoord"
|
||||
OpName %dp "dp"
|
||||
OpName %i "i"
|
||||
OpName %p "p"
|
||||
OpName %tileId "tileId"
|
||||
OpName %offset "offset"
|
||||
OpName %main "main"
|
||||
OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
|
||||
OpDecorate %LightsBuffer Block
|
||||
OpMemberDecorate %LightsBuffer 0 Offset 0
|
||||
OpMemberDecorate %LightData 0 Offset 0
|
||||
OpMemberDecorate %LightData 1 Offset 16
|
||||
OpMemberDecorate %LightData 2 Offset 28
|
||||
OpDecorate %_runtimearr_LightData ArrayStride 32
|
||||
OpDecorate %lightsBuffer DescriptorSet 0
|
||||
OpDecorate %lightsBuffer Binding 0
|
||||
OpDecorate %Tiles Block
|
||||
OpMemberDecorate %Tiles 0 Offset 0
|
||||
OpMemberDecorate %TileLightIdData 0 Offset 0
|
||||
OpMemberDecorate %TileLightIdData 1 Offset 4
|
||||
OpDecorate %_arr_uint_uint_64 ArrayStride 4
|
||||
OpDecorate %_arr_TileLightIdData_uint_4 ArrayStride 260
|
||||
OpDecorate %tileLightId DescriptorSet 1
|
||||
OpDecorate %tileLightId Binding 0
|
||||
OpDecorate %Config Block
|
||||
OpMemberDecorate %Config 0 Offset 0
|
||||
OpMemberDecorate %Config 1 Offset 4
|
||||
OpMemberDecorate %Config 2 Offset 8
|
||||
OpMemberDecorate %Config 3 Offset 12
|
||||
OpMemberDecorate %Config 4 Offset 16
|
||||
OpMemberDecorate %Config 5 Offset 20
|
||||
OpDecorate %config NonWritable
|
||||
OpDecorate %config DescriptorSet 2
|
||||
OpDecorate %config Binding 0
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 16
|
||||
OpMemberDecorate %Uniforms 2 Offset 32
|
||||
OpMemberDecorate %Uniforms 2 ColMajor
|
||||
OpMemberDecorate %Uniforms 2 MatrixStride 16
|
||||
OpMemberDecorate %Uniforms 3 Offset 96
|
||||
OpMemberDecorate %Uniforms 3 ColMajor
|
||||
OpMemberDecorate %Uniforms 3 MatrixStride 16
|
||||
OpMemberDecorate %Uniforms 4 Offset 160
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 3
|
||||
OpDecorate %uniforms Binding 0
|
||||
OpDecorate %_arr_v4float_uint_6 ArrayStride 16
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%v3float = OpTypeVector %float 3
|
||||
%LightData = OpTypeStruct %v4float %v3float %float
|
||||
%_runtimearr_LightData = OpTypeRuntimeArray %LightData
|
||||
%LightsBuffer = OpTypeStruct %_runtimearr_LightData
|
||||
%_ptr_StorageBuffer_LightsBuffer = OpTypePointer StorageBuffer %LightsBuffer
|
||||
%lightsBuffer = OpVariable %_ptr_StorageBuffer_LightsBuffer StorageBuffer
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_uint_uint_64 = OpTypeArray %uint %uint_64
|
||||
%TileLightIdData = OpTypeStruct %uint %_arr_uint_uint_64
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%_arr_TileLightIdData_uint_4 = OpTypeArray %TileLightIdData %uint_4
|
||||
%Tiles = OpTypeStruct %_arr_TileLightIdData_uint_4
|
||||
%_ptr_StorageBuffer_Tiles = OpTypePointer StorageBuffer %Tiles
|
||||
%tileLightId = OpVariable %_ptr_StorageBuffer_Tiles StorageBuffer
|
||||
%Config = OpTypeStruct %uint %uint %uint %uint %uint %uint
|
||||
%_ptr_Uniform_Config = OpTypePointer Uniform %Config
|
||||
%config = OpVariable %_ptr_Uniform_Config Uniform
|
||||
%mat4v4float = OpTypeMatrix %v4float 4
|
||||
%Uniforms = OpTypeStruct %v4float %v4float %mat4v4float %mat4v4float %v4float
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%28 = OpTypeFunction %void %v3uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%36 = OpConstantNull %uint
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%bool = OpTypeBool
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
|
||||
%float_0_100000001 = OpConstant %float 0.100000001
|
||||
%float_0_00100000005 = OpConstant %float 0.00100000005
|
||||
%float_64 = OpConstant %float 64
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
|
||||
%_ptr_Function_mat4v4float = OpTypePointer Function %mat4v4float
|
||||
%87 = OpConstantNull %mat4v4float
|
||||
%int = OpTypeInt 32 1
|
||||
%int_3 = OpConstant %int 3
|
||||
%int_2 = OpConstant %int 2
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%float_n1 = OpConstant %float -1
|
||||
%101 = OpConstantNull %float
|
||||
%float_1 = OpConstant %float 1
|
||||
%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
|
||||
%_ptr_Function_v4float = OpTypePointer Function %v4float
|
||||
%117 = OpConstantNull %v4float
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%float_0 = OpConstant %float 0
|
||||
%uint_6 = OpConstant %uint 6
|
||||
%_arr_v4float_uint_6 = OpTypeArray %v4float %uint_6
|
||||
%_ptr_Function__arr_v4float_uint_6 = OpTypePointer Function %_arr_v4float_uint_6
|
||||
%156 = OpConstantNull %_arr_v4float_uint_6
|
||||
%int_4 = OpConstant %int 4
|
||||
%int_5 = OpConstant %int 5
|
||||
%int_16 = OpConstant %int 16
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%170 = OpConstantNull %int
|
||||
%v2int = OpTypeVector %int 2
|
||||
%_ptr_Function_v2int = OpTypePointer Function %v2int
|
||||
%198 = OpConstantNull %v2int
|
||||
%float_2 = OpConstant %float 2
|
||||
%v2float = OpTypeVector %float 2
|
||||
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
|
||||
%209 = OpConstantComposite %v2float %float_1 %float_1
|
||||
%_ptr_Function_v2float = OpTypePointer Function %v2float
|
||||
%213 = OpConstantNull %v2float
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
|
||||
%_ptr_StorageBuffer_uint_0 = OpTypePointer StorageBuffer %uint
|
||||
%412 = OpTypeFunction %void
|
||||
%main_inner = OpFunction %void None %28
|
||||
%GlobalInvocationID = OpFunctionParameter %v3uint
|
||||
%32 = OpLabel
|
||||
%index = OpVariable %_ptr_Function_uint Function %36
|
||||
%M = OpVariable %_ptr_Function_mat4v4float Function %87
|
||||
%viewNear = OpVariable %_ptr_Function_float Function %101
|
||||
%viewFar = OpVariable %_ptr_Function_float Function %101
|
||||
%lightPos = OpVariable %_ptr_Function_v4float Function %117
|
||||
%127 = OpVariable %_ptr_Function_v4float Function %117
|
||||
%lightRadius = OpVariable %_ptr_Function_float Function %101
|
||||
%boxMin = OpVariable %_ptr_Function_v4float Function %117
|
||||
%boxMax = OpVariable %_ptr_Function_v4float Function %117
|
||||
%frustumPlanes = OpVariable %_ptr_Function__arr_v4float_uint_6 Function %156
|
||||
%y = OpVariable %_ptr_Function_int Function %170
|
||||
%x = OpVariable %_ptr_Function_int Function %170
|
||||
%tilePixel0Idx = OpVariable %_ptr_Function_v2int Function %198
|
||||
%floorCoord = OpVariable %_ptr_Function_v2float Function %213
|
||||
%ceilCoord = OpVariable %_ptr_Function_v2float Function %213
|
||||
%viewFloorCoord = OpVariable %_ptr_Function_v2float Function %213
|
||||
%viewCeilCoord = OpVariable %_ptr_Function_v2float Function %213
|
||||
%dp = OpVariable %_ptr_Function_float Function %101
|
||||
%i = OpVariable %_ptr_Function_uint Function %36
|
||||
%p = OpVariable %_ptr_Function_v4float Function %117
|
||||
%tileId = OpVariable %_ptr_Function_uint Function %36
|
||||
%offset = OpVariable %_ptr_Function_uint Function %36
|
||||
%33 = OpCompositeExtract %uint %GlobalInvocationID 0
|
||||
OpStore %index %33
|
||||
%37 = OpLoad %uint %index
|
||||
%40 = OpAccessChain %_ptr_Uniform_uint %config %uint_0
|
||||
%41 = OpLoad %uint %40
|
||||
%42 = OpUGreaterThanEqual %bool %37 %41
|
||||
OpSelectionMerge %44 None
|
||||
OpBranchConditional %42 %45 %44
|
||||
%45 = OpLabel
|
||||
OpReturn
|
||||
%44 = OpLabel
|
||||
%46 = OpLoad %uint %index
|
||||
%49 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %46 %uint_0 %uint_1
|
||||
%50 = OpLoad %uint %index
|
||||
%51 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %50 %uint_0 %uint_1
|
||||
%52 = OpLoad %float %51
|
||||
%54 = OpFSub %float %52 %float_0_100000001
|
||||
%57 = OpLoad %uint %index
|
||||
%56 = OpConvertUToF %float %57
|
||||
%62 = OpLoad %uint %index
|
||||
%61 = OpConvertUToF %float %62
|
||||
%63 = OpFDiv %float %61 %float_64
|
||||
%59 = OpExtInst %float %60 Floor %63
|
||||
%64 = OpFMul %float %float_64 %59
|
||||
%65 = OpFSub %float %56 %64
|
||||
%66 = OpFMul %float %float_0_00100000005 %65
|
||||
%67 = OpFAdd %float %54 %66
|
||||
OpStore %49 %67
|
||||
%68 = OpLoad %uint %index
|
||||
%69 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %68 %uint_0 %uint_1
|
||||
%70 = OpLoad %float %69
|
||||
%72 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_1
|
||||
%73 = OpLoad %float %72
|
||||
%74 = OpFOrdLessThan %bool %70 %73
|
||||
OpSelectionMerge %75 None
|
||||
OpBranchConditional %74 %76 %75
|
||||
%76 = OpLabel
|
||||
%77 = OpLoad %uint %index
|
||||
%78 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %77 %uint_0 %uint_1
|
||||
%79 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_1 %uint_1
|
||||
%80 = OpLoad %float %79
|
||||
OpStore %78 %80
|
||||
OpBranch %75
|
||||
%75 = OpLabel
|
||||
%83 = OpAccessChain %_ptr_Uniform_mat4v4float %uniforms %uint_3
|
||||
%84 = OpLoad %mat4v4float %83
|
||||
OpStore %M %84
|
||||
%93 = OpAccessChain %_ptr_Function_float %M %int_3 %int_2
|
||||
%94 = OpLoad %float %93
|
||||
%88 = OpFNegate %float %94
|
||||
%96 = OpAccessChain %_ptr_Function_float %M %int_2 %int_2
|
||||
%97 = OpLoad %float %96
|
||||
%98 = OpFAdd %float %float_n1 %97
|
||||
%99 = OpFDiv %float %88 %98
|
||||
OpStore %viewNear %99
|
||||
%103 = OpAccessChain %_ptr_Function_float %M %int_3 %int_2
|
||||
%104 = OpLoad %float %103
|
||||
%102 = OpFNegate %float %104
|
||||
%106 = OpAccessChain %_ptr_Function_float %M %int_2 %int_2
|
||||
%107 = OpLoad %float %106
|
||||
%108 = OpFAdd %float %float_1 %107
|
||||
%109 = OpFDiv %float %102 %108
|
||||
OpStore %viewFar %109
|
||||
%111 = OpLoad %uint %index
|
||||
%113 = OpAccessChain %_ptr_StorageBuffer_v4float %lightsBuffer %uint_0 %111 %uint_0
|
||||
%114 = OpLoad %v4float %113
|
||||
OpStore %lightPos %114
|
||||
%119 = OpAccessChain %_ptr_Uniform_mat4v4float %uniforms %uint_2
|
||||
%120 = OpLoad %mat4v4float %119
|
||||
%121 = OpLoad %v4float %lightPos
|
||||
%122 = OpMatrixTimesVector %v4float %120 %121
|
||||
OpStore %lightPos %122
|
||||
%123 = OpLoad %v4float %lightPos
|
||||
%124 = OpAccessChain %_ptr_Function_float %lightPos %uint_3
|
||||
%125 = OpLoad %float %124
|
||||
%128 = OpCompositeConstruct %v4float %125 %125 %125 %125
|
||||
%126 = OpFDiv %v4float %123 %128
|
||||
OpStore %lightPos %126
|
||||
%129 = OpLoad %uint %index
|
||||
%130 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %129 %uint_2
|
||||
%131 = OpLoad %float %130
|
||||
OpStore %lightRadius %131
|
||||
%133 = OpLoad %v4float %lightPos
|
||||
%134 = OpLoad %float %lightRadius
|
||||
%135 = OpCompositeConstruct %v3float %134 %134 %134
|
||||
%136 = OpCompositeExtract %float %135 0
|
||||
%137 = OpCompositeExtract %float %135 1
|
||||
%138 = OpCompositeExtract %float %135 2
|
||||
%140 = OpCompositeConstruct %v4float %136 %137 %138 %float_0
|
||||
%141 = OpFSub %v4float %133 %140
|
||||
OpStore %boxMin %141
|
||||
%143 = OpLoad %v4float %lightPos
|
||||
%144 = OpLoad %float %lightRadius
|
||||
%145 = OpCompositeConstruct %v3float %144 %144 %144
|
||||
%146 = OpCompositeExtract %float %145 0
|
||||
%147 = OpCompositeExtract %float %145 1
|
||||
%148 = OpCompositeExtract %float %145 2
|
||||
%149 = OpCompositeConstruct %v4float %146 %147 %148 %float_0
|
||||
%150 = OpFAdd %v4float %143 %149
|
||||
OpStore %boxMax %150
|
||||
%158 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_4
|
||||
%159 = OpLoad %float %viewNear
|
||||
%160 = OpCompositeConstruct %v4float %float_0 %float_0 %float_n1 %159
|
||||
OpStore %158 %160
|
||||
%162 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_5
|
||||
%164 = OpLoad %float %viewFar
|
||||
%163 = OpFNegate %float %164
|
||||
%165 = OpCompositeConstruct %v4float %float_0 %float_0 %float_1 %163
|
||||
OpStore %162 %165
|
||||
OpStore %y %int_0
|
||||
OpBranch %171
|
||||
%171 = OpLabel
|
||||
OpLoopMerge %172 %173 None
|
||||
OpBranch %174
|
||||
%174 = OpLabel
|
||||
%176 = OpLoad %int %y
|
||||
%177 = OpSLessThan %bool %176 %int_2
|
||||
%175 = OpLogicalNot %bool %177
|
||||
OpSelectionMerge %178 None
|
||||
OpBranchConditional %175 %179 %178
|
||||
%179 = OpLabel
|
||||
OpBranch %172
|
||||
%178 = OpLabel
|
||||
OpStore %x %int_0
|
||||
OpBranch %181
|
||||
%181 = OpLabel
|
||||
OpLoopMerge %182 %183 None
|
||||
OpBranch %184
|
||||
%184 = OpLabel
|
||||
%186 = OpLoad %int %x
|
||||
%187 = OpSLessThan %bool %186 %int_2
|
||||
%185 = OpLogicalNot %bool %187
|
||||
OpSelectionMerge %188 None
|
||||
OpBranchConditional %185 %189 %188
|
||||
%189 = OpLabel
|
||||
OpBranch %182
|
||||
%188 = OpLabel
|
||||
%191 = OpLoad %int %x
|
||||
%192 = OpIMul %int %191 %int_16
|
||||
%193 = OpLoad %int %y
|
||||
%194 = OpIMul %int %193 %int_16
|
||||
%195 = OpCompositeConstruct %v2int %192 %194
|
||||
OpStore %tilePixel0Idx %195
|
||||
%202 = OpLoad %v2int %tilePixel0Idx
|
||||
%200 = OpConvertSToF %v2float %202
|
||||
%203 = OpVectorTimesScalar %v2float %200 %float_2
|
||||
%205 = OpAccessChain %_ptr_Uniform_v4float %uniforms %uint_4
|
||||
%206 = OpLoad %v4float %205
|
||||
%207 = OpVectorShuffle %v2float %206 %206 0 1
|
||||
%208 = OpFDiv %v2float %203 %207
|
||||
%210 = OpFSub %v2float %208 %209
|
||||
OpStore %floorCoord %210
|
||||
%215 = OpLoad %v2int %tilePixel0Idx
|
||||
%216 = OpCompositeConstruct %v2int %int_16 %int_16
|
||||
%217 = OpIAdd %v2int %215 %216
|
||||
%214 = OpConvertSToF %v2float %217
|
||||
%218 = OpVectorTimesScalar %v2float %214 %float_2
|
||||
%219 = OpAccessChain %_ptr_Uniform_v4float %uniforms %uint_4
|
||||
%220 = OpLoad %v4float %219
|
||||
%221 = OpVectorShuffle %v2float %220 %220 0 1
|
||||
%222 = OpFDiv %v2float %218 %221
|
||||
%223 = OpFSub %v2float %222 %209
|
||||
OpStore %ceilCoord %223
|
||||
%226 = OpLoad %float %viewNear
|
||||
%225 = OpFNegate %float %226
|
||||
%227 = OpAccessChain %_ptr_Function_float %floorCoord %uint_0
|
||||
%228 = OpLoad %float %227
|
||||
%229 = OpFMul %float %225 %228
|
||||
%230 = OpAccessChain %_ptr_Function_float %M %int_2 %int_0
|
||||
%231 = OpLoad %float %230
|
||||
%232 = OpLoad %float %viewNear
|
||||
%233 = OpFMul %float %231 %232
|
||||
%234 = OpFSub %float %229 %233
|
||||
%235 = OpAccessChain %_ptr_Function_float %M %int_0 %int_0
|
||||
%236 = OpLoad %float %235
|
||||
%237 = OpFDiv %float %234 %236
|
||||
%239 = OpLoad %float %viewNear
|
||||
%238 = OpFNegate %float %239
|
||||
%240 = OpAccessChain %_ptr_Function_float %floorCoord %uint_1
|
||||
%241 = OpLoad %float %240
|
||||
%242 = OpFMul %float %238 %241
|
||||
%244 = OpAccessChain %_ptr_Function_float %M %int_2 %int_1
|
||||
%245 = OpLoad %float %244
|
||||
%246 = OpLoad %float %viewNear
|
||||
%247 = OpFMul %float %245 %246
|
||||
%248 = OpFSub %float %242 %247
|
||||
%249 = OpAccessChain %_ptr_Function_float %M %int_1 %int_1
|
||||
%250 = OpLoad %float %249
|
||||
%251 = OpFDiv %float %248 %250
|
||||
%252 = OpCompositeConstruct %v2float %237 %251
|
||||
OpStore %viewFloorCoord %252
|
||||
%255 = OpLoad %float %viewNear
|
||||
%254 = OpFNegate %float %255
|
||||
%256 = OpAccessChain %_ptr_Function_float %ceilCoord %uint_0
|
||||
%257 = OpLoad %float %256
|
||||
%258 = OpFMul %float %254 %257
|
||||
%259 = OpAccessChain %_ptr_Function_float %M %int_2 %int_0
|
||||
%260 = OpLoad %float %259
|
||||
%261 = OpLoad %float %viewNear
|
||||
%262 = OpFMul %float %260 %261
|
||||
%263 = OpFSub %float %258 %262
|
||||
%264 = OpAccessChain %_ptr_Function_float %M %int_0 %int_0
|
||||
%265 = OpLoad %float %264
|
||||
%266 = OpFDiv %float %263 %265
|
||||
%268 = OpLoad %float %viewNear
|
||||
%267 = OpFNegate %float %268
|
||||
%269 = OpAccessChain %_ptr_Function_float %ceilCoord %uint_1
|
||||
%270 = OpLoad %float %269
|
||||
%271 = OpFMul %float %267 %270
|
||||
%272 = OpAccessChain %_ptr_Function_float %M %int_2 %int_1
|
||||
%273 = OpLoad %float %272
|
||||
%274 = OpLoad %float %viewNear
|
||||
%275 = OpFMul %float %273 %274
|
||||
%276 = OpFSub %float %271 %275
|
||||
%277 = OpAccessChain %_ptr_Function_float %M %int_1 %int_1
|
||||
%278 = OpLoad %float %277
|
||||
%279 = OpFDiv %float %276 %278
|
||||
%280 = OpCompositeConstruct %v2float %266 %279
|
||||
OpStore %viewCeilCoord %280
|
||||
%282 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_0
|
||||
%284 = OpAccessChain %_ptr_Function_float %viewFloorCoord %uint_0
|
||||
%285 = OpLoad %float %284
|
||||
%283 = OpFNegate %float %285
|
||||
%286 = OpLoad %float %viewNear
|
||||
%287 = OpFDiv %float %283 %286
|
||||
%288 = OpCompositeConstruct %v4float %float_1 %float_0 %287 %float_0
|
||||
OpStore %282 %288
|
||||
%289 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_1
|
||||
%290 = OpAccessChain %_ptr_Function_float %viewCeilCoord %uint_0
|
||||
%291 = OpLoad %float %290
|
||||
%292 = OpLoad %float %viewNear
|
||||
%293 = OpFDiv %float %291 %292
|
||||
%294 = OpCompositeConstruct %v4float %float_n1 %float_0 %293 %float_0
|
||||
OpStore %289 %294
|
||||
%295 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_2
|
||||
%297 = OpAccessChain %_ptr_Function_float %viewFloorCoord %uint_1
|
||||
%298 = OpLoad %float %297
|
||||
%296 = OpFNegate %float %298
|
||||
%299 = OpLoad %float %viewNear
|
||||
%300 = OpFDiv %float %296 %299
|
||||
%301 = OpCompositeConstruct %v4float %float_0 %float_1 %300 %float_0
|
||||
OpStore %295 %301
|
||||
%302 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_3
|
||||
%303 = OpAccessChain %_ptr_Function_float %viewCeilCoord %uint_1
|
||||
%304 = OpLoad %float %303
|
||||
%305 = OpLoad %float %viewNear
|
||||
%306 = OpFDiv %float %304 %305
|
||||
%307 = OpCompositeConstruct %v4float %float_0 %float_n1 %306 %float_0
|
||||
OpStore %302 %307
|
||||
OpStore %dp %float_0
|
||||
OpStore %i %uint_0
|
||||
OpBranch %310
|
||||
%310 = OpLabel
|
||||
OpLoopMerge %311 %312 None
|
||||
OpBranch %313
|
||||
%313 = OpLabel
|
||||
%315 = OpLoad %uint %i
|
||||
%316 = OpULessThan %bool %315 %uint_6
|
||||
%314 = OpLogicalNot %bool %316
|
||||
OpSelectionMerge %317 None
|
||||
OpBranchConditional %314 %318 %317
|
||||
%318 = OpLabel
|
||||
OpBranch %311
|
||||
%317 = OpLabel
|
||||
%320 = OpLoad %uint %i
|
||||
%321 = OpAccessChain %_ptr_Function_float %frustumPlanes %320 %uint_0
|
||||
%322 = OpLoad %float %321
|
||||
%323 = OpFOrdGreaterThan %bool %322 %float_0
|
||||
OpSelectionMerge %324 None
|
||||
OpBranchConditional %323 %325 %326
|
||||
%325 = OpLabel
|
||||
%327 = OpAccessChain %_ptr_Function_float %p %uint_0
|
||||
%328 = OpAccessChain %_ptr_Function_float %boxMax %uint_0
|
||||
%329 = OpLoad %float %328
|
||||
OpStore %327 %329
|
||||
OpBranch %324
|
||||
%326 = OpLabel
|
||||
%330 = OpAccessChain %_ptr_Function_float %p %uint_0
|
||||
%331 = OpAccessChain %_ptr_Function_float %boxMin %uint_0
|
||||
%332 = OpLoad %float %331
|
||||
OpStore %330 %332
|
||||
OpBranch %324
|
||||
%324 = OpLabel
|
||||
%333 = OpLoad %uint %i
|
||||
%334 = OpAccessChain %_ptr_Function_float %frustumPlanes %333 %uint_1
|
||||
%335 = OpLoad %float %334
|
||||
%336 = OpFOrdGreaterThan %bool %335 %float_0
|
||||
OpSelectionMerge %337 None
|
||||
OpBranchConditional %336 %338 %339
|
||||
%338 = OpLabel
|
||||
%340 = OpAccessChain %_ptr_Function_float %p %uint_1
|
||||
%341 = OpAccessChain %_ptr_Function_float %boxMax %uint_1
|
||||
%342 = OpLoad %float %341
|
||||
OpStore %340 %342
|
||||
OpBranch %337
|
||||
%339 = OpLabel
|
||||
%343 = OpAccessChain %_ptr_Function_float %p %uint_1
|
||||
%344 = OpAccessChain %_ptr_Function_float %boxMin %uint_1
|
||||
%345 = OpLoad %float %344
|
||||
OpStore %343 %345
|
||||
OpBranch %337
|
||||
%337 = OpLabel
|
||||
%346 = OpLoad %uint %i
|
||||
%347 = OpAccessChain %_ptr_Function_float %frustumPlanes %346 %uint_2
|
||||
%348 = OpLoad %float %347
|
||||
%349 = OpFOrdGreaterThan %bool %348 %float_0
|
||||
OpSelectionMerge %350 None
|
||||
OpBranchConditional %349 %351 %352
|
||||
%351 = OpLabel
|
||||
%353 = OpAccessChain %_ptr_Function_float %p %uint_2
|
||||
%354 = OpAccessChain %_ptr_Function_float %boxMax %uint_2
|
||||
%355 = OpLoad %float %354
|
||||
OpStore %353 %355
|
||||
OpBranch %350
|
||||
%352 = OpLabel
|
||||
%356 = OpAccessChain %_ptr_Function_float %p %uint_2
|
||||
%357 = OpAccessChain %_ptr_Function_float %boxMin %uint_2
|
||||
%358 = OpLoad %float %357
|
||||
OpStore %356 %358
|
||||
OpBranch %350
|
||||
%350 = OpLabel
|
||||
%359 = OpAccessChain %_ptr_Function_float %p %uint_3
|
||||
OpStore %359 %float_1
|
||||
%360 = OpLoad %float %dp
|
||||
%363 = OpLoad %v4float %p
|
||||
%364 = OpLoad %uint %i
|
||||
%365 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %364
|
||||
%366 = OpLoad %v4float %365
|
||||
%362 = OpDot %float %363 %366
|
||||
%361 = OpExtInst %float %60 NMin %float_0 %362
|
||||
%367 = OpFAdd %float %360 %361
|
||||
OpStore %dp %367
|
||||
OpBranch %312
|
||||
%312 = OpLabel
|
||||
%368 = OpLoad %uint %i
|
||||
%369 = OpIAdd %uint %368 %uint_1
|
||||
OpStore %i %369
|
||||
OpBranch %310
|
||||
%311 = OpLabel
|
||||
%370 = OpLoad %float %dp
|
||||
%371 = OpFOrdGreaterThanEqual %bool %370 %float_0
|
||||
OpSelectionMerge %372 None
|
||||
OpBranchConditional %371 %373 %372
|
||||
%373 = OpLabel
|
||||
%375 = OpLoad %int %x
|
||||
%376 = OpLoad %int %y
|
||||
%377 = OpIMul %int %376 %int_2
|
||||
%378 = OpIAdd %int %375 %377
|
||||
%374 = OpBitcast %uint %378
|
||||
OpStore %tileId %374
|
||||
%380 = OpLoad %uint %tileId
|
||||
%381 = OpULessThan %bool %380 %uint_0
|
||||
OpSelectionMerge %382 None
|
||||
OpBranchConditional %381 %382 %383
|
||||
%383 = OpLabel
|
||||
%384 = OpLoad %uint %tileId
|
||||
%385 = OpAccessChain %_ptr_Uniform_uint %config %uint_1
|
||||
%386 = OpLoad %uint %385
|
||||
%387 = OpUGreaterThanEqual %bool %384 %386
|
||||
OpBranch %382
|
||||
%382 = OpLabel
|
||||
%388 = OpPhi %bool %381 %373 %387 %383
|
||||
OpSelectionMerge %389 None
|
||||
OpBranchConditional %388 %390 %389
|
||||
%390 = OpLabel
|
||||
OpBranch %183
|
||||
%389 = OpLabel
|
||||
%393 = OpLoad %uint %tileId
|
||||
%395 = OpAccessChain %_ptr_StorageBuffer_uint %tileLightId %uint_0 %393 %uint_0
|
||||
%391 = OpAtomicIAdd %uint %395 %uint_1 %uint_0 %uint_1
|
||||
OpStore %offset %391
|
||||
%397 = OpLoad %uint %offset
|
||||
%398 = OpAccessChain %_ptr_Uniform_uint %config %uint_4
|
||||
%399 = OpLoad %uint %398
|
||||
%400 = OpUGreaterThanEqual %bool %397 %399
|
||||
OpSelectionMerge %401 None
|
||||
OpBranchConditional %400 %402 %401
|
||||
%402 = OpLabel
|
||||
OpBranch %183
|
||||
%401 = OpLabel
|
||||
%403 = OpLoad %uint %tileId
|
||||
%404 = OpLoad %uint %offset
|
||||
%406 = OpAccessChain %_ptr_StorageBuffer_uint_0 %tileLightId %uint_0 %403 %uint_1 %404
|
||||
%407 = OpCompositeExtract %uint %GlobalInvocationID 0
|
||||
OpStore %406 %407
|
||||
OpBranch %372
|
||||
%372 = OpLabel
|
||||
OpBranch %183
|
||||
%183 = OpLabel
|
||||
%408 = OpLoad %int %x
|
||||
%409 = OpIAdd %int %408 %int_1
|
||||
OpStore %x %409
|
||||
OpBranch %181
|
||||
%182 = OpLabel
|
||||
OpBranch %173
|
||||
%173 = OpLabel
|
||||
%410 = OpLoad %int %y
|
||||
%411 = OpIAdd %int %410 %int_1
|
||||
OpStore %y %411
|
||||
OpBranch %171
|
||||
%172 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %412
|
||||
%414 = OpLabel
|
||||
%416 = OpLoad %v3uint %GlobalInvocationID_1
|
||||
%415 = OpFunctionCall %void %main_inner %416
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,119 @@
|
|||
struct LightData {
|
||||
position : vec4<f32>;
|
||||
color : vec3<f32>;
|
||||
radius : f32;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct LightsBuffer {
|
||||
lights : array<LightData>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read_write> lightsBuffer : LightsBuffer;
|
||||
|
||||
struct TileLightIdData {
|
||||
count : atomic<u32>;
|
||||
lightId : array<u32, 64>;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct Tiles {
|
||||
data : array<TileLightIdData, 4>;
|
||||
};
|
||||
|
||||
[[group(1), binding(0)]] var<storage, read_write> tileLightId : Tiles;
|
||||
|
||||
[[block]]
|
||||
struct Config {
|
||||
numLights : u32;
|
||||
numTiles : u32;
|
||||
tileCountX : u32;
|
||||
tileCountY : u32;
|
||||
numTileLightSlot : u32;
|
||||
tileSize : u32;
|
||||
};
|
||||
|
||||
[[group(2), binding(0)]] var<uniform> config : Config;
|
||||
|
||||
[[block]]
|
||||
struct Uniforms {
|
||||
min : vec4<f32>;
|
||||
max : vec4<f32>;
|
||||
viewMatrix : mat4x4<f32>;
|
||||
projectionMatrix : mat4x4<f32>;
|
||||
fullScreenSize : vec4<f32>;
|
||||
};
|
||||
|
||||
[[group(3), binding(0)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(64, 1, 1)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
var index = GlobalInvocationID.x;
|
||||
if ((index >= config.numLights)) {
|
||||
return;
|
||||
}
|
||||
lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001) + (0.001 * (f32(index) - (64.0 * floor((f32(index) / 64.0))))));
|
||||
if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) {
|
||||
lightsBuffer.lights[index].position.y = uniforms.max.y;
|
||||
}
|
||||
var M : mat4x4<f32> = uniforms.projectionMatrix;
|
||||
var viewNear : f32 = (-(M[3][2]) / (-1.0 + M[2][2]));
|
||||
var viewFar : f32 = (-(M[3][2]) / (1.0 + M[2][2]));
|
||||
var lightPos = lightsBuffer.lights[index].position;
|
||||
lightPos = (uniforms.viewMatrix * lightPos);
|
||||
lightPos = (lightPos / lightPos.w);
|
||||
var lightRadius : f32 = lightsBuffer.lights[index].radius;
|
||||
var boxMin : vec4<f32> = (lightPos - vec4<f32>(vec3<f32>(lightRadius), 0.0));
|
||||
var boxMax : vec4<f32> = (lightPos + vec4<f32>(vec3<f32>(lightRadius), 0.0));
|
||||
var frustumPlanes : array<vec4<f32>, 6>;
|
||||
frustumPlanes[4] = vec4<f32>(0.0, 0.0, -1.0, viewNear);
|
||||
frustumPlanes[5] = vec4<f32>(0.0, 0.0, 1.0, -(viewFar));
|
||||
let TILE_SIZE : i32 = 16;
|
||||
let TILE_COUNT_X : i32 = 2;
|
||||
let TILE_COUNT_Y : i32 = 2;
|
||||
for(var y : i32 = 0; (y < TILE_COUNT_Y); y = (y + 1)) {
|
||||
for(var x : i32 = 0; (x < TILE_COUNT_X); x = (x + 1)) {
|
||||
var tilePixel0Idx : vec2<i32> = vec2<i32>((x * TILE_SIZE), (y * TILE_SIZE));
|
||||
var floorCoord : vec2<f32> = (((2.0 * vec2<f32>(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - vec2<f32>(1.0));
|
||||
var ceilCoord : vec2<f32> = (((2.0 * vec2<f32>((tilePixel0Idx + vec2<i32>(TILE_SIZE)))) / uniforms.fullScreenSize.xy) - vec2<f32>(1.0));
|
||||
var viewFloorCoord : vec2<f32> = vec2<f32>((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
var viewCeilCoord : vec2<f32> = vec2<f32>((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
frustumPlanes[0] = vec4<f32>(1.0, 0.0, (-(viewFloorCoord.x) / viewNear), 0.0);
|
||||
frustumPlanes[1] = vec4<f32>(-1.0, 0.0, (viewCeilCoord.x / viewNear), 0.0);
|
||||
frustumPlanes[2] = vec4<f32>(0.0, 1.0, (-(viewFloorCoord.y) / viewNear), 0.0);
|
||||
frustumPlanes[3] = vec4<f32>(0.0, -1.0, (viewCeilCoord.y / viewNear), 0.0);
|
||||
var dp : f32 = 0.0;
|
||||
for(var i : u32 = 0u; (i < 6u); i = (i + 1u)) {
|
||||
var p : vec4<f32>;
|
||||
if ((frustumPlanes[i].x > 0.0)) {
|
||||
p.x = boxMax.x;
|
||||
} else {
|
||||
p.x = boxMin.x;
|
||||
}
|
||||
if ((frustumPlanes[i].y > 0.0)) {
|
||||
p.y = boxMax.y;
|
||||
} else {
|
||||
p.y = boxMin.y;
|
||||
}
|
||||
if ((frustumPlanes[i].z > 0.0)) {
|
||||
p.z = boxMax.z;
|
||||
} else {
|
||||
p.z = boxMin.z;
|
||||
}
|
||||
p.w = 1.0;
|
||||
dp = (dp + min(0.0, dot(p, frustumPlanes[i])));
|
||||
}
|
||||
if ((dp >= 0.0)) {
|
||||
var tileId : u32 = u32((x + (y * TILE_COUNT_X)));
|
||||
if (((tileId < 0u) || (tileId >= config.numTiles))) {
|
||||
continue;
|
||||
}
|
||||
var offset : u32 = atomicAdd(&(tileLightId.data[tileId].count), 1u);
|
||||
if ((offset >= config.numTileLightSlot)) {
|
||||
continue;
|
||||
}
|
||||
tileLightId.data[tileId].lightId[offset] = GlobalInvocationID.x;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct Light {
|
||||
/* 0x0000 */ packed_float3 position;
|
||||
/* 0x000c */ int8_t tint_pad[4];
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct LeftOver {
|
||||
/* 0x0000 */ float time;
|
||||
/* 0x0004 */ uint padding;
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct lightingInfo {
|
||||
float3 diffuse;
|
||||
float3 specular;
|
||||
|
|
|
|||
|
|
@ -1,6 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ packed_float3 v;
|
||||
/* 0x000c */ uint i;
|
||||
|
|
|
|||
|
|
@ -0,0 +1,11 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x2<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
[[group(0), binding(0)]] var<uniform> data: S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = data.matrix * data.vector;
|
||||
}
|
||||
|
|
@ -0,0 +1,18 @@
|
|||
cbuffer cbuffer_data : register(b0, space0) {
|
||||
uint4 data[3];
|
||||
};
|
||||
|
||||
float3x2 tint_symbol_2(uint4 buffer[3], uint offset) {
|
||||
const uint scalar_offset = ((offset + 0u)) / 4;
|
||||
uint4 ubo_load = buffer[scalar_offset / 4];
|
||||
const uint scalar_offset_1 = ((offset + 8u)) / 4;
|
||||
uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
|
||||
const uint scalar_offset_2 = ((offset + 16u)) / 4;
|
||||
uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
|
||||
return float3x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
|
||||
}
|
||||
|
||||
void main() {
|
||||
const float2 x = mul(asfloat(data[2].xyz), tint_symbol_2(data, 0u));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ float3x2 tint_symbol;
|
||||
/* 0x0018 */ int8_t tint_pad[8];
|
||||
/* 0x0020 */ packed_float3 vector;
|
||||
/* 0x002c */ int8_t tint_pad_1[4];
|
||||
};
|
||||
|
||||
fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
|
||||
float2 const x = (data.tint_symbol * data.vector);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,45 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 22
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main"
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "matrix"
|
||||
OpMemberName %S 1 "vector"
|
||||
OpName %data "data"
|
||||
OpName %main "main"
|
||||
OpDecorate %S Block
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 0 ColMajor
|
||||
OpMemberDecorate %S 0 MatrixStride 8
|
||||
OpMemberDecorate %S 1 Offset 32
|
||||
OpDecorate %data NonWritable
|
||||
OpDecorate %data DescriptorSet 0
|
||||
OpDecorate %data Binding 0
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%mat3v2float = OpTypeMatrix %v2float 3
|
||||
%v3float = OpTypeVector %float 3
|
||||
%S = OpTypeStruct %mat3v2float %v3float
|
||||
%_ptr_Uniform_S = OpTypePointer Uniform %S
|
||||
%data = OpVariable %_ptr_Uniform_S Uniform
|
||||
%void = OpTypeVoid
|
||||
%8 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
|
||||
%main = OpFunction %void None %8
|
||||
%11 = OpLabel
|
||||
%15 = OpAccessChain %_ptr_Uniform_mat3v2float %data %uint_0
|
||||
%16 = OpLoad %mat3v2float %15
|
||||
%19 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
|
||||
%20 = OpLoad %v3float %19
|
||||
%21 = OpMatrixTimesVector %v2float %16 %20
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,12 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x2<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> data : S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = (data.matrix * data.vector);
|
||||
}
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
[[group(0), binding(0)]] var<uniform> data: S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = data.matrix * data.vector;
|
||||
}
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
cbuffer cbuffer_data : register(b0, space0) {
|
||||
uint4 data[4];
|
||||
};
|
||||
|
||||
float3x3 tint_symbol_2(uint4 buffer[4], uint offset) {
|
||||
const uint scalar_offset = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_1 = ((offset + 16u)) / 4;
|
||||
const uint scalar_offset_2 = ((offset + 32u)) / 4;
|
||||
return float3x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz));
|
||||
}
|
||||
|
||||
void main() {
|
||||
const float3 x = mul(asfloat(data[3].xyz), tint_symbol_2(data, 0u));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ float3x3 tint_symbol;
|
||||
/* 0x0030 */ packed_float3 vector;
|
||||
/* 0x003c */ int8_t tint_pad[4];
|
||||
};
|
||||
|
||||
fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
|
||||
float3 const x = (data.tint_symbol * data.vector);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,44 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 21
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main"
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "matrix"
|
||||
OpMemberName %S 1 "vector"
|
||||
OpName %data "data"
|
||||
OpName %main "main"
|
||||
OpDecorate %S Block
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 0 ColMajor
|
||||
OpMemberDecorate %S 0 MatrixStride 16
|
||||
OpMemberDecorate %S 1 Offset 48
|
||||
OpDecorate %data NonWritable
|
||||
OpDecorate %data DescriptorSet 0
|
||||
OpDecorate %data Binding 0
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%mat3v3float = OpTypeMatrix %v3float 3
|
||||
%S = OpTypeStruct %mat3v3float %v3float
|
||||
%_ptr_Uniform_S = OpTypePointer Uniform %S
|
||||
%data = OpVariable %_ptr_Uniform_S Uniform
|
||||
%void = OpTypeVoid
|
||||
%7 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_mat3v3float = OpTypePointer Uniform %mat3v3float
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
|
||||
%main = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%14 = OpAccessChain %_ptr_Uniform_mat3v3float %data %uint_0
|
||||
%15 = OpLoad %mat3v3float %14
|
||||
%18 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
|
||||
%19 = OpLoad %v3float %18
|
||||
%20 = OpMatrixTimesVector %v3float %15 %19
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,12 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> data : S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = (data.matrix * data.vector);
|
||||
}
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
[[group(0), binding(0)]] var<uniform> data: S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = data.vector * data.matrix;
|
||||
}
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
cbuffer cbuffer_data : register(b0, space0) {
|
||||
uint4 data[4];
|
||||
};
|
||||
|
||||
float3x3 tint_symbol_3(uint4 buffer[4], uint offset) {
|
||||
const uint scalar_offset = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_1 = ((offset + 16u)) / 4;
|
||||
const uint scalar_offset_2 = ((offset + 32u)) / 4;
|
||||
return float3x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz));
|
||||
}
|
||||
|
||||
void main() {
|
||||
const float3 x = mul(tint_symbol_3(data, 0u), asfloat(data[3].xyz));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ float3x3 tint_symbol;
|
||||
/* 0x0030 */ packed_float3 vector;
|
||||
/* 0x003c */ int8_t tint_pad[4];
|
||||
};
|
||||
|
||||
fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
|
||||
float3 const x = (data.vector * data.tint_symbol);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,44 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 21
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main"
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "matrix"
|
||||
OpMemberName %S 1 "vector"
|
||||
OpName %data "data"
|
||||
OpName %main "main"
|
||||
OpDecorate %S Block
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 0 ColMajor
|
||||
OpMemberDecorate %S 0 MatrixStride 16
|
||||
OpMemberDecorate %S 1 Offset 48
|
||||
OpDecorate %data NonWritable
|
||||
OpDecorate %data DescriptorSet 0
|
||||
OpDecorate %data Binding 0
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%mat3v3float = OpTypeMatrix %v3float 3
|
||||
%S = OpTypeStruct %mat3v3float %v3float
|
||||
%_ptr_Uniform_S = OpTypePointer Uniform %S
|
||||
%data = OpVariable %_ptr_Uniform_S Uniform
|
||||
%void = OpTypeVoid
|
||||
%7 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_mat3v3float = OpTypePointer Uniform %mat3v3float
|
||||
%main = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%14 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
|
||||
%15 = OpLoad %v3float %14
|
||||
%18 = OpAccessChain %_ptr_Uniform_mat3v3float %data %uint_0
|
||||
%19 = OpLoad %mat3v3float %18
|
||||
%20 = OpVectorTimesMatrix %v3float %15 %19
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,12 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat3x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> data : S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = (data.vector * data.matrix);
|
||||
}
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat4x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
[[group(0), binding(0)]] var<uniform> data: S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = data.vector * data.matrix;
|
||||
}
|
||||
|
|
@ -0,0 +1,16 @@
|
|||
cbuffer cbuffer_data : register(b0, space0) {
|
||||
uint4 data[5];
|
||||
};
|
||||
|
||||
float4x3 tint_symbol_3(uint4 buffer[5], uint offset) {
|
||||
const uint scalar_offset = ((offset + 0u)) / 4;
|
||||
const uint scalar_offset_1 = ((offset + 16u)) / 4;
|
||||
const uint scalar_offset_2 = ((offset + 32u)) / 4;
|
||||
const uint scalar_offset_3 = ((offset + 48u)) / 4;
|
||||
return float4x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz), asfloat(buffer[scalar_offset_3 / 4].xyz));
|
||||
}
|
||||
|
||||
void main() {
|
||||
const float4 x = mul(tint_symbol_3(data, 0u), asfloat(data[4].xyz));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline auto 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 auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct S {
|
||||
/* 0x0000 */ float4x3 tint_symbol;
|
||||
/* 0x0040 */ packed_float3 vector;
|
||||
/* 0x004c */ int8_t tint_pad[4];
|
||||
};
|
||||
|
||||
fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
|
||||
float4 const x = (data.vector * data.tint_symbol);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,45 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 22
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main"
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpName %S "S"
|
||||
OpMemberName %S 0 "matrix"
|
||||
OpMemberName %S 1 "vector"
|
||||
OpName %data "data"
|
||||
OpName %main "main"
|
||||
OpDecorate %S Block
|
||||
OpMemberDecorate %S 0 Offset 0
|
||||
OpMemberDecorate %S 0 ColMajor
|
||||
OpMemberDecorate %S 0 MatrixStride 16
|
||||
OpMemberDecorate %S 1 Offset 64
|
||||
OpDecorate %data NonWritable
|
||||
OpDecorate %data DescriptorSet 0
|
||||
OpDecorate %data Binding 0
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%mat4v3float = OpTypeMatrix %v3float 4
|
||||
%S = OpTypeStruct %mat4v3float %v3float
|
||||
%_ptr_Uniform_S = OpTypePointer Uniform %S
|
||||
%data = OpVariable %_ptr_Uniform_S Uniform
|
||||
%void = OpTypeVoid
|
||||
%7 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_mat4v3float = OpTypePointer Uniform %mat4v3float
|
||||
%v4float = OpTypeVector %float 4
|
||||
%main = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%14 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
|
||||
%15 = OpLoad %v3float %14
|
||||
%18 = OpAccessChain %_ptr_Uniform_mat4v3float %data %uint_0
|
||||
%19 = OpLoad %mat4v3float %18
|
||||
%20 = OpVectorTimesMatrix %v4float %15 %19
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,12 @@
|
|||
[[block]]
|
||||
struct S {
|
||||
matrix : mat4x3<f32>;
|
||||
vector : vec3<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> data : S;
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn main() {
|
||||
let x = (data.vector * data.matrix);
|
||||
}
|
||||
Loading…
Reference in New Issue