<div dir="ltr"><br><div class="gmail_extra"><br><br><div class="gmail_quote">On Wed, Jul 31, 2013 at 11:37 PM, Dylan Baker <span dir="ltr"><<a href="mailto:baker.dylan.c@gmail.com" target="_blank">baker.dylan.c@gmail.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr"><div>I have some python comments for you, they're mainly style type comments, but some of them should help your code readability a lot.<br>
</div><div class="gmail_extra"><br><br><div class="gmail_quote"><div><div class="h5">
On Wed, Jul 31, 2013 at 7:16 PM, Tom Stellard <span dir="ltr"><<a href="mailto:tom@stellard.net" target="_blank">tom@stellard.net</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
From: Tom Stellard <<a href="mailto:thomas.stellard@amd.com" target="_blank">thomas.stellard@amd.com</a>><br>
<br>
---<br>
generated_tests/CMakeLists.txt | 4 +<br>
generated_tests/cl/store/store-kernels-global.inc | 21 ++++++<br>
generated_tests/cl/store/store-kernels-local.inc | 14 ++++<br>
generated_tests/generate-cl-store-tests.py | 90 +++++++++++++++++++++++<br>
tests/all_cl.tests | 3 +<br>
5 files changed, 132 insertions(+)<br>
create mode 100644 generated_tests/cl/store/store-kernels-global.inc<br>
create mode 100644 generated_tests/cl/store/store-kernels-local.inc<br>
create mode 100644 generated_tests/generate-cl-store-tests.py<br>
<br>
diff --git a/generated_tests/CMakeLists.txt b/generated_tests/CMakeLists.txt<br>
index db3734f..6ac5a9c 100644<br>
--- a/generated_tests/CMakeLists.txt<br>
+++ b/generated_tests/CMakeLists.txt<br>
@@ -54,6 +54,9 @@ piglit_make_generated_tests(<br>
piglit_make_generated_tests(<br>
builtin_cl_int_tests.list<br>
generate-cl-int-builtins.py)<br>
+piglit_make_generated_tests(<br>
+ cl_store_tests.list<br>
+ generate-cl-store-tests.py)<br>
<br>
# Add a "gen-tests" target that can be used to generate all the<br>
# tests without doing any other compilation.<br>
@@ -62,6 +65,7 @@ add_custom_target(gen-tests ALL<br>
builtin_uniform_tests.list<br>
constant_array_size_tests.list<br>
builtin_cl_int_tests.list<br>
+ cl_store_tests.list<br>
interpolation_tests.list<br>
non-lvalue_tests.list<br>
texture_query_lod_tests.list<br>
diff --git a/generated_tests/cl/store/store-kernels-global.inc b/generated_tests/cl/store/store-kernels-global.inc<br>
new file mode 100644<br>
index 0000000..b6220d0<br>
--- /dev/null<br>
+++ b/generated_tests/cl/store/store-kernels-global.inc<br>
@@ -0,0 +1,21 @@<br>
+typedef TYPE type_t;<br>
+<br>
+#if TYPE == double<br>
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable<br>
+#endif<br>
+<br>
+kernel void store_global(global type_t *out, global type_t *in) {<br>
+ out[0] = in[0];<br>
+ out[1] = in[1];<br>
+ out[2] = in[2];<br>
+ out[3] = in[3];<br>
+ out[4] = in[4];<br>
+ out[5] = in[5];<br>
+ out[6] = in[6];<br>
+ out[7] = in[7];<br>
+}<br>
+<br>
+kernel void store_global_wi(global type_t *out, global type_t *in) {<br>
+ size_t id = get_global_id(0);<br>
+ out[id] = in[id];<br>
+}<br>
diff --git a/generated_tests/cl/store/store-kernels-local.inc b/generated_tests/cl/store/store-kernels-local.inc<br>
new file mode 100644<br>
index 0000000..7d70d13<br>
--- /dev/null<br>
+++ b/generated_tests/cl/store/store-kernels-local.inc<br>
@@ -0,0 +1,14 @@<br>
+typedef TYPE type_t;<br>
+<br>
+#if TYPE == double<br>
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable<br>
+#endif<br>
+<br>
+kernel void store_local(global type_t *out, global type_t *in) {<br>
+ local type_t local_data[8];<br>
+ size_t id = get_local_id(0);<br>
+ size_t store_index = (id + 1) % 8;<br>
+ local_data[store_index] = store_index;<br>
+ barrier(CLK_LOCAL_MEM_FENCE);<br>
+ out[id] = local_data[id];<br>
+}<br>
diff --git a/generated_tests/generate-cl-store-tests.py b/generated_tests/generate-cl-store-tests.py<br>
new file mode 100644<br>
index 0000000..488aa62<br>
--- /dev/null<br>
+++ b/generated_tests/generate-cl-store-tests.py<br>
@@ -0,0 +1,90 @@<br>
+#!/usr/bin/env python<br>
+#<br>
+# Copyright 2013 Advanced Micro Devices, Inc.<br>
+#<br>
+# Permission is hereby granted, free of charge, to any person obtaining a<br>
+# copy of this software and associated documentation files (the "Software"),<br>
+# to deal in the Software without restriction, including without limitation<br>
+# the rights to use, copy, modify, merge, publish, distribute, sublicense,<br>
+# and/or sell copies of the Software, and to permit persons to whom the<br>
+# Software is furnished to do so, subject to the following conditions:<br>
+#<br>
+# The above copyright notice and this permission notice (including the next<br>
+# paragraph) shall be included in all copies or substantial portions of the<br>
+# Software.<br>
+#<br>
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR<br>
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,<br>
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL<br>
+# THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER<br>
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,<br>
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE<br>
+# SOFTWARE.<br>
+#<br>
+# Authors: Tom Stellard <<a href="mailto:thomas.stellard@amd.com" target="_blank">thomas.stellard@amd.com</a>><br>
+#<br>
+#<br>
+<br>
+import os<br>
+<br>
+TYPES = ['char', 'uchar', 'short', 'ushort', 'int', 'uint', 'long', 'ulong', 'float', 'double']<br>
+VEC_SIZES = ['', '2', '4', '8', '16']<br>
+<br>
+dirName = os.path.join("cl", "store")<br>
+if not os.path.exists(dirName):<br>
+ os.makedirs(dirName)<br>
+<br>
+<br>
+def gen_array(size):<br>
+ items = size * 8<br>
+ array = list(xrange(items))<br>
+ return ' '.join(map(str,array))<br></blockquote></div></div><div>In python comma is always followed by a space.<br><br></div><div>Also, have you considered a simpler list comprehension? They are lazy, and thus use less memory, besides not creating two useless objects. ie:<br>
</div><div>def gen_array(size):<br></div><div> return ' '.join([str(i) for i in xrange(size * 8)])<br></div><div><br>In python all top level function and class definitions should have exactly two blank lines proceeding and following them: there should be two, not one blank line here<br>
</div><div class="im"><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
+<br>
+def print_config(f, type_name, addr_space):<br>
+ f.write( \<br>
+ '[config]\n' + \<br>
+ 'name: Store (' + type_name + ')\n' + \<br>
+ 'program_source_file: store-kernels-' + addr_space + '.inc\n' + \<br>
+ 'build_options: -D TYPE=' + type_name + '\n' + \<br>
+ 'dimensions: 1\n')<br></blockquote><div><br></div></div><div>and here <br></div><div class="im"><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
+<br>
+def begin_test(type_name, addr_space):<br>
+ fileName = os.path.join(dirName, 'store-' + type_name + '-' + addr_space + '.program_test')<br>
+ print(fileName)<br>
+ f = open(fileName, 'w')<br>
+ print_config(f, type_name, addr_space)<br>
+ return f <br></blockquote></div></div></div></div></blockquote><div><br></div><div>sorry, one more thing. The above function is double indented.<br> <br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="ltr"><div class="gmail_extra"><div class="gmail_quote"><div class="im"><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"></blockquote><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
+<br>
+<br>
+for t in TYPES:<br>
+ for s in VEC_SIZES:<br>
+ if s == '':<br>
+ size = 1<br></blockquote><div> </div></div><div>I only see the one use of VEC_SIZES. Is there a reason element 0 is set to '' and then you have an if statement here to change it to 1?<br><br></div><div class="im">
<blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
+ else:<br>
+ size = int(s)<br>
+ type_name = t + s<br>
+ f = begin_test(type_name, 'global')<br>
+ f.write( \<br>
+ '[test]\n' + \<br>
+ 'name: global address space\n' + \<br>
+ 'global_size: 1 0 0\n' + \<br>
+ 'kernel_name: store_global\n' + \<br>
+ 'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \<br>
+ 'arg_in: 1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \<br>
+ '[test]\n' + \<br>
+ 'name: global address space work items\n' + \<br>
+ 'global_size: 8 0 0\n' + \<br>
+ 'kernel_name: store_global_wi\n' + \<br>
+ 'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \<br>
+ 'arg_in: 1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n')<br></blockquote><div><br></div></div><div>In python do not use \ inside of parenthesis, python will concatenate those lines.<br>
</div><div>However, that said I would use a triple quote and format() here to make this code more readable. I'm not familiar with CL, so I may not get the CL syntax perfect, but something like:<br></div><div>f.write("""<br>
</div><div>[test]<br></div><div>name: global address space<br>global_size: 1 0 0<br></div><div>kernel_name: store_global<br></div><div>arg_out: 0 buffer {type_name}[8] {gen_array}<br></div><div>arg_in: 1 buffer {type_name}[8] {gen_array}<br>
<br></div><div>[test]<br></div><div class="im"><div>name: global address space work items<br></div></div><div>global_size: 8 0 0<br></div><div>kernel_name: store_global_wi<br></div><div>arg_out: 0 buffer {type_name}[8] {gen_array}<br>
</div>arg_in: 1 buffer {type_name}[8] {gen_array}<br>
</div><div class="gmail_quote">""".format(type_name=type_name, gen_array=gen_array(size)))<br><br></div><div class="gmail_quote">if you import textwrap you can use textwrap.dedent() to indent the text in the for/if block and still have it print at the root of your text document.<br>
</div><div class="gmail_quote"><br></div><div class="gmail_quote">and for the next one also.<br></div><div><div class="h5"><div class="gmail_quote"><br><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
+ f.close()<br>
+ f = begin_test(type_name, 'local')<br>
+ f.write( \<br>
+ '[test]\n' + \<br>
+ 'name: local address space\n' + \<br>
+ 'global_size: 8 0 0\n' + \<br>
+ 'local_size: 8 0 0\n' + \<br>
+ 'kernel_name: store_local\n' + \<br>
+ 'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \<br>
+ 'arg_in: 1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n')<br>
+ f.close() <br></blockquote><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
diff --git a/tests/all_cl.tests b/tests/all_cl.tests<br>
index b722511..efb04e4 100644<br>
--- a/tests/all_cl.tests<br>
+++ b/tests/all_cl.tests<br>
@@ -110,3 +110,6 @@ add_program_test_dir(program_execute, 'tests/cl/program/execute')<br>
program_execute_builtin = Group()<br>
program["Execute"]["Builtin"] = program_execute_builtin<br>
add_program_test_dir(program_execute_builtin, 'generated_tests/cl/builtin/int')<br>
+program_execute_store = Group()<br>
+program["Execute"]["Store"] = program_execute_store<br>
+add_program_test_dir(program_execute_store, 'generated_tests/cl/store')<br>
<span><font color="#888888">--<br>
1.7.11.4<br>
<br>
_______________________________________________<br>
Piglit mailing list<br>
<a href="mailto:Piglit@lists.freedesktop.org" target="_blank">Piglit@lists.freedesktop.org</a><br>
<a href="http://lists.freedesktop.org/mailman/listinfo/piglit" target="_blank">http://lists.freedesktop.org/mailman/listinfo/piglit</a><br>
</font></span></blockquote></div><br></div></div></div></div>
</blockquote></div><br></div></div>