mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-10 05:57:51 +00:00
tint: Deprecated module-scope 'let' for 'const'
Enable the parsing of 'const'. Warn on use of module-scope 'let', and automatically replace with 'const'. Fixed: tint:1580 Change-Id: I214aabca80686dc6b60ae21a7a57fbfb4898ea83 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/93786 Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
committed by
Dawn LUCI CQ
parent
03f88e6f49
commit
c64ca23d94
@@ -1,2 +1,2 @@
|
||||
let
|
||||
const
|
||||
H=1;
|
||||
|
||||
@@ -4,4 +4,3 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
const int H = 1;
|
||||
|
||||
@@ -3,4 +3,3 @@ void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
static const int H = 1;
|
||||
|
||||
@@ -1,5 +1,3 @@
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
constant int H = 1;
|
||||
|
||||
|
||||
@@ -1,19 +1,16 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 7
|
||||
; Bound: 5
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
|
||||
OpExecutionMode %unused_entry_point LocalSize 1 1 1
|
||||
OpName %H "H"
|
||||
OpName %unused_entry_point "unused_entry_point"
|
||||
%int = OpTypeInt 32 1
|
||||
%H = OpConstant %int 1
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%unused_entry_point = OpFunction %void None %3
|
||||
%6 = OpLabel
|
||||
%1 = OpTypeFunction %void
|
||||
%unused_entry_point = OpFunction %void None %1
|
||||
%4 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
||||
@@ -1 +1 @@
|
||||
let H = 1;
|
||||
const H = 1;
|
||||
|
||||
@@ -84,7 +84,7 @@ void tint_symbol_2(uvec3 GlobalInvocationID) {
|
||||
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = (x_1 + 1)) {
|
||||
ivec2 tilePixel0Idx = ivec2((x_1 * TILE_SIZE), (y_1 * TILE_SIZE));
|
||||
vec2 floorCoord = (((2.0f * vec2(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - vec2(1.0f));
|
||||
vec2 ceilCoord = (((2.0f * vec2((tilePixel0Idx + ivec2(16)))) / uniforms.fullScreenSize.xy) - vec2(1.0f));
|
||||
vec2 ceilCoord = (((2.0f * vec2((tilePixel0Idx + ivec2(TILE_SIZE)))) / uniforms.fullScreenSize.xy) - vec2(1.0f));
|
||||
vec2 viewFloorCoord = vec2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
vec2 viewCeilCoord = vec2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
frustumPlanes[0] = vec4(1.0f, 0.0f, (-(viewFloorCoord.x) / viewNear), 0.0f);
|
||||
|
||||
@@ -58,7 +58,7 @@ void main_inner(uint3 GlobalInvocationID) {
|
||||
[loop] 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) - (1.0f).xx);
|
||||
float2 ceilCoord = (((2.0f * float2((tilePixel0Idx + (16).xx))) / asfloat(uniforms[10]).xy) - (1.0f).xx);
|
||||
float2 ceilCoord = (((2.0f * float2((tilePixel0Idx + int2((TILE_SIZE).xx)))) / asfloat(uniforms[10]).xy) - (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);
|
||||
|
||||
@@ -88,7 +88,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const ti
|
||||
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)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
|
||||
float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(16)))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
|
||||
float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
|
||||
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);
|
||||
|
||||
@@ -175,7 +175,6 @@
|
||||
%207 = OpConstantComposite %v2float %float_1 %float_1
|
||||
%_ptr_Function_v2float = OpTypePointer Function %v2float
|
||||
%211 = OpConstantNull %v2float
|
||||
%214 = OpConstantComposite %v2int %int_16 %int_16
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
|
||||
%_ptr_StorageBuffer_uint_0 = OpTypePointer StorageBuffer %uint
|
||||
@@ -358,6 +357,7 @@
|
||||
%208 = OpFSub %v2float %206 %207
|
||||
OpStore %floorCoord %208
|
||||
%213 = OpLoad %v2int %tilePixel0Idx
|
||||
%214 = OpCompositeConstruct %v2int %int_16 %int_16
|
||||
%215 = OpIAdd %v2int %213 %214
|
||||
%212 = OpConvertSToF %v2float %215
|
||||
%216 = OpVectorTimesScalar %v2float %212 %float_2
|
||||
|
||||
@@ -6,7 +6,7 @@ void unused_entry_point() {
|
||||
}
|
||||
void f() {
|
||||
mat4 m = mat4(vec4(1.0f), vec4(1.0f), vec4(1.0f), vec4(1.0f));
|
||||
vec4 v1 = vec4(1.0f);
|
||||
float a = 1.0f;
|
||||
vec4 v1 = m[0];
|
||||
float a = v1[0];
|
||||
}
|
||||
|
||||
|
||||
@@ -5,6 +5,6 @@ void unused_entry_point() {
|
||||
|
||||
void f() {
|
||||
const float4x4 m = float4x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
|
||||
const float4 v1 = (1.0f).xxxx;
|
||||
const float a = 1.0f;
|
||||
const float4 v1 = m[0];
|
||||
const float a = v1[0];
|
||||
}
|
||||
|
||||
@@ -3,7 +3,7 @@
|
||||
using namespace metal;
|
||||
void f() {
|
||||
float4x4 const m = float4x4(float4(1.0f), float4(1.0f), float4(1.0f), float4(1.0f));
|
||||
float4 const v1 = float4(1.0f);
|
||||
float const a = 1.0f;
|
||||
float4 const v1 = m[0];
|
||||
float const a = v1[0];
|
||||
}
|
||||
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 13
|
||||
; Bound: 17
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
@@ -17,11 +17,15 @@
|
||||
%float_1 = OpConstant %float 1
|
||||
%11 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%12 = OpConstantComposite %mat4v4float %11 %11 %11 %11
|
||||
%int = OpTypeInt 32 1
|
||||
%14 = OpConstantNull %int
|
||||
%unused_entry_point = OpFunction %void None %1
|
||||
%4 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %1
|
||||
%6 = OpLabel
|
||||
%15 = OpCompositeExtract %v4float %12 0
|
||||
%16 = OpCompositeExtract %float %15 0
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
||||
@@ -2,7 +2,7 @@
|
||||
values : array<f32>,
|
||||
};
|
||||
|
||||
let width : u32 = 128u;
|
||||
const width : u32 = 128u;
|
||||
|
||||
@group(0) @binding(0) var tex : texture_depth_2d;
|
||||
@group(0) @binding(1) var<storage, read_write> result : Result;
|
||||
|
||||
@@ -1,12 +1,11 @@
|
||||
#version 310 es
|
||||
|
||||
const uint width = 128u;
|
||||
layout(binding = 1, std430) buffer Result_1 {
|
||||
float values[];
|
||||
} result;
|
||||
uniform highp sampler2D tex_1;
|
||||
void tint_symbol(uvec3 GlobalInvocationId) {
|
||||
result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = texelFetch(tex_1, ivec2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0).x;
|
||||
result.values[((GlobalInvocationId.y * 128u) + GlobalInvocationId.x)] = texelFetch(tex_1, ivec2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0).x;
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
@@ -1,5 +1,3 @@
|
||||
static const uint width = 128u;
|
||||
|
||||
Texture2D tex : register(t0, space0);
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
|
||||
@@ -8,7 +6,7 @@ struct tint_symbol_1 {
|
||||
};
|
||||
|
||||
void main_inner(uint3 GlobalInvocationId) {
|
||||
result.Store((4u * ((GlobalInvocationId.y * width) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
|
||||
result.Store((4u * ((GlobalInvocationId.y * 128u) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
|
||||
@@ -18,10 +18,8 @@ struct Result {
|
||||
/* 0x0000 */ tint_array<float, 1> values;
|
||||
};
|
||||
|
||||
constant uint width = 128u;
|
||||
|
||||
void tint_symbol_inner(uint3 GlobalInvocationId, device Result* const tint_symbol_1, depth2d<float, access::sample> tint_symbol_2) {
|
||||
(*(tint_symbol_1)).values[((GlobalInvocationId[1] * width) + GlobalInvocationId[0])] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId[0]), int(GlobalInvocationId[1]))), 0);
|
||||
(*(tint_symbol_1)).values[((GlobalInvocationId[1] * 128u) + GlobalInvocationId[0])] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId[0]), int(GlobalInvocationId[1]))), 0);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(device Result* tint_symbol_3 [[buffer(0)]], depth2d<float, access::sample> tint_symbol_4 [[texture(0)]], uint3 GlobalInvocationId [[thread_position_in_grid]]) {
|
||||
|
||||
@@ -8,7 +8,6 @@
|
||||
OpEntryPoint GLCompute %main "main" %GlobalInvocationId_1
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %GlobalInvocationId_1 "GlobalInvocationId_1"
|
||||
OpName %width "width"
|
||||
OpName %tex "tex"
|
||||
OpName %Result "Result"
|
||||
OpMemberName %Result 0 "values"
|
||||
@@ -28,33 +27,33 @@
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%GlobalInvocationId_1 = OpVariable %_ptr_Input_v3uint Input
|
||||
%width = OpConstant %uint 128
|
||||
%float = OpTypeFloat 32
|
||||
%8 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||
%_ptr_UniformConstant_8 = OpTypePointer UniformConstant %8
|
||||
%tex = OpVariable %_ptr_UniformConstant_8 UniformConstant
|
||||
%7 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||
%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
|
||||
%tex = OpVariable %_ptr_UniformConstant_7 UniformConstant
|
||||
%_runtimearr_float = OpTypeRuntimeArray %float
|
||||
%Result = OpTypeStruct %_runtimearr_float
|
||||
%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
|
||||
%result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
|
||||
%void = OpTypeVoid
|
||||
%14 = OpTypeFunction %void %v3uint
|
||||
%13 = OpTypeFunction %void %v3uint
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_128 = OpConstant %uint 128
|
||||
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
|
||||
%v4float = OpTypeVector %float 4
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%37 = OpConstantNull %int
|
||||
%38 = OpTypeFunction %void
|
||||
%main_inner = OpFunction %void None %14
|
||||
%main_inner = OpFunction %void None %13
|
||||
%GlobalInvocationId = OpFunctionParameter %v3uint
|
||||
%18 = OpLabel
|
||||
%20 = OpCompositeExtract %uint %GlobalInvocationId 1
|
||||
%21 = OpIMul %uint %20 %width
|
||||
%17 = OpLabel
|
||||
%19 = OpCompositeExtract %uint %GlobalInvocationId 1
|
||||
%21 = OpIMul %uint %19 %uint_128
|
||||
%22 = OpCompositeExtract %uint %GlobalInvocationId 0
|
||||
%23 = OpIAdd %uint %21 %22
|
||||
%25 = OpAccessChain %_ptr_StorageBuffer_float %result %uint_0 %23
|
||||
%29 = OpLoad %8 %tex
|
||||
%29 = OpLoad %7 %tex
|
||||
%33 = OpCompositeExtract %uint %GlobalInvocationId 0
|
||||
%32 = OpBitcast %int %33
|
||||
%35 = OpCompositeExtract %uint %GlobalInvocationId 1
|
||||
|
||||
@@ -2,7 +2,7 @@ struct Result {
|
||||
values : array<f32>,
|
||||
}
|
||||
|
||||
let width : u32 = 128u;
|
||||
const width : u32 = 128u;
|
||||
|
||||
@group(0) @binding(0) var tex : texture_depth_2d;
|
||||
|
||||
|
||||
@@ -38,11 +38,11 @@ fn mm_write(row : u32, col : u32, value : f32) {
|
||||
}
|
||||
}
|
||||
|
||||
let RowPerThread : u32 = 4u;
|
||||
let ColPerThread : u32 = 4u;
|
||||
let TileAOuter : u32 = 64u;
|
||||
let TileBOuter : u32 = 64u;
|
||||
let TileInner : u32 = 64u;
|
||||
const RowPerThread : u32 = 4u;
|
||||
const ColPerThread : u32 = 4u;
|
||||
const TileAOuter : u32 = 64u;
|
||||
const TileBOuter : u32 = 64u;
|
||||
const TileInner : u32 = 64u;
|
||||
var<workgroup> mm_Asub : array<array<f32, 64>, 64>;
|
||||
var<workgroup> mm_Bsub : array<array<f32, 64>, 64>;
|
||||
@compute @workgroup_size(16, 16, 1)
|
||||
|
||||
@@ -56,9 +56,6 @@ void mm_write(uint row, uint col, float value) {
|
||||
}
|
||||
}
|
||||
|
||||
const uint RowPerThread = 4u;
|
||||
const uint ColPerThread = 4u;
|
||||
const uint TileInner = 64u;
|
||||
shared float mm_Asub[64][64];
|
||||
shared float mm_Bsub[64][64];
|
||||
void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
||||
@@ -71,32 +68,32 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
}
|
||||
barrier();
|
||||
uint tileRow = (local_id.y * RowPerThread);
|
||||
uint tileCol = (local_id.x * ColPerThread);
|
||||
uint globalRow = (global_id.y * RowPerThread);
|
||||
uint globalCol = (global_id.x * ColPerThread);
|
||||
uint numTiles = (((uniforms.dimInner - 1u) / TileInner) + 1u);
|
||||
uint tileRow = (local_id.y * 4u);
|
||||
uint tileCol = (local_id.x * 4u);
|
||||
uint globalRow = (global_id.y * 4u);
|
||||
uint globalCol = (global_id.x * 4u);
|
||||
uint numTiles = (((uniforms.dimInner - 1u) / 64u) + 1u);
|
||||
float acc[16] = float[16](0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||
float ACached = 0.0f;
|
||||
float BCached[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
|
||||
{
|
||||
for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
}
|
||||
uint ColPerThreadA = (TileInner / 16u);
|
||||
uint ColPerThreadA = (64u / 16u);
|
||||
uint tileColA = (local_id.x * ColPerThreadA);
|
||||
uint RowPerThreadB = (TileInner / 16u);
|
||||
uint RowPerThreadB = (64u / 16u);
|
||||
uint tileRowB = (local_id.y * RowPerThreadB);
|
||||
{
|
||||
for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
uint inputRow = (tileRow + innerRow);
|
||||
uint inputCol = (tileColA + innerCol);
|
||||
float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
|
||||
float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
|
||||
mm_Asub[inputRow][inputCol] = tint_symbol_1;
|
||||
}
|
||||
}
|
||||
@@ -105,10 +102,10 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint inputRow = (tileRowB + innerRow);
|
||||
uint inputCol = (tileCol + innerCol);
|
||||
float tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
|
||||
float tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
|
||||
mm_Bsub[innerCol][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
@@ -116,18 +113,18 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
barrier();
|
||||
{
|
||||
for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
{
|
||||
for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
|
||||
}
|
||||
}
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
ACached = mm_Asub[(tileRow + innerRow)][k];
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * 4u) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
@@ -139,10 +136,10 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
}
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * 4u) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -40,12 +40,6 @@ void mm_write(uint row, uint col, float value) {
|
||||
}
|
||||
}
|
||||
|
||||
static const uint RowPerThread = 4u;
|
||||
static const uint ColPerThread = 4u;
|
||||
static const uint TileAOuter = 64u;
|
||||
static const uint TileBOuter = 64u;
|
||||
static const uint TileInner = 64u;
|
||||
|
||||
groupshared float mm_Asub[64][64];
|
||||
groupshared float mm_Bsub[64][64];
|
||||
|
||||
@@ -65,32 +59,32 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
const uint tileRow = (local_id.y * RowPerThread);
|
||||
const uint tileCol = (local_id.x * ColPerThread);
|
||||
const uint globalRow = (global_id.y * RowPerThread);
|
||||
const uint globalCol = (global_id.x * ColPerThread);
|
||||
const uint numTiles = (((uniforms[0].y - 1u) / TileInner) + 1u);
|
||||
const uint tileRow = (local_id.y * 4u);
|
||||
const uint tileCol = (local_id.x * 4u);
|
||||
const uint globalRow = (global_id.y * 4u);
|
||||
const uint globalCol = (global_id.x * 4u);
|
||||
const uint numTiles = (((uniforms[0].y - 1u) / 64u) + 1u);
|
||||
float acc[16] = (float[16])0;
|
||||
float ACached = 0.0f;
|
||||
float BCached[4] = (float[4])0;
|
||||
{
|
||||
[loop] for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
[loop] for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
}
|
||||
const uint ColPerThreadA = (TileInner / 16u);
|
||||
const uint ColPerThreadA = (64u / 16u);
|
||||
const uint tileColA = (local_id.x * ColPerThreadA);
|
||||
const uint RowPerThreadB = (TileInner / 16u);
|
||||
const uint RowPerThreadB = (64u / 16u);
|
||||
const uint tileRowB = (local_id.y * RowPerThreadB);
|
||||
{
|
||||
[loop] for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
const uint inputRow = (tileRow + innerRow);
|
||||
const uint inputCol = (tileColA + innerCol);
|
||||
const float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
|
||||
const float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
|
||||
mm_Asub[inputRow][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
@@ -99,10 +93,10 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
const uint inputRow = (tileRowB + innerRow);
|
||||
const uint inputCol = (tileCol + innerCol);
|
||||
const float tint_symbol_3 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
|
||||
const float tint_symbol_3 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
|
||||
mm_Bsub[innerCol][inputCol] = tint_symbol_3;
|
||||
}
|
||||
}
|
||||
@@ -110,18 +104,18 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
{
|
||||
[loop] for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
[loop] for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
{
|
||||
[loop] for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
[loop] for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
|
||||
}
|
||||
}
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
ACached = mm_Asub[(tileRow + innerRow)][k];
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * 4u) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
@@ -133,10 +127,10 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
||||
}
|
||||
}
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * 4u) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -47,16 +47,6 @@ void mm_write(uint row, uint col, float value, const constant Uniforms* const ti
|
||||
}
|
||||
}
|
||||
|
||||
constant uint RowPerThread = 4u;
|
||||
|
||||
constant uint ColPerThread = 4u;
|
||||
|
||||
constant uint TileAOuter = 64u;
|
||||
|
||||
constant uint TileBOuter = 64u;
|
||||
|
||||
constant uint TileInner = 64u;
|
||||
|
||||
void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_9, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_10, const constant Uniforms* const tint_symbol_11, const device Matrix* const tint_symbol_12, const device Matrix* const tint_symbol_13, device Matrix* const tint_symbol_14) {
|
||||
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
|
||||
uint const i = (idx / 64u);
|
||||
@@ -65,56 +55,56 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
|
||||
(*(tint_symbol_10))[i][i_1] = 0.0f;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
uint const tileRow = (local_id[1] * RowPerThread);
|
||||
uint const tileCol = (local_id[0] * ColPerThread);
|
||||
uint const globalRow = (global_id[1] * RowPerThread);
|
||||
uint const globalCol = (global_id[0] * ColPerThread);
|
||||
uint const numTiles = ((((*(tint_symbol_11)).dimInner - 1u) / TileInner) + 1u);
|
||||
uint const tileRow = (local_id[1] * 4u);
|
||||
uint const tileCol = (local_id[0] * 4u);
|
||||
uint const globalRow = (global_id[1] * 4u);
|
||||
uint const globalCol = (global_id[0] * 4u);
|
||||
uint const numTiles = ((((*(tint_symbol_11)).dimInner - 1u) / 64u) + 1u);
|
||||
tint_array<float, 16> acc = {};
|
||||
float ACached = 0.0f;
|
||||
tint_array<float, 4> BCached = {};
|
||||
for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
uint const ColPerThreadA = (TileInner / 16u);
|
||||
uint const ColPerThreadA = (64u / 16u);
|
||||
uint const tileColA = (local_id[0] * ColPerThreadA);
|
||||
uint const RowPerThreadB = (TileInner / 16u);
|
||||
uint const RowPerThreadB = (64u / 16u);
|
||||
uint const tileRowB = (local_id[1] * RowPerThreadB);
|
||||
for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
uint const inputRow = (tileRow + innerRow);
|
||||
uint const inputCol = (tileColA + innerCol);
|
||||
float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol), tint_symbol_11, tint_symbol_12);
|
||||
float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_11, tint_symbol_12);
|
||||
(*(tint_symbol_9))[inputRow][inputCol] = tint_symbol_1;
|
||||
}
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint const inputRow = (tileRowB + innerRow);
|
||||
uint const inputCol = (tileCol + innerCol);
|
||||
float const tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
|
||||
float const tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
|
||||
(*(tint_symbol_10))[innerCol][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
BCached[inner] = (*(tint_symbol_10))[k][(tileCol + inner)];
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
ACached = (*(tint_symbol_9))[(tileRow + innerRow)][k];
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * ColPerThread) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * 4u) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * ColPerThread) + innerCol);
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * 4u) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_11, tint_symbol_14);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -20,11 +20,6 @@
|
||||
OpMemberName %Uniforms 1 "dimInner"
|
||||
OpMemberName %Uniforms 2 "dimBOuter"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %RowPerThread "RowPerThread"
|
||||
OpName %RowPerThread "ColPerThread"
|
||||
OpName %TileAOuter "TileAOuter"
|
||||
OpName %TileAOuter "TileBOuter"
|
||||
OpName %TileAOuter "TileInner"
|
||||
OpName %mm_Asub "mm_Asub"
|
||||
OpName %mm_Bsub "mm_Bsub"
|
||||
OpName %mm_readA "mm_readA"
|
||||
@@ -80,10 +75,10 @@
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 0
|
||||
OpDecorate %uniforms Binding 3
|
||||
OpDecorate %_arr_float_TileAOuter ArrayStride 4
|
||||
OpDecorate %_arr__arr_float_TileAOuter_TileAOuter ArrayStride 256
|
||||
OpDecorate %_arr_float_uint_64 ArrayStride 4
|
||||
OpDecorate %_arr__arr_float_uint_64_uint_64 ArrayStride 256
|
||||
OpDecorate %_arr_float_uint_16 ArrayStride 4
|
||||
OpDecorate %_arr_float_RowPerThread ArrayStride 4
|
||||
OpDecorate %_arr_float_uint_4 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
@@ -101,199 +96,199 @@
|
||||
%Uniforms = OpTypeStruct %uint %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%RowPerThread = OpConstant %uint 4
|
||||
%TileAOuter = OpConstant %uint 64
|
||||
%_arr_float_TileAOuter = OpTypeArray %float %TileAOuter
|
||||
%_arr__arr_float_TileAOuter_TileAOuter = OpTypeArray %_arr_float_TileAOuter %TileAOuter
|
||||
%_ptr_Workgroup__arr__arr_float_TileAOuter_TileAOuter = OpTypePointer Workgroup %_arr__arr_float_TileAOuter_TileAOuter
|
||||
%mm_Asub = OpVariable %_ptr_Workgroup__arr__arr_float_TileAOuter_TileAOuter Workgroup
|
||||
%mm_Bsub = OpVariable %_ptr_Workgroup__arr__arr_float_TileAOuter_TileAOuter Workgroup
|
||||
%25 = OpTypeFunction %float %uint %uint
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_float_uint_64 = OpTypeArray %float %uint_64
|
||||
%_arr__arr_float_uint_64_uint_64 = OpTypeArray %_arr_float_uint_64 %uint_64
|
||||
%_ptr_Workgroup__arr__arr_float_uint_64_uint_64 = OpTypePointer Workgroup %_arr__arr_float_uint_64_uint_64
|
||||
%mm_Asub = OpVariable %_ptr_Workgroup__arr__arr_float_uint_64_uint_64 Workgroup
|
||||
%mm_Bsub = OpVariable %_ptr_Workgroup__arr__arr_float_uint_64_uint_64 Workgroup
|
||||
%24 = OpTypeFunction %float %uint %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
|
||||
%52 = OpConstantNull %float
|
||||
%51 = OpConstantNull %float
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%void = OpTypeVoid
|
||||
%75 = OpTypeFunction %void %uint %uint %float
|
||||
%98 = OpTypeFunction %void %v3uint %v3uint %uint
|
||||
%74 = OpTypeFunction %void %uint %uint %float
|
||||
%97 = OpTypeFunction %void %v3uint %v3uint %uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%106 = OpConstantNull %uint
|
||||
%105 = OpConstantNull %uint
|
||||
%uint_4096 = OpConstant %uint 4096
|
||||
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
|
||||
%uint_256 = OpConstant %uint 256
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%uint_16 = OpConstant %uint 16
|
||||
%_arr_float_uint_16 = OpTypeArray %float %uint_16
|
||||
%_ptr_Function__arr_float_uint_16 = OpTypePointer Function %_arr_float_uint_16
|
||||
%146 = OpConstantNull %_arr_float_uint_16
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%_arr_float_RowPerThread = OpTypeArray %float %RowPerThread
|
||||
%_ptr_Function__arr_float_RowPerThread = OpTypePointer Function %_arr_float_RowPerThread
|
||||
%152 = OpConstantNull %_arr_float_RowPerThread
|
||||
%_arr_float_uint_4 = OpTypeArray %float %uint_4
|
||||
%_ptr_Function__arr_float_uint_4 = OpTypePointer Function %_arr_float_uint_4
|
||||
%152 = OpConstantNull %_arr_float_uint_4
|
||||
%367 = OpTypeFunction %void
|
||||
%mm_readA = OpFunction %float None %25
|
||||
%mm_readA = OpFunction %float None %24
|
||||
%row = OpFunctionParameter %uint
|
||||
%col = OpFunctionParameter %uint
|
||||
%29 = OpLabel
|
||||
%32 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%33 = OpLoad %uint %32
|
||||
%34 = OpULessThan %bool %row %33
|
||||
OpSelectionMerge %36 None
|
||||
OpBranchConditional %34 %37 %36
|
||||
%37 = OpLabel
|
||||
%39 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%40 = OpLoad %uint %39
|
||||
%41 = OpULessThan %bool %col %40
|
||||
OpBranch %36
|
||||
%28 = OpLabel
|
||||
%31 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%32 = OpLoad %uint %31
|
||||
%33 = OpULessThan %bool %row %32
|
||||
OpSelectionMerge %35 None
|
||||
OpBranchConditional %33 %36 %35
|
||||
%36 = OpLabel
|
||||
%42 = OpPhi %bool %34 %29 %41 %37
|
||||
OpSelectionMerge %43 None
|
||||
OpBranchConditional %42 %44 %43
|
||||
%44 = OpLabel
|
||||
%45 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%46 = OpLoad %uint %45
|
||||
%47 = OpIMul %uint %row %46
|
||||
%48 = OpIAdd %uint %47 %col
|
||||
%50 = OpAccessChain %_ptr_StorageBuffer_float %firstMatrix %uint_0 %48
|
||||
%51 = OpLoad %float %50
|
||||
OpReturnValue %51
|
||||
%38 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%39 = OpLoad %uint %38
|
||||
%40 = OpULessThan %bool %col %39
|
||||
OpBranch %35
|
||||
%35 = OpLabel
|
||||
%41 = OpPhi %bool %33 %28 %40 %36
|
||||
OpSelectionMerge %42 None
|
||||
OpBranchConditional %41 %43 %42
|
||||
%43 = OpLabel
|
||||
OpReturnValue %52
|
||||
%44 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%45 = OpLoad %uint %44
|
||||
%46 = OpIMul %uint %row %45
|
||||
%47 = OpIAdd %uint %46 %col
|
||||
%49 = OpAccessChain %_ptr_StorageBuffer_float %firstMatrix %uint_0 %47
|
||||
%50 = OpLoad %float %49
|
||||
OpReturnValue %50
|
||||
%42 = OpLabel
|
||||
OpReturnValue %51
|
||||
OpFunctionEnd
|
||||
%mm_readB = OpFunction %float None %25
|
||||
%mm_readB = OpFunction %float None %24
|
||||
%row_0 = OpFunctionParameter %uint
|
||||
%col_0 = OpFunctionParameter %uint
|
||||
%56 = OpLabel
|
||||
%57 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%58 = OpLoad %uint %57
|
||||
%59 = OpULessThan %bool %row_0 %58
|
||||
OpSelectionMerge %60 None
|
||||
OpBranchConditional %59 %61 %60
|
||||
%61 = OpLabel
|
||||
%63 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%64 = OpLoad %uint %63
|
||||
%65 = OpULessThan %bool %col_0 %64
|
||||
OpBranch %60
|
||||
%55 = OpLabel
|
||||
%56 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%57 = OpLoad %uint %56
|
||||
%58 = OpULessThan %bool %row_0 %57
|
||||
OpSelectionMerge %59 None
|
||||
OpBranchConditional %58 %60 %59
|
||||
%60 = OpLabel
|
||||
%66 = OpPhi %bool %59 %56 %65 %61
|
||||
OpSelectionMerge %67 None
|
||||
OpBranchConditional %66 %68 %67
|
||||
%68 = OpLabel
|
||||
%69 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%70 = OpLoad %uint %69
|
||||
%71 = OpIMul %uint %row_0 %70
|
||||
%72 = OpIAdd %uint %71 %col_0
|
||||
%73 = OpAccessChain %_ptr_StorageBuffer_float %secondMatrix %uint_0 %72
|
||||
%74 = OpLoad %float %73
|
||||
OpReturnValue %74
|
||||
%62 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%63 = OpLoad %uint %62
|
||||
%64 = OpULessThan %bool %col_0 %63
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
%65 = OpPhi %bool %58 %55 %64 %60
|
||||
OpSelectionMerge %66 None
|
||||
OpBranchConditional %65 %67 %66
|
||||
%67 = OpLabel
|
||||
OpReturnValue %52
|
||||
%68 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%69 = OpLoad %uint %68
|
||||
%70 = OpIMul %uint %row_0 %69
|
||||
%71 = OpIAdd %uint %70 %col_0
|
||||
%72 = OpAccessChain %_ptr_StorageBuffer_float %secondMatrix %uint_0 %71
|
||||
%73 = OpLoad %float %72
|
||||
OpReturnValue %73
|
||||
%66 = OpLabel
|
||||
OpReturnValue %51
|
||||
OpFunctionEnd
|
||||
%mm_write = OpFunction %void None %75
|
||||
%mm_write = OpFunction %void None %74
|
||||
%row_1 = OpFunctionParameter %uint
|
||||
%col_1 = OpFunctionParameter %uint
|
||||
%value = OpFunctionParameter %float
|
||||
%81 = OpLabel
|
||||
%82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%83 = OpLoad %uint %82
|
||||
%84 = OpULessThan %bool %row_1 %83
|
||||
OpSelectionMerge %85 None
|
||||
OpBranchConditional %84 %86 %85
|
||||
%86 = OpLabel
|
||||
%87 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%88 = OpLoad %uint %87
|
||||
%89 = OpULessThan %bool %col_1 %88
|
||||
OpBranch %85
|
||||
%80 = OpLabel
|
||||
%81 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%82 = OpLoad %uint %81
|
||||
%83 = OpULessThan %bool %row_1 %82
|
||||
OpSelectionMerge %84 None
|
||||
OpBranchConditional %83 %85 %84
|
||||
%85 = OpLabel
|
||||
%90 = OpPhi %bool %84 %81 %89 %86
|
||||
OpSelectionMerge %91 None
|
||||
OpBranchConditional %90 %92 %91
|
||||
%92 = OpLabel
|
||||
%93 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%94 = OpLoad %uint %93
|
||||
%95 = OpIMul %uint %row_1 %94
|
||||
%96 = OpIAdd %uint %col_1 %95
|
||||
%97 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %96
|
||||
OpStore %97 %value
|
||||
OpBranch %91
|
||||
%86 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%87 = OpLoad %uint %86
|
||||
%88 = OpULessThan %bool %col_1 %87
|
||||
OpBranch %84
|
||||
%84 = OpLabel
|
||||
%89 = OpPhi %bool %83 %80 %88 %85
|
||||
OpSelectionMerge %90 None
|
||||
OpBranchConditional %89 %91 %90
|
||||
%91 = OpLabel
|
||||
%92 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2
|
||||
%93 = OpLoad %uint %92
|
||||
%94 = OpIMul %uint %row_1 %93
|
||||
%95 = OpIAdd %uint %col_1 %94
|
||||
%96 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %95
|
||||
OpStore %96 %value
|
||||
OpBranch %90
|
||||
%90 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main_inner = OpFunction %void None %98
|
||||
%main_inner = OpFunction %void None %97
|
||||
%local_id = OpFunctionParameter %v3uint
|
||||
%global_id = OpFunctionParameter %v3uint
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%103 = OpLabel
|
||||
%idx = OpVariable %_ptr_Function_uint Function %106
|
||||
%102 = OpLabel
|
||||
%idx = OpVariable %_ptr_Function_uint Function %105
|
||||
%acc = OpVariable %_ptr_Function__arr_float_uint_16 Function %146
|
||||
%ACached = OpVariable %_ptr_Function_float Function %52
|
||||
%BCached = OpVariable %_ptr_Function__arr_float_RowPerThread Function %152
|
||||
%index = OpVariable %_ptr_Function_uint Function %106
|
||||
%t = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerRow = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerCol = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerRow_0 = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerCol_0 = OpVariable %_ptr_Function_uint Function %106
|
||||
%k = OpVariable %_ptr_Function_uint Function %106
|
||||
%inner = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerRow_1 = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerCol_1 = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerRow_2 = OpVariable %_ptr_Function_uint Function %106
|
||||
%innerCol_2 = OpVariable %_ptr_Function_uint Function %106
|
||||
%ACached = OpVariable %_ptr_Function_float Function %51
|
||||
%BCached = OpVariable %_ptr_Function__arr_float_uint_4 Function %152
|
||||
%index = OpVariable %_ptr_Function_uint Function %105
|
||||
%t = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerRow = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerCol = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerRow_0 = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerCol_0 = OpVariable %_ptr_Function_uint Function %105
|
||||
%k = OpVariable %_ptr_Function_uint Function %105
|
||||
%inner = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerRow_1 = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerCol_1 = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerRow_2 = OpVariable %_ptr_Function_uint Function %105
|
||||
%innerCol_2 = OpVariable %_ptr_Function_uint Function %105
|
||||
OpStore %idx %local_invocation_index
|
||||
OpBranch %107
|
||||
%107 = OpLabel
|
||||
OpLoopMerge %108 %109 None
|
||||
OpBranch %110
|
||||
%110 = OpLabel
|
||||
%112 = OpLoad %uint %idx
|
||||
%114 = OpULessThan %bool %112 %uint_4096
|
||||
%111 = OpLogicalNot %bool %114
|
||||
OpSelectionMerge %115 None
|
||||
OpBranchConditional %111 %116 %115
|
||||
%116 = OpLabel
|
||||
OpBranch %108
|
||||
%115 = OpLabel
|
||||
%117 = OpLoad %uint %idx
|
||||
%118 = OpUDiv %uint %117 %TileAOuter
|
||||
%119 = OpLoad %uint %idx
|
||||
%120 = OpUMod %uint %119 %TileAOuter
|
||||
%122 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %118 %120
|
||||
OpStore %122 %52
|
||||
%123 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %118 %120
|
||||
OpStore %123 %52
|
||||
OpBranch %106
|
||||
%106 = OpLabel
|
||||
OpLoopMerge %107 %108 None
|
||||
OpBranch %109
|
||||
%109 = OpLabel
|
||||
%124 = OpLoad %uint %idx
|
||||
%126 = OpIAdd %uint %124 %uint_256
|
||||
OpStore %idx %126
|
||||
%111 = OpLoad %uint %idx
|
||||
%113 = OpULessThan %bool %111 %uint_4096
|
||||
%110 = OpLogicalNot %bool %113
|
||||
OpSelectionMerge %114 None
|
||||
OpBranchConditional %110 %115 %114
|
||||
%115 = OpLabel
|
||||
OpBranch %107
|
||||
%114 = OpLabel
|
||||
%116 = OpLoad %uint %idx
|
||||
%117 = OpUDiv %uint %116 %uint_64
|
||||
%118 = OpLoad %uint %idx
|
||||
%119 = OpUMod %uint %118 %uint_64
|
||||
%121 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %117 %119
|
||||
OpStore %121 %51
|
||||
%122 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %117 %119
|
||||
OpStore %122 %51
|
||||
OpBranch %108
|
||||
%108 = OpLabel
|
||||
%123 = OpLoad %uint %idx
|
||||
%125 = OpIAdd %uint %123 %uint_256
|
||||
OpStore %idx %125
|
||||
OpBranch %106
|
||||
%107 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%129 = OpCompositeExtract %uint %local_id 1
|
||||
%130 = OpIMul %uint %129 %RowPerThread
|
||||
%128 = OpCompositeExtract %uint %local_id 1
|
||||
%130 = OpIMul %uint %128 %uint_4
|
||||
%131 = OpCompositeExtract %uint %local_id 0
|
||||
%132 = OpIMul %uint %131 %RowPerThread
|
||||
%132 = OpIMul %uint %131 %uint_4
|
||||
%133 = OpCompositeExtract %uint %global_id 1
|
||||
%134 = OpIMul %uint %133 %RowPerThread
|
||||
%134 = OpIMul %uint %133 %uint_4
|
||||
%135 = OpCompositeExtract %uint %global_id 0
|
||||
%136 = OpIMul %uint %135 %RowPerThread
|
||||
%136 = OpIMul %uint %135 %uint_4
|
||||
%137 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%138 = OpLoad %uint %137
|
||||
%139 = OpISub %uint %138 %uint_1
|
||||
%140 = OpUDiv %uint %139 %TileAOuter
|
||||
%140 = OpUDiv %uint %139 %uint_64
|
||||
%141 = OpIAdd %uint %140 %uint_1
|
||||
OpStore %index %106
|
||||
OpStore %index %105
|
||||
OpBranch %154
|
||||
%154 = OpLabel
|
||||
OpLoopMerge %155 %156 None
|
||||
OpBranch %157
|
||||
%157 = OpLabel
|
||||
%159 = OpLoad %uint %index
|
||||
%160 = OpIMul %uint %RowPerThread %RowPerThread
|
||||
%160 = OpIMul %uint %uint_4 %uint_4
|
||||
%161 = OpULessThan %bool %159 %160
|
||||
%158 = OpLogicalNot %bool %161
|
||||
OpSelectionMerge %162 None
|
||||
@@ -303,7 +298,7 @@
|
||||
%162 = OpLabel
|
||||
%164 = OpLoad %uint %index
|
||||
%165 = OpAccessChain %_ptr_Function_float %acc %164
|
||||
OpStore %165 %52
|
||||
OpStore %165 %51
|
||||
OpBranch %156
|
||||
%156 = OpLabel
|
||||
%166 = OpLoad %uint %index
|
||||
@@ -311,13 +306,13 @@
|
||||
OpStore %index %167
|
||||
OpBranch %154
|
||||
%155 = OpLabel
|
||||
%168 = OpUDiv %uint %TileAOuter %uint_16
|
||||
%168 = OpUDiv %uint %uint_64 %uint_16
|
||||
%169 = OpCompositeExtract %uint %local_id 0
|
||||
%170 = OpIMul %uint %169 %168
|
||||
%171 = OpUDiv %uint %TileAOuter %uint_16
|
||||
%171 = OpUDiv %uint %uint_64 %uint_16
|
||||
%172 = OpCompositeExtract %uint %local_id 1
|
||||
%173 = OpIMul %uint %172 %171
|
||||
OpStore %t %106
|
||||
OpStore %t %105
|
||||
OpBranch %175
|
||||
%175 = OpLabel
|
||||
OpLoopMerge %176 %177 None
|
||||
@@ -331,21 +326,21 @@
|
||||
%183 = OpLabel
|
||||
OpBranch %176
|
||||
%182 = OpLabel
|
||||
OpStore %innerRow %106
|
||||
OpStore %innerRow %105
|
||||
OpBranch %185
|
||||
%185 = OpLabel
|
||||
OpLoopMerge %186 %187 None
|
||||
OpBranch %188
|
||||
%188 = OpLabel
|
||||
%190 = OpLoad %uint %innerRow
|
||||
%191 = OpULessThan %bool %190 %RowPerThread
|
||||
%191 = OpULessThan %bool %190 %uint_4
|
||||
%189 = OpLogicalNot %bool %191
|
||||
OpSelectionMerge %192 None
|
||||
OpBranchConditional %189 %193 %192
|
||||
%193 = OpLabel
|
||||
OpBranch %186
|
||||
%192 = OpLabel
|
||||
OpStore %innerCol %106
|
||||
OpStore %innerCol %105
|
||||
OpBranch %195
|
||||
%195 = OpLabel
|
||||
OpLoopMerge %196 %197 None
|
||||
@@ -366,7 +361,7 @@
|
||||
%209 = OpLoad %uint %innerRow
|
||||
%210 = OpIAdd %uint %134 %209
|
||||
%211 = OpLoad %uint %t
|
||||
%212 = OpIMul %uint %211 %TileAOuter
|
||||
%212 = OpIMul %uint %211 %uint_64
|
||||
%213 = OpIAdd %uint %212 %207
|
||||
%208 = OpFunctionCall %float %mm_readA %210 %213
|
||||
%214 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %205 %207
|
||||
@@ -385,7 +380,7 @@
|
||||
OpStore %innerRow %218
|
||||
OpBranch %185
|
||||
%186 = OpLabel
|
||||
OpStore %innerRow_0 %106
|
||||
OpStore %innerRow_0 %105
|
||||
OpBranch %220
|
||||
%220 = OpLabel
|
||||
OpLoopMerge %221 %222 None
|
||||
@@ -399,14 +394,14 @@
|
||||
%228 = OpLabel
|
||||
OpBranch %221
|
||||
%227 = OpLabel
|
||||
OpStore %innerCol_0 %106
|
||||
OpStore %innerCol_0 %105
|
||||
OpBranch %230
|
||||
%230 = OpLabel
|
||||
OpLoopMerge %231 %232 None
|
||||
OpBranch %233
|
||||
%233 = OpLabel
|
||||
%235 = OpLoad %uint %innerCol_0
|
||||
%236 = OpULessThan %bool %235 %RowPerThread
|
||||
%236 = OpULessThan %bool %235 %uint_4
|
||||
%234 = OpLogicalNot %bool %236
|
||||
OpSelectionMerge %237 None
|
||||
OpBranchConditional %234 %238 %237
|
||||
@@ -418,7 +413,7 @@
|
||||
%241 = OpLoad %uint %innerCol_0
|
||||
%242 = OpIAdd %uint %132 %241
|
||||
%244 = OpLoad %uint %t
|
||||
%245 = OpIMul %uint %244 %TileAOuter
|
||||
%245 = OpIMul %uint %244 %uint_64
|
||||
%246 = OpIAdd %uint %245 %240
|
||||
%247 = OpLoad %uint %innerCol_0
|
||||
%248 = OpIAdd %uint %136 %247
|
||||
@@ -441,28 +436,28 @@
|
||||
OpBranch %220
|
||||
%221 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
OpStore %k %106
|
||||
OpStore %k %105
|
||||
OpBranch %257
|
||||
%257 = OpLabel
|
||||
OpLoopMerge %258 %259 None
|
||||
OpBranch %260
|
||||
%260 = OpLabel
|
||||
%262 = OpLoad %uint %k
|
||||
%263 = OpULessThan %bool %262 %TileAOuter
|
||||
%263 = OpULessThan %bool %262 %uint_64
|
||||
%261 = OpLogicalNot %bool %263
|
||||
OpSelectionMerge %264 None
|
||||
OpBranchConditional %261 %265 %264
|
||||
%265 = OpLabel
|
||||
OpBranch %258
|
||||
%264 = OpLabel
|
||||
OpStore %inner %106
|
||||
OpStore %inner %105
|
||||
OpBranch %267
|
||||
%267 = OpLabel
|
||||
OpLoopMerge %268 %269 None
|
||||
OpBranch %270
|
||||
%270 = OpLabel
|
||||
%272 = OpLoad %uint %inner
|
||||
%273 = OpULessThan %bool %272 %RowPerThread
|
||||
%273 = OpULessThan %bool %272 %uint_4
|
||||
%271 = OpLogicalNot %bool %273
|
||||
OpSelectionMerge %274 None
|
||||
OpBranchConditional %271 %275 %274
|
||||
@@ -484,14 +479,14 @@
|
||||
OpStore %inner %284
|
||||
OpBranch %267
|
||||
%268 = OpLabel
|
||||
OpStore %innerRow_1 %106
|
||||
OpStore %innerRow_1 %105
|
||||
OpBranch %286
|
||||
%286 = OpLabel
|
||||
OpLoopMerge %287 %288 None
|
||||
OpBranch %289
|
||||
%289 = OpLabel
|
||||
%291 = OpLoad %uint %innerRow_1
|
||||
%292 = OpULessThan %bool %291 %RowPerThread
|
||||
%292 = OpULessThan %bool %291 %uint_4
|
||||
%290 = OpLogicalNot %bool %292
|
||||
OpSelectionMerge %293 None
|
||||
OpBranchConditional %290 %294 %293
|
||||
@@ -504,14 +499,14 @@
|
||||
%298 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %296 %297
|
||||
%299 = OpLoad %float %298
|
||||
OpStore %ACached %299
|
||||
OpStore %innerCol_1 %106
|
||||
OpStore %innerCol_1 %105
|
||||
OpBranch %301
|
||||
%301 = OpLabel
|
||||
OpLoopMerge %302 %303 None
|
||||
OpBranch %304
|
||||
%304 = OpLabel
|
||||
%306 = OpLoad %uint %innerCol_1
|
||||
%307 = OpULessThan %bool %306 %RowPerThread
|
||||
%307 = OpULessThan %bool %306 %uint_4
|
||||
%305 = OpLogicalNot %bool %307
|
||||
OpSelectionMerge %308 None
|
||||
OpBranchConditional %305 %309 %308
|
||||
@@ -519,7 +514,7 @@
|
||||
OpBranch %302
|
||||
%308 = OpLabel
|
||||
%310 = OpLoad %uint %innerRow_1
|
||||
%311 = OpIMul %uint %310 %RowPerThread
|
||||
%311 = OpIMul %uint %310 %uint_4
|
||||
%312 = OpLoad %uint %innerCol_1
|
||||
%313 = OpIAdd %uint %311 %312
|
||||
%314 = OpAccessChain %_ptr_Function_float %acc %313
|
||||
@@ -561,28 +556,28 @@
|
||||
OpStore %t %331
|
||||
OpBranch %175
|
||||
%176 = OpLabel
|
||||
OpStore %innerRow_2 %106
|
||||
OpStore %innerRow_2 %105
|
||||
OpBranch %333
|
||||
%333 = OpLabel
|
||||
OpLoopMerge %334 %335 None
|
||||
OpBranch %336
|
||||
%336 = OpLabel
|
||||
%338 = OpLoad %uint %innerRow_2
|
||||
%339 = OpULessThan %bool %338 %RowPerThread
|
||||
%339 = OpULessThan %bool %338 %uint_4
|
||||
%337 = OpLogicalNot %bool %339
|
||||
OpSelectionMerge %340 None
|
||||
OpBranchConditional %337 %341 %340
|
||||
%341 = OpLabel
|
||||
OpBranch %334
|
||||
%340 = OpLabel
|
||||
OpStore %innerCol_2 %106
|
||||
OpStore %innerCol_2 %105
|
||||
OpBranch %343
|
||||
%343 = OpLabel
|
||||
OpLoopMerge %344 %345 None
|
||||
OpBranch %346
|
||||
%346 = OpLabel
|
||||
%348 = OpLoad %uint %innerCol_2
|
||||
%349 = OpULessThan %bool %348 %RowPerThread
|
||||
%349 = OpULessThan %bool %348 %uint_4
|
||||
%347 = OpLogicalNot %bool %349
|
||||
OpSelectionMerge %350 None
|
||||
OpBranchConditional %347 %351 %350
|
||||
@@ -590,7 +585,7 @@
|
||||
OpBranch %344
|
||||
%350 = OpLabel
|
||||
%352 = OpLoad %uint %innerRow_2
|
||||
%353 = OpIMul %uint %352 %RowPerThread
|
||||
%353 = OpIMul %uint %352 %uint_4
|
||||
%354 = OpLoad %uint %innerCol_2
|
||||
%355 = OpIAdd %uint %353 %354
|
||||
%357 = OpLoad %uint %innerRow_2
|
||||
|
||||
@@ -39,15 +39,15 @@ fn mm_write(row : u32, col : u32, value : f32) {
|
||||
}
|
||||
}
|
||||
|
||||
let RowPerThread : u32 = 4u;
|
||||
const RowPerThread : u32 = 4u;
|
||||
|
||||
let ColPerThread : u32 = 4u;
|
||||
const ColPerThread : u32 = 4u;
|
||||
|
||||
let TileAOuter : u32 = 64u;
|
||||
const TileAOuter : u32 = 64u;
|
||||
|
||||
let TileBOuter : u32 = 64u;
|
||||
const TileBOuter : u32 = 64u;
|
||||
|
||||
let TileInner : u32 = 64u;
|
||||
const TileInner : u32 = 64u;
|
||||
|
||||
var<workgroup> mm_Asub : array<array<f32, 64>, 64>;
|
||||
|
||||
|
||||
Reference in New Issue
Block a user