diff --git a/benchmarks/opencl/NonUniformBarrier.cl b/benchmarks/opencl/NonUniformBarrier.cl index 3c04e25b93..517c64ddae 100644 --- a/benchmarks/opencl/NonUniformBarrier.cl +++ b/benchmarks/opencl/NonUniformBarrier.cl @@ -1,6 +1,8 @@ __kernel void non_uniform_barrier() { - if(get_global_id(0) == 0) { + // PASS liveness if each WG has a single tread + // FAIL otherwise + if(get_local_id(0) == 0) { barrier(CLK_GLOBAL_MEM_FENCE); } diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/basic/SpirvLivenessTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/basic/SpirvLivenessTest.java index 56e54edbb4..997b504334 100644 --- a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/basic/SpirvLivenessTest.java +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/basic/SpirvLivenessTest.java @@ -48,7 +48,8 @@ public SpirvLivenessTest(String file, int bound, Result expected) { @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") public static Iterable data() throws IOException { return Arrays.asList(new Object[][]{ - {"NonUniformBarrier.spv.dis", 1, PASS}, + {"NonUniformBarrier-1.spv.dis", 1, PASS}, + {"NonUniformBarrier-2.spv.dis", 1, FAIL}, }); } diff --git a/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-1.spv.dis b/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-1.spv.dis new file mode 100644 index 0000000000..deada79b1f --- /dev/null +++ b/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-1.spv.dis @@ -0,0 +1,51 @@ +; @Config: 1, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 34 +; Schema: 0 + OpCapability Shader + %28 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %13 "non_uniform_barrier" %gl_LocalInvocationID %10 + OpSource OpenCL_C 300 + %29 = OpString "non_uniform_barrier" + %30 = OpString "__kernel" + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %5 SpecId 0 + OpDecorate %6 SpecId 1 + OpDecorate %7 SpecId 2 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %5 = OpSpecConstant %uint 1 + %6 = OpSpecConstant %uint 1 + %7 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %5 %6 %7 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint + %void = OpTypeVoid + %12 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_72 = OpConstant %uint 72 + %uint_1 = OpConstant %uint 1 +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %10 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %13 = OpFunction %void None %12 + %14 = OpLabel + %17 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %18 = OpLoad %uint %17 Aligned 16 + %20 = OpIEqual %bool %18 %uint_0 + OpSelectionMerge %27 None + OpBranchConditional %20 %23 %27 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_72 + OpBranch %27 + %27 = OpLabel + OpReturn + OpFunctionEnd + %31 = OpExtInst %void %28 Kernel %13 %29 %uint_0 %uint_0 %30 + %33 = OpExtInst %void %28 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 diff --git a/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-2.spv.dis b/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-2.spv.dis new file mode 100644 index 0000000000..b812230927 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier-2.spv.dis @@ -0,0 +1,51 @@ +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 34 +; Schema: 0 + OpCapability Shader + %28 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %13 "non_uniform_barrier" %gl_LocalInvocationID %10 + OpSource OpenCL_C 300 + %29 = OpString "non_uniform_barrier" + %30 = OpString "__kernel" + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %5 SpecId 0 + OpDecorate %6 SpecId 1 + OpDecorate %7 SpecId 2 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %5 = OpSpecConstant %uint 1 + %6 = OpSpecConstant %uint 1 + %7 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %5 %6 %7 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint + %void = OpTypeVoid + %12 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_72 = OpConstant %uint 72 + %uint_1 = OpConstant %uint 1 +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %10 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %13 = OpFunction %void None %12 + %14 = OpLabel + %17 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %18 = OpLoad %uint %17 Aligned 16 + %20 = OpIEqual %bool %18 %uint_0 + OpSelectionMerge %27 None + OpBranchConditional %20 %23 %27 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_72 + OpBranch %27 + %27 = OpLabel + OpReturn + OpFunctionEnd + %31 = OpExtInst %void %28 Kernel %13 %29 %uint_0 %uint_0 %30 + %33 = OpExtInst %void %28 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 diff --git a/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier.spv.dis b/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier.spv.dis deleted file mode 100644 index 363ee8b0f2..0000000000 --- a/dartagnan/src/test/resources/spirv/basic/NonUniformBarrier.spv.dis +++ /dev/null @@ -1,62 +0,0 @@ -; @Config: 2, 1, 1 -; SPIR-V -; Version: 1.6 -; Generator: Google Clspv; 0 -; Bound: 43 -; Schema: 0 - OpCapability Shader - %35 = OpExtInstImport "NonSemantic.ClspvReflection.5" - OpMemoryModel Logical Vulkan - OpEntryPoint GLCompute %16 "non_uniform_barrier" %gl_GlobalInvocationID %13 %5 - OpSource OpenCL_C 300 - %36 = OpString "non_uniform_barrier" - %37 = OpString "__kernel" - OpMemberDecorate %_struct_3 0 Offset 0 - OpDecorate %_struct_3 Block - OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId - OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize - OpDecorate %8 SpecId 0 - OpDecorate %9 SpecId 1 - OpDecorate %10 SpecId 2 - %uint = OpTypeInt 32 0 - %v3uint = OpTypeVector %uint 3 - %_struct_3 = OpTypeStruct %v3uint -%_ptr_PushConstant__struct_3 = OpTypePointer PushConstant %_struct_3 -%_ptr_Input_v3uint = OpTypePointer Input %v3uint - %8 = OpSpecConstant %uint 1 - %9 = OpSpecConstant %uint 1 - %10 = OpSpecConstant %uint 1 -%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %8 %9 %10 -%_ptr_Private_v3uint = OpTypePointer Private %v3uint - %void = OpTypeVoid - %15 = OpTypeFunction %void -%_ptr_Input_uint = OpTypePointer Input %uint - %uint_0 = OpConstant %uint 0 -%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint - %bool = OpTypeBool - %uint_2 = OpConstant %uint 2 - %uint_72 = OpConstant %uint 72 - %uint_12 = OpConstant %uint 12 - %uint_1 = OpConstant %uint 1 - %5 = OpVariable %_ptr_PushConstant__struct_3 PushConstant -%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input - %13 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize - %16 = OpFunction %void None %15 - %17 = OpLabel - %20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 - %21 = OpLoad %uint %20 Aligned 16 - %23 = OpAccessChain %_ptr_PushConstant_uint %5 %uint_0 %uint_0 - %24 = OpLoad %uint %23 Aligned 16 - %25 = OpISub %uint %uint_0 %24 - %27 = OpIEqual %bool %21 %25 - OpSelectionMerge %34 None - OpBranchConditional %27 %30 %34 - %30 = OpLabel - OpControlBarrier %uint_2 %uint_2 %uint_72 - OpBranch %34 - %34 = OpLabel - OpReturn - OpFunctionEnd - %40 = OpExtInst %void %35 PushConstantRegionOffset %uint_0 %uint_12 - %38 = OpExtInst %void %35 Kernel %16 %36 %uint_0 %uint_0 %37 - %42 = OpExtInst %void %35 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2