[Piglit] [PATCH 1/2] cl: Add suport for builtins with multiple outputs
Jan Vesely
jan.vesely at rutgers.edu
Fri Feb 5 17:34:20 UTC 2016
On Thu, 2016-02-04 at 22:15 -0600, Aaron Watry wrote:
> Type-o in the subject (Suport).
>
> On Wed, Feb 3, 2016 at 6:37 PM, Jan Vesely <jan.vesely at rutgers.edu>
> wrote:
>
> > Reuses parts of Tom's fract test patch.
> > There is no change in generated tests with single output
> >
> > Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> > ---
> > generated_tests/genclbuiltins.py | 108
> > ++++++++++++++++++++++++++++-----------
> > 1 file changed, 79 insertions(+), 29 deletions(-)
> >
> > diff --git a/generated_tests/genclbuiltins.py
> > b/generated_tests/genclbuiltins.py
> > index 24bf249..7df55c9 100644
> > --- a/generated_tests/genclbuiltins.py
> > +++ b/generated_tests/genclbuiltins.py
> > @@ -109,13 +109,24 @@ B = {
> >
> > # vecSizes has the layout [in0width, ..., inNwidth] where outType
> > width is
> > # assumed to match the width of the first input
> > -def gen_kernel(f, fnName, inTypes, outType, vecSizes, typePrefix):
> > +def gen_kernel(f, fnName, inTypes, outTypes, vecSizes, typePrefix,
> > outLoc='private'):
> > f.write('kernel void test_' + typePrefix + str(vecSizes[0]) +
> > '_' +
> > fnName
> > - + '_' + inTypes[0]+'(global '+outType+'* out')
> > + + (('_' + outLoc) if len(outTypes) > 1 else '')
> > + + '_' + inTypes[0] + '(global ' + outTypes[0] + '*
> > out')
> > + for arg in range(1, len(outTypes)):
> > + f.write(', global '+outTypes[arg]+'* out'+str(arg))
> > for arg in range(0, len(inTypes)):
> > f.write(', global '+inTypes[arg]+'* in'+str(arg))
> > f.write('){\n')
> >
> > + for arg in range(1, len(outTypes)):
> > + f.write(' ' + outLoc + ' ' + outTypes[arg] + ('' if
> > vecSizes[0]
> > == 1 else str(vecSizes[0])));
> > + if (outLoc == 'global'):
> > + f.write(' *tmp' + str(arg) + ' = out' + str(arg) +
> > ';\n');
> > + else:
> > + #FIXME: This assumes WG size of 1
> >
>
> Was this FIXME left here intentionally for whoever needs that feature
> later? Cause yeah, the generated code for local memory tests with
> non-1 WG
> sizes is definitely broken.
>
> Do you plan on tackling that, or do you want to punt on it and just
> assume
> WG size =1 for now?
The FIXME is intentional. I don't expect anyone to ever need WG > 1 for
these tests and supporting it properly looked like to much effort for
little gain:
We'd need something like
local type tmpX[get_local_size(0)];
which is not allowed for obvious reasons.
we could to do:
local type tmpX[MAX_LOCAL_SIZE_X];
but we'd need to add that macro to cl-program-tester,
regards,
Jan
PS: I fixed the typos and whitespace locally
>
>
> > + f.write(' tmp' + str(arg) + ';\n');
> > +
> > suffix = ';'
> > if (vecSizes[0] == 1):
> > f.write(' out[get_global_id(0)] = ')
> > @@ -134,88 +145,113 @@ def gen_kernel(f, fnName, inTypes, outType,
> > vecSizes, typePrefix):
> > f.write('in'+str(arg)+'[get_global_id(0)]')
> > else:
> > f.write('vload'+str(vecSizes[arg])+'(get_global_id(0),
> > in'+str(arg)+')')
> >
>
> Might be nice to have a blank line here.
>
>
> > + for arg in range(1, len(outTypes)):
> > + if (outLoc == 'global'):
> > + f.write(', &tmp' + str(arg) + '[get_global_id(0)]')
> > + else:
> > + f.write(', &tmp' + str(arg))
> > +
> > + f.write(suffix + '\n')
> > +
> >
>
> And one less here.
>
> Let me know about the local WG size thing. Otherwise, this looks ok
> with
> the spelling/formatting changes. I'm by no means a python expert,
> but this
> looks functional, and I've already re-used it to add frexp tests.
>
> --Aaron
>
>
> > +
> > + if (outLoc != 'global'):
> > + for arg in range(1, len(outTypes)):
> > + if (vecSizes[0] == 1):
> > + f.write(' out'+ str(arg) +'[get_global_id(0)] =
> > tmp' +
> > str(arg))
> > + else:
> > + f.write(' vstore' + str(vecSizes[0]) + '(tmp' +
> > str(arg)
> > + ', get_global_id(0), out' + str(arg) + ')')
> > + f.write(';\n')
> >
> > - f.write(suffix+'\n}\n\n')
> > + f.write('}\n\n')
> >
> >
> >
> > -def gen_kernel_1_arg(f, fnName, inType, outType):
> > +def gen_kernel_1_arg(f, fnName, inType, outTypes, loc =
> > 'private'):
> > for vecSize in ALL_WIDTHS:
> > - gen_kernel(f, fnName, [inType], outType, [vecSize], '')
> > + gen_kernel(f, fnName, [inType], outTypes, [vecSize], '',
> > loc)
> >
> >
> > # 2 argument kernel with input types that match their vector size
> > -def gen_kernel_2_arg_same_size(f, fnName, inTypes, outType):
> > +def gen_kernel_2_arg_same_size(f, fnName, inTypes, outTypes):
> > for vecSize in ALL_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType, [vecSize,
> > vecSize],
> > + gen_kernel(f, fnName, inTypes, outTypes, [vecSize,
> > vecSize],
> > '')
> >
> >
> > # 2 argument kernel with 1 vector and one scalar input argument
> > -def gen_kernel_2_arg_mixed_size(f, fnName, inTypes, outType):
> > +def gen_kernel_2_arg_mixed_size(f, fnName, inTypes, outTypes):
> > for vecSize in VEC_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType, [vecSize, 1],
> > 'tss_')
> > + gen_kernel(f, fnName, inTypes, outTypes, [vecSize, 1],
> > 'tss_')
> >
> >
> > # 2 argument kernel with 1 vector and one scalar input argument
> > with
> > multiple
> > # input data types
> > -def gen_kernel_2_arg_mixed_sign(f, fnName, inTypes, outType):
> > +def gen_kernel_2_arg_mixed_sign(f, fnName, inTypes, outTypes):
> > for vecSize in ALL_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType, [vecSize,
> > vecSize],
> > + gen_kernel(f, fnName, inTypes, outTypes, [vecSize,
> > vecSize],
> > '')
> >
> >
> > # 3-argument built-in functions
> >
> >
> > -def gen_kernel_3_arg_same_type(f, fnName, inTypes, outType):
> > +def gen_kernel_3_arg_same_type(f, fnName, inTypes, outTypes):
> > for vecSize in ALL_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType,
> > + gen_kernel(f, fnName, inTypes, outTypes,
> > [vecSize, vecSize, vecSize], ''
> > )
> >
> > -def gen_kernel_3_arg_mixed_size_tss(f, fnName, inTypes, outType):
> > +def gen_kernel_3_arg_mixed_size_tss(f, fnName, inTypes, outTypes):
> > for vecSize in VEC_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType,
> > + gen_kernel(f, fnName, inTypes, outTypes,
> > [vecSize, 1, 1], 'tss_')
> >
> > -def gen_kernel_3_arg_mixed_size_tts(f, fnName, inTypes, outType):
> > +def gen_kernel_3_arg_mixed_size_tts(f, fnName, inTypes, outTypes):
> > for vecSize in VEC_WIDTHS:
> > - gen_kernel(f, fnName, inTypes, outType,
> > + gen_kernel(f, fnName, inTypes, outTypes,
> > [vecSize, vecSize, 1], 'tts_')
> >
> >
> > def generate_kernels(f, dataType, fnName, fnDef):
> > +
> > argTypes = getArgTypes(dataType, fnDef['arg_types'])
> >
> > # For len(argTypes), remember that this includes the output
> > arg
> > if (len(argTypes) == 2):
> > - gen_kernel_1_arg(f, fnName, argTypes[1], argTypes[0])
> > + gen_kernel_1_arg(f, fnName, argTypes[1], [argTypes[0]])
> > return
> >
> > if (len(argTypes) == 3 and not fnName is 'upsample'):
> > - gen_kernel_2_arg_same_size(f, fnName,
> > - [argTypes[1], argTypes[2]],
> > argTypes[0])
> > + if (getNumOutArgs(fnDef) == 2):
> > + gen_kernel_1_arg(f, fnName,
> > + argTypes[2], [argTypes[0],
> > argTypes[1]],
> > 'private')
> > + gen_kernel_1_arg(f, fnName,
> > + argTypes[2], [argTypes[0],
> > argTypes[1]],
> > 'local')
> > + gen_kernel_1_arg(f, fnName,
> > + argTypes[2], [argTypes[0],
> > argTypes[1]],
> > 'global')
> > + else:
> > + gen_kernel_2_arg_same_size(f, fnName,
> > + [argTypes[1], argTypes[2]],
> > [argTypes[0]])
> > if (fnDef['function_type'] is 'tss'):
> > gen_kernel_2_arg_mixed_size(f, fnName,
> > - [argTypes[1], argTypes[2]],
> > argTypes[0])
> > + [argTypes[1], argTypes[2]],
> > [argTypes[0]])
> > return
> >
> > if (len(argTypes) == 4):
> > gen_kernel_3_arg_same_type(f, fnName,
> > - [argTypes[1], argTypes[2], argTypes[3]],
> > argTypes[0])
> > + [argTypes[1], argTypes[2], argTypes[3]],
> > [argTypes[0]])
> > if (fnDef['function_type'] is 'tss'):
> > gen_kernel_3_arg_mixed_size_tss(f, fnName,
> > - [argTypes[1], argTypes[2], argTypes[3]],
> > argTypes[0])
> > + [argTypes[1], argTypes[2], argTypes[3]],
> > [argTypes[0]])
> > if (fnDef['function_type'] is 'tts'):
> > gen_kernel_3_arg_mixed_size_tts(f, fnName,
> > - [argTypes[1], argTypes[2], argTypes[3]],
> > argTypes[0])
> > + [argTypes[1], argTypes[2], argTypes[3]],
> > [argTypes[0]])
> > return
> >
> > if (fnName is 'upsample'):
> > gen_kernel_2_arg_mixed_sign(f, fnName,
> > [argTypes[1], argTypes[2]],
> > - argTypes[0])
> > + [argTypes[0]])
> > return
> >
> > def getValue(type, val, isVector):
> > @@ -301,6 +337,15 @@ def getArgTypes(baseType, argTypes):
> > def isFloatType(t):
> > return t not in U
> >
> > +def getNumOutArgs(functionDef):
> > + if 'num_out_args' in functionDef:
> > + return functionDef['num_out_args']
> > + else:
> > + return 1
> > +
> > +def isOutArg(functionDef, argIdx):
> > + return argIdx < getNumOutArgs(functionDef)
> > +
> > # Print a test with all-vector inputs/outputs and/or mixed
> > vector/scalar
> > args
> > def print_test(f, fnName, argType, functionDef, tests, numTests,
> > vecSize,
> > fntype):
> > # If the test allows mixed vector/scalar arguments, handle the
> > case
> > with
> > @@ -332,7 +377,7 @@ def print_test(f, fnName, argType, functionDef,
> > tests,
> > numTests, vecSize, fntype
> > for arg in range(0, argCount):
> > argInOut = ''
> > argVal = getStrVal(argType, tests[arg], (vecSize > 1))
> > - if arg == 0:
> > + if isOutArg(functionDef, arg):
> > argInOut = 'arg_out: '
> > else:
> > argInOut = 'arg_in: '
> > @@ -345,7 +390,7 @@ def print_test(f, fnName, argType, functionDef,
> > tests,
> > numTests, vecSize, fntype
> > '[' + str(numTests * vecSize) + '] ' +
> > ''.join(map(lambda x: (x + ' ') * vecSize,
> > argVal.split()))
> > )
> > - if arg == 0:
> > + if isOutArg(functionDef, arg) :
> > f.write(' tolerance {0} '.format(tolerance))
> > # Use ulp tolerance for float types
> > if isFloatType(argTypes[arg]):
> > @@ -418,8 +463,13 @@ def gen(types, minVersions, functions,
> > testDefs,
> > dirName):
> > sizes = sorted(VEC_WIDTHS)
> > sizes.insert(0, 1) # Add 1-wide scalar to the vector
> > widths
> > for vecSize in sizes:
> > - print_test(f, fnName, dataType, functionDef,
> > tests,
> > - numTests, vecSize, fnType)
> > + if (getNumOutArgs(functionDef) == 1):
> > + print_test(f, fnName, dataType, functionDef,
> > tests,
> > + numTests, vecSize, fnType)
> > + else:
> > + for loc in ['_private', '_local', '_global']:
> > + print_test(f, fnName + loc, dataType,
> > functionDef, tests,
> > + numTests, vecSize, fnType)
> >
> > # Terminate the header section
> > f.write('!*/\n\n')
> > --
> > 2.5.0
> >
> > _______________________________________________
> > Piglit mailing list
> > Piglit at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/piglit
> >
--
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <http://lists.freedesktop.org/archives/piglit/attachments/20160205/8e406815/attachment-0001.sig>
More information about the Piglit
mailing list