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 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 -#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 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 -T* createAndRegister(const char *name) +template 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(#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(#name, version); \ + int test_##name(cl_device_id deviceID, cl_context context, \ + cl_command_queue queue, int num_elements) -std::vector 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 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 00000000..a7a47186 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv32 differ 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 00000000..62cdf072 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_double_simple.spv64 differ 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 00000000..045abe52 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv32 differ 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 00000000..701527d8 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_false_simple.spv64 differ 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 00000000..5ac5fdfb Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv32 differ 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 00000000..ac4f4440 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_float_simple.spv64 differ 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 00000000..b775b80d Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv32 differ 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 00000000..31079495 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_half_simple.spv64 differ 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 00000000..f60e2835 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv32 differ 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 00000000..4d96c8e8 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_true_simple.spv64 differ 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 00000000..94ce6666 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv32 differ 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 00000000..6dc98167 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uchar_simple.spv64 differ 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 00000000..a7734d98 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv32 differ 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 00000000..6bb0eef7 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_uint_simple.spv64 differ 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 00000000..eb1469f8 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv32 differ 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 00000000..e37d3226 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ulong_simple.spv64 differ 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 00000000..c81e309a Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv32 differ 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 00000000..8f707428 Binary files /dev/null and b/test_conformance/spirv_new/spirv_bin/op_spec_constant_ushort_simple.spv64 differ 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 +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) +{ + 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 +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(deviceID, context, queue, name, init_buffer, + spec_constant_value, final_value, false); + err |= run_case(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(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(deviceID, context, queue, + "op_spec_constant_false_simple", + init_value, 1, final_value); +}