mirror of
https://github.com/gfx-rs/wgpu.git
synced 2026-04-22 03:02:01 -04:00
[wgsl] test @workgroup_size attribute with constants
This commit is contained in:
teoxoy
@@ -1,7 +1,10 @@
|
||||
@group(0) @binding(0) var<storage, read_write> out: vec4<i32>;
|
||||
@group(0) @binding(1) var<storage, read_write> out2: i32;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
const TWO: u32 = 2u;
|
||||
const THREE: i32 = 3i;
|
||||
|
||||
@compute @workgroup_size(TWO, THREE, TWO - 1u)
|
||||
fn main() {
|
||||
swizzle_of_compose();
|
||||
index_of_compose();
|
||||
|
||||
@@ -3,8 +3,10 @@
|
||||
precision highp float;
|
||||
precision highp int;
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
layout(local_size_x = 2, local_size_y = 3, local_size_z = 1) in;
|
||||
|
||||
const uint TWO = 2u;
|
||||
const int THREE = 3;
|
||||
const int FOUR = 4;
|
||||
const int FOUR_ALIAS = 4;
|
||||
const int TEST_CONSTANT_ADDITION = 8;
|
||||
|
||||
@@ -1,3 +1,5 @@
|
||||
static const uint TWO = 2u;
|
||||
static const int THREE = 3;
|
||||
static const int FOUR = 4;
|
||||
static const int FOUR_ALIAS = 4;
|
||||
static const int TEST_CONSTANT_ADDITION = 8;
|
||||
@@ -58,7 +60,7 @@ void compose_of_constant()
|
||||
return;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
[numthreads(2, 3, 1)]
|
||||
void main()
|
||||
{
|
||||
swizzle_of_compose();
|
||||
|
||||
@@ -4,6 +4,8 @@
|
||||
|
||||
using metal::uint;
|
||||
|
||||
constant uint TWO = 2u;
|
||||
constant int THREE = 3;
|
||||
constant int FOUR = 4;
|
||||
constant int FOUR_ALIAS = 4;
|
||||
constant int TEST_CONSTANT_ADDITION = 8;
|
||||
|
||||
@@ -1,127 +1,128 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 82
|
||||
; Bound: 83
|
||||
OpCapability Shader
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %72 "main"
|
||||
OpExecutionMode %72 LocalSize 1 1 1
|
||||
OpDecorate %7 DescriptorSet 0
|
||||
OpDecorate %7 Binding 0
|
||||
OpDecorate %8 Block
|
||||
OpMemberDecorate %8 0 Offset 0
|
||||
OpEntryPoint GLCompute %73 "main"
|
||||
OpExecutionMode %73 LocalSize 2 3 1
|
||||
OpDecorate %10 DescriptorSet 0
|
||||
OpDecorate %10 Binding 1
|
||||
OpDecorate %10 Binding 0
|
||||
OpDecorate %11 Block
|
||||
OpMemberDecorate %11 0 Offset 0
|
||||
OpDecorate %13 DescriptorSet 0
|
||||
OpDecorate %13 Binding 1
|
||||
OpDecorate %14 Block
|
||||
OpMemberDecorate %14 0 Offset 0
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeInt 32 1
|
||||
%3 = OpTypeVector %4 4
|
||||
%5 = OpConstant %4 4
|
||||
%6 = OpConstant %4 8
|
||||
%8 = OpTypeStruct %3
|
||||
%9 = OpTypePointer StorageBuffer %8
|
||||
%7 = OpVariable %9 StorageBuffer
|
||||
%11 = OpTypeStruct %4
|
||||
%5 = OpTypeInt 32 0
|
||||
%6 = OpConstant %5 2
|
||||
%7 = OpConstant %4 3
|
||||
%8 = OpConstant %4 4
|
||||
%9 = OpConstant %4 8
|
||||
%11 = OpTypeStruct %3
|
||||
%12 = OpTypePointer StorageBuffer %11
|
||||
%10 = OpVariable %12 StorageBuffer
|
||||
%15 = OpTypeFunction %2
|
||||
%16 = OpTypePointer StorageBuffer %3
|
||||
%18 = OpTypeInt 32 0
|
||||
%17 = OpConstant %18 0
|
||||
%20 = OpConstant %4 1
|
||||
%21 = OpConstant %4 2
|
||||
%22 = OpConstant %4 3
|
||||
%23 = OpConstantComposite %3 %5 %22 %21 %20
|
||||
%27 = OpTypePointer StorageBuffer %4
|
||||
%35 = OpConstant %4 6
|
||||
%42 = OpConstant %4 30
|
||||
%43 = OpConstant %4 70
|
||||
%45 = OpTypePointer Function %4
|
||||
%47 = OpConstantNull %4
|
||||
%49 = OpConstantNull %4
|
||||
%64 = OpConstant %4 -4
|
||||
%65 = OpConstantComposite %3 %64 %64 %64 %64
|
||||
%14 = OpFunction %2 None %15
|
||||
%13 = OpLabel
|
||||
%19 = OpAccessChain %16 %7 %17
|
||||
OpBranch %24
|
||||
%24 = OpLabel
|
||||
OpStore %19 %23
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%26 = OpFunction %2 None %15
|
||||
%14 = OpTypeStruct %4
|
||||
%15 = OpTypePointer StorageBuffer %14
|
||||
%13 = OpVariable %15 StorageBuffer
|
||||
%18 = OpTypeFunction %2
|
||||
%19 = OpTypePointer StorageBuffer %3
|
||||
%20 = OpConstant %5 0
|
||||
%22 = OpConstant %4 1
|
||||
%23 = OpConstant %4 2
|
||||
%24 = OpConstantComposite %3 %8 %7 %23 %22
|
||||
%28 = OpTypePointer StorageBuffer %4
|
||||
%36 = OpConstant %4 6
|
||||
%43 = OpConstant %4 30
|
||||
%44 = OpConstant %4 70
|
||||
%46 = OpTypePointer Function %4
|
||||
%48 = OpConstantNull %4
|
||||
%50 = OpConstantNull %4
|
||||
%65 = OpConstant %4 -4
|
||||
%66 = OpConstantComposite %3 %65 %65 %65 %65
|
||||
%17 = OpFunction %2 None %18
|
||||
%16 = OpLabel
|
||||
%21 = OpAccessChain %19 %10 %20
|
||||
OpBranch %25
|
||||
%25 = OpLabel
|
||||
%28 = OpAccessChain %27 %10 %17
|
||||
OpBranch %29
|
||||
%29 = OpLabel
|
||||
%30 = OpLoad %4 %28
|
||||
%31 = OpIAdd %4 %30 %21
|
||||
OpStore %28 %31
|
||||
OpStore %21 %24
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%33 = OpFunction %2 None %15
|
||||
%32 = OpLabel
|
||||
%34 = OpAccessChain %27 %10 %17
|
||||
OpBranch %36
|
||||
%36 = OpLabel
|
||||
%37 = OpLoad %4 %34
|
||||
%38 = OpIAdd %4 %37 %35
|
||||
OpStore %34 %38
|
||||
%27 = OpFunction %2 None %18
|
||||
%26 = OpLabel
|
||||
%29 = OpAccessChain %28 %13 %20
|
||||
OpBranch %30
|
||||
%30 = OpLabel
|
||||
%31 = OpLoad %4 %29
|
||||
%32 = OpIAdd %4 %31 %23
|
||||
OpStore %29 %32
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%40 = OpFunction %2 None %15
|
||||
%39 = OpLabel
|
||||
%46 = OpVariable %45 Function %47
|
||||
%50 = OpVariable %45 Function %43
|
||||
%44 = OpVariable %45 Function %42
|
||||
%48 = OpVariable %45 Function %49
|
||||
%41 = OpAccessChain %16 %7 %17
|
||||
OpBranch %51
|
||||
%51 = OpLabel
|
||||
%52 = OpLoad %4 %44
|
||||
OpStore %46 %52
|
||||
%53 = OpLoad %4 %46
|
||||
OpStore %48 %53
|
||||
%54 = OpLoad %4 %44
|
||||
%55 = OpLoad %4 %46
|
||||
%56 = OpLoad %4 %48
|
||||
%57 = OpLoad %4 %50
|
||||
%58 = OpCompositeConstruct %3 %54 %55 %56 %57
|
||||
%59 = OpLoad %3 %41
|
||||
%60 = OpIAdd %3 %59 %58
|
||||
OpStore %41 %60
|
||||
%34 = OpFunction %2 None %18
|
||||
%33 = OpLabel
|
||||
%35 = OpAccessChain %28 %13 %20
|
||||
OpBranch %37
|
||||
%37 = OpLabel
|
||||
%38 = OpLoad %4 %35
|
||||
%39 = OpIAdd %4 %38 %36
|
||||
OpStore %35 %39
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%62 = OpFunction %2 None %15
|
||||
%61 = OpLabel
|
||||
%63 = OpAccessChain %16 %7 %17
|
||||
OpBranch %66
|
||||
%66 = OpLabel
|
||||
OpStore %63 %65
|
||||
%41 = OpFunction %2 None %18
|
||||
%40 = OpLabel
|
||||
%47 = OpVariable %46 Function %48
|
||||
%51 = OpVariable %46 Function %44
|
||||
%45 = OpVariable %46 Function %43
|
||||
%49 = OpVariable %46 Function %50
|
||||
%42 = OpAccessChain %19 %10 %20
|
||||
OpBranch %52
|
||||
%52 = OpLabel
|
||||
%53 = OpLoad %4 %45
|
||||
OpStore %47 %53
|
||||
%54 = OpLoad %4 %47
|
||||
OpStore %49 %54
|
||||
%55 = OpLoad %4 %45
|
||||
%56 = OpLoad %4 %47
|
||||
%57 = OpLoad %4 %49
|
||||
%58 = OpLoad %4 %51
|
||||
%59 = OpCompositeConstruct %3 %55 %56 %57 %58
|
||||
%60 = OpLoad %3 %42
|
||||
%61 = OpIAdd %3 %60 %59
|
||||
OpStore %42 %61
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%68 = OpFunction %2 None %15
|
||||
%63 = OpFunction %2 None %18
|
||||
%62 = OpLabel
|
||||
%64 = OpAccessChain %19 %10 %20
|
||||
OpBranch %67
|
||||
%67 = OpLabel
|
||||
%69 = OpAccessChain %16 %7 %17
|
||||
OpBranch %70
|
||||
%70 = OpLabel
|
||||
OpStore %69 %65
|
||||
OpStore %64 %66
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%72 = OpFunction %2 None %15
|
||||
%69 = OpFunction %2 None %18
|
||||
%68 = OpLabel
|
||||
%70 = OpAccessChain %19 %10 %20
|
||||
OpBranch %71
|
||||
%71 = OpLabel
|
||||
%73 = OpAccessChain %16 %7 %17
|
||||
%74 = OpAccessChain %27 %10 %17
|
||||
OpBranch %75
|
||||
%75 = OpLabel
|
||||
%76 = OpFunctionCall %2 %14
|
||||
%77 = OpFunctionCall %2 %26
|
||||
%78 = OpFunctionCall %2 %33
|
||||
%79 = OpFunctionCall %2 %40
|
||||
%80 = OpFunctionCall %2 %62
|
||||
%81 = OpFunctionCall %2 %68
|
||||
OpStore %70 %66
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%73 = OpFunction %2 None %18
|
||||
%72 = OpLabel
|
||||
%74 = OpAccessChain %19 %10 %20
|
||||
%75 = OpAccessChain %28 %13 %20
|
||||
OpBranch %76
|
||||
%76 = OpLabel
|
||||
%77 = OpFunctionCall %2 %17
|
||||
%78 = OpFunctionCall %2 %27
|
||||
%79 = OpFunctionCall %2 %34
|
||||
%80 = OpFunctionCall %2 %41
|
||||
%81 = OpFunctionCall %2 %63
|
||||
%82 = OpFunctionCall %2 %69
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -1,3 +1,5 @@
|
||||
const TWO: u32 = 2u;
|
||||
const THREE: i32 = 3;
|
||||
const FOUR: i32 = 4;
|
||||
const FOUR_ALIAS: i32 = 4;
|
||||
const TEST_CONSTANT_ADDITION: i32 = 8;
|
||||
@@ -54,7 +56,7 @@ fn compose_of_constant() {
|
||||
return;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1, 1, 1)
|
||||
@compute @workgroup_size(2, 3, 1)
|
||||
fn main() {
|
||||
swizzle_of_compose();
|
||||
index_of_compose();
|
||||
|
||||
Reference in New Issue
Block a user