From 0052af2227c5980e570ceaf9a763819585e19a74 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 27 Feb 2024 09:57:52 -0800 Subject: [PATCH] add tests for cl_khr_expect_assume (#1888) * initial support for cl_khr_expect_assume Tests expect with 64-bit SPIR-V binaries. * add support for assume testing with 64-bit binaries * add 32-bit SPIR-V files * fix formatting * address review comments --- .../spirv_new/spirv_asm/assume.spvasm32 | 37 ++++ .../spirv_new/spirv_asm/assume.spvasm64 | 39 ++++ .../spirv_new/spirv_asm/expect_char.spvasm32 | 85 +++++++++ .../spirv_new/spirv_asm/expect_char.spvasm64 | 86 +++++++++ .../spirv_new/spirv_asm/expect_int.spvasm32 | 83 +++++++++ .../spirv_new/spirv_asm/expect_int.spvasm64 | 85 +++++++++ .../spirv_new/spirv_asm/expect_long.spvasm32 | 85 +++++++++ .../spirv_new/spirv_asm/expect_long.spvasm64 | 84 +++++++++ .../spirv_new/spirv_asm/expect_short.spvasm32 | 85 +++++++++ .../spirv_new/spirv_asm/expect_short.spvasm64 | 86 +++++++++ .../spirv_new/test_cl_khr_expect_assume.cpp | 176 ++++++++++++++++++ 11 files changed, 931 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/assume.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/assume.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_char.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_char.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_long.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_long.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_short.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/expect_short.spvasm64 create mode 100644 test_conformance/spirv_new/test_cl_khr_expect_assume.cpp diff --git a/test_conformance/spirv_new/spirv_asm/assume.spvasm32 b/test_conformance/spirv_new/spirv_asm/assume.spvasm32 new file mode 100644 index 0000000000..bad59c2266 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/assume.spvasm32 @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %test_assume "test_assume" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %void = OpTypeVoid + %bool = OpTypeBool + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_0 = OpConstantNull %uint +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %functype = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%test_assume = OpFunction %void None %functype + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %value = OpFunctionParameter %uint + %13 = OpLabel + %test = OpSGreaterThan %bool %value %uint_0 + OpAssumeTrueKHR %test + %global_id = OpLoad %v3uint %gl_GlobalInvocationID Aligned 32 + %gid_0 = OpCompositeExtract %uint %global_id 0 + %dst_gid_0 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %gid_0 + OpStore %dst_gid_0 %value Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/assume.spvasm64 b/test_conformance/spirv_new/spirv_asm/assume.spvasm64 new file mode 100644 index 0000000000..da33eed248 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/assume.spvasm64 @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %test_assume "test_assume" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %void = OpTypeVoid + %bool = OpTypeBool + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %v3ulong = OpTypeVector %ulong 3 + %uint_0 = OpConstantNull %uint +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %functype = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input +%test_assume = OpFunction %void None %functype + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %value = OpFunctionParameter %uint + %13 = OpLabel + %test = OpSGreaterThan %bool %value %uint_0 + OpAssumeTrueKHR %test + %global_id = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %gid_0 = OpCompositeExtract %ulong %global_id 0 + %dst_gid_0 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %gid_0 + OpStore %dst_gid_0 %value Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_char.spvasm32 b/test_conformance/spirv_new/spirv_asm/expect_char.spvasm32 new file mode 100644 index 0000000000..496fe08d71 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_char.spvasm32 @@ -0,0 +1,85 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int8 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %expect_char "expect_char" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %uchar = OpTypeInt 8 0 + %uchar2 = OpTypeVector %uchar 2 + %uchar3 = OpTypeVector %uchar 3 + %uchar4 = OpTypeVector %uchar 4 + %uchar8 = OpTypeVector %uchar 8 + %uchar16 = OpTypeVector %uchar 16 + %uint = OpTypeInt 32 0 + %uchar_0 = OpConstantNull %uchar + %uchar2_0 = OpConstantNull %uchar2 + %uchar3_0 = OpConstantNull %uchar3 + %uchar4_0 = OpConstantNull %uchar4 + %uchar8_0 = OpConstantNull %uchar8 + %uchar16_0 = OpConstantNull %uchar16 + %index_1 = OpConstant %uint 1 + %index_2 = OpConstant %uint 2 + %index_3 = OpConstant %uint 3 + %index_4 = OpConstant %uint 4 + %index_5 = OpConstant %uint 5 +%_ptr_CrossWorkgroup_uchar16 = OpTypePointer CrossWorkgroup %uchar16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar16 %uchar +%expect_char = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uchar16 + %value = OpFunctionParameter %uchar + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %uchar2 %value %uchar2_0 0 + ; scalar expect: + ; char v1e = __builtin_expect(value, 0); + ; dst[0] = (char16)(v1e, 0, ...); + %v1e = OpExpectKHR %uchar %value %uchar_0 + %v1v16 = OpCompositeInsert %uchar16 %v1e %uchar16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; char2 v2 = (char2)(value); + ; char2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (char16)(v2e, 0, ...); + %v2 = OpVectorShuffle %uchar2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %uchar2 %v2 %uchar2_0 + %v2v16 = OpVectorShuffle %uchar16 %v2e %uchar2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %uchar3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %uchar3 %v3 %uchar3_0 + %v3v16 = OpVectorShuffle %uchar16 %v3e %uchar2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %uchar4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %uchar4 %v4 %uchar4_0 + %v4v16 = OpVectorShuffle %uchar16 %v4e %uchar2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %uchar8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %uchar8 %v8 %uchar8_0 + %v8v16 = OpVectorShuffle %uchar16 %v8e %uchar2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %uchar16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %uchar16 %v16 %uchar16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_char.spvasm64 b/test_conformance/spirv_new/spirv_asm/expect_char.spvasm64 new file mode 100644 index 0000000000..c54c97fce0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_char.spvasm64 @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int8 + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %expect_char "expect_char" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %uchar = OpTypeInt 8 0 + %uchar2 = OpTypeVector %uchar 2 + %uchar3 = OpTypeVector %uchar 3 + %uchar4 = OpTypeVector %uchar 4 + %uchar8 = OpTypeVector %uchar 8 + %uchar16 = OpTypeVector %uchar 16 + %ulong = OpTypeInt 64 0 + %uchar_0 = OpConstantNull %uchar + %uchar2_0 = OpConstantNull %uchar2 + %uchar3_0 = OpConstantNull %uchar3 + %uchar4_0 = OpConstantNull %uchar4 + %uchar8_0 = OpConstantNull %uchar8 + %uchar16_0 = OpConstantNull %uchar16 + %index_1 = OpConstant %ulong 1 + %index_2 = OpConstant %ulong 2 + %index_3 = OpConstant %ulong 3 + %index_4 = OpConstant %ulong 4 + %index_5 = OpConstant %ulong 5 +%_ptr_CrossWorkgroup_uchar16 = OpTypePointer CrossWorkgroup %uchar16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar16 %uchar +%expect_char = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uchar16 + %value = OpFunctionParameter %uchar + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %uchar2 %value %uchar2_0 0 + ; scalar expect: + ; char v1e = __builtin_expect(value, 0); + ; dst[0] = (char16)(v1e, 0, ...); + %v1e = OpExpectKHR %uchar %value %uchar_0 + %v1v16 = OpCompositeInsert %uchar16 %v1e %uchar16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; char2 v2 = (char2)(value); + ; char2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (char16)(v2e, 0, ...); + %v2 = OpVectorShuffle %uchar2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %uchar2 %v2 %uchar2_0 + %v2v16 = OpVectorShuffle %uchar16 %v2e %uchar2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %uchar3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %uchar3 %v3 %uchar3_0 + %v3v16 = OpVectorShuffle %uchar16 %v3e %uchar2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %uchar4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %uchar4 %v4 %uchar4_0 + %v4v16 = OpVectorShuffle %uchar16 %v4e %uchar2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %uchar8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %uchar8 %v8 %uchar8_0 + %v8v16 = OpVectorShuffle %uchar16 %v8e %uchar2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %uchar16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %uchar16 %v16 %uchar16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/expect_int.spvasm32 new file mode 100644 index 0000000000..3334ae52d1 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_int.spvasm32 @@ -0,0 +1,83 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %expect_int "expect_int" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uint2 = OpTypeVector %uint 2 + %uint3 = OpTypeVector %uint 3 + %uint4 = OpTypeVector %uint 4 + %uint8 = OpTypeVector %uint 8 + %uint16 = OpTypeVector %uint 16 + %uint_0 = OpConstantNull %uint + %uint2_0 = OpConstantNull %uint2 + %uint3_0 = OpConstantNull %uint3 + %uint4_0 = OpConstantNull %uint4 + %uint8_0 = OpConstantNull %uint8 + %uint16_0 = OpConstantNull %uint16 + %index_1 = OpConstant %uint 1 + %index_2 = OpConstant %uint 2 + %index_3 = OpConstant %uint 3 + %index_4 = OpConstant %uint 4 + %index_5 = OpConstant %uint 5 +%_ptr_CrossWorkgroup_uint16 = OpTypePointer CrossWorkgroup %uint16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint16 %uint + %expect_int = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint16 + %value = OpFunctionParameter %uint + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %uint2 %value %uint2_0 0 + ; scalar expect: + ; int v1e = __builtin_expect(value, 0); + ; dst[0] = (int16)(v1e, 0, ...); + %v1e = OpExpectKHR %uint %value %uint_0 + %v1v16 = OpCompositeInsert %uint16 %v1e %uint16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; int2 v2 = (int2)(value); + ; int2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (int16)(v2e, 0, ...); + %v2 = OpVectorShuffle %uint2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %uint2 %v2 %uint2_0 + %v2v16 = OpVectorShuffle %uint16 %v2e %uint2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %uint3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %uint3 %v3 %uint3_0 + %v3v16 = OpVectorShuffle %uint16 %v3e %uint2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %uint4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %uint4 %v4 %uint4_0 + %v4v16 = OpVectorShuffle %uint16 %v4e %uint2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %uint8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %uint8 %v8 %uint8_0 + %v8v16 = OpVectorShuffle %uint16 %v8e %uint2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %uint16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %uint16 %v16 %uint16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/expect_int.spvasm64 new file mode 100644 index 0000000000..9b54bf7908 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_int.spvasm64 @@ -0,0 +1,85 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %expect_int "expect_int" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %uint2 = OpTypeVector %uint 2 + %uint3 = OpTypeVector %uint 3 + %uint4 = OpTypeVector %uint 4 + %uint8 = OpTypeVector %uint 8 + %uint16 = OpTypeVector %uint 16 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstantNull %uint + %uint2_0 = OpConstantNull %uint2 + %uint3_0 = OpConstantNull %uint3 + %uint4_0 = OpConstantNull %uint4 + %uint8_0 = OpConstantNull %uint8 + %uint16_0 = OpConstantNull %uint16 + %index_1 = OpConstant %ulong 1 + %index_2 = OpConstant %ulong 2 + %index_3 = OpConstant %ulong 3 + %index_4 = OpConstant %ulong 4 + %index_5 = OpConstant %ulong 5 +%_ptr_CrossWorkgroup_uint16 = OpTypePointer CrossWorkgroup %uint16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint16 %uint + %expect_int = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint16 + %value = OpFunctionParameter %uint + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %uint2 %value %uint2_0 0 + ; scalar expect: + ; int v1e = __builtin_expect(value, 0); + ; dst[0] = (int16)(v1e, 0, ...); + %v1e = OpExpectKHR %uint %value %uint_0 + %v1v16 = OpCompositeInsert %uint16 %v1e %uint16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; int2 v2 = (int2)(value); + ; int2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (int16)(v2e, 0, ...); + %v2 = OpVectorShuffle %uint2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %uint2 %v2 %uint2_0 + %v2v16 = OpVectorShuffle %uint16 %v2e %uint2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %uint3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %uint3 %v3 %uint3_0 + %v3v16 = OpVectorShuffle %uint16 %v3e %uint2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %uint4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %uint4 %v4 %uint4_0 + %v4v16 = OpVectorShuffle %uint16 %v4e %uint2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %uint8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %uint8 %v8 %uint8_0 + %v8v16 = OpVectorShuffle %uint16 %v8e %uint2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %uint16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %uint16 %v16 %uint16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_long.spvasm32 b/test_conformance/spirv_new/spirv_asm/expect_long.spvasm32 new file mode 100644 index 0000000000..1028aad074 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_long.spvasm32 @@ -0,0 +1,85 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %expect_long "expect_long" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %ulong2 = OpTypeVector %ulong 2 + %ulong3 = OpTypeVector %ulong 3 + %ulong4 = OpTypeVector %ulong 4 + %ulong8 = OpTypeVector %ulong 8 + %ulong16 = OpTypeVector %ulong 16 + %ulong_0 = OpConstantNull %ulong + %ulong2_0 = OpConstantNull %ulong2 + %ulong3_0 = OpConstantNull %ulong3 + %ulong4_0 = OpConstantNull %ulong4 + %ulong8_0 = OpConstantNull %ulong8 + %ulong16_0 = OpConstantNull %ulong16 + %index_1 = OpConstant %uint 1 + %index_2 = OpConstant %uint 2 + %index_3 = OpConstant %uint 3 + %index_4 = OpConstant %uint 4 + %index_5 = OpConstant %uint 5 +%_ptr_CrossWorkgroup_ulong16 = OpTypePointer CrossWorkgroup %ulong16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong16 %ulong +%expect_long = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_ulong16 + %value = OpFunctionParameter %ulong + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %ulong2 %value %ulong2_0 0 + ; scalar expect: + ; long v1e = __builtin_expect(value, 0); + ; dst[0] = (long16)(v1e, 0, ...); + %v1e = OpExpectKHR %ulong %value %ulong_0 + %v1v16 = OpCompositeInsert %ulong16 %v1e %ulong16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; long2 v2 = (long2)(value); + ; long2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (long16)(v2e, 0, ...); + %v2 = OpVectorShuffle %ulong2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %ulong2 %v2 %ulong2_0 + %v2v16 = OpVectorShuffle %ulong16 %v2e %ulong2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %ulong3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %ulong3 %v3 %ulong3_0 + %v3v16 = OpVectorShuffle %ulong16 %v3e %ulong2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %ulong4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %ulong4 %v4 %ulong4_0 + %v4v16 = OpVectorShuffle %ulong16 %v4e %ulong2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %ulong8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %ulong8 %v8 %ulong8_0 + %v8v16 = OpVectorShuffle %ulong16 %v8e %ulong2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %ulong16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %ulong16 %v16 %ulong16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_long.spvasm64 b/test_conformance/spirv_new/spirv_asm/expect_long.spvasm64 new file mode 100644 index 0000000000..4453b5649f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_long.spvasm64 @@ -0,0 +1,84 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %expect_long "expect_long" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %ulong2 = OpTypeVector %ulong 2 + %ulong3 = OpTypeVector %ulong 3 + %ulong4 = OpTypeVector %ulong 4 + %ulong8 = OpTypeVector %ulong 8 + %ulong16 = OpTypeVector %ulong 16 + %ulong_0 = OpConstantNull %ulong + %ulong2_0 = OpConstantNull %ulong2 + %ulong3_0 = OpConstantNull %ulong3 + %ulong4_0 = OpConstantNull %ulong4 + %ulong8_0 = OpConstantNull %ulong8 + %ulong16_0 = OpConstantNull %ulong16 + %index_1 = OpConstant %ulong 1 + %index_2 = OpConstant %ulong 2 + %index_3 = OpConstant %ulong 3 + %index_4 = OpConstant %ulong 4 + %index_5 = OpConstant %ulong 5 +%_ptr_CrossWorkgroup_ulong16 = OpTypePointer CrossWorkgroup %ulong16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong16 %ulong +%expect_long = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_ulong16 + %value = OpFunctionParameter %ulong + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %ulong2 %value %ulong2_0 0 + ; scalar expect: + ; long v1e = __builtin_expect(value, 0); + ; dst[0] = (long16)(v1e, 0, ...); + %v1e = OpExpectKHR %ulong %value %ulong_0 + %v1v16 = OpCompositeInsert %ulong16 %v1e %ulong16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; long2 v2 = (long2)(value); + ; long2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (long16)(v2e, 0, ...); + %v2 = OpVectorShuffle %ulong2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %ulong2 %v2 %ulong2_0 + %v2v16 = OpVectorShuffle %ulong16 %v2e %ulong2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %ulong3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %ulong3 %v3 %ulong3_0 + %v3v16 = OpVectorShuffle %ulong16 %v3e %ulong2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %ulong4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %ulong4 %v4 %ulong4_0 + %v4v16 = OpVectorShuffle %ulong16 %v4e %ulong2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %ulong8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %ulong8 %v8 %ulong8_0 + %v8v16 = OpVectorShuffle %ulong16 %v8e %ulong2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %ulong16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %ulong16 %v16 %ulong16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ulong16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_short.spvasm32 b/test_conformance/spirv_new/spirv_asm/expect_short.spvasm32 new file mode 100644 index 0000000000..c7b008a809 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_short.spvasm32 @@ -0,0 +1,85 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int16 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %expect_short "expect_short" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %ushort = OpTypeInt 16 0 + %ushort2 = OpTypeVector %ushort 2 + %ushort3 = OpTypeVector %ushort 3 + %ushort4 = OpTypeVector %ushort 4 + %ushort8 = OpTypeVector %ushort 8 + %ushort16 = OpTypeVector %ushort 16 + %uint = OpTypeInt 32 0 + %ushort_0 = OpConstantNull %ushort + %ushort2_0 = OpConstantNull %ushort2 + %ushort3_0 = OpConstantNull %ushort3 + %ushort4_0 = OpConstantNull %ushort4 + %ushort8_0 = OpConstantNull %ushort8 + %ushort16_0 = OpConstantNull %ushort16 + %index_1 = OpConstant %uint 1 + %index_2 = OpConstant %uint 2 + %index_3 = OpConstant %uint 3 + %index_4 = OpConstant %uint 4 + %index_5 = OpConstant %uint 5 +%_ptr_CrossWorkgroup_ushort16 = OpTypePointer CrossWorkgroup %ushort16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_ushort16 %ushort + %expect_short = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_ushort16 + %value = OpFunctionParameter %ushort + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %ushort2 %value %ushort2_0 0 + ; scalar expect: + ; short v1e = __builtin_expect(value, 0); + ; dst[0] = (short16)(v1e, 0, ...); + %v1e = OpExpectKHR %ushort %value %ushort_0 + %v1v16 = OpCompositeInsert %ushort16 %v1e %ushort16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; short2 v2 = (short2)(value); + ; short2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (short16)(v2e, 0, ...); + %v2 = OpVectorShuffle %ushort2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %ushort2 %v2 %ushort2_0 + %v2v16 = OpVectorShuffle %ushort16 %v2e %ushort2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %ushort3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %ushort3 %v3 %ushort3_0 + %v3v16 = OpVectorShuffle %ushort16 %v3e %ushort2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %ushort4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %ushort4 %v4 %ushort4_0 + %v4v16 = OpVectorShuffle %ushort16 %v4e %ushort2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %ushort8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %ushort8 %v8 %ushort8_0 + %v8v16 = OpVectorShuffle %ushort16 %v8e %ushort2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %ushort16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %ushort16 %v16 %ushort16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/expect_short.spvasm64 b/test_conformance/spirv_new/spirv_asm/expect_short.spvasm64 new file mode 100644 index 0000000000..b9884b6a43 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/expect_short.spvasm64 @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 58 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Vector16 + OpCapability Int16 + OpCapability Int64 + OpCapability ExpectAssumeKHR + OpExtension "SPV_KHR_expect_assume" + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %expect_short "expect_short" + OpSource OpenCL_C 102000 + OpDecorate %dst FuncParamAttr NoCapture + OpDecorate %dst Alignment 64 + %void = OpTypeVoid + %ushort = OpTypeInt 16 0 + %ushort2 = OpTypeVector %ushort 2 + %ushort3 = OpTypeVector %ushort 3 + %ushort4 = OpTypeVector %ushort 4 + %ushort8 = OpTypeVector %ushort 8 + %ushort16 = OpTypeVector %ushort 16 + %ulong = OpTypeInt 64 0 + %ushort_0 = OpConstantNull %ushort + %ushort2_0 = OpConstantNull %ushort2 + %ushort3_0 = OpConstantNull %ushort3 + %ushort4_0 = OpConstantNull %ushort4 + %ushort8_0 = OpConstantNull %ushort8 + %ushort16_0 = OpConstantNull %ushort16 + %index_1 = OpConstant %ulong 1 + %index_2 = OpConstant %ulong 2 + %index_3 = OpConstant %ulong 3 + %index_4 = OpConstant %ulong 4 + %index_5 = OpConstant %ulong 5 +%_ptr_CrossWorkgroup_ushort16 = OpTypePointer CrossWorkgroup %ushort16 + %6 = OpTypeFunction %void %_ptr_CrossWorkgroup_ushort16 %ushort + %expect_short = OpFunction %void None %6 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_ushort16 + %value = OpFunctionParameter %ushort + %10 = OpLabel + ; setup + %value_vec = OpCompositeInsert %ushort2 %value %ushort2_0 0 + ; scalar expect: + ; short v1e = __builtin_expect(value, 0); + ; dst[0] = (short16)(v1e, 0, ...); + %v1e = OpExpectKHR %ushort %value %ushort_0 + %v1v16 = OpCompositeInsert %ushort16 %v1e %ushort16_0 0 + OpStore %dst %v1v16 Aligned 64 + ; vec2 expect: + ; short2 v2 = (short2)(value); + ; short2 v2e = __builtin_expect(v2, 0); + ; dst[1] = (short16)(v2e, 0, ...); + %v2 = OpVectorShuffle %ushort2 %value_vec %value_vec 0 0 + %v2e = OpExpectKHR %ushort2 %v2 %ushort2_0 + %v2v16 = OpVectorShuffle %ushort16 %v2e %ushort2_0 0 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 + %dst_1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_1 + OpStore %dst_1 %v2v16 Aligned 64 + ; vec3 expect + %v3 = OpVectorShuffle %ushort3 %value_vec %value_vec 0 0 0 + %v3e = OpExpectKHR %ushort3 %v3 %ushort3_0 + %v3v16 = OpVectorShuffle %ushort16 %v3e %ushort2_0 0 1 2 3 3 3 3 3 3 3 3 3 3 3 3 3 + %dst_2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_2 + OpStore %dst_2 %v3v16 Aligned 64 + ; vec4 expect + %v4 = OpVectorShuffle %ushort4 %value_vec %value_vec 0 0 0 0 + %v4e = OpExpectKHR %ushort4 %v4 %ushort4_0 + %v4v16 = OpVectorShuffle %ushort16 %v4e %ushort2_0 0 1 2 3 4 4 4 4 4 4 4 4 4 4 4 4 + %dst_3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_3 + OpStore %dst_3 %v4v16 Aligned 64 + ; vec8 expect + %v8 = OpVectorShuffle %ushort8 %value_vec %value_vec 0 0 0 0 0 0 0 0 + %v8e = OpExpectKHR %ushort8 %v8 %ushort8_0 + %v8v16 = OpVectorShuffle %ushort16 %v8e %ushort2_0 0 1 2 3 4 5 6 7 8 8 8 8 8 8 8 8 + %dst_4 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_4 + OpStore %dst_4 %v8v16 Aligned 64 + ; vec16 expect + %v16 = OpVectorShuffle %ushort16 %value_vec %value_vec 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + %v16e = OpExpectKHR %ushort16 %v16 %ushort16_0 + %dst_5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_ushort16 %dst %index_5 + OpStore %dst_5 %v16e Aligned 64 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_cl_khr_expect_assume.cpp b/test_conformance/spirv_new/test_cl_khr_expect_assume.cpp new file mode 100644 index 0000000000..05c5068a03 --- /dev/null +++ b/test_conformance/spirv_new/test_cl_khr_expect_assume.cpp @@ -0,0 +1,176 @@ +// +// 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 + +template struct TestInfo +{ +}; +template <> struct TestInfo +{ + static constexpr const char* typeName = "char"; + static constexpr const char* testName = "expect_char"; +}; +template <> struct TestInfo +{ + static constexpr const char* typeName = "short"; + static constexpr const char* testName = "expect_short"; +}; +template <> struct TestInfo +{ + static constexpr const char* typeName = "int"; + static constexpr const char* testName = "expect_int"; +}; +template <> struct TestInfo +{ + static constexpr const char* typeName = "long"; + static constexpr const char* testName = "expect_long"; +}; + +template +static int test_expect_type(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + log_info(" testing type %s\n", TestInfo::typeName); + + const T value = 42; + cl_int error = CL_SUCCESS; + + std::vector vecSizes({ 1, 2, 3, 4, 8, 16 }); + std::vector testData; + testData.reserve(16 * vecSizes.size()); + + for (auto v : vecSizes) + { + size_t i; + for (i = 0; i < v; i++) + { + testData.push_back(value); + } + for (; i < 16; i++) + { + testData.push_back(0); + } + } + + clMemWrapper dst = + clCreateBuffer(context, CL_MEM_WRITE_ONLY, testData.size() * sizeof(T), + nullptr, &error); + test_error(error, "Unable to create destination buffer"); + + clProgramWrapper prog; + error = get_program_with_il(prog, device, context, TestInfo::testName); + test_error(error, "Unable to build SPIR-V program"); + + clKernelWrapper kernel = + clCreateKernel(prog, TestInfo::testName, &error); + test_error(error, "Unable to create SPIR-V kernel"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(value), &value); + test_error(error, "Unable to set kernel arguments"); + + size_t global = 1; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + test_error(error, "Unable to enqueue kernel"); + + std::vector resData(testData.size()); + error = + clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, resData.size() * sizeof(T), + resData.data(), 0, NULL, NULL); + test_error(error, "Unable to read destination buffer"); + + if (resData != testData) + { + log_error("Values do not match!\n"); + return TEST_FAIL; + } + + return TEST_PASS; +} + +TEST_SPIRV_FUNC(op_expect) +{ + if (!is_extension_available(deviceID, "cl_khr_expect_assume")) + { + log_info("cl_khr_expect_assume is not supported; skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + int result = TEST_PASS; + + result |= test_expect_type(deviceID, context, queue); + result |= test_expect_type(deviceID, context, queue); + result |= test_expect_type(deviceID, context, queue); + if (gHasLong) + { + result |= test_expect_type(deviceID, context, queue); + } + + return result; +} + +TEST_SPIRV_FUNC(op_assume) +{ + if (!is_extension_available(deviceID, "cl_khr_expect_assume")) + { + log_info("cl_khr_expect_assume is not supported; skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int error = CL_SUCCESS; + + clMemWrapper dst = + clCreateBuffer(context, 0, num_elements * sizeof(cl_int), NULL, &error); + test_error(error, "Unable to create destination buffer"); + + clProgramWrapper prog; + error = get_program_with_il(prog, deviceID, context, "assume"); + test_error(error, "Unable to build SPIR-V program"); + + clKernelWrapper kernel = clCreateKernel(prog, "test_assume", &error); + test_error(error, "Unable to create SPIR-V kernel"); + + const cl_int value = 42; + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(value), &value); + test_error(error, "Unable to set kernel arguments"); + + size_t global = num_elements; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + test_error(error, "Unable to enqueue kernel"); + + std::vector h_dst(num_elements); + error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, + h_dst.size() * sizeof(cl_int), h_dst.data(), 0, + NULL, NULL); + test_error(error, "Unable to read destination buffer"); + + for (int i = 0; i < num_elements; i++) + { + if (h_dst[i] != value) + { + log_error("Values do not match at location %d\n", i); + return TEST_FAIL; + } + } + + return TEST_PASS; +}