<div dir="ltr"><div>This patch is Reviewed-by: Aaron Watry <<a href="mailto:awatry@gmail.com">awatry@gmail.com</a>> </div>The other one I'm not as qualified to review the selected test values, so it is: Acked-By: Aaron Watry <<a href="mailto:awatry@gmail.com">awatry@gmail.com</a>> </div><div class="gmail_extra"> <div class="gmail_quote">On Thu, Feb 18, 2016 at 2:25 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">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 <div class="HOEnZb"><div class="h5"> 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 <<a href="mailto:jan.vesely@rutgers.edu">jan.vesely@rutgers.edu</a>> > > wrote: > > > > > 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 > > > > > > > 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 > > > <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> > > > -- Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu">jan.vesely@rutgers.edu</a>> </div></div></blockquote></div> </div>