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 fb6a06c3..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv32 and /dev/null 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 3c2a9517..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_coherent.spv64 and /dev/null 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 62594fa5..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv32 and /dev/null 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 c56fc97f..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_nonreadable.spv64 and /dev/null 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 5d1e2ac2..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv32 and /dev/null 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 05450370..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_nonwritable.spv64 and /dev/null 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 78c6b05d..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv32 and /dev/null 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 794c0faa..00000000 Binary files a/test_conformance/spirv_new/spirv_bin/decorate_volatile.spv64 and /dev/null differ 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,