[Piglit] [PATCH 1/2] cl: Add suport for builtins with multiple outputs
Jan Vesely
jan.vesely at rutgers.edu
Thu Feb 18 20:25:51 UTC 2016
Hi,
I plan to push both patches by the end of the week unless anyone
objects. still a RB would be nice if anyone can spare some time.
thank you,
Jan
On Fri, 2016-02-05 at 12:34 -0500, Jan Vesely wrote:
> 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: <https://lists.freedesktop.org/archives/piglit/attachments/20160218/3dafdc05/attachment-0001.sig>
More information about the Piglit
mailing list