From f732cd5b7e9ce899341444cb7f984326fea23696 Mon Sep 17 00:00:00 2001 From: Stuart Brady Date: Fri, 25 Sep 2020 14:57:41 +0100 Subject: [PATCH] Remove invalid SPIR-V image decoration tests (#22) (#958) The Volatile, Coherent, NonWritable and NonReadable decorations can only be applied to images if they are storage images, which require the Shader capability. Note that in SPIR-V 1.4 and later, the NonWritable decoration can also be applied to objects in the Private and Function storage classes (although Private also requires the Shader capability), but this was not covered by the existing test for the NonWritable decoration. Signed-off-by: Stuart Brady --- .../spirv_asm/decorate_coherent.spvasm32 | 81 --------- .../spirv_asm/decorate_coherent.spvasm64 | 94 ---------- .../spirv_asm/decorate_nonreadable.spvasm32 | 66 ------- .../spirv_asm/decorate_nonreadable.spvasm64 | 79 -------- .../spirv_asm/decorate_nonwritable.spvasm32 | 81 --------- .../spirv_asm/decorate_nonwritable.spvasm64 | 94 ---------- .../spirv_asm/decorate_volatile.spvasm32 | 81 --------- .../spirv_asm/decorate_volatile.spvasm64 | 94 ---------- .../spirv_bin/decorate_coherent.spv32 | Bin 1396 -> 0 bytes .../spirv_bin/decorate_coherent.spv64 | Bin 1608 -> 0 bytes .../spirv_bin/decorate_nonreadable.spv32 | Bin 1128 -> 0 bytes .../spirv_bin/decorate_nonreadable.spv64 | Bin 1340 -> 0 bytes .../spirv_bin/decorate_nonwritable.spv32 | Bin 1400 -> 0 bytes .../spirv_bin/decorate_nonwritable.spv64 | Bin 1612 -> 0 bytes .../spirv_bin/decorate_volatile.spv32 | Bin 1396 -> 0 bytes .../spirv_bin/decorate_volatile.spv64 | Bin 1608 -> 0 bytes test_conformance/spirv_new/test_decorate.cpp | 168 ------------------ 17 files changed, 838 deletions(-) delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm32 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm64 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm32 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm64 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm32 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm64 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm32 delete mode 100644 test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm64 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_coherent.spv32 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_coherent.spv64 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv32 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv64 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv32 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv64 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_volatile.spv32 delete mode 100644 test_conformance/spirv_new/spirv_bin/decorate_volatile.spv64 diff --git a/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm32 b/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm32 deleted file mode 100644 index 158d7e43..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm32 +++ /dev/null @@ -1,81 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 40 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical32 OpenCL - OpEntryPoint Kernel %2 "decorate_coherent" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %call1 "call1" - OpName %call2 "call2" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit6 "vecinit6" - OpName %TempSampledImage "TempSampledImage" - OpName %call7_old "call7.old" - OpName %arrayidx "arrayidx" - OpDecorate %src Coherent - OpDecorate %19 Constant - %19 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %19 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %uint_16 = OpConstant %uint 16 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3uint = OpTypeVector %uint 3 -%_ptr_Input_v3uint = OpTypePointer Input %v3uint - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %28 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %29 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %28 - %v2uint = OpTypeVector %uint 2 - %31 = OpTypeSampler - %32 = OpTypeSampledImage %28 - %float = OpTypeFloat 32 - %34 = OpConstantSampler %31 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3uint Input - %2 = OpFunction %void None %29 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %28 - %entry = OpLabel - %36 = OpUndef %v2uint - %37 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %uint %37 0 - %38 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %uint %38 1 - %39 = OpLoad %v3uint %__spirv_BuiltInGlobalSize Aligned 0 - %call2 = OpCompositeExtract %uint %39 0 - %mul = OpIMul %uint %call2 %call1 - %add = OpIAdd %uint %mul %call - %vecinit = OpCompositeInsert %v2uint %call1 %36 0 - %vecinit6 = OpCompositeInsert %v2uint %call %vecinit 1 -%TempSampledImage = OpSampledImage %32 %src %34 - %call7_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit6 Lod %float_0 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %add - OpStore %arrayidx %call7_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm64 b/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm64 deleted file mode 100644 index 205a21c3..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm64 +++ /dev/null @@ -1,94 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 47 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability Int64 - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %2 "decorate_coherent" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %conv "conv" - OpName %call1 "call1" - OpName %conv2 "conv2" - OpName %conv3 "conv3" - OpName %call4 "call4" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit10 "vecinit10" - OpName %TempSampledImage "TempSampledImage" - OpName %call11_old "call11.old" - OpName %sext "sext" - OpName %idxprom "idxprom" - OpName %arrayidx "arrayidx" - OpDecorate %src Coherent - OpDecorate %24 Constant - %24 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %24 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %ulong = OpTypeInt 64 0 - %uint_16 = OpConstant %uint 16 - %ulong_32 = OpConstant %ulong 32 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3ulong = OpTypeVector %ulong 3 -%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %35 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %36 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %35 - %v2uint = OpTypeVector %uint 2 - %38 = OpTypeSampler - %39 = OpTypeSampledImage %35 - %float = OpTypeFloat 32 - %41 = OpConstantSampler %38 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3ulong Input - %2 = OpFunction %void None %36 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %35 - %entry = OpLabel - %43 = OpUndef %v2uint - %44 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %ulong %44 0 - %conv = OpUConvert %uint %call - %45 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %ulong %45 1 - %conv2 = OpUConvert %uint %call1 - %conv3 = OpSConvert %ulong %conv2 - %46 = OpLoad %v3ulong %__spirv_BuiltInGlobalSize Aligned 0 - %call4 = OpCompositeExtract %ulong %46 0 - %mul = OpIMul %ulong %conv3 %call4 - %add = OpIAdd %ulong %mul %call - %vecinit = OpCompositeInsert %v2uint %conv2 %43 0 - %vecinit10 = OpCompositeInsert %v2uint %conv %vecinit 1 -%TempSampledImage = OpSampledImage %39 %src %41 - %call11_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit10 Lod %float_0 - %sext = OpShiftLeftLogical %ulong %add %ulong_32 - %idxprom = OpShiftRightArithmetic %ulong %sext %ulong_32 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %idxprom - OpStore %arrayidx %call11_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm32 b/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm32 deleted file mode 100644 index 16c89ff5..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm32 +++ /dev/null @@ -1,66 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 31 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability ImageBasic - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical32 OpenCL - OpEntryPoint Kernel %2 "decorate_nonreadable" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpSource OpenCL_C 100000 - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %call1 "call1" - OpName %call2 "call2" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit3 "vecinit3" - OpName %arrayidx "arrayidx" - OpDecorate %dst NonReadable - OpDecorate %16 Constant - %16 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %src FuncParamAttr NoCapture - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpGroupDecorate %16 %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %v3uint = OpTypeVector %uint 3 -%_ptr_Input_v3uint = OpTypePointer Input %v3uint - %void = OpTypeVoid - %21 = OpTypeImage %void 2D 0 0 0 0 Unknown WriteOnly - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %24 = OpTypeFunction %void %21 %_ptr_CrossWorkgroup_v4uint - %v2uint = OpTypeVector %uint 2 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3uint Input - %2 = OpFunction %void None %24 - %dst = OpFunctionParameter %21 - %src = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %entry = OpLabel - %26 = OpUndef %v2uint - %27 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %uint %27 0 - %28 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %uint %28 1 - %29 = OpLoad %v3uint %__spirv_BuiltInGlobalSize Aligned 0 - %call2 = OpCompositeExtract %uint %29 0 - %mul = OpIMul %uint %call2 %call1 - %add = OpIAdd %uint %mul %call - %vecinit = OpCompositeInsert %v2uint %call1 %26 0 - %vecinit3 = OpCompositeInsert %v2uint %call %vecinit 1 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %src %add - %30 = OpLoad %v4uint %arrayidx Aligned 16 - OpImageWrite %dst %vecinit3 %30 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm64 b/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm64 deleted file mode 100644 index cd5a0065..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm64 +++ /dev/null @@ -1,79 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 38 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability Int64 - OpCapability ImageBasic - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %2 "decorate_nonreadable" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpSource OpenCL_C 100000 - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %conv "conv" - OpName %call1 "call1" - OpName %conv2 "conv2" - OpName %conv3 "conv3" - OpName %call4 "call4" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit7 "vecinit7" - OpName %sext "sext" - OpName %idxprom "idxprom" - OpName %arrayidx "arrayidx" - OpDecorate %dst NonReadable - OpDecorate %21 Constant - %21 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %src FuncParamAttr NoCapture - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpGroupDecorate %21 %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %ulong = OpTypeInt 64 0 - %uint = OpTypeInt 32 0 - %ulong_32 = OpConstant %ulong 32 - %v3ulong = OpTypeVector %ulong 3 -%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong - %void = OpTypeVoid - %28 = OpTypeImage %void 2D 0 0 0 0 Unknown WriteOnly - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %31 = OpTypeFunction %void %28 %_ptr_CrossWorkgroup_v4uint - %v2uint = OpTypeVector %uint 2 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3ulong Input - %2 = OpFunction %void None %31 - %dst = OpFunctionParameter %28 - %src = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %entry = OpLabel - %33 = OpUndef %v2uint - %34 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %ulong %34 0 - %conv = OpUConvert %uint %call - %35 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %ulong %35 1 - %conv2 = OpUConvert %uint %call1 - %conv3 = OpSConvert %ulong %conv2 - %36 = OpLoad %v3ulong %__spirv_BuiltInGlobalSize Aligned 0 - %call4 = OpCompositeExtract %ulong %36 0 - %mul = OpIMul %ulong %conv3 %call4 - %add = OpIAdd %ulong %mul %call - %vecinit = OpCompositeInsert %v2uint %conv2 %33 0 - %vecinit7 = OpCompositeInsert %v2uint %conv %vecinit 1 - %sext = OpShiftLeftLogical %ulong %add %ulong_32 - %idxprom = OpShiftRightArithmetic %ulong %sext %ulong_32 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %src %idxprom - %37 = OpLoad %v4uint %arrayidx Aligned 16 - OpImageWrite %dst %vecinit7 %37 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm32 b/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm32 deleted file mode 100644 index 64c378de..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm32 +++ /dev/null @@ -1,81 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 40 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical32 OpenCL - OpEntryPoint Kernel %2 "decorate_nonwritable" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %call1 "call1" - OpName %call2 "call2" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit6 "vecinit6" - OpName %TempSampledImage "TempSampledImage" - OpName %call7_old "call7.old" - OpName %arrayidx "arrayidx" - OpDecorate %src NonWritable - OpDecorate %19 Constant - %19 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %19 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %uint_16 = OpConstant %uint 16 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3uint = OpTypeVector %uint 3 -%_ptr_Input_v3uint = OpTypePointer Input %v3uint - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %28 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %29 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %28 - %v2uint = OpTypeVector %uint 2 - %31 = OpTypeSampler - %32 = OpTypeSampledImage %28 - %float = OpTypeFloat 32 - %34 = OpConstantSampler %31 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3uint Input - %2 = OpFunction %void None %29 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %28 - %entry = OpLabel - %36 = OpUndef %v2uint - %37 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %uint %37 0 - %38 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %uint %38 1 - %39 = OpLoad %v3uint %__spirv_BuiltInGlobalSize Aligned 0 - %call2 = OpCompositeExtract %uint %39 0 - %mul = OpIMul %uint %call2 %call1 - %add = OpIAdd %uint %mul %call - %vecinit = OpCompositeInsert %v2uint %call1 %36 0 - %vecinit6 = OpCompositeInsert %v2uint %call %vecinit 1 -%TempSampledImage = OpSampledImage %32 %src %34 - %call7_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit6 Lod %float_0 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %add - OpStore %arrayidx %call7_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm64 b/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm64 deleted file mode 100644 index 7df1989a..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm64 +++ /dev/null @@ -1,94 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 47 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability Int64 - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %2 "decorate_nonwritable" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %conv "conv" - OpName %call1 "call1" - OpName %conv2 "conv2" - OpName %conv3 "conv3" - OpName %call4 "call4" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit10 "vecinit10" - OpName %TempSampledImage "TempSampledImage" - OpName %call11_old "call11.old" - OpName %sext "sext" - OpName %idxprom "idxprom" - OpName %arrayidx "arrayidx" - OpDecorate %src NonWritable - OpDecorate %24 Constant - %24 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %24 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %ulong = OpTypeInt 64 0 - %uint_16 = OpConstant %uint 16 - %ulong_32 = OpConstant %ulong 32 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3ulong = OpTypeVector %ulong 3 -%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %35 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %36 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %35 - %v2uint = OpTypeVector %uint 2 - %38 = OpTypeSampler - %39 = OpTypeSampledImage %35 - %float = OpTypeFloat 32 - %41 = OpConstantSampler %38 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3ulong Input - %2 = OpFunction %void None %36 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %35 - %entry = OpLabel - %43 = OpUndef %v2uint - %44 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %ulong %44 0 - %conv = OpUConvert %uint %call - %45 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %ulong %45 1 - %conv2 = OpUConvert %uint %call1 - %conv3 = OpSConvert %ulong %conv2 - %46 = OpLoad %v3ulong %__spirv_BuiltInGlobalSize Aligned 0 - %call4 = OpCompositeExtract %ulong %46 0 - %mul = OpIMul %ulong %conv3 %call4 - %add = OpIAdd %ulong %mul %call - %vecinit = OpCompositeInsert %v2uint %conv2 %43 0 - %vecinit10 = OpCompositeInsert %v2uint %conv %vecinit 1 -%TempSampledImage = OpSampledImage %39 %src %41 - %call11_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit10 Lod %float_0 - %sext = OpShiftLeftLogical %ulong %add %ulong_32 - %idxprom = OpShiftRightArithmetic %ulong %sext %ulong_32 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %idxprom - OpStore %arrayidx %call11_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm32 b/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm32 deleted file mode 100644 index f6eb3577..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm32 +++ /dev/null @@ -1,81 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 40 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical32 OpenCL - OpEntryPoint Kernel %2 "decorate_volatile" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %call1 "call1" - OpName %call2 "call2" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit6 "vecinit6" - OpName %TempSampledImage "TempSampledImage" - OpName %call7_old "call7.old" - OpName %arrayidx "arrayidx" - OpDecorate %src Volatile - OpDecorate %19 Constant - %19 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %19 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %uint_16 = OpConstant %uint 16 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3uint = OpTypeVector %uint 3 -%_ptr_Input_v3uint = OpTypePointer Input %v3uint - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %28 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %29 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %28 - %v2uint = OpTypeVector %uint 2 - %31 = OpTypeSampler - %32 = OpTypeSampledImage %28 - %float = OpTypeFloat 32 - %34 = OpConstantSampler %31 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3uint Input - %2 = OpFunction %void None %29 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %28 - %entry = OpLabel - %36 = OpUndef %v2uint - %37 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %uint %37 0 - %38 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %uint %38 1 - %39 = OpLoad %v3uint %__spirv_BuiltInGlobalSize Aligned 0 - %call2 = OpCompositeExtract %uint %39 0 - %mul = OpIMul %uint %call2 %call1 - %add = OpIAdd %uint %mul %call - %vecinit = OpCompositeInsert %v2uint %call1 %36 0 - %vecinit6 = OpCompositeInsert %v2uint %call %vecinit 1 -%TempSampledImage = OpSampledImage %32 %src %34 - %call7_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit6 Lod %float_0 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %add - OpStore %arrayidx %call7_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm64 b/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm64 deleted file mode 100644 index ef2623b7..00000000 --- a/test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm64 +++ /dev/null @@ -1,94 +0,0 @@ -; SPIR-V -; Version: 1.0 -; Generator: Khronos SPIR-V Tools Assembler; 0 -; Bound: 47 -; Schema: 0 - OpCapability Addresses - OpCapability Linkage - OpCapability Kernel - OpCapability Int64 - OpCapability ImageBasic - OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" - OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %2 "decorate_volatile" %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - OpName %sampler "sampler" - OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" - OpName %__spirv_BuiltInGlobalSize "__spirv_BuiltInGlobalSize" - OpName %dst "dst" - OpName %src "src" - OpName %entry "entry" - OpName %call "call" - OpName %conv "conv" - OpName %call1 "call1" - OpName %conv2 "conv2" - OpName %conv3 "conv3" - OpName %call4 "call4" - OpName %mul "mul" - OpName %add "add" - OpName %vecinit "vecinit" - OpName %vecinit10 "vecinit10" - OpName %TempSampledImage "TempSampledImage" - OpName %call11_old "call11.old" - OpName %sext "sext" - OpName %idxprom "idxprom" - OpName %arrayidx "arrayidx" - OpDecorate %src Volatile - OpDecorate %24 Constant - %24 = OpDecorationGroup - OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId - OpDecorate %__spirv_BuiltInGlobalSize BuiltIn GlobalSize - OpDecorate %dst FuncParamAttr NoCapture - OpDecorate %sampler LinkageAttributes "sampler" Export - OpDecorate %__spirv_BuiltInGlobalSize LinkageAttributes "__spirv_BuiltInGlobalSize" Import - OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import - OpDecorate %sampler Alignment 4 - OpGroupDecorate %24 %sampler %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInGlobalSize - %uint = OpTypeInt 32 0 - %ulong = OpTypeInt 64 0 - %uint_16 = OpConstant %uint 16 - %ulong_32 = OpConstant %ulong 32 -%_ptr_Input_uint = OpTypePointer Input %uint -%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint - %v3ulong = OpTypeVector %ulong 3 -%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong - %void = OpTypeVoid - %v4uint = OpTypeVector %uint 4 -%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint - %35 = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly - %36 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %35 - %v2uint = OpTypeVector %uint 2 - %38 = OpTypeSampler - %39 = OpTypeSampledImage %35 - %float = OpTypeFloat 32 - %41 = OpConstantSampler %38 None 0 Nearest - %float_0 = OpConstant %float 0 - %sampler = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_16 -%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input -%__spirv_BuiltInGlobalSize = OpVariable %_ptr_Input_v3ulong Input - %2 = OpFunction %void None %36 - %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint - %src = OpFunctionParameter %35 - %entry = OpLabel - %43 = OpUndef %v2uint - %44 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call = OpCompositeExtract %ulong %44 0 - %conv = OpUConvert %uint %call - %45 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 0 - %call1 = OpCompositeExtract %ulong %45 1 - %conv2 = OpUConvert %uint %call1 - %conv3 = OpSConvert %ulong %conv2 - %46 = OpLoad %v3ulong %__spirv_BuiltInGlobalSize Aligned 0 - %call4 = OpCompositeExtract %ulong %46 0 - %mul = OpIMul %ulong %conv3 %call4 - %add = OpIAdd %ulong %mul %call - %vecinit = OpCompositeInsert %v2uint %conv2 %43 0 - %vecinit10 = OpCompositeInsert %v2uint %conv %vecinit 1 -%TempSampledImage = OpSampledImage %39 %src %41 - %call11_old = OpImageSampleExplicitLod %v4uint %TempSampledImage %vecinit10 Lod %float_0 - %sext = OpShiftLeftLogical %ulong %add %ulong_32 - %idxprom = OpShiftRightArithmetic %ulong %sext %ulong_32 - %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %idxprom - OpStore %arrayidx %call11_old Aligned 16 - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv32 b/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv32 deleted file mode 100644 index fb6a06c37d47da2363968f84dda2211f70c90b0e..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1396 zcmaKr+iMe15XL9n#AsuywYIgbm&AHSE$Bll6n#)ILdxtq0D|<$N>eo7pqL+*aGHVPE@+ zr{#3`>M+f!LEg`bFJU&wr$rQ2X^{_fov;^Nj#IG*>5s%=0>>ksRKf?URCkq;^roGc zE;i+n!!?!7EDEzs!ky21`>D4(-hSrnz%ELT4l|dR-h^@N;ufV(lPJwo?rff236I_FBAeCOvMch1Wf>l&@po{2U(1@;mnB!V2kJO_ zU-i*rIA8ZR{l>oKZT`(uzvFt0|9^ev74>bw-FYVV#Q&Sc%>Msl z#l!EdSim1H^&45y32&~e&3feR73uV`E1TNMH^ViJzOb;@bbih2H+@4AEU{;;o+57Q5FGKukej}fo0-_qw=21+DtnS(R;AyT>}o&J z|3GeOA79iC=8kvn=`L!8Ba5Ew2}6C<+BjxY6ZVeKaDxX6uDe3Rl{bM7A78%-Kri?K&=8nFsMKf(Gj#eE#;b`5_IY-Zm)=kxB zcs-2L=dW6$A~a^+c*R$#>S~IMR1F{u)^y&*BJ2h5S98She!L z_=-6xUTE?I>sL`WWBNgoh**0KIZMZ-Odqhfo&A)v-*qzm!T+9<=aJ=q-^!j@5FP9# z7RPDvK^WRPi{i&oKTcyXR?V5ea+~>7g?dT!T{IZ(+CGKdLGX#=S%2nXakX1nqQ-Ex zW$8xIKCQO$ifA0}5A$qbV^$>xc^(`{iBHG%20i>aQDVB9AbdwTnHfE};AG~9yy0Z# zeN~%zcGPv~4efHaVx_O!nW6vR+_;?`GiPfa>wD-vOkyVY>6YTb_fjap50>)IEz5FV zvaiiTUK4$apw)Etb@9xpsc7t++t+6K^0_EV%}K6`xr9Lrc~y4$LL7DR_eJqXzofo4 zMDg1cB?i3sZHa<+qOY0D;_)T+xYba^6&dj(o)~g74|4Mk*Yv$6x-LZ5Me$h?Pk(*w zZOsh(P3?ad{lKT`lsl@$S;5GnPdoDQML8=tOPMnvKah7r6dWwDK9o0Dxi|3O;eGy+ z7jJV)I`A{g$MT|Y+uGD9^Sn`BdOYiFBIb**aAw{w`ZBg9{-{{)Fg9?OvAk{cxpI-8 zw~emMbz3}VJB))jkB#-+VdLEskL{b&aj(#W@2$?pIr$GzgY&XVJ^usFf*p+Sm4o^5 zPoQ7S>w~yWQT_$!i2)lk`{`=(PoOUM_=nT+#?hmDAszGJAHYnoJ&_KXyyPSPw@^IN F{sW2&i^2c^ diff --git a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv32 b/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv32 deleted file mode 100644 index 62594fa5421ed99c73de14a9815550315f686960..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1128 zcmaKq%Ss$U6o${-yvAG97?Y?Yf{=|611bbJ2GUFjgd`6zl-<2(XuGAS9ipo&#OLra zvl0BiqB;wiSWtcL|M~0m>9qHz&01zQZ9DpMR;+E^AbUYh1-T$OZ#`e}{W2{+e>)h} ziP@rcysxeJk}}9gPtv%oT%8V!vZzv*xbr-fTiND&bflr*-?fA~@ActubeUD-;n`J| z*MlO;%X60xig6jcIxCBT+9`L|*H6s-mEF*!Cp7pcquT3m8db6Mu3l@F7Ik&)^SU-8 z8M{1}@bTGTe+c$mun)c65mU2^tK8$hknWPi*D;@Q8fQgD&z|P;&;BUY`<5hKRk`ad z`He3MJU4^BEcr@zpb4T!;>(Qa`7YS}A8h^^ao)DC%X=e5h`-l;Za)0=$QC^Ie{-MA zotm#ghgb6t*JZEB?(1?kWWSWeM^0b%swB1eSM{zn$xWNHbxCw$lMa6Hmg3%m_U z-im)$+G|O0z#~Qt^Q4A7{3II=@By>2!NWch2KD4}z&lnpc=SOBmpT7c7W>nZj~;Jy zPZi;9Z*}R>;QUZE9M~^noFzUc2?dZ!nZwodEd>Httz(;|vh>vXOGv7_K z^3~h>O`DpXwZ5{=)2brKLr=4$s$H8+%Brk0m%6Pkdcf>-W{Lhp807O#INJ35^A%4QDiG(OHJ)xr$fYwTU_w&Z!a>^nZ{>~h%;?6t7KXXmkXShpq9 z+sQh33-~x+Lv6eV^6`BMHr@qx>ieMHB5~dachE1Vv-3{q1wS?Z3>dr>aM+C>!Nz;R kj_tK<%#=67T=7vq|MZRsJN1K`s2lzRcj5c3o1TjP0R9tpz5oCK diff --git a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv32 b/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv32 deleted file mode 100644 index 5d1e2ac2b6d5fb5158327653fad3d275c591f172..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1400 zcmaKr+iMd+6vj_3YqYVgwYIgbm&AHSE$Bll6n#)IL%i{0edo+~&dfCGTP?GinYHYZ{;VadThq(H%L`u4dwET=ZGp@DFpje~ z@3$vqWVT=p=c_BeCPwiNb~t8pfFKl#B%fvUkJnv_mg(>5uF(wlZs zdYqNTvBNc0&1@K^sf3HqdHb2S+unZe?1o*I937>uF7pYa$mK0apT@%^OV~5eJpS1` zQsrsc70J)>XuR+Gj{2kUtD6h>CB@*eyPc=AJ_EZdUsx33aT0x(YCZKuFW0@kEV-($ zuc@Q=)E_g3^9^rvw%E75&0Oy5WBxr&3CUfk&O)mEZ8u~5|C=*6sb}+U&vUV-&Tp18 zJO3xjhu=G~fInQSJF=n(_pYOl_1N1h(wSpNHoa4ChHKiIbQboS;>-!1n&8(Z!Q)@n z%r+#!ZAyZJ3pi#67yRosXE!Cml6&6jD&v-d;K(P3+RUBWyo(L}JCd7X)RhFYD*cXR zN8eMO2Xaf_$wlvA?t15*_M%rfvY5%9F!V>SwG%cyVebeHH+Zn%x+^qXxfAH{;eP!V zig!|z5B~S`y%dVMad+6Gs`F8G>5q3rj8$PDNIwh98zTm<3j0txTu;~K?h<4D@_X?H z@l)rg_wg3-!{d{|IJRLA>bs_bNipmoHNt0(Y0~25i=XNyZW+bZPaR>2A-bqwC?GYr)Nd$7Pts+ zdU5vR^=7|}%%*KjQOxZbA@E0z;!a+KW!%p4>`Rf9;cgmBEws6ufxhJIhuvNp7vk&A z)^7KENik?Y-%HZ6l{M3RH%wdEAn$}_l4mWY6n4VJI1;;)d?it!Xw(w*OULW*)-O8Z zYj#q6oR!6a!_^ectP`fGh_g>3=h>i==>_&JZ$IVj_q{Zhoq3!+K1-z)O2So(iDffnwtwfxZp4~*L&p4 zUZ&62w3%m9U5DP(u4XG%`nvN&|G&9$Lz_0`avqC4bRUL3!~1kw{@m$Ts)G8drMh$T zvYeOP&}Jboi9SQnTKD#4@yuym-q<;JLz~qQ=AtM$hp}pQ34<2$qU`hqKk|Zah=NBy zr@pR;f?E~EhkC)SiBj)SKWSIQgT?o_wIYwJGJ?Y&A7V2PV)GW4^t~>+tcq@kf>{tx ze;e8xni=+6+W#*40cPDRx0Q>tQX`8#Z41*7<*d|M$(#xKq0lR$)WM?GM?zC8_XeJN zxP!lh;%$ydNBzw5iBR-yLz^5`oVSWgk0-oM#6GJQ&dmEoU-@o|KgyRojEy=gU*0zQ zLa~U?+eTN#x+9*m9j=2nkB#-+*T%cYAKN#t<6fbszIQqs=j1;?4$jM}^!yJv3wCP! zpcu@Le**ntULVD+it;Z&j}Nsmv!6aE{{-@KkAHX_ZyY_km(np0{te6o+f(U~iAy~E Kf2)ee+J68|8I2JD diff --git a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv32 b/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv32 deleted file mode 100644 index 78c6b05db247cd97d91b9ee6468b305d3bc4516a..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1396 zcmaKr+iTQ76voe9;%e7cYi(;=FT2(&YC#`bq3DBxVMP$Fh;K{W1PNr5C7UkRCm;Ny zeG&YAvlA($bzpLC-#PQ0lbKd?yJOZcvyMH~m$ht7YkL`ZdCtp4FRw~=EpVCdCrSSL z-OjX%&6cd?d`-nK*@BUW9w(!sjH+ZfE3&9cvqZj@Hs=KTa?UgwPqL(x-gU9za5_oL z+3?k2npJ~5%!;p3HppkiNEM2FpyPzS;A)(TJxG6WPM~V=h^Lj))v!&=k@U8mm!9NR zdE{_yRWln!StjA)3*LU}?XI_d=su2gwC zc3JXMGM*f`zT?3-`UYKxUzUW&-cFIt`wZ-gd{J3OM``>WT&TY2W!38|l55%nO&mQ` zf6N%p*S*cWv2S^sx!luc{-LIXFflmxdW zNe*1VF*~^6-_V`DAqke;vsPajHx&d&J~`B8?$qW@Z0g&S+)|f)Nib{DZ%g*HpXh%e zx3!NidIxjIJNI=Ky~2^jO!kGLKYDE(v*`(YS7^Awg9X<;q2bD%K!*?a>$gz6lZJfo zzpMR1DCWl9;f%V@2i2uN-VrfYjlD1ZG%Rn77`$ri1L<%*S(m#@jP=W(#T&#=ouA&v zTf`5KkKzh%-Xs|KvuZzY6C80kd{&LVB=?H`ng17Q{D0mhesbW(KX@ZNbK`w7JI;P7 PA2xV!(18g z4hL~Q>OS9#lcJNhlk8KFbkb4QQ-&<$Kk zB{#HL$g85y5VSVEeN8-b+Eg@l&fU;vHRN+ql$zsQRl9^i3wcF$`a&Fa@o$LYkA6{o zZHnTzE=mk|@!Jpu?^r)$SH$B>>~X89h^sQySgoQKne$kh)E%8Uia)+^jvyA0!qtBI# z{Jd>+Wv)BoIon|zym@S_?;acPo_K8EypDT?9(?a~HqOa^fEt{aRqFX4a2D)f{Gc4n zkADLFVqPD`t&8$6Ku-+VnAuNXlebM>?(q+=OSo diff --git a/test_conformance/spirv_new/test_decorate.cpp b/test_conformance/spirv_new/test_decorate.cpp index 766a6b68..ccd74315 100644 --- a/test_conformance/spirv_new/test_decorate.cpp +++ b/test_conformance/spirv_new/test_decorate.cpp @@ -310,174 +310,6 @@ TEST_SATURATED_CONVERSION(float, uint, ushort) TEST_SATURATED_CONVERSION(double, long, int) TEST_SATURATED_CONVERSION(double, ulong, uint) -int test_image_decorate(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - const char *name) -{ - const int width = 4096; - const int height = 4096; - - std::vector src(width * height); - RandomSeed seed(gRandomSeed); - - for (auto &val : src) { - val = genrand(seed); - } - - cl_image_format imageFormat; - imageFormat.image_channel_data_type = CL_UNSIGNED_INT32; - imageFormat.image_channel_order = CL_RGBA; - - cl_image_desc desc; - desc.image_type = CL_MEM_OBJECT_IMAGE2D; - desc.image_width = width; - desc.image_height = height; - desc.image_depth = 0; - desc.image_array_size = 0; - desc.image_row_pitch = 0; // Automatically calculate the pitch - desc.image_slice_pitch = 0; // Not needed for 2D - desc.num_mip_levels = 0; - desc.num_samples = 0; - desc.mem_object = NULL; - - cl_int err = CL_SUCCESS; - - clMemWrapper srcImage = clCreateImage(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, - &imageFormat, &desc, - &src[0], &err); - SPIRV_CHECK_ERROR(err, "Failed to create image object"); - - size_t bytes = src.size() * sizeof(cl_uint4); - clMemWrapper dstBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); - SPIRV_CHECK_ERROR(err, "Failed to create image object"); - - clProgramWrapper prog; - err = get_program_with_il(prog, deviceID, context, name); - SPIRV_CHECK_ERROR(err, "Failed to build program"); - - clKernelWrapper kernel = clCreateKernel(prog, name, &err); - SPIRV_CHECK_ERROR(err, "Failed to create spv kernel"); - - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dstBuffer); - SPIRV_CHECK_ERROR(err, "Failed to set arg 2 of the kernel"); - - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &srcImage); - SPIRV_CHECK_ERROR(err, "Failed to set arg 1 of the kernel"); - - size_t global[] = {width, height}; - err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, NULL, 0, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to enqueue kernel"); - - std::vector dst(src.size()); - err = clEnqueueReadBuffer(queue, dstBuffer, CL_TRUE, 0, bytes, &dst[0], 0, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to copy data back to host"); - - for (int j = 0; j < height; j++) { - for (int i = 0; i < width; i++) { - int srcIdx = j * width + i; - int dstIdx = i * height + j; - if (dst[dstIdx] != src[srcIdx]) { - log_error("Values do not match at location (%d, %d) of src\n", i, j); - } - } - } - - return 0; -} - -#define TEST_SPIRV_IMAGE_DECORATE(type) \ - TEST_SPIRV_FUNC(decorate_##type) \ - { \ - PASSIVE_REQUIRE_IMAGE_SUPPORT(deviceID); \ - return test_image_decorate(deviceID, context, queue, \ - "decorate_" #type); \ - } \ - -TEST_SPIRV_IMAGE_DECORATE(volatile) -TEST_SPIRV_IMAGE_DECORATE(coherent) -TEST_SPIRV_IMAGE_DECORATE(nonwritable) - -TEST_SPIRV_FUNC(decorate_nonreadable) -{ - PASSIVE_REQUIRE_IMAGE_SUPPORT(deviceID); - const char *name = "decorate_nonreadable"; - const int width = 4096; - const int height = 4096; - cl_int err = CL_SUCCESS; - - std::vector src(width * height); - RandomSeed seed(gRandomSeed); - - for (auto &val : src) { - val = genrand(seed); - } - - size_t bytes = src.size() * sizeof(cl_uint4); - clMemWrapper srcBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); - SPIRV_CHECK_ERROR(err, "Failed to create image object"); - - err = clEnqueueWriteBuffer(queue, srcBuffer, CL_TRUE, 0, bytes, &src[0], 0, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to copy data back to host"); - - cl_image_format imageFormat; - imageFormat.image_channel_data_type = CL_UNSIGNED_INT32; - imageFormat.image_channel_order = CL_RGBA; - - cl_image_desc desc; - desc.image_type = CL_MEM_OBJECT_IMAGE2D; - desc.image_width = height; - desc.image_height = width; - desc.image_depth = 0; - desc.image_array_size = 0; - desc.image_row_pitch = 0; // Automatically calculate the pitch - desc.image_slice_pitch = 0; // Not needed for 2D - desc.num_mip_levels = 0; - desc.num_samples = 0; - desc.mem_object = NULL; - - - clMemWrapper dstImage = clCreateImage(context, CL_MEM_WRITE_ONLY, - &imageFormat, &desc, - NULL, &err); - SPIRV_CHECK_ERROR(err, "Failed to create image object"); - - clProgramWrapper prog; - err = get_program_with_il(prog, deviceID, context, name); - SPIRV_CHECK_ERROR(err, "Failed to build program"); - - clKernelWrapper kernel = clCreateKernel(prog, name, &err); - SPIRV_CHECK_ERROR(err, "Failed to create spv kernel"); - - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dstImage); - SPIRV_CHECK_ERROR(err, "Failed to set arg 2 of the kernel"); - - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &srcBuffer); - SPIRV_CHECK_ERROR(err, "Failed to set arg 1 of the kernel"); - - size_t global[] = {width, height}; - err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, NULL, 0, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to enqueue kernel"); - - std::vector dst(src.size()); - size_t origin[] = {0, 0, 0}; - size_t region[] = {height, width, 1}; - err = clEnqueueReadImage(queue, dstImage, CL_TRUE, origin, region, 0, 0, &dst[0], 0, NULL, NULL); - - - for (int j = 0; j < height; j++) { - for (int i = 0; i < width; i++) { - int srcIdx = j * width + i; - int dstIdx = i * height + j; - if (dst[dstIdx] != src[srcIdx]) { - log_error("Values do not match at location (%d, %d) of src\n", i, j); - } - } - } - - return 0; -} - template int test_fp_rounding(cl_device_id deviceID, cl_context context,