[Beignet] [PATCH v2] Add test cases generator.

Yi Sun yi.sun at intel.com
Fri Nov 29 23:11:13 PST 2013


v1:
File utest_generator.py contain the base class and function for generating
File utest_math_gen.py can generate most math function for all the gentype
utest_math_gen.py can be run during cmake.

v2:
1. Put all the generated unit test cases to folder utest/generated.
2. Delete all generated folder while involve make clean.
3. At the top of the generated test cases, add some comments
4. Insert of defined FLT_ULP(0.000001) as the ulp unit, caculate the float ulp before using it.
5. Add several math functions' test case.

Signed-off-by: Yi Sun <yi.sun at intel.com>
Signed-off-by: YangweiX Shui <yangwei.shui at intel.com>
---
 utests/CMakeLists.txt     |    8 +
 utests/utest_generator.py |  308 +++++++++++++++++++++++
 utests/utest_helper.cpp   |   13 +
 utests/utest_helper.hpp   |    1 +
 utests/utest_math_gen.py  |  599 +++++++++++++++++++++++++++++++++++++++++++++
 5 files changed, 929 insertions(+), 0 deletions(-)
 create mode 100644 utests/utest_generator.py
 create mode 100755 utests/utest_math_gen.py

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 5e0bc19..f858e9e 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -1,10 +1,18 @@
 INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}
                     ${CMAKE_CURRENT_SOURCE_DIR}/../include)
 
