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>
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 158d7e4..0000000
--- 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 205a21c..0000000
--- 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 16c89ff..0000000
--- 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 cd5a006..0000000
--- 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 64c378d..0000000
--- 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 7df1989..0000000
--- 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 f6eb357..0000000
--- 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 ef2623b..0000000
--- 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 fb6a06c..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv32
+++ /dev/null
Binary files differ
diff --git a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv64 b/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv64
deleted file mode 100644
index 3c2a951..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv64
+++ /dev/null
Binary files differ
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 62594fa..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv32
+++ /dev/null
Binary files differ
diff --git a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv64 b/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv64
deleted file mode 100644
index c56fc97..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv64
+++ /dev/null
Binary files differ
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 5d1e2ac..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv32
+++ /dev/null
Binary files differ
diff --git a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv64 b/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv64
deleted file mode 100644
index 0545037..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv64
+++ /dev/null
Binary files differ
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 78c6b05..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv32
+++ /dev/null
Binary files differ
diff --git a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv64 b/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv64
deleted file mode 100644
index 794c0fa..0000000
--- a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv64
+++ /dev/null
Binary files differ
diff --git a/test_conformance/spirv_new/test_decorate.cpp b/test_conformance/spirv_new/test_decorate.cpp
index 766a6b6..ccd7431 100644
--- a/test_conformance/spirv_new/test_decorate.cpp
+++ b/test_conformance/spirv_new/test_decorate.cpp
@@ -310,174 +310,6 @@
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,
