mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
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
This commit is contained in:
Grzegorz Wawiorko
committed by
GitHub
GitHub
parent
afa2fcca96
commit
41cd9c6d98
@@ -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
|
||||
{
|
||||
|
||||
@@ -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);
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
159
test_conformance/spirv_new/test_op_spec_constant.cpp
Normal file
159
test_conformance/spirv_new/test_op_spec_constant.cpp
Normal file
@@ -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);
|
||||
}
|
||||
Reference in New Issue
Block a user