<div dir="ltr">Type-o in the subject (Suport). <div><div class="gmail_extra"> <div class="gmail_quote">On Wed, Feb 3, 2016 at 6:37 PM, Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu" target="_blank">jan.vesely@rutgers.edu</a>> wrote: <blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Reuses parts of Tom's fract test patch. There is no change in generated tests with single output Signed-off-by: Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu">jan.vesely@rutgers.edu</a>> --- 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 </blockquote><div> </div><div>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. </div><div>Do you plan on tackling that, or do you want to punt on it and just assume WG size =1 for now? </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> + 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)+')') </blockquote><div> </div><div>Might be nice to have a blank line here. </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> + 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') + </blockquote><div> </div><div>And one less here. </div><div>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. </div><div>--Aaron </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> + + 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 <a href="mailto:Piglit@lists.freedesktop.org">Piglit@lists.freedesktop.org</a> <a href="http://lists.freedesktop.org/mailman/listinfo/piglit" rel="noreferrer" target="_blank">http://lists.freedesktop.org/mailman/listinfo/piglit</a> </blockquote></div> </div></div></div>