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

Sun, Yi yi.sun at intel.com
Thu Dec 26 22:51:26 PST 2013


Anyone else can reproduce this issue?
I have tried it on 3 machines, but all following command are passed.
make -j(2, 4, 8, 16, 32)


Thanks
  --Sun, Yi

> -----Original Message-----
> From: Zou, Nanhai
> Sent: Friday, December 27, 2013 2:38 PM
> To: Zhigang Gong; Sun, Yi
> Cc: Song, Ruiling; beignet at lists.freedesktop.org; Shui, YangweiX
> Subject: RE: [Beignet] [PATCH v5] Add test cases generator.
> 
> Hi Yi,
> 	It seems that this patch breaks make -jx
> 
> Can you have a check?
> 
> Thanks
> Zou Nanhai
> 
> -----Original Message-----
> From: beignet-bounces at lists.freedesktop.org
> [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of Zhigang Gong
> Sent: Wednesday, December 25, 2013 2:01 PM
> To: Sun, Yi
> Cc: Song, Ruiling; beignet at lists.freedesktop.org; Shui, YangweiX
> Subject: Re: [Beignet] [PATCH v5] Add test cases generator.
> 
> Pushed, thanks.
> 
> On Tue, Dec 24, 2013 at 11:15:18AM +0800, Yi Sun wrote:
> >     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. Instead of defined FLT_ULP(0.000001) as the ulp unit, caculate the
> float ulp before using it.
> >     5. Add several math functions' test case.
> >
> >     v3:
> >     1. Refine the calculation for float, and calculate each float got from cpu
> function.
> >
> >     v4:
> >     Refine the calculation for float.
> >
> >     Following fucntions test cases fail with input 0, 1 or 3.14:
> > 	builtin_atan2_float
> > 	builtin_atanh_float
> > 	builtin_rootn_float
> >     builtin_cos_float
> >     builtin_cospi_float
> >     builtin_erf_float
> >     builtin_erfc_float
> >     builtin_mad_float
> >     builtin_nextafter_float
> >     builtin_pown_float
> >     builtin_powr_float
> >     builtin_rint_float
> >     builtin_sinpi_float
> >     builtin_tan_float
> >     builtin_tanpi_float
> >
> >     v5:
> >     remove case builtin_mad_float
> >
> > 	todo:
> > 	atan2pi
> > 	fmax
> > 	fmin
> > 	sincos
> >
> > Signed-off-by: Yi Sun <yi.sun at intel.com>
> > Signed-off-by: Yangwei Shui <yangweix.shui at intel.com>
> > ---
> >  utests/CMakeLists.txt     |    9 +
> >  utests/utest_generator.py |  374 +++++++++++++++++++++++++++++++
> >  utests/utest_helper.cpp   |   30 +++
> >  utests/utest_helper.hpp   |    6 +
> >  utests/utest_math_gen.py  |  531
> > +++++++++++++++++++++++++++++++++++++++++++++
> >  5 files changed, 950 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..836c80d 100644
> > --- a/utests/CMakeLists.txt
> > +++ b/utests/CMakeLists.txt
> > @@ -1,10 +1,19 @@
> >  INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}
> >                      ${CMAKE_CURRENT_SOURCE_DIR}/../include)
> >
> > +EXEC_PROGRAM(mkdir ${CMAKE_CURRENT_SOURCE_DIR} ARGS
> generated -p)
> > +##### Math Function Part:
> > +EXEC_PROGRAM(python ${CMAKE_CURRENT_SOURCE_DIR} ARGS
> > +utest_math_gen.py OUTPUT_VARIABLE GEN_MATH_STRING) string(REGEX
> > +REPLACE " " ";" ADDMATHFUNC ${GEN_MATH_STRING}) string(REGEX
> REPLACE
> > +" " "\n" NAMEMATHLIST ${GEN_MATH_STRING}) MESSAGE(STATUS
> "Generated
> > +Builtin Math Functions:\n" ${NAMEMATHLIST})
> > +set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES
> > +generated/)
> > +
> >  link_directories (${LLVM_LIBRARY_DIR})  set (utests_sources
> >    cl_create_kernel.cpp
> >    utest_error.c
> > +  ${ADDMATHFUNC}
> >    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..626ac96
> > --- /dev/null
> > +++ b/utests/utest_generator.py
> > @@ -0,0 +1,374 @@
> > +#!/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':'%.20f','int':'%d','double':'%lf','uint':'%d','
> > +string':'%s'}
> > +
> > +
> > +def ulpUnit(ulpSize):
> > +  return re.findall(r"([a-zA-Z_]+)",ulpSize)[0]
> > +
> > +def ulpNum(ulpSize):
> > +  return re.findall(r"([0-9]+)",ulpSize)[0]
> > +
> > +def udebug(ulpSize,returnType):
> > +  #ulpUnit=re.findall(r"([a-zA-Z_]+)",ulpSize)[0]
> > +  #ulpNum=re.findall(r"([0-9]+)",ulpSize)[0]
> > +  text='''
> > +    static const char* INFORNAN;
> > +    static %s ULPSIZE;
> > +
> > +    if (isinf(cpu_data[index])){
> > +      INFORNAN="INF";
> > +    }
> > +    else if (isnan(cpu_data[index])){
> > +      INFORNAN="NAN";
> > +    }
> > +    else{
> > +      ULPSIZE=cl_%s(cpu_data[index]) * %s;
> > +    }
> > +
> > +#if udebug
> > +    if (isinf(cpu_data[index])){
> > +      if (isinf(gpu_data[index]))
> > +        printf("%s expect:%s\\n", log, INFORNAN);
> > +      else
> > +        printf_c("%s expect:%s\\n", log, INFORNAN);
> > +      }
> > +    else if (isnan(cpu_data[index])){
> > +      if (isnan(gpu_data[index]))
> > +        printf("%s expect:%s\\n", log, INFORNAN);
> > +      else
> > +        printf_c("%s expect:%s\\n", log, INFORNAN);
> > +      }
> > +    else if (diff <= ULPSIZE){
> > +      printf("%s expect:%s\\n", log, ULPSIZE);
> > +      }
> > +    else
> > +      printf_c("%s expect:%s\\n", log, ULPSIZE); #else
> > +    if (isinf(cpu_data[index])){
> > +      sprintf(log, "%s expect:%s\\n", log, INFORNAN);
> > +      OCL_ASSERTM(isinf(gpu_data[index]),log);
> > +      }
> > +    else if (isnan(cpu_data[index])){
> > +      sprintf(log, "%s expect:%s\\n", log, INFORNAN);
> > +      OCL_ASSERTM(isnan(gpu_data[index]),log);
> > +      }
> > +    else{
> > +      sprintf(log, "%s expect:%s\\n", log, ULPSIZE);
> > +      OCL_ASSERTM(fabs(gpu_data[index]-cpu_data[index]) <= ULPSIZE,
> log);
> > +      }
> > +#endif
> > +  }
> > +}\n'''%(returnType,\
> > +        ulpUnit(ulpSize),ulpNum(ulpSize),\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['%s'%(returnType)],\
> > +        paraTypeList['string'],paraTypeList['%s'%(returnType)],\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['string'],\
> > +        paraTypeList['string'],paraTypeList['%s'%(returnType)])
> > +
> > +  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,cpuFuncName,inputType,outputType,values,ulp,
> cpu_func=''):
> > +    self.funcName = name
> > +    self.cpuFuncName = cpuFuncName
> > +    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 <stdio.h>
> > +#include <math.h>
> > +#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 returnVector(self,index):
> > +    returnVector=re.findall(r"[0-9]+",self.outputtype[index])
> > +    if returnVector:
> > +      returnVector=returnVector[0]
> > +    else:
> > +      returnVector=1
> > +    return returnVector
> > +
> > +  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]);" ]
> > +    self.cpplines += [ "const int vector =
> > +%s;\n"%(self.argvector(self.inputtype.__len__()-1,index)) ]
> > +
> > +#####Cpu Function
> > +  def GenCpuCompilerMath(self,index):
> > +    #namesuffix=self.inputtype[0][index]
> > +    defline='static void cpu_compiler_math(%s *dst, '%(self.retType(index))
> > +    cpufunargs='('
> > +    funcline = ['{']
> > +    vectorargs=[]
> > +
> > +    if (self.returnVector(index) == 1 and self.argvector(0,index) != 1):
> > +      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 ','
> > +        vectorargs.append('(')
> > +      for i in range(0,self.values.__len__()):
> > +        for j in range(0,self.vector):
> > +          vectorargs += "x%d%d"%(i+1,j+1)
> > +          vectorargs += ( j == self.vector-1 ) and ');' or ','
> > +          funcline += ["  const %s x%d%d =
> > + *(src%d+%d);"%(self.argtype(i,index),i+1,j+1,i+1,j)]
> > +
> > +      return 0
> > +
> > +    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.cpuFuncName, 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;
> > +  %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.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=''
> > +    if (self.returnVector(index) == 1 and self.argvector(0,index) != 1):
> > +      funccompare='''
> > +  for (index = 0; index < count_input/vector; index++)  {
> > +    cpu_compiler_math( cpu_data + index, '''
> > +    else:
> > +      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\","%(self.outputNumFormat(index),self.outputNumFormat(index),self.o
> utputNumFormat(index))#,self.outputNumFormat(index))
> > +    funcsprintfb += " gpu_data[index], cpu_data[index],
> > + diff);"#%(ulpUnit(self.ulp),ulpNum(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));'
> > +    funcdiff = "    diff = fabs((gpu_data[index]-cpu_data[index]));"
> > +    funcline += [ funcdiff ]
> > +    funcline += [ funcsprintfa + funcsprintfb ]
> > +
> > +    self.cpplines += funcline
> > +
> > +    self.cpplines += [ udebug(self.ulp,self.retType(index)) ]
> > +    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.returnVector(index)) == 1):
> > +      clLine += [ '  dst[i] = ret;' ]
> > +    else:
> > +      for i in range(0,int(self.returnVector(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]
> > +
> > +      #Parameters:
> > +      self.GenInputValues(i)
> > +
> > +      #cpu function generator:
> > +      self.cpplines += [self.cpufunc]
> > +
> > +      #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..8b9ae5f 100644
> > --- a/utests/utest_helper.cpp
> > +++ b/utests/utest_helper.cpp
> > @@ -641,3 +641,33 @@ int cl_check_image(const int *img, int w, int h,
> const char *bmp)
> >    return (float(discrepancy) / float(n) > max_error_ratio) ? 0 : 1; }
> >
> > +typedef struct
> > +{
> > +  unsigned int mantissa:23;
> > +  unsigned int exponent:8;
> > +  unsigned int sign:1;
> > +} FLOAT;
> > +
> > +typedef union
> > +{
> > +  float f;
> > +  unsigned int i;
> > +  FLOAT spliter;
> > +} SF;
> > +
> > +const float cl_FLT_ULP(float float_number) {
> > +  SF floatBin, ulpBin;
> > +  floatBin.f = float_number;
> > +
> > +  ulpBin.spliter.sign = floatBin.spliter.sign;
> > + ulpBin.spliter.exponent = floatBin.spliter.exponent;
> > + ulpBin.spliter.mantissa = 0x1;
> > +
> > +  return ulpBin.f;
> > +}
> > +
> > +const int cl_INT_ULP(int int_number)
> > +{
> > +  return 0;
> > +}
> > diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp index
> > 29a21d5..79e7417 100644
> > --- a/utests/utest_helper.hpp
> > +++ b/utests/utest_helper.hpp
> > @@ -216,5 +216,11 @@ extern void cl_write_bmp(const int *data, int
> > width, int height, const char *fil
> >  /* Check data from img against bmp file located at "bmp" */  extern
> > int cl_check_image(const int *img, int w, int h, const char *bmp);
> >
> > +/* Calculator ULP of each FLOAT value */ extern const float
> > +cl_FLT_ULP(float float_number);
> > +
> > +/* Calculator ULP of each INT value */ extern const int
> > +cl_INT_ULP(int int_number);
> > +
> >  #endif /* __UTEST_HELPER_HPP__ */
> >
> > diff --git a/utests/utest_math_gen.py b/utests/utest_math_gen.py new
> > file mode 100755 index 0000000..7a4b678
> > --- /dev/null
> > +++ b/utests/utest_math_gen.py
> > @@ -0,0 +1,531 @@
> > +#!/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]
> > +
> > +#func:
> > +#    gpufuncName
> > +#    cpuFuncName
> > +#    fileName: 'builtin_'+name
> > +#    inputtype: a 2-D list because there're more than one input data
> > +#    outputtype: a list
> > +#    values
> > +#    ulp
> > +
> > +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']
> > +  acosUtests =
> > +func('acos','acos',[acos_input_type],acos_output_type,[acos_input_val
> > +ues],'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']
> > +  acoshUtests =
> > + func('acosh','acosh',[acosh_input_type],acosh_output_type,[acosh_inp
> > + ut_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']
> > +  acospi_cpu_func='''
> > +static float acospi(float x){
> > +  return acos(x)/M_PI;
> > +} '''
> > +  acospiUtests =
> > +func('acospi','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']
> > +  asinUtests =
> > + func('asin','asin',[asin_input_type],asin_output_type,[asin_input_va
> > + lues],'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']
> > +  asinhUtests =
> > + func('asinh','asinh',[asinh_input_type],asinh_output_type,[asinh_inp
> > + ut_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']
> > +  asinpi_cpu_func='''
> > +static float asinpi(float x){
> > +  return asin(x)/M_PI;
> > +} '''
> > +  asinpiUtests =
> > +func('asinpi','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']
> > +  atanUtests =
> > + func('atan','atan',[atan_input_type],atan_output_type,[atan_input_va
> > + lues],'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']
> > +  atan2Utests =
> > + func('atan2','atan2',[atan2_input_type1,atan2_input_type2],atan2_out
> > + put_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']
> > +  atanhUtests =
> > + func('atanh','atanh',[atanh_input_type],atanh_output_type,[atanh_inp
> > + ut_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']
> > +  atanpi_cpu_func='''
> > +static float atanpi(float x){
> > +  return atan(x)/M_PI;
> > +} '''
> > +  atanpiUtests =
> > +func('atanpi','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_i
> > +nput_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']
> > +#  atan2pi_cpu_func='''
> > +#static float atan2pi(float y, float x){ #  return atan2(y,x)/M_PI;
> > +#} '''
> > +#  atan2piUtests =
> > +func('atan2pi','atan2pi',[atan2pi_input_type1,atan2pi_input_type2],at
> > +an2pi_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']
> > +  cbrtUtests =
> > + func('cbrt','cbrt',[cbrt_input_type],cbrt_output_type,[cbrt_input_va
> > + lues],'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']
> > +  ceilUtests =
> > + func('ceil','ceil',[ceil_input_type],ceil_output_type,[ceil_input_va
> > + lues],'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(copysi
> > + gn_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']
> > +  copysignUtests =
> > + func('copysign','copysign',[copysign_input_type1,copysign_input_type
> > + 2],copysign_output_type,[copysign_input_values1,copysign_input_value
> > + s2],'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']
> > +  cosUtests =
> > + func('cos','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']
> > +  coshUtests =
> > + func('cosh','cosh',[cosh_input_type],cosh_output_type,[cosh_input_va
> > + lues],'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']
> > +  cospi_cpu_func='''
> > +static float cospi(float x){
> > +  return cos(M_PI * x);
> > +} '''
> > +  cospiUtests =
> > +func('cospi','cospi',[cospi_input_type],cospi_output_type,[cospi_inpu
> > +t_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']
> > +  erfUtests =
> > + func('erf','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']
> > +  erfcUtests =
> > + func('erfc','erfc',[erfc_input_type],erfc_output_type,[erfc_input_va
> > + lues],'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']
> > +  expUtests =
> > + func('exp','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']
> > +  exp2Utests =
> > + func('exp2','exp2',[exp2_input_type],exp2_output_type,[exp2_input_va
> > + lues],'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']
> > +  exp10Utests =
> > + func('exp10','exp10',[exp10_input_type],exp10_output_type,[exp10_inp
> > + ut_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']
> > +  expm1Utests =
> > + func('expm1','expm1',[expm1_input_type],expm1_output_type,[expm1_inp
> > + ut_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']
> > +  fabsUtests =
> > + func('fabs','fabs',[fabs_input_type],fabs_output_type,[fabs_input_va
> > + lues],'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_val
> > + ues1,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']
> > +  fdimUtests =
> > + func('fdim','fdim',[fdim_input_type1,fdim_input_type2],fdim_output_t
> > + ype,[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']
> > +  floorUtests =
> > + func('floor','floor',[floor_input_type],floor_output_type,[floor_inp
> > + ut_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=gene3ValuesLoo
> > + p(fma_input_values1,fma_input_values2,fma_input_values3,fma_base_val
> > + ues)
> > +  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']
> > +  fmaUtests =
> > + func('fma','fma',[fma_input_type1,fma_input_type2,fma_input_type3],f
> > + ma_output_type,[fma_input_values1,fma_input_values2,fma_input_values
> > + 3],'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_val
> > + ues1,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']
> > +  fmaxUtests =
> > + func('fmax','fmax',[fmax_input_type1,fmax_input_type2],fmax_output_t
> > + ype,[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=gene2ValuesLo
> >
> +op(fmax_gentypef_input_values1,fmax_gentypef_input_values2,fmax_genty
> > +pef_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','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_val
> > + ues1,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']
> > +  fminUtests =
> > + func('fmin','fmin',[fmin_input_type1,fmin_input_type2],fmin_output_t
> > + ype,[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=gene2ValuesLo
> > +op(fmin_gentypef_input_values1,fmin_gentypef_input_values2,fmin_genty
> > +pef_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','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_valu
> > +es1,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']
> > +  fmodUtests =
> > +func('fmod','fmod',[fmod_input_type1,fmod_input_type2],fmod_output_ty
> > +pe,[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']
> > +  hypotUtests =
> > + func('hypot','hypot',[hypot_input_type1,hypot_input_type2],hypot_out
> > + put_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']
> > +  ilogbUtests =
> > + func('ilogb','ilogb',[ilogb_input_type],ilogb_output_type,[ilogb_inp
> > + ut_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']
> > +  lgammaUtests =
> > +
> func('lgamma','lgamma',[lgamma_input_type],lgamma_output_type,[lgamm
> > + a_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']
> > +  logUtests =
> > + func('log','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']
> > +  log2Utests =
> > + func('log2','log2',[log2_input_type],log2_output_type,[log2_input_va
> > + lues],'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']
> > +  log10Utests =
> > + func('log10','log10',[log10_input_type],log10_output_type,[log10_inp
> > + ut_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']
> > +  log1pUtests =
> > + func('log1p','log1p',[log1p_input_type],log1p_output_type,[log1p_inp
> > + ut_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']
> > +  logbUtests =
> > + func('logb','logb',[logb_input_type],logb_output_type,[logb_input_va
> > + lues],'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_
> inpu
> > +t_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']
> > +  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',[maxmag_input_type1,maxmag_input_type2],max
> mag
> > +_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_in
> pu
> > +t_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']
> > +  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',[minmag_input_type1,minmag_input_type2],minma
> g
> > +_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']
> > +#  nanUtests =
> > +func('nan','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(next
> > + after_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']
> > +  nextafterUtests =
> > + func('nextafter','nextafter',[nextafter_input_type1,nextafter_input_
> > + type2],nextafter_output_type,[nextafter_input_values1,nextafter_inpu
> > + t_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_values
> > + 1,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']
> > +  powUtests =
> > + func('pow','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']
> > +  pown_cpu_func='''
> > +static float pown(float x, int y){
> > +    return pow(x,y);
> > +} '''
> > +  pownUtests =
> > +func('pown','pown',[pown_input_type1,pown_input_type2],pown_output_ty
> > +pe,[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']
> > +  powr_cpu_func='''
> > +static float powr(float x, int y){
> > +    return pow(x,y);
> > +} '''
> > +  powrUtests =
> > +func('powr','powr',[powr_input_type1,powr_input_type2],powr_output_ty
> > +pe,[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(rema
> > + inder_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']
> > +  remainderUtests =
> > + func('remainder','remainder',[remainder_input_type1,remainder_input_
> > + type2],remainder_output_type,[remainder_input_values1,remainder_inpu
> > + t_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']
> > +  rintUtests =
> > + func('rint','rint',[rint_input_type],rint_output_type,[rint_input_va
> > + lues],'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']
> > +  rootn_cpu_func='''
> > +static float rootn(float x, int y){
> > +    return pow(x,1.0/y);
> > +} '''
> > +  rootnUtests =
> > +func('rootn','rootn',[rootn_input_type1,rootn_input_type2],rootn_outp
> > +ut_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']
> > +  roundUtests =
> > + func('round','round',[round_input_type],round_output_type,[round_inp
> > + ut_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']
> > +  rsqrt_cpu_func='''
> > +static float rsqrt(float x)
> > +{ return 1/sqrt(x);} '''
> > +  rsqrtUtests =
> > +func('rsqrt','rsqrt',[rsqrt_input_type],rsqrt_output_type,[rsqrt_inpu
> > +t_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']
> > +  sinUtests =
> > + func('sin','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',[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']
> > +  sinhUtests =
> > + func('sinh','sinh',[sinh_input_type],sinh_output_type,[sinh_input_va
> > + lues],'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']
> > +  sinpi_cpu_func='''
> > +static float sinpi(float x){
> > +  return sin(M_PI*x);
> > +} '''
> > +  sinpiUtests =
> > +func('sinpi','sinpi',[sinpi_input_type],sinpi_output_type,[sinpi_inpu
> > +t_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']
> > +  sqrtUtests =
> > + func('sqrt','sqrt',[sqrt_input_type],sqrt_output_type,[sqrt_input_va
> > + lues],'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']
> > +  tanUtests =
> > + func('tan','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']
> > +  tanhUtests =
> > + func('tanh','tanh',[tanh_input_type],tanh_output_type,[tanh_input_va
> > + lues],'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']
> > +  tanpi_cpu_func='''
> > +static float tanpi(float x){
> > +  return tan(M_PI*x);
> > +} '''
> > +  tanpiUtests =
> > +func('tanpi','tanpi',[tanpi_input_type],tanpi_output_type,[tanpi_inpu
> > +t_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']
> > +  tgammaUtests =
> > +
> func('tgamma','tgamma',[tgamma_input_type],tgamma_output_type,[tgamm
> > + a_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']
> > +  truncUtests =
> > + func('trunc','trunc',[trunc_input_type],trunc_output_type,[trunc_inp
> > + ut_values],'0 * FLT_ULP')
> > +
> > +if __name__ == "__main__":
> > +  main()
> > --
> > 1.7.6.4
> >
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list