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 <stuart.brady@arm.com>

test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm32

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_coherent.spvasm64

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm32

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_nonreadable.spvasm64

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm32

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_nonwritable.spvasm64

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm32

@@ -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

test_conformance/spirv_new/spirv_asm/decorate_volatile.spvasm64

@@ -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

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<cl_uint4> src(width * height);
-    RandomSeed seed(gRandomSeed);
-
-    for (auto &val : src) {
-        val = genrand<cl_uint4>(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<cl_uint4> 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<cl_uint4> src(width * height);
-    RandomSeed seed(gRandomSeed);
-
-    for (auto &val : src) {
-        val = genrand<cl_uint4>(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<cl_uint4> 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<typename Ti, typename To>
 int test_fp_rounding(cl_device_id deviceID,
                      cl_context context,