[Piglit] [PATCH 1/2] cl: Add suport for builtins with multiple outputs
Aaron Watry
awatry at gmail.com
Fri Feb 5 05:15:40 CET 2016
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?
> + 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
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/piglit/attachments/20160204/8e1b53fd/attachment-0001.html>
More information about the Piglit
mailing list