kbz_8/Vulkan-CTS-bin
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

yes

Browse Source
This commit is contained in:
kbz_8
2026-05-06 23:44:13 +02:00
parent a5f787d80f
commit fd7038b853
1574 changed files with 365439 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
vulkan/amber/compute/atomic_barrier_sum_small.amber
+58
View File
@@ -0,0 +1,58 @@
#!amber
# Copyright 2024 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# https://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# The purpose of this test is to check the behavior of workgroup atomics in
# combination with barriers.
# This is known to cause issues on the mobile devices given the CTS failures
# for webgpu
# See: crbug.com/42241359
SHADER compute workgroup_shared_atomic_shader GLSL
#version 430
layout(local_size_x = 3, local_size_y = 1, local_size_z = 1) in;
layout(set = 0, binding = 0) buffer block1 {
uint ssbo_write[];
};
shared uint wg_shared;
void main() {
if( gl_LocalInvocationID.x == 0){
atomicExchange(wg_shared, 7);
}
barrier();
atomicAdd(wg_shared,1);
barrier();
if(gl_LocalInvocationID.x == 0) {
ssbo_write[0] = atomicExchange(wg_shared,0);
}
}
END
BUFFER buf_uint DATA_TYPE uint32 SIZE 16 FILL 123
PIPELINE compute pipeline
ATTACH workgroup_shared_atomic_shader
BIND BUFFER buf_uint AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT buf_uint IDX 0 EQ 10
vulkan/amber/compute/branch_past_barrier.amber
+98
View File
@@ -0,0 +1,98 @@
#!amber
# Test barriers insude control-flow which is sometimes uniform and sometimes non-uniform. This is
# valid provided the barrier is never executed when the control-flow is non-uniform. To do this we
# have 2 conditions, which generate 3 possible paths. The first workgroup chooses non-uniformly
# between the two paths which don't reach the barrier, while the second workgroup uniformly chooses
# the path with the barrier. The barrier path does some simple operations through shared memory in
# order to prevent compilers removing the barrier altogether.
#
# TODO: On implementations that use a subgroup size of 128 the barrier is still equivalent to a subgroup
# barrier, so the difference between uniform and non-uniform control won't make a difference. We should
# force the minimum subgroup size or something.
SHADER compute compute_shader GLSL
#version 310 es
layout(local_size_x = 128) in;
layout(set=0, binding = 0) buffer B {
uint a1[64];
uint a2[64];
uint x;
} b;
layout(set=0, binding=1) buffer C {
uint y[256];
} c;
shared uint s[128];
void main() {
uint l = b.x;
/* To save us having enormous buffers, force each workgroup to consume only 32 slots */
uint a_idx = (gl_LocalInvocationIndex & 31u) + 32u * gl_WorkGroupID.x;
if (b.a1[a_idx] == 0u) {
l = 2u;
} else {
if (b.a2[a_idx] == 0u) {
l = 3u;
} else {
/* In order for the barrier to be required, reverse the values of a2 through shared memory.
If the subgroup is as large as the workgroup here then the barrier can still be removed,
so the test should make sure that doesn't happen. */
s[gl_LocalInvocationIndex] = l + b.a2[a_idx];
barrier();
l = s[127u - gl_LocalInvocationIndex];
}
}
c.y[gl_GlobalInvocationID.x] = l;
}
END
BUFFER input_buf DATA_TYPE uint32 DATA
# First 32 values are used by WG 0 to select the non-uniform path. Second 32 are uniform for WG 1.
0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
0 0 0 0 2 0 0 0 7 0 3 0 9 0 1 0 0 0 5 0 0 0 0 0 8 0 2 0 7 0 0 0 1 1 1 1 2 1 1 1 7 1 3 1 9 1 1 1 1 1 5 1 1 1 1 1 8 1 2 1 7 1 1 1
# b.x at the end. A random number.
42
END
BUFFER res_buf DATA_TYPE uint32 DATA
# Initialise result to a value not stored in the test so that every lane must write the correct value to pass.
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
END
BUFFER reference_buffer DATA_TYPE uint32 DATA
# First workgroup hits either the l = 2 or l = 3 paths non-uniformly.
2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3
2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3
2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3
2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3 2 3
# Second workgroup has uniform control flow past the barrier. Result is 42 + b.a2, reversed through shared memory.
43 43 43 49 43 44 43 50 43 43 43 43 43 47 43 43 43 43 43 51 43 45 43 49 43 43 43 44 43 43 43 43
43 43 43 49 43 44 43 50 43 43 43 43 43 47 43 43 43 43 43 51 43 45 43 49 43 43 43 44 43 43 43 43
43 43 43 49 43 44 43 50 43 43 43 43 43 47 43 43 43 43 43 51 43 45 43 49 43 43 43 44 43 43 43 43
43 43 43 49 43 44 43 50 43 43 43 43 43 47 43 43 43 43 43 51 43 45 43 49 43 43 43 44 43 43 43 43
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER input_buf AS storage DESCRIPTOR_SET 0 BINDING 0
BIND BUFFER res_buf AS storage DESCRIPTOR_SET 0 BINDING 1
END
# Spawn 2 workgroups, one with non-uniform control flow, the other with a barrier.
RUN pipeline 2 1 1
EXPECT res_buf EQ_BUFFER reference_buffer
vulkan/amber/compute/pk-immediate.amber
+1269
View File
File diff suppressed because it is too large Load Diff
vulkan/amber/compute/pkadd-immediate.amber
+77
View File
@@ -0,0 +1,77 @@
#!amber
# Copyright (c) 2024 The Khronos Group Inc.
# Copyright (c) 2024 AMD
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
#
# Immediate/inline arguments to packed 16-bit operations might handled
# incorrectly in the compiler backend.
DEVICE_FEATURE Storage16BitFeatures.storageBuffer16BitAccess
DEVICE_FEATURE shaderInt16
SHADER compute compute_shader GLSL
#version 460
#extension GL_EXT_shader_16bit_storage : require
#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
struct Output {
i16vec2 nw;
i16vec2 n;
i16vec2 ne;
i16vec2 w;
i16vec2 e;
i16vec2 sw;
i16vec2 s;
i16vec2 se;
};
layout(set = 0, binding = 0) buffer Result {
Output arr[];
} result;
void main() {
ivec2 id32 = ivec2(gl_LocalInvocationID.xy);
i16vec2 id16 = i16vec2(gl_LocalInvocationID.xy);
int idx = id32.y * 8 + id32.x;
result.arr[idx].nw = id16 + i16vec2(-1s, 1s);
result.arr[idx].n = id16 + i16vec2( 0s, 1s);
result.arr[idx].ne = id16 + i16vec2( 1s, 1s);
result.arr[idx].w = id16 + i16vec2(-1s, 0s);
result.arr[idx].e = id16 + i16vec2( 1s, 0s);
result.arr[idx].sw = id16 + i16vec2(-1s, -1s);
result.arr[idx].s = id16 + i16vec2( 0s, -1s);
result.arr[idx].se = id16 + i16vec2( 1s, -1s);
}
END
BUFFER result DATA_TYPE int16 SIZE 1024 FILL 0
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
# Amber bug and/or misleading behavior: index must be given in bytes
EXPECT result IDX 0 EQ -1 1 0 1 1 1 -1 0 1 0 -1 -1 0 -1 1 -1
EXPECT result IDX 32 EQ 0 1 1 1 2 1 0 0 2 0 0 -1 1 -1 2 -1
EXPECT result IDX 64 EQ 1 1 2 1 3 1 1 0 3 0 1 -1 2 -1 3 -1
EXPECT result IDX 256 EQ -1 2 0 2 1 2 -1 1 1 1 -1 0 0 0 1 0
EXPECT result IDX 288 EQ 0 2 1 2 2 2 0 1 2 1 0 0 1 0 2 0
vulkan/amber/compute/vec2_nclamp_nan_component.amber
+118
View File
@@ -0,0 +1,118 @@
#!amber
# Copyright 2025 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# https://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
# The purpose of this test is to check the nan behavior of NClamp
# NClamp will be called on vec2(nan, 1.0) with the low=0.0 and high = 1.0
# The y component of the vec2 will be checked and should remain 1.0
#
# crbug.com/407109052
#
# fn compute() -> f32 {
# var k = 0.0;
# var x = vec2(k/0.0, 1.0);
# var p = clamp(x, vec2(0), vec2(1));
# return p.y;
# }
#
# @group(0) @binding(0) var<storage, read_write> _shader_output: array<f32>;
#
# @compute @workgroup_size(1) fn _computeSomething() {
# _shader_output[0] = compute();
# }
#
SHADER compute dawn_entry_point SPIRV-ASM
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 1
; Bound: 39
; Schema: 0
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%21 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %_computeSomething "main"
OpExecutionMode %_computeSomething LocalSize 1 1 1
OpMemberName %_shader_output_block_tint_explicit_layout 0 "inner"
OpName %_shader_output_block_tint_explicit_layout "_shader_output_block_tint_explicit_layout"
OpName %compute "compute"
OpName %k "k"
OpName %x "x"
OpName %p "p"
OpName %_computeSomething "_computeSomething"
OpDecorate %_runtimearr_float ArrayStride 4
OpMemberDecorate %_shader_output_block_tint_explicit_layout 0 Offset 0
OpDecorate %_shader_output_block_tint_explicit_layout Block
OpDecorate %1 DescriptorSet 0
OpDecorate %1 Binding 0
OpDecorate %1 Coherent
%float = OpTypeFloat 32
%_runtimearr_float = OpTypeRuntimeArray %float
%_shader_output_block_tint_explicit_layout = OpTypeStruct %_runtimearr_float
%_ptr_StorageBuffer__shader_output_block_tint_explicit_layout = OpTypePointer StorageBuffer %_shader_output_block_tint_explicit_layout
%1 = OpVariable %_ptr_StorageBuffer__shader_output_block_tint_explicit_layout StorageBuffer
%7 = OpTypeFunction %float
%_ptr_Function_float = OpTypePointer Function %float
%float_0 = OpConstant %float 0
%v2float = OpTypeVector %float 2
%float_1 = OpConstant %float 1
%_ptr_Function_v2float = OpTypePointer Function %v2float
%22 = OpConstantNull %v2float
%23 = OpConstantComposite %v2float %float_1 %float_1
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%31 = OpTypeFunction %void
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%compute = OpFunction %float None %7
%8 = OpLabel
%k = OpVariable %_ptr_Function_float Function
%x = OpVariable %_ptr_Function_v2float Function
%p = OpVariable %_ptr_Function_v2float Function
OpStore %k %float_0
%12 = OpLoad %float %k None
%13 = OpFDiv %float %12 %float_0
%15 = OpCompositeConstruct %v2float %13 %float_1
OpStore %x %15
%19 = OpLoad %v2float %x None
%20 = OpExtInst %v2float %21 NClamp %19 %22 %23
OpStore %p %20
%25 = OpAccessChain %_ptr_Function_float %p %uint_1
%28 = OpLoad %float %25 None
OpReturnValue %28
OpFunctionEnd
%_computeSomething = OpFunction %void None %31
%32 = OpLabel
%33 = OpAccessChain %_ptr_StorageBuffer_float %1 %uint_0 %int_0
%38 = OpFunctionCall %float %compute
OpStore %33 %38 None
OpReturn
OpFunctionEnd
END
BUFFER buf_float DATA_TYPE float SIZE 1000 FILL 777.0
PIPELINE compute pipeline
ATTACH dawn_entry_point
BIND BUFFER buf_float AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT buf_float IDX 0 EQ 1.0
vulkan/amber/compute/webgl_spirv_loop.amber
+89
View File
@@ -0,0 +1,89 @@
#!amber
SHADER compute compute_shader SPIRV-ASM
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 36
; Schema: 0
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %flow_block "flow_block"
OpMemberName %flow_block 0 "inner"
OpName %flow "flow"
OpName %main "main"
OpName %LOOP_COUNTER "LOOP_COUNTER"
OpDecorate %flow_block Block
OpMemberDecorate %flow_block 0 Offset 0
OpDecorate %_arr_uint_uint_2 ArrayStride 4
OpDecorate %flow DescriptorSet 0
OpDecorate %flow Binding 0
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
%flow_block = OpTypeStruct %_arr_uint_uint_2
%_ptr_StorageBuffer_flow_block = OpTypePointer StorageBuffer %flow_block
%flow = OpVariable %_ptr_StorageBuffer_flow_block StorageBuffer
%void = OpTypeVoid
%7 = OpTypeFunction %void
%11 = OpConstantNull %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%uint_1 = OpConstant %uint 1
%bool = OpTypeBool
%main = OpFunction %void None %7
%10 = OpLabel
%LOOP_COUNTER = OpVariable %_ptr_Function_uint Function %11
OpStore %LOOP_COUNTER %11
OpBranch %14
%14 = OpLabel
OpLoopMerge %15 %16 None
OpBranch %17
%17 = OpLabel
%20 = OpAccessChain %_ptr_StorageBuffer_uint %flow %uint_0 %11
%21 = OpAccessChain %_ptr_StorageBuffer_uint %flow %uint_0 %11
%22 = OpLoad %uint %21
%24 = OpIAdd %uint %22 %uint_1
OpStore %20 %24
%25 = OpLoad %uint %LOOP_COUNTER
%26 = OpUGreaterThanEqual %bool %25 %uint_1
OpSelectionMerge %28 None
OpBranchConditional %26 %29 %28
%29 = OpLabel
%30 = OpAccessChain %_ptr_StorageBuffer_uint %flow %uint_0 %uint_1
%31 = OpAccessChain %_ptr_StorageBuffer_uint %flow %uint_0 %uint_1
%32 = OpLoad %uint %31
%33 = OpIAdd %uint %32 %uint_1
OpStore %30 %33
OpBranch %15
%28 = OpLabel
%34 = OpLoad %uint %LOOP_COUNTER
%35 = OpIAdd %uint %34 %uint_1
OpStore %LOOP_COUNTER %35
OpBranch %16
%16 = OpLabel
OpBranch %14
%15 = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER buf0 DATA_TYPE uint32 DATA 0 0 END
BUFFER expected0 DATA_TYPE uint32 DATA 2 1 END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER buf0 AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
#EXPECT buf0 IDX 0 EQ 2
#EXPECT buf0 IDX 4 EQ 1
EXPECT buf0 EQ_BUFFER expected0
vulkan/amber/compute/workgroup_memory_explicit_layout/copy_memory_basic.amber
+124
View File
@@ -0,0 +1,124 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %input_buffer_0 %input_buffer_1 %output_buffer %wg %half_wg_0 %half_wg_1
OpExecutionMode %main LocalSize 1 1 1
OpDecorate %buffer_type Block
OpDecorate %half_buffer_type Block
OpDecorate %other_half_buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpDecorate %array ArrayStride 4
OpMemberDecorate %half_buffer_type 0 Offset 0
OpMemberDecorate %other_half_buffer_type 0 Offset 64
OpDecorate %half_array ArrayStride 4
OpDecorate %input_buffer_0 DescriptorSet 0
OpDecorate %input_buffer_0 Binding 0
OpDecorate %input_buffer_1 DescriptorSet 0
OpDecorate %input_buffer_1 Binding 1
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 2
OpDecorate %wg Aliased
OpDecorate %half_wg_0 Aliased
OpDecorate %half_wg_1 Aliased
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_16 = OpConstant %uint 16
%const_uint_32 = OpConstant %uint 32
%const_uint_64 = OpConstant %uint 64
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%half_array = OpTypeArray %uint %const_uint_16
%half_array_wg_ptr = OpTypePointer Workgroup %half_array
%half_array_sb_ptr = OpTypePointer StorageBuffer %half_array
%buffer_type = OpTypeStruct %array
%half_buffer_type = OpTypeStruct %half_array
%other_half_buffer_type = OpTypeStruct %half_array
%input_buffer_0_ptr = OpTypePointer StorageBuffer %half_buffer_type
%input_buffer_1_ptr = OpTypePointer StorageBuffer %half_buffer_type
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%half_wg_0_ptr = OpTypePointer Workgroup %half_buffer_type
%half_wg_1_ptr = OpTypePointer Workgroup %other_half_buffer_type
;;; Workgroup has a block covering the entire memory and another two
;;; blocks covering each half of the former.
%wg = OpVariable %wg_ptr Workgroup
%half_wg_0 = OpVariable %half_wg_0_ptr Workgroup
%half_wg_1 = OpVariable %half_wg_1_ptr Workgroup
%input_buffer_0 = OpVariable %input_buffer_0_ptr StorageBuffer
%input_buffer_1 = OpVariable %input_buffer_1_ptr StorageBuffer
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Copy first input into half of the workgroup memory. Because the
;;; whole types match here, we can afford to copy the Variables
;;; directly.
OpCopyMemory %half_wg_0 %input_buffer_0
;;; Then copy the second input into the other half of Workgroup
;;; memory, use OpAccessChain to extract the array pointers (what we
;;; are copying), since the whole types don't match (different Offsets
;;; in the block struct).
%a = OpAccessChain %half_array_wg_ptr %half_wg_1 %const_uint_0
%b = OpAccessChain %half_array_sb_ptr %input_buffer_1 %const_uint_0
OpCopyMemory %a %b
;;; The two halves of Workgroup memory were filled, now copy the large
;;; Workgroup block that alias the halves into the output.
OpCopyMemory %output_buffer %wg
OpReturn
OpFunctionEnd
END
BUFFER input_buffer_0 DATA_TYPE uint32 SIZE 16 SERIES_FROM 1 INC_BY 1
BUFFER input_buffer_1 DATA_TYPE uint32 SIZE 16 SERIES_FROM 17 INC_BY 1
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER expected_buffer DATA_TYPE uint32 SIZE 32 SERIES_FROM 1 INC_BY 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER input_buffer_0 AS storage DESCRIPTOR_SET 0 BINDING 0
BIND BUFFER input_buffer_1 AS storage DESCRIPTOR_SET 0 BINDING 1
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 2
END
RUN pipeline 1 1 1
EXPECT output_buffer EQ_BUFFER expected_buffer
vulkan/amber/compute/workgroup_memory_explicit_layout/copy_memory_two_invocations.amber
+180
View File
@@ -0,0 +1,180 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %index %push_constants %input_buffer_0 %input_buffer_1 %output_buffer %wg %half_wg_0 %half_wg_1
OpExecutionMode %main LocalSize 32 4 1
OpDecorate %index BuiltIn LocalInvocationIndex
OpDecorate %buffer_type Block
OpDecorate %half_buffer_type Block
OpDecorate %other_half_buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpDecorate %array ArrayStride 4
OpMemberDecorate %half_buffer_type 0 Offset 0
OpMemberDecorate %other_half_buffer_type 0 Offset 64
OpDecorate %half_array ArrayStride 4
OpDecorate %input_buffer_0 DescriptorSet 0
OpDecorate %input_buffer_0 Binding 0
OpDecorate %input_buffer_1 DescriptorSet 0
OpDecorate %input_buffer_1 Binding 1
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 2
OpDecorate %wg Aliased
OpDecorate %half_wg_0 Aliased
OpDecorate %half_wg_1 Aliased
OpMemberDecorate %push_constants_type 0 Offset 0
OpMemberDecorate %push_constants_type 1 Offset 4
OpDecorate %push_constants_type Block
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_1 = OpConstant %uint 1
%const_uint_2 = OpConstant %uint 2
%const_uint_16 = OpConstant %uint 16
%const_uint_32 = OpConstant %uint 32
%const_uint_64 = OpConstant %uint 64
%const_uint_128 = OpConstant %uint 128
%uint_input_ptr = OpTypePointer Input %uint
%uint_pc_ptr = OpTypePointer PushConstant %uint
%bool = OpTypeBool
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%half_array = OpTypeArray %uint %const_uint_16
%half_array_wg_ptr = OpTypePointer Workgroup %half_array
%half_array_sb_ptr = OpTypePointer StorageBuffer %half_array
%buffer_type = OpTypeStruct %array
%half_buffer_type = OpTypeStruct %half_array
%other_half_buffer_type = OpTypeStruct %half_array
%input_buffer_0_ptr = OpTypePointer StorageBuffer %half_buffer_type
%input_buffer_1_ptr = OpTypePointer StorageBuffer %half_buffer_type
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%half_wg_0_ptr = OpTypePointer Workgroup %half_buffer_type
%half_wg_1_ptr = OpTypePointer Workgroup %other_half_buffer_type
%push_constants_type = OpTypeStruct %uint %uint
%push_constants_type_ptr = OpTypePointer PushConstant %push_constants_type
;;; Workgroup has a block covering the entire memory and another two
;;; blocks covering each half of the former.
%wg = OpVariable %wg_ptr Workgroup
%half_wg_0 = OpVariable %half_wg_0_ptr Workgroup
%half_wg_1 = OpVariable %half_wg_1_ptr Workgroup
%input_buffer_0 = OpVariable %input_buffer_0_ptr StorageBuffer
%input_buffer_1 = OpVariable %input_buffer_1_ptr StorageBuffer
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%index = OpVariable %uint_input_ptr Input
%push_constants = OpVariable %push_constants_type_ptr PushConstant
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Look up in the push constants values a and b that will be use to select
;;; which invocation will do the work.
%index_val = OpLoad %uint %index
%first_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_0
%second_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_1
%first = OpLoad %uint %first_ptr
%second = OpLoad %uint %second_ptr
%is_first = OpIEqual %bool %index_val %first
%is_second = OpIEqual %bool %index_val %second
;;; Copy first input into half of the workgroup memory. Because the
;;; whole types match here, we can afford to copy the Variables
;;; directly. Only the first invocation specified will do that.
OpSelectionMerge %after_first_copy None
OpBranchConditional %is_first %first_copy %after_first_copy
%first_copy = OpLabel
OpCopyMemory %half_wg_0 %input_buffer_0
OpBranch %after_first_copy
%after_first_copy = OpLabel
;;; Then copy the second input into the other half of Workgroup
;;; memory, use OpAccessChain to extract the array pointers (what we
;;; are copying), since the whole types don't match (different Offsets
;;; in the block struct). Only the second invocation specified will
;;; do that.
OpSelectionMerge %after_second_copy None
OpBranchConditional %is_second %second_copy %after_second_copy
%second_copy = OpLabel
%a = OpAccessChain %half_array_wg_ptr %half_wg_1 %const_uint_0
%b = OpAccessChain %half_array_sb_ptr %input_buffer_1 %const_uint_0
OpCopyMemory %a %b
OpBranch %after_second_copy
%after_second_copy = OpLabel
OpControlBarrier %const_uint_2 %const_uint_2 %const_uint_0
;;; The two halves of Workgroup memory were filled, now copy the large
;;; Workgroup block that alias the halves into the output. Only the first
;;; invocation specified will do that.
OpSelectionMerge %after_third_copy None
OpBranchConditional %is_first %third_copy %after_third_copy
%third_copy = OpLabel
OpCopyMemory %output_buffer %wg
OpBranch %after_third_copy
%after_third_copy = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER input_buffer_0 DATA_TYPE uint32 SIZE 16 SERIES_FROM 1 INC_BY 1
BUFFER input_buffer_1 DATA_TYPE uint32 SIZE 16 SERIES_FROM 17 INC_BY 1
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER expected_buffer DATA_TYPE uint32 SIZE 32 SERIES_FROM 1 INC_BY 1
BUFFER const_buf DATA_TYPE uint32 DATA
30 90
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER const_buf AS push_constant
BIND BUFFER input_buffer_0 AS storage DESCRIPTOR_SET 0 BINDING 0
BIND BUFFER input_buffer_1 AS storage DESCRIPTOR_SET 0 BINDING 1
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 2
END
RUN pipeline 4 1 1
EXPECT output_buffer EQ_BUFFER expected_buffer
vulkan/amber/compute/workgroup_memory_explicit_layout/copy_memory_variable_pointers.amber
+202
View File
@@ -0,0 +1,202 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_FEATURE VariablePointerFeatures.variablePointers
DEVICE_EXTENSION VK_EXT_descriptor_indexing
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpCapability VariablePointers
OpCapability ShaderNonUniform
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpExtension "SPV_EXT_descriptor_indexing"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %index %push_constants %input_buffer_0 %input_buffer_1 %output_buffer %wg %half_wg_0 %half_wg_1
OpExecutionMode %main LocalSize 32 4 1
OpDecorate %index BuiltIn LocalInvocationIndex
OpDecorate %buffer_type Block
OpDecorate %half_buffer_type Block
OpDecorate %other_half_buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpDecorate %array ArrayStride 4
OpMemberDecorate %half_buffer_type 0 Offset 0
OpMemberDecorate %other_half_buffer_type 0 Offset 64
OpDecorate %half_array ArrayStride 4
OpDecorate %input_buffer_0 DescriptorSet 0
OpDecorate %input_buffer_0 Binding 0
OpDecorate %input_buffer_1 DescriptorSet 0
OpDecorate %input_buffer_1 Binding 1
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 2
OpDecorate %wg Aliased
OpDecorate %half_wg_0 Aliased
OpDecorate %half_wg_1 Aliased
OpMemberDecorate %push_constants_type 0 Offset 0
OpMemberDecorate %push_constants_type 1 Offset 4
OpDecorate %push_constants_type Block
OpDecorate %src NonUniform
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_1 = OpConstant %uint 1
%const_uint_2 = OpConstant %uint 2
%const_uint_16 = OpConstant %uint 16
%const_uint_32 = OpConstant %uint 32
%const_uint_64 = OpConstant %uint 64
%const_uint_128 = OpConstant %uint 128
%uint_input_ptr = OpTypePointer Input %uint
%uint_pc_ptr = OpTypePointer PushConstant %uint
%bool = OpTypeBool
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%half_array = OpTypeArray %uint %const_uint_16
%half_array_wg_ptr = OpTypePointer Workgroup %half_array
%half_array_sb_ptr = OpTypePointer StorageBuffer %half_array
%buffer_type = OpTypeStruct %array
%half_buffer_type = OpTypeStruct %half_array
%other_half_buffer_type = OpTypeStruct %half_array
%input_buffer_0_ptr = OpTypePointer StorageBuffer %half_buffer_type
%input_buffer_1_ptr = OpTypePointer StorageBuffer %half_buffer_type
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%half_wg_0_ptr = OpTypePointer Workgroup %half_buffer_type
%half_wg_1_ptr = OpTypePointer Workgroup %other_half_buffer_type
%push_constants_type = OpTypeStruct %uint %uint
%push_constants_type_ptr = OpTypePointer PushConstant %push_constants_type
%null_src = OpConstantNull %half_array_sb_ptr
;;; Workgroup has a block covering the entire memory and another two
;;; blocks covering each half of the former.
%wg = OpVariable %wg_ptr Workgroup
%half_wg_0 = OpVariable %half_wg_0_ptr Workgroup
%half_wg_1 = OpVariable %half_wg_1_ptr Workgroup
%input_buffer_0 = OpVariable %input_buffer_0_ptr StorageBuffer
%input_buffer_1 = OpVariable %input_buffer_1_ptr StorageBuffer
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%index = OpVariable %uint_input_ptr Input
%push_constants = OpVariable %push_constants_type_ptr PushConstant
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Look up in the push constants values a and b that will be use to select
;;; which invocation will do the work.
%index_val = OpLoad %uint %index
%first_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_0
%second_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_1
%first = OpLoad %uint %first_ptr
%second = OpLoad %uint %second_ptr
%is_first = OpIEqual %bool %index_val %first
%is_second = OpIEqual %bool %index_val %second
%is_first_or_second = OpLogicalOr %bool %is_first %is_second
;;; Set the values for %dst using Select.
%ptr_to_array_wg_0 = OpAccessChain %half_array_wg_ptr %half_wg_0 %const_uint_0
%ptr_to_array_wg_1 = OpAccessChain %half_array_wg_ptr %half_wg_1 %const_uint_0
%dst = OpSelect %half_array_wg_ptr %is_first %ptr_to_array_wg_0 %ptr_to_array_wg_1
;;; Set the values for %src using OpPhi.
OpSelectionMerge %after_src_choice None
OpBranchConditional %is_first %src_for_first %next
%src_for_first = OpLabel
%first_src = OpAccessChain %half_array_sb_ptr %input_buffer_0 %const_uint_0
OpBranch %after_src_choice
%next = OpLabel
OpSelectionMerge %merge_inner None
OpBranchConditional %is_second %src_for_second %src_is_null
%src_for_second = OpLabel
%second_src = OpAccessChain %half_array_sb_ptr %input_buffer_1 %const_uint_0
OpBranch %merge_inner
%src_is_null = OpLabel
OpBranch %merge_inner
%merge_inner = OpLabel
%inner_src = OpPhi %half_array_sb_ptr %second_src %src_for_second %null_src %src_is_null
OpBranch %after_src_choice
%after_src_choice = OpLabel
%src = OpPhi %half_array_sb_ptr %first_src %src_for_first %inner_src %merge_inner
;;; The first and second invocations identified will copy from src to
;;; dst.
OpSelectionMerge %after_var_copy None
OpBranchConditional %is_first_or_second %var_copy %after_var_copy
%var_copy = OpLabel
OpCopyMemory %dst %src
OpBranch %after_var_copy
%after_var_copy = OpLabel
OpControlBarrier %const_uint_2 %const_uint_2 %const_uint_0
;;; The two halves of Workgroup memory were filled, now copy the large
;;; Workgroup block that alias the halves into the output. Only the first
;;; invocation specified will do that.
OpSelectionMerge %after_output_copy None
OpBranchConditional %is_first %output_copy %after_output_copy
%output_copy = OpLabel
OpCopyMemory %output_buffer %wg
OpBranch %after_output_copy
%after_output_copy = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER input_buffer_0 DATA_TYPE uint32 SIZE 16 SERIES_FROM 1 INC_BY 1
BUFFER input_buffer_1 DATA_TYPE uint32 SIZE 16 SERIES_FROM 17 INC_BY 1
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER expected_buffer DATA_TYPE uint32 SIZE 32 SERIES_FROM 1 INC_BY 1
BUFFER const_buf DATA_TYPE uint32 DATA
30 90
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER const_buf AS push_constant
BIND BUFFER input_buffer_0 AS storage DESCRIPTOR_SET 0 BINDING 0
BIND BUFFER input_buffer_1 AS storage DESCRIPTOR_SET 0 BINDING 1
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 2
END
RUN pipeline 4 1 1
EXPECT output_buffer EQ_BUFFER expected_buffer
vulkan/amber/compute/workgroup_memory_explicit_layout/zero_ext_block.amber
+106
View File
@@ -0,0 +1,106 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %index %push_constants %output_buffer %wg
OpExecutionMode %main LocalSize 32 4 1
OpDecorate %index BuiltIn LocalInvocationIndex
OpDecorate %buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpDecorate %array ArrayStride 4
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 1
OpDecorate %wg Aliased
OpMemberDecorate %push_constants_type 0 Offset 0
OpDecorate %push_constants_type Block
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_32 = OpConstant %uint 32
%uint_input_ptr = OpTypePointer Input %uint
%uint_pc_ptr = OpTypePointer PushConstant %uint
%bool = OpTypeBool
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%buffer_type = OpTypeStruct %array
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%push_constants_type = OpTypeStruct %uint
%push_constants_type_ptr = OpTypePointer PushConstant %push_constants_type
%null_buffer = OpConstantNull %buffer_type
;;; Workgroup variable has a "null initializer".
%wg = OpVariable %wg_ptr Workgroup %null_buffer
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%index = OpVariable %uint_input_ptr Input
%push_constants = OpVariable %push_constants_type_ptr PushConstant
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Look up in the push constant to select which invocation will do
;;; the work.
%index_val = OpLoad %uint %index
%worker_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_0
%worker = OpLoad %uint %worker_ptr
%is_worker = OpIEqual %bool %index_val %worker
OpSelectionMerge %after_copy None
OpBranchConditional %is_worker %copy %after_copy
%copy = OpLabel
OpCopyMemory %output_buffer %wg
OpBranch %after_copy
%after_copy = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER expected_buffer DATA_TYPE uint32 SIZE 32 FILL 0
BUFFER const_buf DATA_TYPE uint32 DATA
30
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER const_buf AS push_constant
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 1
END
RUN pipeline 4 1 1
EXPECT output_buffer EQ_BUFFER expected_buffer
vulkan/amber/compute/workgroup_memory_explicit_layout/zero_ext_block_with_offset.amber
+128
View File
@@ -0,0 +1,128 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %index %push_constants %output_buffer %wg %offset_wg
OpExecutionMode %main LocalSize 32 4 1
OpDecorate %index BuiltIn LocalInvocationIndex
OpDecorate %buffer_type Block
OpDecorate %offset_buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpMemberDecorate %offset_buffer_type 0 Offset 64
OpDecorate %array ArrayStride 4
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 1
OpDecorate %wg Aliased
OpDecorate %offset_wg Aliased
OpMemberDecorate %push_constants_type 0 Offset 0
OpDecorate %push_constants_type Block
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_32 = OpConstant %uint 32
%uint_input_ptr = OpTypePointer Input %uint
%uint_pc_ptr = OpTypePointer PushConstant %uint
%bool = OpTypeBool
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%buffer_type = OpTypeStruct %array
%offset_buffer_type = OpTypeStruct %array
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%offset_wg_ptr = OpTypePointer Workgroup %offset_buffer_type
%push_constants_type = OpTypeStruct %uint
%push_constants_type_ptr = OpTypePointer PushConstant %push_constants_type
%null_offset_buffer = OpConstantNull %offset_buffer_type
;;; One variable has a "null initializer" and the other not.
%wg = OpVariable %wg_ptr Workgroup
%offset_wg = OpVariable %offset_wg_ptr Workgroup %null_offset_buffer
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%index = OpVariable %uint_input_ptr Input
%push_constants = OpVariable %push_constants_type_ptr PushConstant
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Look up in the push constant to select which invocation will do
;;; the work.
%index_val = OpLoad %uint %index
%worker_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_0
%worker = OpLoad %uint %worker_ptr
%is_worker = OpIEqual %bool %index_val %worker
OpSelectionMerge %after_copy None
OpBranchConditional %is_worker %copy %after_copy
%copy = OpLabel
;;; Copying variable to output. Because the aliasing at least half of
;;; it should be zero initialized.
OpCopyMemory %output_buffer %wg
OpBranch %after_copy
%after_copy = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER const_buf DATA_TYPE uint32 DATA
30
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER const_buf AS push_constant
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 1
END
RUN pipeline 4 1 1
EXPECT output_buffer IDX 64 EQ 0
EXPECT output_buffer IDX 68 EQ 0
EXPECT output_buffer IDX 72 EQ 0
EXPECT output_buffer IDX 76 EQ 0
EXPECT output_buffer IDX 80 EQ 0
EXPECT output_buffer IDX 84 EQ 0
EXPECT output_buffer IDX 88 EQ 0
EXPECT output_buffer IDX 92 EQ 0
EXPECT output_buffer IDX 96 EQ 0
EXPECT output_buffer IDX 100 EQ 0
EXPECT output_buffer IDX 104 EQ 0
EXPECT output_buffer IDX 108 EQ 0
EXPECT output_buffer IDX 112 EQ 0
EXPECT output_buffer IDX 116 EQ 0
EXPECT output_buffer IDX 120 EQ 0
EXPECT output_buffer IDX 124 EQ 0
vulkan/amber/compute/workgroup_memory_explicit_layout/zero_ext_other_block.amber
+111
View File
@@ -0,0 +1,111 @@
#!amber
# Copyright 2021 Intel Corporation
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_spirv_1_4
DEVICE_EXTENSION VK_KHR_workgroup_memory_explicit_layout
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader SPIRV-ASM
OpCapability Shader
OpCapability WorkgroupMemoryExplicitLayoutKHR
OpExtension "SPV_KHR_workgroup_memory_explicit_layout"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %index %push_constants %output_buffer %wg %other_wg
OpExecutionMode %main LocalSize 32 4 1
OpDecorate %index BuiltIn LocalInvocationIndex
OpDecorate %buffer_type Block
OpMemberDecorate %buffer_type 0 Offset 0
OpDecorate %array ArrayStride 4
OpDecorate %output_buffer DescriptorSet 0
OpDecorate %output_buffer Binding 1
OpDecorate %wg Aliased
OpDecorate %other_wg Aliased
OpMemberDecorate %push_constants_type 0 Offset 0
OpDecorate %push_constants_type Block
%uint = OpTypeInt 32 0
%const_uint_0 = OpConstant %uint 0
%const_uint_32 = OpConstant %uint 32
%uint_input_ptr = OpTypePointer Input %uint
%uint_pc_ptr = OpTypePointer PushConstant %uint
%bool = OpTypeBool
%void = OpTypeVoid
%main_type = OpTypeFunction %void
%array = OpTypeArray %uint %const_uint_32
%buffer_type = OpTypeStruct %array
%output_buffer_ptr = OpTypePointer StorageBuffer %buffer_type
%wg_ptr = OpTypePointer Workgroup %buffer_type
%push_constants_type = OpTypeStruct %uint
%push_constants_type_ptr = OpTypePointer PushConstant %push_constants_type
%null_buffer = OpConstantNull %buffer_type
;;; One variable has a "null initializer" and the other not.
%wg = OpVariable %wg_ptr Workgroup %null_buffer
%other_wg = OpVariable %wg_ptr Workgroup
%output_buffer = OpVariable %output_buffer_ptr StorageBuffer
%index = OpVariable %uint_input_ptr Input
%push_constants = OpVariable %push_constants_type_ptr PushConstant
%main = OpFunction %void None %main_type
%entry = OpLabel
;;; Look up in the push constant to select which invocation will do
;;; the work.
%index_val = OpLoad %uint %index
%worker_ptr = OpAccessChain %uint_pc_ptr %push_constants %const_uint_0
%worker = OpLoad %uint %worker_ptr
%is_worker = OpIEqual %bool %index_val %worker
OpSelectionMerge %after_copy None
OpBranchConditional %is_worker %copy %after_copy
%copy = OpLabel
;;; Copying the other variable to output. Because the aliasing, it
;;; should be zero initalized too.
OpCopyMemory %output_buffer %other_wg
OpBranch %after_copy
%after_copy = OpLabel
OpReturn
OpFunctionEnd
END
BUFFER output_buffer DATA_TYPE uint32 SIZE 32 FILL 99
BUFFER expected_buffer DATA_TYPE uint32 SIZE 32 FILL 0
BUFFER const_buf DATA_TYPE uint32 DATA
30
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER const_buf AS push_constant
BIND BUFFER output_buffer AS storage DESCRIPTOR_SET 0 BINDING 1
END
RUN pipeline 4 1 1
EXPECT output_buffer EQ_BUFFER expected_buffer
vulkan/amber/compute/write_ssbo_array.amber
+53
View File
@@ -0,0 +1,53 @@
#!amber
# Copyright 2020 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
SHADER compute compute_shader GLSL
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0) buffer block0
{
int data;
} ssbo_array[2];
void main()
{
ssbo_array[0].data++;
ssbo_array[0].data++;
ssbo_array[0].data++;
ssbo_array[1].data++;
ssbo_array[1].data++;
}
END
BUFFER buf0 DATA_TYPE int32 DATA
0
END
BUFFER buf1 DATA_TYPE int32 DATA
0
END
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER_ARRAY buf0 buf1 AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT buf0 IDX 0 EQ 3
EXPECT buf1 IDX 0 EQ 2
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_128.amber
+55
View File
@@ -0,0 +1,55 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 128) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx = gl_LocalInvocationID.x;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_2x8x8.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 2, local_size_y = 8, local_size_z = 8) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_4x4x8.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 4, local_size_y = 4, local_size_z = 8) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_4x8x4.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 4, local_size_y = 8, local_size_z = 4) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_8x2x8.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 8, local_size_y = 2, local_size_z = 8) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_8x4x4.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 8, local_size_y = 4, local_size_z = 4) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer
vulkan/amber/compute/zero_initialize_workgroup_memory/workgroup_size_8x8x2.amber
+59
View File
@@ -0,0 +1,59 @@
#!amber
# Copyright 2022 Google LLC
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
DEVICE_EXTENSION VK_KHR_zero_initialize_workgroup_memory
SHADER compute compute_shader GLSL TARGET_ENV spv1.3
#version 450
#extension GL_EXT_null_initializer : enable
layout(local_size_x = 8, local_size_y = 8, local_size_z = 2) in;
layout(set = 0, binding = 0) buffer A
{
uint a[];
} a;
shared uint wg_mem = {};
void main()
{
uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y;
uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x;
uint idx_x = gl_LocalInvocationID.x;
uint idx = idx_x + idx_y + idx_z;
if (gl_LocalInvocationIndex == 0)
{
wg_mem = 1;
}
barrier();
a.a[idx] = wg_mem;
}
END
BUFFER result_buffer DATA_TYPE uint32 SIZE 128 FILL 99
BUFFER reference_buffer DATA_TYPE uint32 SIZE 128 FILL 1
PIPELINE compute pipeline
ATTACH compute_shader
BIND BUFFER result_buffer AS storage DESCRIPTOR_SET 0 BINDING 0
END
RUN pipeline 1 1 1
EXPECT result_buffer EQ_BUFFER reference_buffer