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

11 files changed, 931 insertions, 0 deletions

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 00000000..bad59c22
--- /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 00000000..da33eed2
--- /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 00000000..496fe08d
--- /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 00000000..c54c97fc
--- /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 00000000..3334ae52
--- /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 00000000..9b54bf79
--- /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 00000000..1028aad0
--- /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 00000000..4453b564
--- /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 00000000..c7b008a8
--- /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 00000000..b9884b6a
--- /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 00000000..05c5068a
--- /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 <vector>
+
+template <typename T> struct TestInfo
+{
+};
+template <> struct TestInfo<cl_char>
+{
+    static constexpr const char* typeName = "char";
+    static constexpr const char* testName = "expect_char";
+};
+template <> struct TestInfo<cl_short>
+{
+    static constexpr const char* typeName = "short";
+    static constexpr const char* testName = "expect_short";
+};
+template <> struct TestInfo<cl_int>
+{
+    static constexpr const char* typeName = "int";
+    static constexpr const char* testName = "expect_int";
+};
+template <> struct TestInfo<cl_long>
+{
+    static constexpr const char* typeName = "long";
+    static constexpr const char* testName = "expect_long";
+};
+
+template <typename T>
+static int test_expect_type(cl_device_id device, cl_context context,
+                            cl_command_queue queue)
+{
+    log_info("    testing type %s\n", TestInfo<T>::typeName);
+
+    const T value = 42;
+    cl_int error = CL_SUCCESS;
+
+    std::vector<size_t> vecSizes({ 1, 2, 3, 4, 8, 16 });
+    std::vector<T> 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<T>::testName);
+    test_error(error, "Unable to build SPIR-V program");
+
+    clKernelWrapper kernel =
+        clCreateKernel(prog, TestInfo<T>::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<T> 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<cl_char>(deviceID, context, queue);
+    result |= test_expect_type<cl_short>(deviceID, context, queue);
+    result |= test_expect_type<cl_int>(deviceID, context, queue);
+    if (gHasLong)
+    {
+        result |= test_expect_type<cl_long>(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<cl_int> 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;
+}