[Piglit] [PATCH 1/2] cl: Add suport for builtins with multiple outputs
Jan Vesely
jan.vesely at rutgers.edu
Tue Feb 23 16:06:37 UTC 2016
On Mon, 2016-02-22 at 20:44 -0600, Aaron Watry wrote:
> This patch is
> Reviewed-by: Aaron Watry <awatry at gmail.com>
>
> The other one I'm not as qualified to review the selected test
> values, so
> it is:
> Acked-By: Aaron Watry <awatry at gmail.com>
pushed. thank you,
Jan
>
> On Thu, Feb 18, 2016 at 2:25 PM, Jan Vesely <jan.vesely at rutgers.edu>
> wrote:
>
> > 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>
> >
--
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/20160223/8c9ad4b3/attachment.sig>
More information about the Piglit
mailing list