From 41cd9c6d9823c15bdd09d3910f580ea39c8b2867 Mon Sep 17 00:00:00 2001
From: Grzegorz Wawiorko <grzegorz.wawiorko@intel.com>
Date: Thu, 20 Aug 2020 01:39:06 +0200
Subject: [PATCH] New spirv_new test case OpSpecConstant  (#719)

* New spirv_new test case OpSpecConstant OpSpecConstantTrue OpSpecConstantfalse

* Register test case with minimal OpenCL version

* Specialization constant - add files uint/int cases

* Specialization constants - fix second build program

* Fix changes clang format.

* spirv_new - Make functions visible outside of main

* Fix clang format issues

* Fix int/uint 32 bit cases

* Fix minimal required version for test_op_spec_constant

* Remove not needed OpName. Update binaries

* Fix code format

* Remove op_spec_constant_int cases in spirv_new tests

* op_spec_constant - add simplified spirv files

* no redundant OpUConvert OpSConvert
* no redundant OpName
* instead of buffers scalar variable

* op_spec_constant - refactor to address review issues

* avoid using program that has already kernel program attached
* remove not used buffer
* Simplified test case - instead of buffers use scalar variable

* spirv_new - remove spec_const duplicated cases (singed versions)

* spirv_new - set clSetProgramSpecializationConstant before clBuildProgram

* Test spirv_new - fix spec const after rebase

* Test spirv_new - fix spec const set min version for bool type tests
---
 test_conformance/spirv_new/main.cpp           |  25 ++-
 test_conformance/spirv_new/procs.h            | 106 ++++++------
 .../op_spec_constant_double_simple.spvasm32   |  36 ++++
 .../op_spec_constant_double_simple.spvasm64   |  36 ++++
 .../op_spec_constant_false_simple.spvasm32    |  57 +++++++
 .../op_spec_constant_false_simple.spvasm64    |  57 +++++++
 .../op_spec_constant_float_simple.spvasm32    |  35 ++++
 .../op_spec_constant_float_simple.spvasm64    |  35 ++++
 .../op_spec_constant_half_simple.spvasm32     |  36 ++++
 .../op_spec_constant_half_simple.spvasm64     |  36 ++++
 .../op_spec_constant_true_simple.spvasm32     |  57 +++++++
 .../op_spec_constant_true_simple.spvasm64     |  57 +++++++
 .../op_spec_constant_uchar_simple.spvasm32    |  39 +++++
 .../op_spec_constant_uchar_simple.spvasm64    |  39 +++++
 .../op_spec_constant_uint_simple.spvasm32     |  35 ++++
 .../op_spec_constant_uint_simple.spvasm64     |  35 ++++
 .../op_spec_constant_ulong_simple.spvasm32    |  36 ++++
 .../op_spec_constant_ulong_simple.spvasm64    |  36 ++++
 .../op_spec_constant_ushort_simple.spvasm32   |  38 +++++
 .../op_spec_constant_ushort_simple.spvasm64   |  38 +++++
 .../op_spec_constant_double_simple.spv32      | Bin 0 -> 560 bytes
 .../op_spec_constant_double_simple.spv64      | Bin 0 -> 560 bytes
 .../op_spec_constant_false_simple.spv32       | Bin 0 -> 836 bytes
 .../op_spec_constant_false_simple.spv64       | Bin 0 -> 836 bytes
 .../op_spec_constant_float_simple.spv32       | Bin 0 -> 548 bytes
 .../op_spec_constant_float_simple.spv64       | Bin 0 -> 548 bytes
 .../op_spec_constant_half_simple.spv32        | Bin 0 -> 556 bytes
 .../op_spec_constant_half_simple.spv64        | Bin 0 -> 556 bytes
 .../op_spec_constant_true_simple.spv32        | Bin 0 -> 836 bytes
 .../op_spec_constant_true_simple.spv64        | Bin 0 -> 836 bytes
 .../op_spec_constant_uchar_simple.spv32       | Bin 0 -> 628 bytes
 .../op_spec_constant_uchar_simple.spv64       | Bin 0 -> 628 bytes
 .../op_spec_constant_uint_simple.spv32        | Bin 0 -> 552 bytes
 .../op_spec_constant_uint_simple.spv64        | Bin 0 -> 552 bytes
 .../op_spec_constant_ulong_simple.spv32       | Bin 0 -> 564 bytes
 .../op_spec_constant_ulong_simple.spv64       | Bin 0 -> 564 bytes
 .../op_spec_constant_ushort_simple.spv32      | Bin 0 -> 616 bytes
 .../op_spec_constant_ushort_simple.spv64      | Bin 0 -> 616 bytes
 .../spirv_new/test_op_spec_constant.cpp       | 159 ++++++++++++++++++
 39 files changed, 969 insertions(+), 59 deletions(-)
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32
 create mode 100644 test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv64
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32
 create mode 100644 test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64
 create mode 100644 test_conformance/spirv_new/test_op_spec_constant.cpp

diff --git a/test_conformance/spirv_new/main.cpp b/test_conformance/spirv_new/main.cpp
index cef4677a..68a1e02a 100644
--- a/test_conformance/spirv_new/main.cpp
+++ b/test_conformance/spirv_new/main.cpp
@@ -74,7 +74,8 @@ size_t spirvTestsRegistry::getNumTests()
     return testDefinitions.size();
 }
 