+EXEC_PROGRAM(mkdir ${CMAKE_CURRENT_SOURCE_DIR} ARGS generated -p)
+EXEC_PROGRAM(python ${CMAKE_CURRENT_SOURCE_DIR} ARGS utest_math_gen.py OUTPUT_VARIABLE GEN_STRING)
+string(REGEX REPLACE " " ";" ADDFUNC ${GEN_STRING})
+string(REGEX REPLACE " " "\n" NAMELIST ${GEN_STRING})
+MESSAGE(STATUS "Generated Builtin Math Functions:\n" ${NAMELIST})
+set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES generated/)
+
 link_directories (${LLVM_LIBRARY_DIR})
 set (utests_sources
   cl_create_kernel.cpp
   utest_error.c
+  ${ADDFUNC}
   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..3a52d3d
--- /dev/null
+++ b/utests/utest_generator.py
@@ -0,0 +1,308 @@
+#!/usr/bin/python
+import os,sys,re
+
+FLT_MAX_POSI='0x1.fffffep127f'
+FLT_MIN_NEGA='-0x1.fffffep127f'
+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 += [j]
+
+  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='''/*
+This file is generated by utest_generator.py.
+Usually you need NOT modify this file manually.
+But when any bug occured, you can change the value of udebug from 0 to 1,
+which can print more values and information to assist debuging the issue.
+*/
+
+#include "utest_helper.hpp"
+#include <cmath>
+#include <algorithm>
+
+#define udebug 0
+#define FLT_MAX 0x1.fffffep127f
+#define FLT_MIN 0x1.0p-126f
+#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("generated/%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("generated/%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_helper.cpp b/utests/utest_helper.cpp
index 65af727..8155300 100644
--- a/utests/utest_helper.cpp
+++ b/utests/utest_helper.cpp
@@ -641,3 +641,16 @@ int cl_check_image(const int *img, int w, int h, const char *bmp)
   return (float(discrepancy) / float(n) > max_error_ratio) ? 0 : 1;
 }
 
+static const float cl_float_ulp()
+{
+  float x = 1.0, ulp = 0.1, delta = 0.1;
+
+  while( x != x + delta)
+  {
+    ulp = delta;
+    delta *= 0.1;
+  }
+  return ulp;
+}
+
+const float FLT_ULP = cl_float_ulp();
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index 29a21d5..4898045 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -172,6 +172,7 @@ extern cl_mem buf[MAX_BUFFER_N];
 extern void* buf_data[MAX_BUFFER_N];
 extern size_t globals[3];
 extern size_t locals[3];
+extern const float FLT_ULP;
 
 enum {
   SOURCE = 0,
diff --git a/utests/utest_math_gen.py b/utests/utest_math_gen.py
new file mode 100755
index 0000000..204481b
--- /dev/null
+++ b/utests/utest_math_gen.py
@@ -0,0 +1,599 @@
+#!/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 acospi(gentype x)
+  acospi_input_values = base_input_values
+  acospi_input_type = ['float','float2','float4','float8','float16']
+  acospi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype acospi(gentype x)
+  acospi_cpu_func='''
+static float acospi(float x){
+  return acos(x)/M_PI;
+} '''
+  acospiUtests = func('acospi',[acospi_input_type],acospi_output_type,[acospi_input_values],'4 * FLT_ULP',acospi_cpu_func)
+  
+  ##### 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 asinpi(gentype x)
+  asinpi_input_values = base_input_values
+  asinpi_input_type = ['float','float2','float4','float8','float16']
+  asinpi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype asinpi(gentype x)
+  asinpi_cpu_func='''
+static float asinpi(float x){
+  return asin(x)/M_PI;
+} '''
+  asinpiUtests = func('asinpi',[asinpi_input_type],asinpi_output_type,[asinpi_input_values],'4 * FLT_ULP',asinpi_cpu_func)
+  
+  ##### 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 atanpi(gentype x)
+  atanpi_input_values = base_input_values
+  atanpi_input_type = ['float','float2','float4','float8','float16']
+  atanpi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype atanpi(gentype x)
+  atanpi_cpu_func='''
+static float atanpi(float x){
+  return atan(x)/M_PI;
+} '''
+  atanpiUtests = func('atanpi',[atanpi_input_type],atanpi_output_type,[atanpi_input_values],'4 * FLT_ULP',atanpi_cpu_func)
+  
+#  ##### gentype atan2pi(gentype y, gentype x)
+#  atan2pi_base_values = base_input_values
+#  atan2pi_input_values1 = []
+#  atan2pi_input_values2 = []
+#  atan2pi_input_values1,atan2pi_input_values2=gene2ValuesLoop(atan2pi_input_values1,atan2pi_input_values2,atan2pi_base_values)
+#  atan2pi_input_type1 = ['float','float2','float4','float8','float16']
+#  atan2pi_input_type2 = ['float','float2','float4','float8','float16']
+#  atan2pi_output_type = ['float','float2','float4','float8','float16']
+#  ##### gentype atan2pi(gentype y, gentype x)
+#  atan2pi_cpu_func='''
+#static float atan2pi(float y, float x){
+#  return atan2(y,x)/M_PI;
+#} '''
+#  atan2piUtests = func('atan2pi',[atan2pi_input_type1,atan2pi_input_type2],atan2pi_output_type,[atan2pi_input_values1,atan2pi_input_values2],'6 * FLT_ULP',atan2pi_cpu_func)
+  
+  ##### 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 cospi(gentype x)
+  cospi_input_values = base_input_values
+  cospi_input_type = ['float','float2','float4','float8','float16']
+  cospi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype cospi(gentype x)
+  cospi_cpu_func='''
+static float cospi(float x){
+  return cos(M_PI * x);
+} '''
+  cospiUtests = func('cospi',[cospi_input_type],cospi_output_type,[cospi_input_values],'2 * FLT_ULP',cospi_cpu_func)
+  
+  ##### 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 lgamma(gentype x)
+  lgamma_input_values = base_input_values
+  lgamma_input_type = ['float','float2','float4','float8','float16']
+  lgamma_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype lgamma(gentype x)
+  lgammaUtests = func('lgamma',[lgamma_input_type],lgamma_output_type,[lgamma_input_values],'4 * FLT_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)
+  mad_cpu_func='''
+static float mad(float a, float b, float c){
+  return a * b + 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',mad_cpu_func)
+  
+  ##### 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)
+  maxmag_cpu_func='''
+static float maxmag(float x, float y){
+  if(fabs(x) > fabs(y))
+    return x;
+  else if (fabs(x) < fabs(y))
+    return y;
+  else
+    return fmax(x,y);
+} '''
+  maxmagUtests = func('maxmag',[maxmag_input_type1,maxmag_input_type2],maxmag_output_type,[maxmag_input_values1,maxmag_input_values2],'0 * FLT_ULP',maxmag_cpu_func)
+  
+  ##### 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)
+  minmag_cpu_func='''
+static float minmag(float x, float y){
+  if(fabs(x) < fabs(y))
+    return x;
+  else if (fabs(x) > fabs(y))
+    return y;
+  else
+    return fmin(x,y);
+} '''
+  minmagUtests = func('minmag',[minmag_input_type1,minmag_input_type2],minmag_output_type,[minmag_input_values1,minmag_input_values2],'0 * FLT_ULP',minmag_cpu_func)
+  
+  ##### 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)
+  pown_cpu_func='''
+static float pown(float x, int y){
+    return pow(x,y);
+} '''
+  pownUtests = func('pown',[pown_input_type1,pown_input_type2],pown_output_type,[pown_input_values1,pown_input_values2],'16 * FLT_ULP', pown_cpu_func)
+  
+  ##### 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']
+  ##### gentype powr(gentype x, gentype y)
+  powr_cpu_func='''
+static float powr(float x, int y){
+    return pow(x,y);
+} '''
+  powrUtests = func('powr',[powr_input_type1,powr_input_type2],powr_output_type,[powr_input_values1,powr_input_values2],'16 * FLT_ULP', powr_cpu_func)
+  
+  ##### 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)
+  rootn_cpu_func='''
+static float rootn(float x, int y){
+    return pow(x,1.0/y);
+} '''
+  rootnUtests = func('rootn',[rootn_input_type1,rootn_input_type2],rootn_output_type,[rootn_input_values1,rootn_input_values2],'4 * FLT_ULP',rootn_cpu_func)
+  
+  ##### 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 sinpi(gentype x)
+  sinpi_input_values = base_input_values
+  sinpi_input_type = ['float','float2','float4','float8','float16']
+  sinpi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype sinpi(gentype x)
+  sinpi_cpu_func='''
+static float sinpi(float x){
+  return sin(M_PI*x);
+} '''
+  sinpiUtests = func('sinpi',[sinpi_input_type],sinpi_output_type,[sinpi_input_values],'4 * FLT_ULP',sinpi_cpu_func)
+  
+  ##### 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 tanpi(gentype x)
+  tanpi_input_values = base_input_values
+  tanpi_input_type = ['float','float2','float4','float8','float16']
+  tanpi_output_type = ['float','float2','float4','float8','float16']
+  ##### gentype tanpi(gentype x)
+  tanpi_cpu_func='''
+static float tanpi(float x){
+  return tan(M_PI*x);
+} '''
+  tanpiUtests = func('tanpi',[tanpi_input_type],tanpi_output_type,[tanpi_input_values],'4 * FLT_ULP',tanpi_cpu_func)
+  
+  ##### 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



More information about the Beignet mailing list