ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

add SPIR-V 1.4 testing for Image Operands and Loop Controls (#2034)

Browse Source 
This PR adds targeted testing for SPIR-V 1.4 features. Specifically,
this PR adds testing for:

1. The new Image Operand for SignExtend and ZeroExtend.
2. The new Loop Controls MinIterations, MaxIterations,
IterationMultiple, PeelCount, and PartialCount.
This commit is contained in:
Ben Ashbaugh
2024-10-22 09:47:31 -07:00
committed by GitHub
parent d1fe1ec252
commit 94257defea
17 changed files with 850 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
test_conformance/spirv_new/CMakeLists.txt
View File

@@ -27,6 +27,7 @@ set(${MODULE_NAME}_SOURCES
test_op_vector_extract.cpp
test_op_vector_insert.cpp
test_op_vector_times_scalar.cpp
test_spirv_14.cpp
)
set(TEST_HARNESS_SOURCES

41
test_conformance/spirv_new/spirvInfo.hpp Normal file
View File

@@ -0,0 +1,41 @@
//
// Copyright (c) 2024 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#pragma once
#include "harness/compat.h"
#include <string>
extern bool gVersionSkip;
static bool is_spirv_version_supported(cl_device_id deviceID,
const char* version)
{
std::string ilVersions = get_device_il_version_string(deviceID);
if (gVersionSkip)
{
log_info(" Skipping version check for %s.\n", version);
return true;
}
else if (ilVersions.find(version) == std::string::npos)
{
return false;
}
return true;
}

37
test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 Normal file
View File

@@ -0,0 +1,37 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 31
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability ImageBasic
OpCapability LiteralSampler
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %read_image_test "read_image_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%index_0 = OpConstant %uint 0
%void = OpTypeVoid
%v4uint = OpTypeVector %uint 4
%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint
%image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly
%7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t
%sampler_t = OpTypeSampler
%sampledimage_t = OpTypeSampledImage %image2d_t
%v2uint = OpTypeVector %uint 2
%float = OpTypeFloat 32
%sampler = OpConstantSampler %sampler_t None 0 Nearest
%coord_0_0 = OpConstantNull %v2uint
%float_0 = OpConstant %float 0
%read_image_test = OpFunction %void None %7
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint
%image = OpFunctionParameter %image2d_t
%entry = OpLabel
%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler
%call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|SignExtend %float_0
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0
OpStore %arrayidx %call Aligned 16
OpReturn
OpFunctionEnd

39
test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 Normal file
View File

@@ -0,0 +1,39 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 31
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpCapability ImageBasic
OpCapability LiteralSampler
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %read_image_test "read_image_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%index_0 = OpConstant %ulong 0
%void = OpTypeVoid
%v4uint = OpTypeVector %uint 4
%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint
%image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly
%7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t
%sampler_t = OpTypeSampler
%sampledimage_t = OpTypeSampledImage %image2d_t
%v2uint = OpTypeVector %uint 2
%float = OpTypeFloat 32
%sampler = OpConstantSampler %sampler_t None 0 Nearest
%coord_0_0 = OpConstantNull %v2uint
%float_0 = OpConstant %float 0
%read_image_test = OpFunction %void None %7
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint
%image = OpFunctionParameter %image2d_t
%entry = OpLabel
%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler
%call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|SignExtend %float_0
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0
OpStore %arrayidx %call Aligned 16
OpReturn
OpFunctionEnd

37
test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 Normal file
View File

@@ -0,0 +1,37 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 31
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability ImageBasic
OpCapability LiteralSampler
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %read_image_test "read_image_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%index_0 = OpConstant %uint 0
%void = OpTypeVoid
%v4uint = OpTypeVector %uint 4
%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint
%image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly
%7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t
%sampler_t = OpTypeSampler
%sampledimage_t = OpTypeSampledImage %image2d_t
%v2uint = OpTypeVector %uint 2
%float = OpTypeFloat 32
%sampler = OpConstantSampler %sampler_t None 0 Nearest
%coord_0_0 = OpConstantNull %v2uint
%float_0 = OpConstant %float 0
%read_image_test = OpFunction %void None %7
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint
%image = OpFunctionParameter %image2d_t
%entry = OpLabel
%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler
%call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|ZeroExtend %float_0
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0
OpStore %arrayidx %call Aligned 16
OpReturn
OpFunctionEnd

39
test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 Normal file
View File

@@ -0,0 +1,39 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 31
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpCapability ImageBasic
OpCapability LiteralSampler
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %read_image_test "read_image_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%index_0 = OpConstant %ulong 0
%void = OpTypeVoid
%v4uint = OpTypeVector %uint 4
%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint
%image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly
%7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t
%sampler_t = OpTypeSampler
%sampledimage_t = OpTypeSampledImage %image2d_t
%v2uint = OpTypeVector %uint 2
%float = OpTypeFloat 32
%sampler = OpConstantSampler %sampler_t None 0 Nearest
%coord_0_0 = OpConstantNull %v2uint
%float_0 = OpConstant %float 0
%read_image_test = OpFunction %void None %7
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint
%image = OpFunctionParameter %image2d_t
%entry = OpLabel
%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler
%call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|ZeroExtend %float_0
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0
OpStore %arrayidx %call Aligned 16
OpReturn
OpFunctionEnd

43
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 Normal file
View File

@@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc IterationMultiple 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

46
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 Normal file
View File

@@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%ulong_0 = OpConstant %ulong 0
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc IterationMultiple 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

43
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 Normal file
View File

@@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc MaxIterations 16
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

46
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 Normal file
View File

@@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%ulong_0 = OpConstant %ulong 0
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc MaxIterations 16
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

43
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 Normal file
View File

@@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc MinIterations 4
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

46
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 Normal file
View File

@@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%ulong_0 = OpConstant %ulong 0
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc MinIterations 4
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

43
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 Normal file
View File

@@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc PartialCount 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

46
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 Normal file
View File

@@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%ulong_0 = OpConstant %ulong 0
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc PartialCount 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

43
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 Normal file
View File

@@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc PeelCount 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

46
test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 Normal file
View File

@@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 37
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %loop_control_test "loop_control_test"
OpSource OpenCL_C 102000
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%ulong_0 = OpConstant %ulong 0
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint
%bool = OpTypeBool
%loop_control_test = OpFunction %void None %5
%dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%count = OpFunctionParameter %uint
%value = OpFunctionParameter %uint
%entry = OpLabel
OpBranch %for_cond
%for_cond = OpLabel
%sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc
%i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc
%cmp = OpSLessThan %bool %i_0 %count
OpLoopMerge %for_cond_cleanup %for_inc PeelCount 2
OpBranchConditional %cmp %for_body %for_cond_cleanup
%for_cond_cleanup = OpLabel
OpBranch %for_end
%for_body = OpLabel
%add = OpIAdd %uint %sum_0 %value
OpBranch %for_inc
%for_inc = OpLabel
%inc = OpIAdd %uint %i_0 %uint_1
OpBranch %for_cond
%for_end = OpLabel
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0
OpStore %arrayidx %sum_0 Aligned 4
OpReturn
OpFunctionEnd

211
test_conformance/spirv_new/test_spirv_14.cpp Normal file
View File

@@ -0,0 +1,211 @@
//
// Copyright (c) 2024 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "testBase.h"
#include "spirvInfo.hpp"
#include "types.hpp"
#include <string>
#include <vector>
static int test_image_operand_helper(cl_device_id deviceID, cl_context context,
cl_command_queue queue, bool signExtend)
{
const char* filename = signExtend ? "spv1.4/image_operand_signextend"
: "spv1.4/image_operand_zeroextend";
cl_image_format image_format = {
CL_RGBA,
signExtend ? CL_SIGNED_INT8 : CL_UNSIGNED_INT8,
};
cl_int error = CL_SUCCESS;
std::vector<cl_uchar> imgData({ 0x1, 0x80, 0xFF, 0x0 });
std::vector<cl_uint> expected;
for (auto v : imgData)
{
if (signExtend)
{
expected.push_back((cl_int)(cl_char)v);
}
else
{
expected.push_back(v);
}
}
clProgramWrapper prog;
error = get_program_with_il(prog, deviceID, context, filename);
SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
clKernelWrapper kernel = clCreateKernel(prog, "read_image_test", &error);
SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
std::vector<cl_uint> h_dst({ 0, 0, 0, 0 });
clMemWrapper dst =
clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
h_dst.size() * sizeof(cl_uint), h_dst.data(), &error);
SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
clMemWrapper src =
clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
&image_format, 1, 1, 0, imgData.data(), &error);
SPIRV_CHECK_ERROR(error, "Failed to create src image");
error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
error |= clSetKernelArg(kernel, 1, sizeof(src), &src);
SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
size_t global = 1;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0,
NULL, NULL);
SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0,
h_dst.size() * sizeof(cl_uint), h_dst.data(), 0,
NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
if (h_dst != expected)
{
log_error("Mismatch! Got: %u, %u, %u, %u, Wanted: %u, %u, %u, %u\n",
h_dst[0], h_dst[1], h_dst[2], h_dst[3], expected[0],
expected[1], expected[2], expected[3]);
return TEST_FAIL;
}
return TEST_PASS;
}
TEST_SPIRV_FUNC(spirv14_image_operand_signextend)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_image_operand_helper(deviceID, context, queue, true);
}
TEST_SPIRV_FUNC(spirv14_image_operand_zeroextend)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_image_operand_helper(deviceID, context, queue, false);
}
static int test_loop_control_helper(cl_device_id deviceID, cl_context context,
cl_command_queue queue,
const char* filename)
{
const int count = 10;
const int value = 5;
cl_int error = CL_SUCCESS;
clProgramWrapper prog;
std::string full_filename = "spv1.4/" + std::string(filename);
error = get_program_with_il(prog, deviceID, context, full_filename.c_str());
SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
clKernelWrapper kernel = clCreateKernel(prog, "loop_control_test", &error);
SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
int h_dst = 0;
clMemWrapper dst = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
sizeof(h_dst), &h_dst, &error);
SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
error |= clSetKernelArg(kernel, 1, sizeof(count), &count);
error |= clSetKernelArg(kernel, 2, sizeof(value), &value);
SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
size_t global = 1;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0,
NULL, NULL);
SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(h_dst), &h_dst,
0, NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
if (h_dst != count * value)
{
log_error("Mismatch! Got: %i, Wanted: %i\n", h_dst, count * value);
return TEST_FAIL;
}
return TEST_PASS;
}
TEST_SPIRV_FUNC(spirv14_loop_control_miniterations)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_loop_control_helper(deviceID, context, queue,
"loop_control_miniterations");
}
TEST_SPIRV_FUNC(spirv14_loop_control_maxiterations)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_loop_control_helper(deviceID, context, queue,
"loop_control_maxiterations");
}
TEST_SPIRV_FUNC(spirv14_loop_control_iterationmultiple)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_loop_control_helper(deviceID, context, queue,
"loop_control_iterationmultiple");
}
TEST_SPIRV_FUNC(spirv14_loop_control_peelcount)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_loop_control_helper(deviceID, context, queue,
"loop_control_peelcount");
}
TEST_SPIRV_FUNC(spirv14_loop_control_partialcount)
{
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
{
log_info("SPIR-V 1.4 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}
return test_loop_control_helper(deviceID, context, queue,
"loop_control_partialcount");
}