-void spirvTestsRegistry::addTestClass(baseTestClass *test, const char *testName)
+void spirvTestsRegistry::addTestClass(baseTestClass *test, const char *testName,
+                                      Version version)
 {
 
     testClasses.push_back(test);
@@ -135,20 +136,21 @@ static int offline_get_program_with_il(clProgramWrapper &prog,
     return err;
 }
 
-int get_program_with_il(clProgramWrapper &prog,
-                        const cl_device_id deviceID,
-                        const cl_context context,
-                        const char *prog_name)
+int get_program_with_il(clProgramWrapper &prog, const cl_device_id deviceID,
+                        const cl_context context, const char *prog_name,
+                        spec_const spec_const_def)
 {
     cl_int err = 0;
-    if (gCompilationMode == kBinary) {
+    if (gCompilationMode == kBinary)
+    {
         return offline_get_program_with_il(prog, deviceID, context, prog_name);
     }
 
     std::vector<unsigned char> buffer_vec = readSPIRV(prog_name);
 
     int file_bytes = buffer_vec.size();
-    if (file_bytes == 0) {
+    if (file_bytes == 0)
+    {
         log_error("File %s not found\n", prog_name);
         return -1;
     }
@@ -159,6 +161,15 @@ int get_program_with_il(clProgramWrapper &prog,
         prog = clCreateProgramWithIL(context, buffer, file_bytes, &err);
         SPIRV_CHECK_ERROR(
             err, "Failed to create program with clCreateProgramWithIL");
+
+        if (spec_const_def.spec_value != NULL)
+        {
+            err = clSetProgramSpecializationConstant(
+                prog, spec_const_def.spec_id, spec_const_def.spec_size,
+                spec_const_def.spec_value);
+            SPIRV_CHECK_ERROR(
+                err, "Failed to run clSetProgramSpecializationConstant");
+        }
     }
     else
     {
diff --git a/test_conformance/spirv_new/procs.h b/test_conformance/spirv_new/procs.h
index aa44cdda..fd697a43 100644
--- a/test_conformance/spirv_new/procs.h
+++ b/test_conformance/spirv_new/procs.h
@@ -1,14 +1,16 @@
 /******************************************************************
 Copyright (c) 2016 The Khronos Group Inc. All Rights Reserved.
 
-This code is protected by copyright laws and contains material proprietary to the Khronos Group, Inc.
-This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not be disclosed in whole or in part to
-third parties, and may not be reproduced, republished, distributed, transmitted, displayed,
-broadcast or otherwise exploited in any manner without the express prior written permission
-of Khronos Group. The receipt or possession of this code does not convey any rights to reproduce,
-disclose, or distribute its contents, or to manufacture, use, or sell anything that it may describe,
-in whole or in part other than under the terms of the Khronos Adopters Agreement
-or Khronos Conformance Test Source License Agreement as executed between Khronos and the recipient.
+This code is protected by copyright laws and contains material proprietary to
+the Khronos Group, Inc. This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not
+be disclosed in whole or in part to third parties, and may not be reproduced,
+republished, distributed, transmitted, displayed, broadcast or otherwise
+exploited in any manner without the express prior written permission of Khronos
+Group. The receipt or possession of this code does not convey any rights to
+reproduce, disclose, or distribute its contents, or to manufacture, use, or sell
+anything that it may describe, in whole or in part other than under the terms of
+the Khronos Adopters Agreement or Khronos Conformance Test Source License
+Agreement as executed between Khronos and the recipient.
 ******************************************************************/
 
 #pragma once
@@ -24,16 +26,17 @@ or Khronos Conformance Test Source License Agreement as executed between Khronos
 
 #include <vector>
 
-#define SPIRV_CHECK_ERROR(err, fmt, ...) do {               \
-        if (err == CL_SUCCESS) break;                       \
-        log_error("%s(%d): Error %d\n" fmt "\n",            \
-                  __FILE__, __LINE__, err, ##__VA_ARGS__);  \
-        return -1;                                          \
-    } while(0)
+#define SPIRV_CHECK_ERROR(err, fmt, ...)                                       \
+    do                                                                         \
+    {                                                                          \
+        if (err == CL_SUCCESS) break;                                          \
+        log_error("%s(%d): Error %d\n" fmt "\n", __FILE__, __LINE__, err,      \
+                  ##__VA_ARGS__);                                              \
+        return -1;                                                             \
+    } while (0)
 
 
-class baseTestClass
-{
+class baseTestClass {
 public:
     baseTestClass() {}
     virtual basefn getFunction() = 0;
@@ -45,54 +48,53 @@ private:
     std::vector<test_definition> testDefinitions;
 
 public:
-
-    static spirvTestsRegistry& getInstance();
+    static spirvTestsRegistry &getInstance();
 
     test_definition *getTestDefinitions();
 
     size_t getNumTests();
 
-    void addTestClass(baseTestClass *test, const char *testName);
+    void addTestClass(baseTestClass *test, const char *testName,
+                      Version version);
     spirvTestsRegistry() {}
 };
 
-template<typename T>
-T* createAndRegister(const char *name)
+template <typename T> T *createAndRegister(const char *name, Version version)
 {
     T *testClass = new T();
-    spirvTestsRegistry::getInstance().addTestClass((baseTestClass *)testClass, name);
+    spirvTestsRegistry::getInstance().addTestClass((baseTestClass *)testClass,
+                                                   name, version);
     return testClass;
 }
 
-#define TEST_SPIRV_FUNC(name)                           \
-    extern int test_##name(cl_device_id deviceID,       \
-                           cl_context context,          \
-                           cl_command_queue queue,      \
-                           int num_elements);           \
-    class test_##name##_class  : public baseTestClass   \
-    {                                                   \
-    private:                                            \
-        basefn fn;                                      \
-                                                        \
-    public:                                             \
-    test_##name##_class() : fn(test_##name)             \
-        {                                               \
-        }                                               \
-        basefn getFunction()                            \
-        {                                               \
-            return fn;                                  \
-        }                                               \
-    };                                                  \
-    test_##name##_class *var_##name =                   \
-        createAndRegister<test_##name##_class>(#name);  \
-    int test_##name(cl_device_id deviceID,              \
-                    cl_context context,                 \
-                    cl_command_queue queue,             \
-                    int num_elements)
+#define TEST_SPIRV_FUNC_VERSION(name, version)                                 \
+    extern int test_##name(cl_device_id deviceID, cl_context context,          \
+                           cl_command_queue queue, int num_elements);          \
+    class test_##name##_class : public baseTestClass {                         \
+    private:                                                                   \
+        basefn fn;                                                             \
+                                                                               \
+    public:                                                                    \
+        test_##name##_class(): fn(test_##name) {}                              \
+        basefn getFunction() { return fn; }                                    \
+    };                                                                         \
+    test_##name##_class *var_##name =                                          \
+        createAndRegister<test_##name##_class>(#name, version);                \
+    int test_##name(cl_device_id deviceID, cl_context context,                 \
+                    cl_command_queue queue, int num_elements)
 
-std::vector<unsigned char> readSPIRV(const char *file_name);
+#define TEST_SPIRV_FUNC(name) TEST_SPIRV_FUNC_VERSION(name, Version(2, 1))
 
-int get_program_with_il(clProgramWrapper &prog,
-                        const cl_device_id deviceID,
-                        const cl_context context,
-                        const char *prog_name);
+struct spec_const
+{
+    spec_const(cl_int id = 0, size_t sizet = 0, const void *value = NULL)
+        : spec_id(id), spec_size(sizet), spec_value(value){};
+    cl_int spec_id;
+    size_t spec_size;
+    const void *spec_value;
+};
+
+int get_program_with_il(clProgramWrapper &prog, const cl_device_id deviceID,
+                        const cl_context context, const char *prog_name,
+                        spec_const spec_const_def = spec_const());
+std::vector<unsigned char> readSPIRV(const char *file_name);
\ No newline at end of file
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm32
new file mode 100644
index 00000000..00cd989a
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm32
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Float64
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.double*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 8
+               OpDecorate %double_0 SpecId 101
+     %double = OpTypeFloat 64
+%_ptr_UniformConstant_double = OpTypePointer UniformConstant %double
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_double = OpTypePointer CrossWorkgroup %double
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_double
+   %double_0 = OpSpecConstant %double 0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_double UniformConstant %double_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_double
+      %entry = OpLabel
+         %12 = OpLoad %double %output_value Aligned 8
+         %13 = OpLoad %double %spec_const_kernel_spec_constant_value Aligned 8
+        %add = OpFAdd %double %12 %13
+               OpStore %output_value %add Aligned 8
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm64
new file mode 100644
index 00000000..b50501b8
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_double_simple.spvasm64
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Float64
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.double*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 8
+               OpDecorate %double_0 SpecId 101
+     %double = OpTypeFloat 64
+%_ptr_UniformConstant_double = OpTypePointer UniformConstant %double
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_double = OpTypePointer CrossWorkgroup %double
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_double
+   %double_0 = OpSpecConstant %double 0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_double UniformConstant %double_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_double
+      %entry = OpLabel
+         %12 = OpLoad %double %output_value Aligned 8
+         %13 = OpLoad %double %spec_const_kernel_spec_constant_value Aligned 8
+        %add = OpFAdd %double %12 %13
+               OpStore %output_value %add Aligned 8
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32
new file mode 100644
index 00000000..36be62fd
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm32
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 28
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Linkage
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %11 "spec_const_kernel"
+         %27 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %test_value "test_value"
+               OpName %entry "entry"
+               OpName %output_value "output_value"
+               OpName %entry_0 "entry"
+               OpName %if_then "if.then"
+               OpName %if_end "if.end"
+               OpDecorate %test_value FuncParamAttr Zext
+               OpDecorate %test_value LinkageAttributes "test_value" Export
+               OpDecorate %add NoSignedWrap
+               OpDecorate %false SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpConstant %uchar 0
+    %uchar_1 = OpConstant %uchar 1
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+          %3 = OpTypeFunction %bool
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+         %10 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+      %false = OpSpecConstantFalse %bool
+ %test_value = OpFunction %bool None %3
+      %entry = OpLabel
+               OpReturnValue %false
+               OpFunctionEnd
+         %11 = OpFunction %void None %10
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+    %entry_0 = OpLabel
+       %call = OpFunctionCall %bool %test_value
+   %frombool = OpSelect %uchar %call %uchar_1 %uchar_0
+     %tobool = OpINotEqual %bool %frombool %uchar_0
+               OpBranchConditional %tobool %if_then %if_end
+    %if_then = OpLabel
+         %21 = OpLoad %uchar %output_value Aligned 1
+       %conv = OpUConvert %uint %21
+        %add = OpIAdd %uint %conv %uint_1
+      %conv1 = OpUConvert %uchar %add
+               OpStore %output_value %conv1 Aligned 1
+               OpBranch %if_end
+     %if_end = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64
new file mode 100644
index 00000000..c9f53443
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_false_simple.spvasm64
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 28
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Linkage
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %11 "spec_const_kernel"
+         %27 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %test_value "test_value"
+               OpName %entry "entry"
+               OpName %output_value "output_value"
+               OpName %entry_0 "entry"
+               OpName %if_then "if.then"
+               OpName %if_end "if.end"
+               OpDecorate %test_value FuncParamAttr Zext
+               OpDecorate %test_value LinkageAttributes "test_value" Export
+               OpDecorate %add NoSignedWrap
+               OpDecorate %false SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpConstant %uchar 0
+    %uchar_1 = OpConstant %uchar 1
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+          %3 = OpTypeFunction %bool
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+         %10 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+      %false = OpSpecConstantFalse %bool
+ %test_value = OpFunction %bool None %3
+      %entry = OpLabel
+               OpReturnValue %false
+               OpFunctionEnd
+         %11 = OpFunction %void None %10
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+    %entry_0 = OpLabel
+       %call = OpFunctionCall %bool %test_value
+   %frombool = OpSelect %uchar %call %uchar_1 %uchar_0
+     %tobool = OpINotEqual %bool %frombool %uchar_0
+               OpBranchConditional %tobool %if_then %if_end
+    %if_then = OpLabel
+         %21 = OpLoad %uchar %output_value Aligned 1
+       %conv = OpUConvert %uint %21
+        %add = OpIAdd %uint %conv %uint_1
+      %conv1 = OpUConvert %uchar %add
+               OpStore %output_value %conv1 Aligned 1
+               OpBranch %if_end
+     %if_end = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm32
new file mode 100644
index 00000000..b3163af1
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm32
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.float*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 4
+               OpDecorate %float_0 SpecId 101
+      %float = OpTypeFloat 32
+%_ptr_UniformConstant_float = OpTypePointer UniformConstant %float
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_float
+    %float_0 = OpSpecConstant %float 0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_float UniformConstant %float_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_float
+      %entry = OpLabel
+         %12 = OpLoad %float %output_value Aligned 4
+         %13 = OpLoad %float %spec_const_kernel_spec_constant_value Aligned 4
+        %add = OpFAdd %float %12 %13
+               OpStore %output_value %add Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm64
new file mode 100644
index 00000000..006cdc7d
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_float_simple.spvasm64
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.float*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 4
+               OpDecorate %float_0 SpecId 101
+      %float = OpTypeFloat 32
+%_ptr_UniformConstant_float = OpTypePointer UniformConstant %float
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_float
+    %float_0 = OpSpecConstant %float 0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_float UniformConstant %float_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_float
+      %entry = OpLabel
+         %12 = OpLoad %float %output_value Aligned 4
+         %13 = OpLoad %float %spec_const_kernel_spec_constant_value Aligned 4
+        %add = OpFAdd %float %12 %13
+               OpStore %output_value %add Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm32
new file mode 100644
index 00000000..2195ebe8
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm32
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Float16Buffer
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.half*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 2
+               OpDecorate %half_0x0p_0 SpecId 101
+       %half = OpTypeFloat 16
+%_ptr_UniformConstant_half = OpTypePointer UniformConstant %half
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_half = OpTypePointer CrossWorkgroup %half
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_half
+%half_0x0p_0 = OpSpecConstant %half 0x0p+0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_half UniformConstant %half_0x0p_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_half
+      %entry = OpLabel
+         %12 = OpLoad %half %output_value Aligned 2
+         %13 = OpLoad %half %spec_const_kernel_spec_constant_value Aligned 2
+        %add = OpFAdd %half %12 %13
+               OpStore %output_value %add Aligned 2
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm64
new file mode 100644
index 00000000..47dc418a
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_half_simple.spvasm64
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Float16Buffer
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.half*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 2
+               OpDecorate %half_0x0p_0 SpecId 101
+       %half = OpTypeFloat 16
+%_ptr_UniformConstant_half = OpTypePointer UniformConstant %half
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_half = OpTypePointer CrossWorkgroup %half
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_half
+%half_0x0p_0 = OpSpecConstant %half 0x0p+0
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_half UniformConstant %half_0x0p_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_half
+      %entry = OpLabel
+         %12 = OpLoad %half %output_value Aligned 2
+         %13 = OpLoad %half %spec_const_kernel_spec_constant_value Aligned 2
+        %add = OpFAdd %half %12 %13
+               OpStore %output_value %add Aligned 2
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32
new file mode 100644
index 00000000..5e182e1a
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm32
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 28
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Linkage
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %11 "spec_const_kernel"
+         %27 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %test_value "test_value"
+               OpName %entry "entry"
+               OpName %output_value "output_value"
+               OpName %entry_0 "entry"
+               OpName %if_then "if.then"
+               OpName %if_end "if.end"
+               OpDecorate %test_value FuncParamAttr Zext
+               OpDecorate %test_value LinkageAttributes "test_value" Export
+               OpDecorate %add NoSignedWrap
+               OpDecorate %true SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpConstant %uchar 0
+    %uchar_1 = OpConstant %uchar 1
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+          %3 = OpTypeFunction %bool
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+         %10 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+       %true = OpSpecConstantTrue %bool
+ %test_value = OpFunction %bool None %3
+      %entry = OpLabel
+               OpReturnValue %true
+               OpFunctionEnd
+         %11 = OpFunction %void None %10
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+    %entry_0 = OpLabel
+       %call = OpFunctionCall %bool %test_value
+   %frombool = OpSelect %uchar %call %uchar_1 %uchar_0
+     %tobool = OpINotEqual %bool %frombool %uchar_0
+               OpBranchConditional %tobool %if_end %if_then
+    %if_then = OpLabel
+         %21 = OpLoad %uchar %output_value Aligned 1
+       %conv = OpUConvert %uint %21
+        %add = OpIAdd %uint %conv %uint_1
+      %conv1 = OpUConvert %uchar %add
+               OpStore %output_value %conv1 Aligned 1
+               OpBranch %if_end
+     %if_end = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64
new file mode 100644
index 00000000..720121f6
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_true_simple.spvasm64
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 28
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Linkage
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %11 "spec_const_kernel"
+         %27 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %test_value "test_value"
+               OpName %entry "entry"
+               OpName %output_value "output_value"
+               OpName %entry_0 "entry"
+               OpName %if_then "if.then"
+               OpName %if_end "if.end"
+               OpDecorate %test_value FuncParamAttr Zext
+               OpDecorate %test_value LinkageAttributes "test_value" Export
+               OpDecorate %add NoSignedWrap
+               OpDecorate %true SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpConstant %uchar 0
+    %uchar_1 = OpConstant %uchar 1
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+          %3 = OpTypeFunction %bool
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+         %10 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+       %true = OpSpecConstantTrue %bool
+ %test_value = OpFunction %bool None %3
+      %entry = OpLabel
+               OpReturnValue %true
+               OpFunctionEnd
+         %11 = OpFunction %void None %10
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+    %entry_0 = OpLabel
+       %call = OpFunctionCall %bool %test_value
+   %frombool = OpSelect %uchar %call %uchar_1 %uchar_0
+     %tobool = OpINotEqual %bool %frombool %uchar_0
+               OpBranchConditional %tobool %if_end %if_then
+    %if_then = OpLabel
+         %21 = OpLoad %uchar %output_value Aligned 1
+       %conv = OpUConvert %uint %21
+        %add = OpIAdd %uint %conv %uint_1
+      %conv1 = OpUConvert %uchar %add
+               OpStore %output_value %conv1 Aligned 1
+               OpBranch %if_end
+     %if_end = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32
new file mode 100644
index 00000000..af778e37
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm32
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 20
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %19 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 1
+               OpDecorate %add NoSignedWrap
+               OpDecorate %uchar_0 SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpSpecConstant %uchar 0
+%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_uchar UniformConstant %uchar_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+      %entry = OpLabel
+         %12 = OpLoad %uchar %output_value Aligned 1
+         %15 = OpLoad %uchar %spec_const_kernel_spec_constant_value Aligned 1
+        %add = OpIAdd %uchar %12 %15
+               OpStore %output_value %add Aligned 1
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64
new file mode 100644
index 00000000..a05d7b7a
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uchar_simple.spvasm64
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 20
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %19 = OpString "kernel_arg_type.spec_const_kernel.uchar*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 1
+               OpDecorate %add NoSignedWrap
+               OpDecorate %uchar_0 SpecId 101
+      %uchar = OpTypeInt 8 0
+       %uint = OpTypeInt 32 0
+    %uchar_0 = OpSpecConstant %uchar 0
+%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_uchar UniformConstant %uchar_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+      %entry = OpLabel
+         %12 = OpLoad %uchar %output_value Aligned 1
+         %15 = OpLoad %uchar %spec_const_kernel_spec_constant_value Aligned 1
+        %add = OpIAdd %uchar %12 %15
+               OpStore %output_value %add Aligned 1
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm32
new file mode 100644
index 00000000..7bf0f12c
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm32
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.uint*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 4
+               OpDecorate %uint_0 SpecId 101
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpSpecConstant %uint 0
+%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uint
+      %entry = OpLabel
+         %12 = OpLoad %uint %output_value Aligned 4
+         %13 = OpLoad %uint %spec_const_kernel_spec_constant_value Aligned 4
+        %add = OpIAdd %uint %12 %13
+               OpStore %output_value %add Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm64
new file mode 100644
index 00000000..a73bf240
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_uint_simple.spvasm64
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.uint*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 4
+               OpDecorate %uint_0 SpecId 101
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpSpecConstant %uint 0
+%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_uint UniformConstant %uint_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_uint
+      %entry = OpLabel
+         %12 = OpLoad %uint %output_value Aligned 4
+         %13 = OpLoad %uint %spec_const_kernel_spec_constant_value Aligned 4
+        %add = OpIAdd %uint %12 %13
+               OpStore %output_value %add Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm32
new file mode 100644
index 00000000..c4e2ef79
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm32
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.ulong*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 8
+               OpDecorate %ulong_0 SpecId 101
+      %ulong = OpTypeInt 64 0
+    %ulong_0 = OpSpecConstant %ulong 0
+%_ptr_UniformConstant_ulong = OpTypePointer UniformConstant %ulong
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_ulong UniformConstant %ulong_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
+      %entry = OpLabel
+         %12 = OpLoad %ulong %output_value Aligned 8
+         %13 = OpLoad %ulong %spec_const_kernel_spec_constant_value Aligned 8
+        %add = OpIAdd %ulong %12 %13
+               OpStore %output_value %add Aligned 8
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm64
new file mode 100644
index 00000000..3ec3a8ff
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ulong_simple.spvasm64
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 16
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %15 = OpString "kernel_arg_type.spec_const_kernel.ulong*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpName %add "add"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 8
+               OpDecorate %ulong_0 SpecId 101
+      %ulong = OpTypeInt 64 0
+    %ulong_0 = OpSpecConstant %ulong 0
+%_ptr_UniformConstant_ulong = OpTypePointer UniformConstant %ulong
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_ulong UniformConstant %ulong_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
+      %entry = OpLabel
+         %12 = OpLoad %ulong %output_value Aligned 8
+         %13 = OpLoad %ulong %spec_const_kernel_spec_constant_value Aligned 8
+        %add = OpIAdd %ulong %12 %13
+               OpStore %output_value %add Aligned 8
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32
new file mode 100644
index 00000000..75d2e24f
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm32
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 20
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int16
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %19 = OpString "kernel_arg_type.spec_const_kernel.ushort*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 2
+               OpDecorate %add NoSignedWrap
+               OpDecorate %ushort_0 SpecId 101
+     %ushort = OpTypeInt 16 0
+       %uint = OpTypeInt 32 0
+   %ushort_0 = OpSpecConstant %ushort 0
+%_ptr_UniformConstant_ushort = OpTypePointer UniformConstant %ushort
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_ushort = OpTypePointer CrossWorkgroup %ushort
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_ushort
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_ushort UniformConstant %ushort_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_ushort
+      %entry = OpLabel
+         %12 = OpLoad %ushort %output_value Aligned 2
+         %15 = OpLoad %ushort %spec_const_kernel_spec_constant_value Aligned 2
+        %add = OpIAdd %ushort %12 %15
+               OpStore %output_value %add Aligned 2
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64 b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64
new file mode 100644
index 00000000..4f3582eb
--- /dev/null
+++ b/test_conformance/spirv_new/spirv_asm/op_spec_constant_ushort_simple.spvasm64
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 20
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int16
+               OpExtension "SPV_KHR_no_integer_wrap_decoration"
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %9 "spec_const_kernel"
+         %19 = OpString "kernel_arg_type.spec_const_kernel.ushort*,"
+               OpSource OpenCL_C 102000
+               OpName %spec_const_kernel_spec_constant_value "spec_const_kernel.spec_constant_value"
+               OpName %output_value "output_value"
+               OpName %entry "entry"
+               OpDecorate %spec_const_kernel_spec_constant_value Constant
+               OpDecorate %spec_const_kernel_spec_constant_value Alignment 2
+               OpDecorate %add NoSignedWrap
+               OpDecorate %ushort_0 SpecId 101
+     %ushort = OpTypeInt 16 0
+       %uint = OpTypeInt 32 0
+   %ushort_0 = OpSpecConstant %ushort 0
+%_ptr_UniformConstant_ushort = OpTypePointer UniformConstant %ushort
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_ushort = OpTypePointer CrossWorkgroup %ushort
+          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_ushort
+%spec_const_kernel_spec_constant_value = OpVariable %_ptr_UniformConstant_ushort UniformConstant %ushort_0
+          %9 = OpFunction %void None %8
+%output_value = OpFunctionParameter %_ptr_CrossWorkgroup_ushort
+      %entry = OpLabel
+         %12 = OpLoad %ushort %output_value Aligned 2
+         %15 = OpLoad %ushort %spec_const_kernel_spec_constant_value Aligned 2
+        %add = OpIAdd %ushort %12 %15
+               OpStore %output_value %add Aligned 2
+               OpReturn
+               OpFunctionEnd
diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..a7a4718628790bf4b54f45854b4281c5c6fab1e6
GIT binary patch
literal 560
zcmZWl%SyyR5Ufo$M%|Yp>{;<9l7j~oL=imfMeq}bxEVYcl8{V<J^BNFuNT3p*$jKw
z4l~uMu6|6%(PavZ0HiqKEjvYo@qh;d9up@>yz;HF)yr$vxdNCY_FTmH0*8=Ze%IKw
zS=UwP%%^QDTarsLqZ-%}CT~BC+cY-Y(`H58znAvv8r16pnx7H7PyJl``W=&3&V1!%
zZ<&*DnuJ*QuIYC)xv`(qR<7N6UCh~lyeP0@H@$tyv!W-TWAJ6$f1$Re`b%n$h|fId
zRBA^uacA7ZgeZPW)Ewe3_&*~mCjN#I@v0l_ft+fow@k5Ha>~(eWvX$<{-O3SBij8V
lZ`GUgsHgPG5B!3@>IVLM=&$!HNAV|e$`u!Me|+^l&p%m5JxKrn

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..62cdf072986bad9f50387cfd3d9c093354e967a9
GIT binary patch
literal 560
zcmZWlO-sW-6r47VQTwF`JuBWsa`0dUQ3MaY2>yg6Hj4*CvLqX!NB@Ao*Gs{fZ9)%y
z;bnH`&HLCF4X+bm03g8`|FUxoG3xNB!&BlIaZtWHTfV-fjn9D@BB~L5j$=r#d*kfd
ztgEu|=F8TlEyyL9QVr|~lhvQbZ=6jJv}s<oABDZS1@-y>_d8(sNto+Uzh|=2oA0b>
zEpuW{6BDb}yLL~L8-+P->FZ6<MVt-D@*FF6)7uyH6;C>I48Cl47iwFozo7Pn_(H#=
zQah4~yWkeaMDY`%<`93y`;@4d_&Y|#t8Q<P<Wx(&Ws2RCQ;v2kQ;jA2_qG2R(e9u4
lSG^fsJ*8KE&(C>Q-JbvH`s@A5QT&;la>ez!EnodW{|8$$JxTxo

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..045abe523e6a2c2a77830470d1300d4786eb8b2e
GIT binary patch
literal 836
zcmZXS%T5$Q6owBoj3c)}1jO3{LIM%UByLm|vT=b4CK^rHRyspDOd!=wPc?)cPvCnP
zUr<dj{9kuLNUToHx&7xqmzs8Ky=yjMW?g%!U(SNHY%<ACk`I#Hl$_T0?aQG%`1sX@
z+>Ju?Ltna!GArDXALL~gM|m)tu};nWx$xn`r~Wt|nf0t)`&#0ew}%EV`nd1|H^{>{
zx)Wape=1+s=7ctD94;$|E?ySC{~xwr4UV&NYe(32{h1X%CRDvs&k?;?&$H9YTkT;c
zA7Xha*-;1dkXNy&?obU*i@ArJRJUfM@BMh}L(MbQ*gj~keN|#-8@il*b&vf@_kR5*
z*475is)Z(RMmjH`_H0>Mvy#-}V9!aeV(4tAc&{<<?0MO65C{J#*-P54%GjFr@>mjm
zQNBgVRbiJT*#|Y(WG_n+!~a@2;OXy;Q0UCWf$O(;$o(e?=AAkKM?P=Lf$yz&;A79^
zu-9EN!S_TL&Yz0X!%Ffo4>SHsb)TjfzQ4-AW=FgMoapZp$1~}-|0{NKCOkWNu1oAn
hoY=epK6V9$`Y&|x?+F73HnYI<N7`*Y<h8uad;>SVR!jf@

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..701527d85ea3ccea6e30a2d194cfa3f4f92cafe8
GIT binary patch
literal 836
zcmZXS+e#ck5Qa<F&3cGu&G|wR6b~%q#=}K&lM4e0BqWgAW>`DAC_UqBkK&zA;Co13
zPy-SFZ>AL?HS|_j)n9)ddRxOQ9kU@b>)1Q}BC|GZBT2WDewOr_^tisC-+#F;pMSWJ
zyMBoN(3kGK%nEnldwH3~ejdywtX(sID}4C$wL6FhW>eM@yVYmfMh#x<LE(F@mxn=g
z$G!~yM81wq3T@UnTvi^scv1N7f7ou-JIcz9En!>rXI5Mdsd~GfBYLr(Wha%l+QUpf
z#PUMAtq$lRuVPU>q8gkQ^8`1dZq53?yYa|}nrE!Beb8L{s>IG#b&!2^kG-z*w0;w7
zX@h3fY?C*kJ1?O2Y*AUyrKyErPf9OgNH$Y^sxcq!Y1wcP2S1bS1#MSlY*~Amljc4z
z-@NpaunW@cgPO~-7o~~ef3F<y^tUY(cV<H1x)Be#chX>X)d4v2c~b<w9r3`&p2=aa
zdt!p`g$|t8iqgZ&<YOLY{FCawN-=zYm4VHUcmp`Oe@Gm!b$|F@v6EBb*~uFnVwd8?
g<_++%D=^f5tAqceFmPZq3p}^FJ=8=0lb4zA0XmIVO#lD@

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..5ac5fdfbdad833d608d05b2dfa9795bb3fa90fe1
GIT binary patch
literal 548
zcmZWlyH3ME5S;jd19^dvQU}o?<bnc_5Hu8NkobhtDLyE$Y|B0)rQ`?rUK%83j#w$O
z)_HDcXCJq5beRDo02xmCmz^O(vcn_d1ZhybHm-Vp$va;FbHss*7++uv>2&WJw=tW#
z>b&`IZRJXG8D>-i|AevayYbt`<p<jQt*ouTx~4V`AJBY{*nJx2In?hNTY2+o%ib|J
z<uoa=?tRnmX>v)J(^bCR23^eA09zDTv!C9)<XQ8kyXFwe9}~4D)n8J3Ok9PN;Djg>
zcg`(Lh%4qyiJC+F1^pROG4VI_#jCFDflAd;KbiPja?)$BGUeZ~_fYqT5$*kvf7P1v
hsGszzGw_R@Kk#2vNw2pnNAV|e$`v>0e)-~io<DcqJeB|e

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..ac4f444019d5fb202cf9583accefd2d58529cd51
GIT binary patch
literal 548
zcmZWlyH3ME5S+v|4&((wN*zRpkP8YxLeNm8iIh(`S@A)EWuNRbQc8Y+@1;Rv=7^Od
zYweA9ZuW5-4X;yR03gLVPuT^Ai1&Cx93u&ejkDG3I%|Ca%n(t9;B$;1o#NKnt=ZO9
z>&=I4DqE6EF{K*#CrsYF8^3clJJM!vWu5z*TWX{50q%Ri?vpUjvHrm1l{cSx*;(c$
zoF*aGop;@VCKrb}ZRML?&_$dL$cq9i_S2ge^cDZ~uQ`N@M?`H&^%vBh5|_baoDpT>
zF1dv<amkzsQFDmDVmu}4PW&BX@v19(qEdC#PbU7Job=kOO!*J&-Pir$jrRV`Q?+Jv
h^^;z8dVao-d;W_m>GgKy=>CPAa>ez!U%vQ}{s(!+JeL3f

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..b775b80d42a1f08d9046090d932b583d3c8c1329
GIT binary patch
literal 556
zcmZXQ%}T>i5QV2nW2*gAgsv5LB3ZalL3HDy8^K2ip}8s;(vaK;UHSq(uN%Sh#fC0A
zVRG(0bLQ`c!9{E~FtgZ>^~X*uu;CUDwm1@wE$Wmnl`Ef~lg8&}GYdN|P<(D9L(jRX
z+$vqI%f_c4t}b06F1D%EX8%N*)o-caR4&=mChuACdU++O^n7O3_dwkz-CX<n9g~$l
z{mhEiDJRl2k#ODms@>7zhTWX5^!29Gg__MQ%X3?*8@n&$OPx3~2s?4-9e6G3FQh#b
z@@{8}9SJeGQ)P^W@Dm|(z@O_r6%vEL)*YU@*u9k0p)Ur1BMv<?VC3JbcVG8M5#Iel
ke`?L-^h1yL{9I@1_WYMr=-Hhd;*a9Uh3j>{+W1cX2j96o9smFU

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..3107949596aadf1c9c36bebe5d84fdffc4aa205c
GIT binary patch
literal 556
zcmZWm%}T>S7@RhZsrFA1dRDxN<lw;yq8AUn2tLBHG^>JPyCfT-M_<6_^-^#qHuTUh
zd^@|d^Zjj-{#6F_0Ax7jDmz0TgB>32a7Y{>ZI!RV`RA9sjs-A9LN!{SVF>AUu7g{d
z)!Nt5e7efJl3a!f)xbYtZ1ry9Cb;}Sn}68y?fQmN>h%HPr^o8!wy#6|p0PffFIzT_
zF)6!AiR&hYW>1qFw0$}st4*s**c)Jr0!voY>~s2(JMGnoX3|;jpx09VIkm?`z1s!v
zPKYvb=ZqK;#gB>VL;NMr6Qb_K-|{S8d1Vh&Dv#>Q#NUyVUOmVZzhLby?=Nrk?vGrR
kYf4wW?!R>WjGXdz{5Mt7Yj(xx{)wDo#dW$ZC%&iu0pX!M9{>OV

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..f60e283596a500020b4ad22760d38e2e8282b8bb
GIT binary patch
literal 836
zcmZXSOHUL*6ooG{jH4oNP+kj&iH1ietP~fxae;{>5J=osIztH)Np;gx4PnPG@P82h
zpqNNJUw27JtWM3Td+(`pA2sdP>#o^|nRRVNzpQC%*=WK}!Y2ve3ZLscKK$ywd_Hm^
zcY_f9sW07CnHBED_wzD~gFKjxS*KP$D17*|-y6mgvk7b0zLs<*?Wv*5d06<q>*rw@
z-FIIGe=c9w#>F;k94;$QUA!)Q?>}y@>Yruh#+JD4`ZFtjj3|4j?jw4sUS{W&x7tHb
zKE(1`*ii-4kXNy&9#{=egL%XoRkdb=Z@qZtL#^|yv3*cqd&<Pj-e|M-R6X{(_T&6b
ztgQ)}RnyJAF`an<m1pz1^+HH47WTMs6+>$?#U~o`!Jd?j266aP$)3}6mB!XIml+}F
zS@~v#tK!ZHnFl%7WX}tU;osIB_|*4KEY9@ALhF}wxO*>zv!e>&ai2G3q4!=o=rL#R
zFxL-KLhq$EnzxEl!>i<@AA0<m@-CzpzTdim&HR}Yn&58ISk(F7|B9Jhh|f%xw29qF
h6Pq`{$E@Ixe_0#<M{&@=rWbVXbRsADz(4ZR^ItZZR!aZ?

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..4d96c8e819c3f5cd69ed66148525d00afa4a5efb
GIT binary patch
literal 836
zcmZXSOG_L<6os#jlQBMH)R>qo6d}PR46gE6=q3vV14;z9MKe~ML8)$gsuSGh7x+IU
ze=rIL&(~dwU^O+T?!BkZeblr@*E(h+X4bJc`en`8sEs9TCw!9dop3_m!TzD!-TC1{
z?gk<HV_&-8WmdQ&-^<G^4)S0&Y3*8hukhj9_wFzrnax;B>Q?Vr8#i>ZhlTICULJ<g
zo%k~NQ~5eJEw)+Xa9Mfm;(6h_kGS2c*U!rLAH{9epILD+qU`OukLacPD?6>c)gF5C
zA(rRDwkn{8yoyD2&uVxY%mdz-sx=$@?8d$iwa(MV_CbAZD-$z&tIgV0_1Np$5A!#%
zmL_Oc%{BKXb>;<Bo-OLul#pC3>}lZ&hSrA^pJ~i{dsa3Y#Nj_n_JXFXG`6a_JQs4F
zmv3IUBJP5ad608e_M(s&{texMPko=n;!IC0wEjtlyZ?o7zNi9t+~-YM=zW$Bdd!(S
z%=N33(0iec=B1+4u#|lCLyvz`-j^wc?^-vonLl$v6WmN1uXMiqUon$2@tMi8HnBg_
h#O4j~F)KLaf31ywOB^(?=>?rDoybW(@J3#Gz6CjHR!jf@

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..94ce6666a3b474b9c94c0de4161e8be15e7f8e52
GIT binary patch
literal 628
zcmZWmO-lnY5S@0rw)WE^daig7ME2l83!)bf6%-T&uOaM)S}fTm%__b53;exa1m8=$
z6zqg#GVi^a%+5yj<EBVWM4B?hNgGOC8f(9^_9yUdte2-(b$|D)Y_4Y3o2hZ?)9FHG
zCe5Apv)qbwBn}lH3uA8|l7-Ji`Vs}KjyeNr3G4E{C`_u-+%CL&H_n<lu%`5&7FhvG
zyQ%Wa!X#VTq)gwmJ3WIo3QMH;uHp4==zm+kLA3SiL(fZt+!)=&@OkNrasvZvgg%Y+
zZW(kDdK1xECKJ5J869FzaO0hQaFXPxD=|=Vtu@ZK4iN9c6QeP1!;cui&j0Zv#GMrt
z^7b*cE@I+6c=pYDM~M6I<cME^VIKRU@x2RR<T-a5^_S?is`CSe|K%DdYYwnkpL%{*
dh410cx)uHf74`gP=8(Su##~~R?iW?Y*nf})OG^L%

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..6dc981679c8c0154d65576f3d3bcfeda65cb2ece
GIT binary patch
literal 628
zcmZWmO-lnY5S@0rw)WE^daig7ME2l83!)bf6%-T&uOaM)S}fTm%__b53;exa1m8=$
z6zqg#GVi^a%+5yj<EBVWM4B?hNgGOC8f(9^_9yUdte2-(b$|D)Y_4Y3o2hZ?)9FHG
zCe5Apv)qbwBn}lH3uA8|l7-Ji`Vyh4ggua!urA|8VN#XmcHz~#an{U%HKhl&$O=%}
zO_g63CfU*^W%{Pw=^3<9SR%!D4X<}Y|J(WvqODgSdR`jj#^@%7&r4sF8yHw4^l7Yj
z%b<(Un~2UbnczLn=n#8?8}ICclO#u7iGh-9t#Q6}fOr?47>#ioe#8KF{*NCa?yRVg
zw~wiH5fkshvv1ZrLfnTZNBjy5^Vk=S?_B^R&$-j6zeKN9ogXm#FV{F(bAZkI)bqP4
dd=Gclt?)0XsOL8`hx`pN<`S!Pzo;_C{sWQ@OH2R&

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..a7734d984d17ea826eb7b0e7091003a96e6adc73
GIT binary patch
literal 552
zcmZWlO-sW-5PeM=Q|*^U^jz^Kl7j~oie5bQQt&4%p;<&M+mLL8-uwmrUN3_0Z4-Lv
zgw30I^JaE-;^-y?dH_<K^DmnrLchfm;t)wtyp`6yuCm$}zyxvNBIc(UKsxz#Y1d}G
zb+tF2wsLkuF2$H?U{9F5dN+Po+U!W1H6PC3-cd@uKA`;R@%B-e=U9JWa_7xgzG*CL
z5>As4w~a5G15K_U=CsaNyP%6X8;}<Tmb^`GzT#QZldnO9iU&k(N&PPwpAzRhvimlc
zMmA@Bws(W>f}0t({D`R8RDVs~n5dli1v9GCoig=%Ag8=`E0g|_vvz%cnbGc_`B%>g
ikNQimJ{>=8<BtEPN_xFtHOjw`Q?0m8_s3U1@%#hJQanEZ

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..6bb0eef7ef8799a2fa10fa931304c2794545d28c
GIT binary patch
literal 552
zcmZWlO-sW-5PfYLQ|*@`daigA$-#pPMK2zD5&Q{DXciI6c1bouZ~g**uNT4hwh29S
z!sP9|c{4jZN&h+ndH^z<@-I6>AA=T;h$Ey?@#d|2dCluk0#hVZA^HqMNT;~=c4bx@
zR|oTHD`(f_GEArjc7!RacN4bW<_Frm`EcRpmQw2V0sg1Q+sAR9L;ap9Trgk7y0NTD
zIZaC3G{HA}n%p4HX<ewcQI~KwpeRc$d7IvRNni4$vj!0=9ul=B^}k?zOq|nYcWo?<
zY|i*(cZ|AoZf4Z-W1?nL{S|c+qH^LF%&1Ox%GB?kobuYOO!^1T+V%NmM!SFFUp=RE
h^_O0KI)2v19sf<0^m@N)lz%3tT5+B3kFS2D{{ziFJU{>d

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..eb1469f85a8d88500cd21c162ff137795b76b2f5
GIT binary patch
literal 564
zcmZWl%SyyR5Ufo$W_4eR=vnb5l7j~oL=imfMeq}buo?DXn1o~^?9DImd%Xx&jR||$
z4n5sdRo&C+IJ!!K5r7oOJY^?{Fz(|)AE%5{BtiLJTle~wwY~rrhyxapUt$93mAAEB
zn|0+{Z$50}Y)LG|oNC}7FnO~vep}mYPn&gR<u=zh)W+ch>hBR>pM|;h^*bbY-hAd|
zXIYzYnuM|HeBJF}V&gEUb-vjKUCh~lyeP2Zd;0ft`ihxy%`xfkirSX?pOHOee4)#3
z*;#jK2D_l<sK1Ci<#x17;WI|fC;XE4Iion?k7QJ*_hssNM@+o-E0g}7^A3G}$Y}ph
mJk@hSSAXeMKj2H|>O0_HR7tNLs7CxVG1Uqibie%b2l^kaBRul}

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..e37d3226af44070f1693598c01aec14a9e48e30a
GIT binary patch
literal 564
zcmZWlO-sW-5PfYLQ|*@`dRDxN<lvzSq6i*(5&Q`uG)oVLO-MFEZ~g**ua|=Ftyy~L
zgxR<A=FRM6qTx*f3;-lJ<taPE5ThQCdYlj^h`sWoHumj3ZCwG(5m62A3ydMX;<h&H
zY+c#bWuK<8rX-hOMm6wHWO=j6+_pC9fi~^R%5HA&sEzyw)IS5hKJ{}Q>i0};UG|li
zondXvX<}m4xw_la<VJo@V_mcLx`?v@c~M}+_w?_V^c8QqYmRXr3$-ovzhL&5_)3@E
zv$JApf?ZK_(g(_&b359l_$g8IiN9t%BPu8Ui5b=DeVKYL$tkb>%A|kbyrIu8Gur<%
mPxYMB)n9tm2Y$i3`Ud`+D(STY)hPc$PPO8KZpS}=r2hf06Fl?)

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32
new file mode 100644
index 0000000000000000000000000000000000000000..c81e309a1f31b113221ac0f7176794711086a4cf
GIT binary patch
literal 616
zcmZWmO-lnY5S@OkTl;AdJy*Pm?7@QyMK2yIC@2bE!?F#vSdt~1ReJLm_<OwwzL)OO
zLnkCNdGqFFcDB_#jYS$F63ZATZ7fY`SAMVZ$M9XO*XK9&aQ~u=RZHXbLOb>4(p=4T
zW?kx+)`;{Z3W`s;Hg}In;pZYlX$7naoso2ex{QllXDYL%@ajW5qgTLU8GtRa0aWT1
z%CB>s>~NE^c(=}<oeTI}BKc1PulGa$yZkMZ8m~UnRjE-Ep_>TamOd}HFtB#$QycHr
z$j|Ts=fYXc5HZHn;X9baAf?0*nd>0tRH=W2`#$^uHjVc#{Dc?SMU`g_ZGwAm6Nb7&
zywI=k0X%!9-U)a^c;<*tkzpPCqEYV>81tMpjr=S0TI>8G!+&vulbR!J>XXl1<LsTf
ZHU142`TS<qFn<e-wZv-PAEt1P{SVP-N}vD$

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64 b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64
new file mode 100644
index 0000000000000000000000000000000000000000..8f7074280a9357028bcd79c17c20988e948303d9
GIT binary patch
literal 616
zcmZWmO-lnY5S@OkTl;AdJy*Pm?7@QyMK2yIC@2bE!?F#vSh69TReJLm_<OwwzL)OO
zLnmY>^WMzMWVY2jjYS$F63ZB0+E|*>uKZr*kKwymug`Dl;r>OLTrG{)3+>dGORbvg
zEO)72=0>C^QBZuc+T1-Rg`bNIr3F<9dn6s9F5|-LOl7$#y!z12=oPS72FMoK04jA0
z<=0jxJGn_&yywoJoeT20MC?xkxA#N;yZKurHC}zDt5TyTLN^h7UV2+@VPNghr#9ZL
z!Ow65=fYXcaN-$Phj);}Af?0*$#rn%RH=W2cpv@%o5r{cKVbm7sCd@UCWw2RW2igC
z4gCrqz_VBCogi-rPmcH$4C~kzje3{B$aB^-=3k-LTIUxG|HTcy)Er?`zy4P>&fck8
Z<KM_)KHtn5^0&ZPORUcO!xXNu{{hovN}&J%

literal 0
HcmV?d00001

diff --git a/test_conformance/spirv_new/test_op_spec_constant.cpp b/test_conformance/spirv_new/test_op_spec_constant.cpp
new file mode 100644
index 00000000..a6feddfd
--- /dev/null
+++ b/test_conformance/spirv_new/test_op_spec_constant.cpp
@@ -0,0 +1,159 @@
+/******************************************************************
+Copyright (c) 2020 The Khronos Group Inc. All Rights Reserved.
+
+This code is protected by copyright laws and contains material proprietary to
+the Khronos Group, Inc. This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not
+be disclosed in whole or in part to third parties, and may not be reproduced,
+republished, distributed, transmitted, displayed, broadcast or otherwise
+exploited in any manner without the express prior written permission of Khronos
+Group. The receipt or possession of this code does not convey any rights to
+reproduce, disclose, or distribute its contents, or to manufacture, use, or sell
+anything that it may describe, in whole or in part other than under the terms of
+the Khronos Adopters Agreement or Khronos Conformance Test Source License
+Agreement as executed between Khronos and the recipient.
+******************************************************************/
+
+#include "testBase.h"
+#include "types.hpp"
+
+
+template <typename T>
+int run_case(cl_device_id deviceID, cl_context context, cl_command_queue queue,
+             const char *name, T init_buffer, T spec_constant_value,
+             T final_value, bool use_spec_constant,
+             bool (*notEqual)(const T &, const T &) = isNotEqual<T>)
+{
+    clProgramWrapper prog;
+    cl_int err = CL_SUCCESS;
+    if (use_spec_constant)
+    {
+        spec_const new_spec_const =
+            spec_const(101, sizeof(T), &spec_constant_value);
+
+        err =
+            get_program_with_il(prog, deviceID, context, name, new_spec_const);
+    }
+    else
+    {
+        err = get_program_with_il(prog, deviceID, context, name);
+    }
+    SPIRV_CHECK_ERROR(err, "Failed to build program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "spec_const_kernel", &err);
+    SPIRV_CHECK_ERROR(err, "Failed to create kernel");
+    size_t bytes = sizeof(T);
+    clMemWrapper output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, bytes,
+                       &init_buffer, &err);
+    SPIRV_CHECK_ERROR(err, "Failed to create output_buffer");
+
+    err = clSetKernelArg(kernel, 0, sizeof(clMemWrapper), &output_buffer);
+    SPIRV_CHECK_ERROR(err, "Failed to set kernel argument output_buffer");
+
+    size_t work_size = 1;
+    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_size, NULL, 0,
+                                 NULL, NULL);
+    SPIRV_CHECK_ERROR(err, "Failed to enqueue kernel");
+    clFinish(queue);
+
+    T device_results = 0;
+    err = clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, 0, bytes,
+                              &device_results, 0, NULL, NULL);
+    SPIRV_CHECK_ERROR(err, "Failed to copy from output_buffer");
+    T reference = 0;
+    use_spec_constant ? reference = final_value : reference = init_buffer;
+    if (device_results != reference)
+    {
+        log_error("Values do not match. Expected %d obtained %d\n", reference,
+                  device_results);
+        err = -1;
+    }
+    return err;
+}
+
+template <typename T>
+int test_spec_constant(cl_device_id deviceID, cl_context context,
+                       cl_command_queue queue, const char *name, T init_buffer,
+                       T spec_constant_value, T final_value)
+{
+    if (std::string(name).find("double") != std::string::npos)
+    {
+        if (!is_extension_available(deviceID, "cl_khr_fp64"))
+        {
+            log_info("Extension cl_khr_fp64 not supported; skipping double "
+                     "tests.\n");
+            return TEST_SKIPPED_ITSELF;
+        }
+    }
+    if (std::string(name).find("half") != std::string::npos)
+    {
+        if (!is_extension_available(deviceID, "cl_khr_fp16"))
+        {
+            log_info("Extension cl_khr_fp16 not supported; skipping half "
+                     "tests.\n");
+            return TEST_SKIPPED_ITSELF;
+        }
+    }
+    cl_int err = CL_SUCCESS;
+    err = run_case<T>(deviceID, context, queue, name, init_buffer,
+                      spec_constant_value, final_value, false);
+    err |= run_case<T>(deviceID, context, queue, name, init_buffer,
+                       spec_constant_value, final_value, true);
+
+    if (err == CL_SUCCESS)
+    {
+        return TEST_PASS;
+    }
+    else
+    {
+        return TEST_FAIL;
+    }
+}
+
+
+#define TEST_SPEC_CONSTANT(NAME, type, init_buffer, spec_constant_value)       \
+    TEST_SPIRV_FUNC_VERSION(op_spec_constant_##NAME##_simple, Version(2, 2))   \
+    {                                                                          \
+        type init_value = init_buffer;                                         \
+        type final_value = init_value + spec_constant_value;                   \
+        return test_spec_constant(                                             \
+            deviceID, context, queue, "op_spec_constant_" #NAME "_simple",     \
+            init_value, (type)spec_constant_value, final_value);               \
+    }
+
+// type name, type, value init, spec constant value
+TEST_SPEC_CONSTANT(uint, cl_uint, 25, 43)
+TEST_SPEC_CONSTANT(uchar, cl_uchar, 19, 4)
+TEST_SPEC_CONSTANT(ushort, cl_ushort, 6000, 3000)
+TEST_SPEC_CONSTANT(ulong, cl_ulong, 9223372036854775000UL, 200)
+TEST_SPEC_CONSTANT(float, cl_float, 1.5, -3.7)
+TEST_SPEC_CONSTANT(half, cl_half, 1, 2)
+TEST_SPEC_CONSTANT(double, cl_double, 14534.53453, 1.53453)
+
+// Boolean tests
+// documenation: 'If a specialization constant is a boolean
+// constant, spec_value should be a pointer to a cl_uchar value'
+
+TEST_SPIRV_FUNC_VERSION(op_spec_constant_true_simple, Version(2, 2))
+{
+    // 1-st ndrange init_value is expected value (no change)
+    // 2-nd ndrange sets spec const to 'false' so value = value + 1
+    cl_uchar value = (cl_uchar)7;
+    cl_uchar init_value = value;
+    cl_uchar final_value = value + 1;
+    return test_spec_constant<cl_uchar>(deviceID, context, queue,
+                                        "op_spec_constant_true_simple",
+                                        init_value, 0, final_value);
+}
+
+TEST_SPIRV_FUNC_VERSION(op_spec_constant_false_simple, Version(2, 2))
+{
+    // 1-st ndrange init_value is expected value (no change)
+    // 2-nd ndrange sets spec const to 'true' so value = value + 1
+    cl_uchar value = (cl_uchar)7;
+    cl_uchar init_value = value;
+    cl_uchar final_value = value + 1;
+    return test_spec_constant<cl_uchar>(deviceID, context, queue,
+                                        "op_spec_constant_false_simple",
+                                        init_value, 1, final_value);
+}