File utest_generator.py contain the base class and function for generating test
cases.
File utest_math_gen.py can generate most math function for all the gentype
(floatX, intX).
utest_math_gen.py can be run during cmake.
Signed-off-by: Yi Sun <yi....@intel.com>
Signed-off-by: Yangwei Shui <yangwei.s...@intel.com>
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 5e0bc19..135ea9b 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -1,10 +1,14 @@
INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}
${CMAKE_CURRENT_SOURCE_DIR}/../include)
+EXEC_PROGRAM(python ${CMAKE_CURRENT_SOURCE_DIR} ARGS utest_math_gen.py
OUTPUT_VARIABLE NEWFUNC)
+MESSAGE(STATUS ${NEWFUNC})
+string(REGEX REPLACE " " ";" NEWFUNC ${NEWFUNC})
link_directories (${LLVM_LIBRARY_DIR})
set (utests_sources
cl_create_kernel.cpp
utest_error.c
+ ${NEWFUNC}
compiler_basic_arithmetic.cpp
compiler_displacement_map_element.cpp
compiler_shader_toy.cpp
diff --git a/utests/utest_generator.py b/utests/utest_generator.py
new file mode 100644
index 0000000..1346fc2
--- /dev/null
+++ b/utests/utest_generator.py
@@ -0,0 +1,304 @@
+#!/usr/bin/python
+import os,sys,re
+
+FLT_MAX_POSI='0x1.fffffep127f'
+FLT_MIN_NEGA='-0x1.fffffep127f'
+FLT_ULP='1.0e-6f'
+INT_ULP='0'
+FLT_MIN_POSI='0x1.0p-126f'
+FLT_MAX_NEGA='-0x1.0p-126f'
+
+paraTypeList={'float':'%f','int':'%d','double':'%lf','uint':'%d'}
+
+def udebug(ulpSize):
+ ulpUnit=re.findall(r"([a-zA-Z_]+)",ulpSize)[0]
+ text='''
+#define ULPSIZE %s
+#define ULPUNIT %s
+#if udebug
+ if (isinf(cpu_data[index]) && isinf(gpu_data[index])){
+ printf(log);
+ }
+ else if (isnan(cpu_data[index]) && isnan(gpu_data[index])){
+ printf(log);
+ }
+ else if ( diff <= ULPSIZE \\
+ && ( fabs(gpu_data[index]) > ULPUNIT || fabs(cpu_data[index]) > ULPUNIT
)){
+ printf(log);
+ }
+ else if ( fabs(cpu_data[index]) <= ULPUNIT && fabs(gpu_data[index]) <=
ULPUNIT){
+ printf(log);
+ }
+ else
+ printf_c(log);
+#else
+ if (isinf(cpu_data[index]))
+ OCL_ASSERTM(isinf(gpu_data[index]),log);
+ else if (isnan(cpu_data[index]))
+ OCL_ASSERTM(isnan(gpu_data[index]),log);
+ else if ( fabs(gpu_data[index]) > ULPUNIT || fabs(cpu_data[index]) >
ULPUNIT)
+ OCL_ASSERTM( diff <= ULPSIZE, log);
+ else
+ OCL_ASSERTM(fabs(gpu_data[index]-cpu_data[index]) <= ULPSIZE, log);
+#endif
+ }
+}\n'''%(ulpSize,ulpUnit)
+
+ return text
+
+def gene2ValuesLoop(values1,values2,inputValues):
+ values2=values2+inputValues*len(inputValues)
+
+ for i in inputValues:
+ for j in range(0,len(inputValues)):
+ values1 += [i]
+
+ return values1,values2
+
+def gene3ValuesLoop(values1,values2,values3,inputValues):
+ for i in inputValues:
+ for j in range(0,len(inputValues)):
+ for k in range(0,len(inputValues)):
+ values1 += [i]
+
+ for i in inputValues:
+ for j in inputValues:
+ for k in range(0,len(inputValues)):
+ values2 += [i]
+
+ values3=inputValues*(len(inputValues)**2)
+ return values1,values2,values3
+
+class func:
+ """ This class will define all needed instance attribute in fundation a c
programing file. """
+
+ def __init__(self,name,inputType,outputType,values,ulp, cpu_func=''):
+ self.funcName = name
+ self.fileName = 'builtin_'+name
+ self.inputtype = inputType
+ self.outputtype = outputType
+ self.values = values
+ self.ulp = ulp
+ self.cpufunc=cpu_func
+ self.cpplines = []
+
+#####cpp file required information:
+ self.Head='''#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+#define udebug 0
+#define FLT_MAX 0x1.fffffep127f
+#define FLT_MIN 0x1.0p-126f
+#define FLT_ULP (1.0e-6f)
+#define INT_ULP 0
+
+#define printf_c(...) \\
+{\\
+ printf("\\033[1m\\033[40;31m");\\
+ printf( __VA_ARGS__ );\\
+ printf("\\033[0m");\\
+}
+'''
+ #########Execute class itself
+ self.geneProcess()
+
+#####Computer vector && argument type:
+ def argtype(self,paraN,index):
+ return re.findall(r"[a-zA-Z_]+",self.inputtype[paraN][index])[0]
+
+ def argvector(self,paraN,index):
+ vector=re.findall(r"[0-9]+",self.inputtype[paraN][index])
+ if vector:
+ vector=vector[0]
+ else:
+ vector=1
+ return vector
+
+ def retType(self,index):
+ return re.findall("[a-zA-Z_]+",self.outputtype[index])[0]
+
+ def inputNumFormat(self,paraN,index):
+ return paraTypeList['%s'%(self.argtype(paraN,index))]
+
+ def outputNumFormat(self,index):
+ return paraTypeList['%s'%(self.retType(index))]
+
+#####Cpu values analyse
+ def GenInputValues(self,index):
+ #namesuffix=self.inputtype[0][index]
+ for i in range(0,self.values.__len__()):
+ self.cpplines += [ "const %s input_data%d[] = {%s};"
%(self.argtype(i,index),i+1,str(self.values[i]).strip('[]').replace('\'','')) ]
+ self.cpplines += [ "const int count_input = sizeof(input_data1) /
sizeof(input_data1[0]);\n" ]
+
+#####Cpu Function
+ def GenCpuCompilerMath(self,index):
+ #namesuffix=self.inputtype[0][index]
+ defline='static void cpu_compiler_math(%s *dst, '%(self.retType(index))
+ cpufunargs='('
+ funcline = ['{']
+
+ for i in range(0,self.values.__len__()):
+ defline += 'const %s *src%d'%(self.argtype(i,index),i+1)
+ defline += ( i == self.values.__len__()-1 ) and ')' or ','
+ cpufunargs += "x%d"%(i+1)
+ cpufunargs += ( i == self.values.__len__()-1 ) and ');' or ','
+ funcline += [" const %s x%d = *src%d;"%(self.argtype(i,index),i+1,i+1)]
+
+ funcline += [ " dst[0] = %s%s"%(self.funcName, cpufunargs) ]
+ funcline += [ '}']
+
+ funcline = [defline] + funcline
+
+ self.cpplines += funcline
+# self.writeCPP( '\n'.join(funcline), 'a', namesuffix)
+
+ def writeCPP(self,content,authority,namesuffix):
+ file_object = open("%s_%s.cpp"%(self.fileName,namesuffix),authority)
+ file_object.writelines(content)
+ file_object.close()
+
+ def writeCL(self,content,authority,namesuffix):
+ file_object =
open(os.getcwd()+"/../kernels/%s_%s.cl"%(self.fileName,namesuffix),authority)
+ file_object.writelines(content)
+ file_object.close()
+
+ def nameForCmake(self,content,namesuffix):
+ print("%s_%s.cpp"%(self.fileName,namesuffix)),
+
+ def utestFunc(self,index):
+ funcLines=[]
+ namesuffix=self.inputtype[0][index]
+ funcline=[]
+ funchead='''
+static void %s_%s(void)
+{
+ int index; const int vector = %s;
+ %s gpu_data[count_input] = {0}, cpu_data[count_input] = {0}, diff=0.0;
+ char log[1024] = {0};
+
+ OCL_CREATE_KERNEL(\"%s_%s\");
+ OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, count_input * sizeof(%s),
NULL);
+
+ globals[0] = count_input;
+ locals[0] = 1;
+ '''%(self.fileName,namesuffix,\
+ self.argvector(self.inputtype.__len__()-1,index),\
+ self.retType(index),\
+ self.fileName, namesuffix,\
+ self.retType(index))
+
+ funcline += [funchead]
+ for i in range(1,self.values.__len__()+1):
+ funcline += [" OCL_CREATE_BUFFER(buf[%d], CL_MEM_READ_WRITE,
count_input * sizeof(%s), NULL);"%(i,self.argtype(i-1,index))]
+ funcline += [" clEnqueueWriteBuffer( queue, buf[%d], CL_TRUE, 0,
count_input * sizeof(%s), input_data%d, 0, NULL,
NULL);"%(i,self.argtype(i-1,index),i)]
+
+ funcline += [" OCL_CREATE_BUFFER(buf[%d], CL_MEM_READ_WRITE, sizeof(int),
NULL);"%(self.inputtype.__len__()+1)]
+ funcline += [" clEnqueueWriteBuffer( queue, buf[%d], CL_TRUE, 0,
sizeof(int), &vector, 0, NULL, NULL);"%(self.inputtype.__len__()+1)]
+
+ #0=output 1=input1 2=input2 ... len+2=output
+ for i in range(0,self.values.__len__()+2):
+ funcline += [" OCL_SET_ARG(%d, sizeof(cl_mem), &buf[%d]);"%(i,i)]
+
+ funcrun='''
+ // Run the kernel:
+ OCL_NDRANGE( 1 );
+ clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(%s) * count_input,
gpu_data, 0, NULL, NULL);
+'''%(self.inputtype.__len__()+1)
+ funcline += [ funcrun ]
+
+ funcsprintfa=' sprintf(log, \"'
+ funcsprintfb=''
+ funccompare='''
+ for (index = 0; index < count_input; index++)
+ {
+ cpu_compiler_math( cpu_data + index,'''
+
+ for i in range(0,self.values.__len__()):
+ funccompare += " input_data%d + index"%(i+1)
+ funccompare += (self.values.__len__() - 1 == i) and ');' or ','
+
+ funcsprintfa += "input_data%d:"%(i+1)
+ funcsprintfa += "%s "%(self.inputNumFormat(i,index))
+ funcsprintfb += " input_data%d[index],"%(i+1)
+
+ funcline += [ funccompare ]
+
+ funcsprintfa += " -> gpu:%s cpu:%s diff:%s
expect:%s\\n\","%(self.outputNumFormat(index),self.outputNumFormat(index),self.outputNumFormat(index),self.outputNumFormat(index))
+ funcsprintfb += " gpu_data[index], cpu_data[index], diff, %s);"%self.ulp
+
+ funcdiff = " diff = fabs((gpu_data[index]-cpu_data[index])"
+ funcdiff += (self.retType(index) == "int") and ');' or
'/(cpu_data[index]>1?cpu_data[index]:1));'
+ funcline += [ funcdiff ]
+ funcline += [ funcsprintfa + funcsprintfb ]
+
+ self.cpplines += funcline
+
+ self.cpplines += [ udebug(self.ulp) ]
+ self.cpplines += [
"MAKE_UTEST_FROM_FUNCTION(%s_%s)"%(self.fileName,namesuffix) ]
+
+ def genCL(self,index):
+ namesuffix=self.inputtype[0][index]
+ clLine = []
+ clhead = '__kernel void %s_%s(__global %s *dst,
'%(self.fileName,namesuffix,self.retType(index))
+ clvalueDef=''
+ clcomputer=''
+ tmp=''
+
+ for i in range(0,self.values.__len__()):
+ clhead += ' __global %s *src%d,'%(self.argtype(i,index),i+1)
+ clvalueDef += ' %s x%d = (%s)
('%(self.inputtype[i][index],i+1,self.inputtype[i][index])
+ tmp = 'src%d[i * (*vector) + '%(i+1)
+ for j in range(0,int(self.argvector(i,index))):
+ clvalueDef += tmp + ((int(self.argvector(i-1,index)) == j+1 ) and
'%d]);\n'%(j) or '%d],'%(j))
+ clcomputer += (self.values.__len__() == i+1) and 'x%d);'%(i+1) or
'x%d,'%(i+1)
+
+ clhead += ' __global int *vector) {\n'
+ clhead += ' int i = get_global_id(0);'
+ clLine += [ clhead ]
+ clLine += [ clvalueDef ]
+ clLine += [ ' %s ret;'%(self.outputtype[index]) ]
+ clLine += [ ' ret = %s('%(self.funcName) + clcomputer ]
+
+ if (int(self.argvector(0,index)) == 1):
+ clLine += [ ' dst[i] = ret;' ]
+ else:
+ for i in range(0,int(self.argvector(0,index))):
+ clLine += [ ' dst[i * (*vector) + %d] = ret[%d];'%(i,i) ]
+ clLine += [ '};' ]
+
+ self.writeCL('\n'.join(clLine),'w',namesuffix)
+
+ def geneProcess(self):
+ for i in range(0,self.inputtype[0].__len__()):
+##########Write Cpp file
+ namesuffix=self.inputtype[0][i]
+ self.cpplines = []
+ #The head:
+ self.cpplines += [self.Head]
+
+ #cpu function generator:
+ self.cpplines += [self.cpufunc]
+
+ #Parameters:
+ self.GenInputValues(i)
+
+ #Cpu function:
+ self.GenCpuCompilerMath(i)
+
+ #utest function
+ self.utestFunc(i)
+
+ #kernel cl
+ self.genCL(i)
+
+ #CMakelists.txt
+ self.nameForCmake(self.fileName,namesuffix)
+
+ self.writeCPP( '\n'.join(self.cpplines) ,'w',namesuffix)
+#########End
+
+#def main():
+#
+#if __name__ == "__main__":
+# main()
diff --git a/utests/utest_math_gen.py b/utests/utest_math_gen.py
new file mode 100755
index 0000000..c60d512
--- /dev/null
+++ b/utests/utest_math_gen.py
@@ -0,0 +1,476 @@
+#!/usr/bin/python
+from utest_generator import *
+import os,sys
+
+#base_input_values = [80, -80, 3.14, -3.14, -0.5, 0.5, 1, -1,
0.0,6,-6,1500.24,-1500.24]
+#extend_input_values =
[FLT_MAX_POSI,FLT_MIN_NEGA,FLT_MIN_POSI,FLT_MAX_NEGA,80, -80, 3.14, -3.14,
-0.5, 0.5, 1, -1, 0.0,6,-6,1500.24,-1500.24]
+base_input_values = [ 0, 1, 3.14]
+def main():
+ ##### gentype acos(gentype)
+ acos_input_values = base_input_values
+ acos_input_type = ['float','float2','float4','float8','float16']
+ acos_output_type = ['float','float2','float4','float8','float16']
+ ###### gentype acos(gentype)
+ acosUtests =
func('acos',[acos_input_type],acos_output_type,[acos_input_values],'4 *
FLT_ULP')
+
+ ##### gentype acosh(gentype)
+ acosh_input_values = base_input_values
+ acosh_input_type = ['float','float2','float4','float8','float16']
+ acosh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype acosh(gentype)
+ acoshUtests =
func('acosh',[acosh_input_type],acosh_output_type,[acosh_input_values],'4 *
FLT_ULP')
+
+ ##### gentype asin(gentype)
+ asin_input_values = base_input_values
+ asin_input_type = ['float','float2','float4','float8','float16']
+ asin_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype asin(gentype)
+ asinUtests =
func('asin',[asin_input_type],asin_output_type,[asin_input_values],'4 *
FLT_ULP')
+
+ ##### gentype asinh(gentype)
+ asinh_input_values = base_input_values
+ asinh_input_type = ['float','float2','float4','float8','float16']
+ asinh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype asinh(gentype)
+ asinhUtests =
func('asinh',[asinh_input_type],asinh_output_type,[asinh_input_values],'4 *
FLT_ULP')
+
+ ##### gentype atan(gentype y_over_x)
+ atan_input_values = base_input_values
+ atan_input_type = ['float','float2','float4','float8','float16']
+ atan_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype atan(gentype y_over_x)
+ atanUtests =
func('atan',[atan_input_type],atan_output_type,[atan_input_values],'5 *
FLT_ULP')
+
+ ##### gentype atan2(gentype y, gentype x)
+ atan2_base_values = base_input_values
+ atan2_input_values1 = []
+ atan2_input_values2 = []
+
atan2_input_values1,atan2_input_values2=gene2ValuesLoop(atan2_input_values1,atan2_input_values2,atan2_base_values)
+ atan2_input_type1 = ['float','float2','float4','float8','float16']
+ atan2_input_type2 = ['float','float2','float4','float8','float16']
+ atan2_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype atan2(gentype y, gentype x)
+ atan2Utests =
func('atan2',[atan2_input_type1,atan2_input_type2],atan2_output_type,[atan2_input_values1,atan2_input_values2],'6
* FLT_ULP')
+
+ ##### gentype atanh(gentype)
+ atanh_input_values = base_input_values
+ atanh_input_type = ['float','float2','float4','float8','float16']
+ atanh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype atanh(gentype)
+ atanhUtests =
func('atanh',[atanh_input_type],atanh_output_type,[atanh_input_values],'5 *
FLT_ULP')
+
+ ##### gentype cbrt(gentype)
+ cbrt_input_values = base_input_values
+ cbrt_input_type = ['float','float2','float4','float8','float16']
+ cbrt_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype cbrt(gentype)
+ cbrtUtests =
func('cbrt',[cbrt_input_type],cbrt_output_type,[cbrt_input_values],'4 *
FLT_ULP')
+
+ ##### gentype ceil(gentype)
+ ceil_input_values = base_input_values
+ ceil_input_type = ['float','float2','float4','float8','float16']
+ ceil_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype ceil(gentype)
+ ceilUtests =
func('ceil',[ceil_input_type],ceil_output_type,[ceil_input_values],'0 *
FLT_ULP')
+
+ ##### gentype copysign(gentype x, gentype y)
+ copysign_base_values = base_input_values
+ copysign_input_values1 = []
+ copysign_input_values2 = []
+
copysign_input_values1,copysign_input_values2=gene2ValuesLoop(copysign_input_values1,copysign_input_values2,copysign_base_values)
+ copysign_input_type1 = ['float','float2','float4','float8','float16']
+ copysign_input_type2 = ['float','float2','float4','float8','float16']
+ copysign_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype copysign(gentype x, gentype y)
+ copysignUtests =
func('copysign',[copysign_input_type1,copysign_input_type2],copysign_output_type,[copysign_input_values1,copysign_input_values2],'0
* FLT_ULP')
+
+ ##### gentype cos(gentype)
+ cos_input_values = base_input_values
+ cos_input_type = ['float','float2','float4','float8','float16']
+ cos_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype cos(gentype)
+ cosUtests =
func('cos',[cos_input_type],cos_output_type,[cos_input_values],'4 * FLT_ULP')
+
+ ##### gentype cosh(gentype)
+ cosh_input_values = base_input_values
+ cosh_input_type = ['float','float2','float4','float8','float16']
+ cosh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype cosh(gentype)
+ coshUtests =
func('cosh',[cosh_input_type],cosh_output_type,[cosh_input_values],'4 *
FLT_ULP')
+
+ ##### gentype erf(gentype)
+ erf_input_values = base_input_values
+ erf_input_type = ['float','float2','float4','float8','float16']
+ erf_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype erf(gentype)
+ erfUtests =
func('erf',[erf_input_type],erf_output_type,[erf_input_values],'16 * FLT_ULP')
+
+ ##### gentype erfc(gentype)
+ erfc_input_values = base_input_values
+ erfc_input_type = ['float','float2','float4','float8','float16']
+ erfc_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype erfc(gentype)
+ erfcUtests =
func('erfc',[erfc_input_type],erfc_output_type,[erfc_input_values],'16 *
FLT_ULP')
+
+ ##### gentype exp(gentype x)
+ exp_input_values = base_input_values
+ exp_input_type = ['float','float2','float4','float8','float16']
+ exp_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype exp(gentype x)
+ expUtests =
func('exp',[exp_input_type],exp_output_type,[exp_input_values],'4 * FLT_ULP')
+
+ ##### gentype exp2(gentype)
+ exp2_input_values = base_input_values
+ exp2_input_type = ['float','float2','float4','float8','float16']
+ exp2_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype exp2(gentype)
+ exp2Utests =
func('exp2',[exp2_input_type],exp2_output_type,[exp2_input_values],'4 *
FLT_ULP')
+
+ ##### gentype exp10(gentype)
+ exp10_input_values = base_input_values
+ exp10_input_type = ['float','float2','float4','float8','float16']
+ exp10_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype exp10(gentype)
+ exp10Utests =
func('exp10',[exp10_input_type],exp10_output_type,[exp10_input_values],'4 *
FLT_ULP')
+
+ ##### gentype expm1(gentype x)
+ expm1_input_values = base_input_values
+ expm1_input_type = ['float','float2','float4','float8','float16']
+ expm1_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype expm1(gentype x)
+ expm1Utests =
func('expm1',[expm1_input_type],expm1_output_type,[expm1_input_values],'4 *
FLT_ULP')
+
+ ##### gentype fabs(gentype)
+ fabs_input_values = base_input_values
+ fabs_input_type = ['float','float2','float4','float8','float16']
+ fabs_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fabs(gentype)
+ fabsUtests =
func('fabs',[fabs_input_type],fabs_output_type,[fabs_input_values],'0 *
FLT_ULP')
+
+ ##### gentype fdim(gentype x, gentype y)
+ fdim_base_values = base_input_values
+ fdim_input_values1 = []
+ fdim_input_values2 = []
+
fdim_input_values1,fdim_input_values2=gene2ValuesLoop(fdim_input_values1,fdim_input_values2,fdim_base_values)
+ fdim_input_type1 = ['float','float2','float4','float8','float16']
+ fdim_input_type2 = ['float','float2','float4','float8','float16']
+ fdim_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fdim(gentype x, gentype y)
+ fdimUtests =
func('fdim',[fdim_input_type1,fdim_input_type2],fdim_output_type,[fdim_input_values1,fdim_input_values2],'0
* FLT_ULP')
+
+ ##### gentype floor(gentype)
+ floor_input_values = base_input_values
+ floor_input_type = ['float','float2','float4','float8','float16']
+ floor_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype floor(gentype)
+ floorUtests =
func('floor',[floor_input_type],floor_output_type,[floor_input_values],'0 *
FLT_ULP')
+
+ ##### gentype fma(gentype a, gentype b, gentype c)
+ fma_base_values = base_input_values
+ fma_input_values1 = []
+ fma_input_values2 = []
+ fma_input_values3 = []
+
fma_input_values1,fma_input_values2,fma_input_values3=gene3ValuesLoop(fma_input_values1,fma_input_values2,fma_input_values3,fma_base_values)
+ fma_input_type1 = ['float','float2','float4','float8','float16']
+ fma_input_type2 = ['float','float2','float4','float8','float16']
+ fma_input_type3 = ['float','float2','float4','float8','float16']
+ fma_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fma(gentype a, gentype b, gentype c)
+ fmaUtests =
func('fma',[fma_input_type1,fma_input_type2,fma_input_type3],fma_output_type,[fma_input_values1,fma_input_values2,fma_input_values3],'0
* FLT_ULP')
+
+ ##### gentype fmax(gentype x, gentype y)
+ fmax_base_values = base_input_values
+ fmax_input_values1 = []
+ fmax_input_values2 = []
+
fmax_input_values1,fmax_input_values2=gene2ValuesLoop(fmax_input_values1,fmax_input_values2,fmax_base_values)
+ fmax_input_type1 = ['float','float2','float4','float8','float16']
+ fmax_input_type2 = ['float','float2','float4','float8','float16']
+ fmax_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fmax(gentype x, gentype y)
+ fmaxUtests =
func('fmax',[fmax_input_type1,fmax_input_type2],fmax_output_type,[fmax_input_values1,fmax_input_values2],'0
* FLT_ULP')
+
+ ##### gentypef fmax(gentypef x, float y)
+# fmax_gentypef_base_values = base_input_values
+# fmax_gentypef_input_values1 = []
+# fmax_gentypef_input_values2 = []
+#
fmax_gentypef_input_values2,fmax_gentypef_input_values1=gene2ValuesLoop(fmax_gentypef_input_values1,fmax_gentypef_input_values2,fmax_gentypef_base_values)
+# fmax_gentypef_input_type1 = ['float','float2','float4','float8','float16']
+# fmax_gentypef_input_type2 = ['float','float','float','float','float']
+# fmax_gentypef_output_type = ['float','float2','float4','float8','float16']
+# ##### gentypef fmax(gentypef x, float y)
+# fmax_gentypefUtests =
func('gentypef_fmax',[fmax_gentypef_input_type1,fmax_gentypef_input_type2],fmax_gentypef_output_type,[fmax_gentypef_input_values1,fmax_gentypef_input_values2],'0
* FLT_ULP')
+
+ ##### gentype fmin(gentype x, gentype y)
+ fmin_base_values = base_input_values
+ fmin_input_values1 = []
+ fmin_input_values2 = []
+
fmin_input_values1,fmin_input_values2=gene2ValuesLoop(fmin_input_values1,fmin_input_values2,fmin_base_values)
+ fmin_input_type1 = ['float','float2','float4','float8','float16']
+ fmin_input_type2 = ['float','float2','float4','float8','float16']
+ fmin_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fmin(gentype x, gentype y)
+ fminUtests =
func('fmin',[fmin_input_type1,fmin_input_type2],fmin_output_type,[fmin_input_values1,fmin_input_values2],'0
* FLT_ULP')
+
+# ##### gentypef fmin(gentypef x, float y)
+# fmin_gentypef_base_values = base_input_values
+# fmin_gentypef_input_values1 = []
+# fmin_gentypef_input_values2 = []
+#
fmin_gentypef_input_values2,fmin_gentypef_input_values1=gene2ValuesLoop(fmin_gentypef_input_values1,fmin_gentypef_input_values2,fmin_gentypef_base_values)
+# fmin_gentypef_input_type1 = ['float','float2','float4','float8','float16']
+# fmin_gentypef_input_type2 = ['float','float','float','float','float']
+# fmin_gentypef_output_type = ['float','float2','float4','float8','float16']
+# ##### gentypef fmin(gentypef x, float y)
+# fmin_gentypefUtests =
func('gentypef_fmin',[fmin_gentypef_input_type1,fmin_gentypef_input_type2],fmin_gentypef_output_type,[fmin_gentypef_input_values1,fmin_gentypef_input_values2],'0
* FLT_ULP')
+#
+ ##### gentype fmod(gentype x, gentype y)
+ fmod_base_values = base_input_values
+ fmod_input_values1 = []
+ fmod_input_values2 = []
+
fmod_input_values1,fmod_input_values2=gene2ValuesLoop(fmod_input_values1,fmod_input_values2,fmod_base_values)
+ fmod_input_type1 = ['float','float2','float4','float8','float16']
+ fmod_input_type2 = ['float','float2','float4','float8','float16']
+ fmod_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype fmod(gentype x, gentype y)
+ fmodUtests =
func('fmod',[fmod_input_type1,fmod_input_type2],fmod_output_type,[fmod_input_values1,fmod_input_values2],'0
* FLT_ULP')
+
+ ##### gentype hypot(gentype x, gentype y)
+ hypot_base_values = base_input_values
+ hypot_input_values1 = []
+ hypot_input_values2 = []
+
hypot_input_values1,hypot_input_values2=gene2ValuesLoop(hypot_input_values1,hypot_input_values2,hypot_base_values)
+ hypot_input_type1 = ['float','float2','float4','float8','float16']
+ hypot_input_type2 = ['float','float2','float4','float8','float16']
+ hypot_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype hypot(gentype x, gentype y)
+ hypotUtests =
func('hypot',[hypot_input_type1,hypot_input_type2],hypot_output_type,[hypot_input_values1,hypot_input_values2],'4
* FLT_ULP')
+
+ ##### intn ilogb(floartn x)
+ ilogb_input_values = base_input_values
+ ilogb_input_type = ['float','float2','float4','float8','float16']
+ ilogb_output_type = ['int','int2','int4','int8','int16']
+ ##### intn ilogb(floatn x)
+ ilogbUtests =
func('ilogb',[ilogb_input_type],ilogb_output_type,[ilogb_input_values],'0 *
INT_ULP')
+
+ ##### gentype log(gentype)
+ log_input_values = base_input_values
+ log_input_type = ['float','float2','float4','float8','float16']
+ log_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype log(gentype)
+ logUtests =
func('log',[log_input_type],log_output_type,[log_input_values],'4 * FLT_ULP')
+
+ ##### gentype log2(gentype)
+ log2_input_values = base_input_values
+ log2_input_type = ['float','float2','float4','float8','float16']
+ log2_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype log2(gentype)
+ log2Utests =
func('log2',[log2_input_type],log2_output_type,[log2_input_values],'4 *
FLT_ULP')
+
+ ##### gentype log10(gentype)
+ log10_input_values = base_input_values
+ log10_input_type = ['float','float2','float4','float8','float16']
+ log10_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype log10(gentype)
+ log10Utests =
func('log10',[log10_input_type],log10_output_type,[log10_input_values],'4 *
FLT_ULP')
+
+ ##### gentype log1p(gentype x)
+ log1p_input_values = base_input_values
+ log1p_input_type = ['float','float2','float4','float8','float16']
+ log1p_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype log1p(gentype x)
+ log1pUtests =
func('log1p',[log1p_input_type],log1p_output_type,[log1p_input_values],'4 *
FLT_ULP')
+
+ ##### gentype logb(gentype x)
+ logb_input_values = base_input_values
+ logb_input_type = ['float','float2','float4','float8','float16']
+ logb_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype logb(gentype x)
+ logbUtests =
func('logb',[logb_input_type],logb_output_type,[logb_input_values],'0 *
FLT_ULP')
+
+# ##### gentype mad(gentype a, gentype b, gentype c)
+# #mad_base_values = base_input_values
+# mad_base_values = base_input_values
+# mad_input_values1 = []
+# mad_input_values2 = []
+# mad_input_values3 = []
+#
mad_input_values1,mad_input_values2,mad_input_values3=gene3ValuesLoop(mad_input_values1,mad_input_values2,mad_input_values3,mad_base_values)
+# mad_input_type1 = ['float','float2','float4','float8','float16']
+# mad_input_type2 = ['float','float2','float4','float8','float16']
+# mad_input_type3 = ['float','float2','float4','float8','float16']
+# mad_output_type = ['float','float2','float4','float8','float16']
+# ##### gentype mad(gentype a, gentype b, gentype c)
+# madUtests =
func('mad',[mad_input_type1,mad_input_type2,mad_input_type3],mad_output_type,[mad_input_values1,mad_input_values2,mad_input_values3],'0
* FLT_ULP')
+
+# ##### gentype maxmag(gentype x, gentype y)
+# maxmag_base_values = base_input_values
+# maxmag_input_values1 = []
+# maxmag_input_values2 = []
+#
maxmag_input_values1,maxmag_input_values2=gene2ValuesLoop(maxmag_input_values1,maxmag_input_values2,maxmag_base_values)
+# maxmag_input_type1 = ['float','float2','float4','float8','float16']
+# maxmag_input_type2 = ['float','float2','float4','float8','float16']
+# maxmag_output_type = ['float','float2','float4','float8','float16']
+# ##### gentype maxmag(gentype x, gentype y)
+# maxmagUtests =
func('maxmag',[maxmag_input_type1,maxmag_input_type2],maxmag_output_type,[maxmag_input_values1,maxmag_input_values2],'0
* FLT_ULP')
+
+# ##### gentype minmag(gentype x, gentype y)
+# minmag_base_values = base_input_values
+# minmag_input_values1 = []
+# minmag_input_values2 = []
+#
minmag_input_values1,minmag_input_values2=gene2ValuesLoop(minmag_input_values1,minmag_input_values2,minmag_base_values)
+# minmag_input_type1 = ['float','float2','float4','float8','float16']
+# minmag_input_type2 = ['float','float2','float4','float8','float16']
+# minmag_output_type = ['float','float2','float4','float8','float16']
+# ##### gentype minmag(gentype x, gentype y)
+# minmagUtests =
func('minmag',[minmag_input_type1,minmag_input_type2],minmag_output_type,[minmag_input_values1,minmag_input_values2],'0
* FLT_ULP')
+
+ ##### floatn nan(uintn nancode)
+ nan_input_values = base_input_values
+ nan_input_type = ['uint','uint2','uint4','uint8','uint16']
+ nan_output_type = ['float','float2','float4','float8','float16']
+ ###### floatn nan(uintn nancode)
+ # nanUtests =
func('nan',[nan_input_type],nan_output_type,[nan_input_values],'0 * FLT_ULP')
+
+ ##### gentype nextafter(gentype x, gentype y)
+ nextafter_base_values = base_input_values
+ nextafter_input_values1 = []
+ nextafter_input_values2 = []
+
nextafter_input_values1,nextafter_input_values2=gene2ValuesLoop(nextafter_input_values1,nextafter_input_values2,nextafter_base_values)
+ nextafter_input_type1 = ['float','float2','float4','float8','float16']
+ nextafter_input_type2 = ['float','float2','float4','float8','float16']
+ nextafter_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype nextafter(gentype x, gentype y)
+ nextafterUtests =
func('nextafter',[nextafter_input_type1,nextafter_input_type2],nextafter_output_type,[nextafter_input_values1,nextafter_input_values2],'0
* FLT_ULP')
+
+ ##### gentype pow(gentype x, gentype y)
+ pow_base_values = base_input_values
+ pow_input_values1 = []
+ pow_input_values2 = []
+
pow_input_values1,pow_input_values2=gene2ValuesLoop(pow_input_values1,pow_input_values2,pow_base_values)
+ pow_input_type1 = ['float','float2','float4','float8','float16']
+ pow_input_type2 = ['float','float2','float4','float8','float16']
+ pow_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype pow(gentype x, gentype y)
+ powUtests =
func('pow',[pow_input_type1,pow_input_type2],pow_output_type,[pow_input_values1,pow_input_values2],'16
* FLT_ULP')
+
+# ##### floatn pown(floatn x, intn y)
+# pown_input_values1 =
[FLT_MAX_POSI,FLT_MIN_NEGA,FLT_MIN_POSI,FLT_MAX_NEGA,80, -80, 3.14, -3.14,
-0.5, 0.5, 1, -1, 0.0,6,-6,1500.24,-1500.24]
+# pown_input_values2 = [-1,-2,-3,4,5,6,7,8,9,10,11,12,13,14,15,16,12]
+# pown_input_type1 = ['float','float2','float4','float8','float16']
+# pown_input_type2 = ['int','int2','int4','int8','int16']
+# pown_output_type = ['float','float2','float4','float8','float16']
+# ##### floatn pown(floatn x, intn y)
+# pownUtests =
func('pown',[pown_input_type1,pown_input_type2],pown_output_type,[pown_input_values1,pown_input_values2],'16
* FLT_ULP')
+# ##### gentype pown(gentype x, gentype y)
+#
+# ##### gentype powr(gentype x, gentype y)
+# powr_input_values1 =
[FLT_MAX_POSI,FLT_MIN_NEGA,FLT_MIN_POSI,FLT_MAX_NEGA,80, -80, 3.14, -3.14,
-0.5, 0.5, 1, -1, 0.0,6,-6,1500.24,-1500.24]
+# powr_input_values2 = [1,2,3.14,4,5,6,7,8,9.889,10,11,12,13,14.33,15,0,12]
+# powr_input_type1 = ['float','float2','float4','float8','float16']
+# powr_input_type2 = ['float','float2','float4','float8','float16']
+# powr_output_type = ['float','float2','float4','float8','float16']
+# powrUtests =
func('powr',[powr_input_type1,powr_input_type2],powr_output_type,[powr_input_values1,powr_input_values2],'16
* FLT_ULP')
+
+ ##### gentype remainder(gentype x, gentype y)
+ remainder_base_values = base_input_values
+ remainder_input_values1 = []
+ remainder_input_values2 = []
+
remainder_input_values1,remainder_input_values2=gene2ValuesLoop(remainder_input_values1,remainder_input_values2,remainder_base_values)
+ remainder_input_type1 = ['float','float2','float4','float8','float16']
+ remainder_input_type2 = ['float','float2','float4','float8','float16']
+ remainder_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype remainder(gentype x, gentype y)
+ remainderUtests =
func('remainder',[remainder_input_type1,remainder_input_type2],remainder_output_type,[remainder_input_values1,remainder_input_values2],'0
* FLT_ULP')
+
+ ##### gentype rint(gentype x)
+ rint_input_values = base_input_values
+ rint_input_type = ['float','float2','float4','float8','float16']
+ rint_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype rint(gentype)
+ rintUtests =
func('rint',[rint_input_type],rint_output_type,[rint_input_values],'0 *
FLT_ULP')
+
+# ##### floatn rootn(floatn x, intn y)
+# rootn_input_values1 =
[FLT_MAX_POSI,FLT_MIN_NEGA,FLT_MIN_POSI,FLT_MAX_NEGA,80, -80, 3.14, -3.14,
-0.5, 0.5, 1, -1, 0.0,6,-6,1500.24,-1500.24,2,3,4]
+# rootn_input_values2 = [-1,-2,-3,2,3,6,7,8,9,2,11,12,13,14,15,16,2,2,2,2]
+# rootn_input_type1 = ['float','float2','float4','float8','float16']
+# rootn_input_type2 = ['int','int2','int4','int8','int16']
+# rootn_output_type = ['float','float2','float4','float8','float16']
+# ##### floatn rootn(floatn x, intn y)
+# rootnUtests =
func('rootn',[rootn_input_type1,rootn_input_type2],rootn_output_type,[rootn_input_values1,rootn_input_values2],'4
* FLT_ULP')
+
+ ##### gentype round(gentype x)
+ round_input_values = base_input_values
+ round_input_type = ['float','float2','float4','float8','float16']
+ round_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype round(gentype x)
+ roundUtests =
func('round',[round_input_type],round_output_type,[round_input_values],'0 *
FLT_ULP')
+
+ ##### gentype rsqrt(gentype)
+ rsqrt_input_values = base_input_values
+ rsqrt_input_type = ['float','float2','float4','float8','float16']
+ rsqrt_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype rsqrt(gentype)
+ rsqrt_cpu_func='''
+static float rsqrt(float x)
+{ return 1/sqrt(x);} '''
+ rsqrtUtests =
func('rsqrt',[rsqrt_input_type],rsqrt_output_type,[rsqrt_input_values],'4 *
FLT_ULP', rsqrt_cpu_func)
+
+
+ ##### gentype sin(gentype)
+ sin_input_values = base_input_values
+ sin_input_type = ['float','float2','float4','float8','float16']
+ sin_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype sin(gentype)
+ sinUtests =
func('sin',[sin_input_type],sin_output_type,[sin_input_values],'4 * FLT_ULP')
+
+# ##### gentype sincos(gentype)
+# sincos_input_values1 =
[FLT_MAX_POSI,FLT_MIN_NEGA,FLT_MIN_POSI,FLT_MAX_NEGA,80, -80, 3.14, -3.14,
-0.5, 0.5, 1, -1, 0.0,6,-6,1500.24,-1500.24]
+# sincos_input_values2 = []
+# sincos_input_type1 = ['float','float2','float4','float8','float16']
+# sincos_input_type2 = ['float','float2','float4','float8','float16']
+# sincos_output_type = ['float','float2','float4','float8','float16']
+# ###### gentype sincos(gentype)
+# # sincosUtests =
func('sincos',[sincos_input_type1,sincos_input_type2],sincos_output_type,[sincos_input_values1,sincos_input_values2],'4
* FLT_ULP')
+
+ ##### gentype sinh(gentype)
+ sinh_input_values = base_input_values
+ sinh_input_type = ['float','float2','float4','float8','float16']
+ sinh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype sinh(gentype)
+ sinhUtests =
func('sinh',[sinh_input_type],sinh_output_type,[sinh_input_values],'4 *
FLT_ULP')
+
+ ##### gentype sqrt(gentype)
+ sqrt_input_values = base_input_values
+ sqrt_input_type = ['float','float2','float4','float8','float16']
+ sqrt_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype sqrt(gentype)
+ sqrtUtests =
func('sqrt',[sqrt_input_type],sqrt_output_type,[sqrt_input_values],'4 *
FLT_ULP')
+
+ ##### gentype tan(gentype)
+ tan_input_values = base_input_values
+ tan_input_type = ['float','float2','float4','float8','float16']
+ tan_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype tan(gentype)
+ tanUtests =
func('tan',[tan_input_type],tan_output_type,[tan_input_values],'5 * FLT_ULP')
+
+ ##### gentype tanh(gentype)
+ tanh_input_values = base_input_values
+ tanh_input_type = ['float','float2','float4','float8','float16']
+ tanh_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype tanh(gentype)
+ tanhUtests =
func('tanh',[tanh_input_type],tanh_output_type,[tanh_input_values],'5 *
FLT_ULP')
+
+ ##### gentype tgamma(gentype)
+ tgamma_input_values = base_input_values
+ tgamma_input_type = ['float','float2','float4','float8','float16']
+ tgamma_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype tgamma(gentype)
+ tgammaUtests =
func('tgamma',[tgamma_input_type],tgamma_output_type,[tgamma_input_values],'16
* FLT_ULP')
+
+ ##### gentype trunc(gentype)
+ trunc_input_values = base_input_values
+ trunc_input_type = ['float','float2','float4','float8','float16']
+ trunc_output_type = ['float','float2','float4','float8','float16']
+ ##### gentype trunc(gentype)
+ truncUtests =
func('trunc',[trunc_input_type],trunc_output_type,[trunc_input_values],'0 *
FLT_ULP')
+
+if __name__ == "__main__":
+ main()
--
1.7.6.4
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
