summaryrefslogtreecommitdiff
path: root/tests/cl
diff options
context:
space:
mode:
authorDylan Baker <baker.dylan.c@gmail.com>2015-02-20 13:38:53 -0800
committerDylan Baker <baker.dylan.c@gmail.com>2015-02-23 15:47:10 -0800
commitff8898df3c78324b2d8945c855a9d2e0034f5525 (patch)
tree2601a4c1560805f36b63d414aadee0b9048eff4e /tests/cl
parenta5712c7d52b4708b7f7b32c71ffc062425c0d7be (diff)
cl tests: move non-generated tests out of the generated tests dir
It really doesn't make sense to have static tests in the generated tests directory, this patch adds them to the tests directory and adds the appropriate code to cl.py for this change Signed-off-by: Dylan Baker <dylanx.c.baker@intel.com> Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Diffstat (limited to 'tests/cl')
-rw-r--r--tests/cl/store/store-kernels-global.inc21
-rw-r--r--tests/cl/store/store-kernels-local.inc14
2 files changed, 35 insertions, 0 deletions
diff --git a/tests/cl/store/store-kernels-global.inc b/tests/cl/store/store-kernels-global.inc
new file mode 100644
index 000000000..b6220d0e3
--- /dev/null
+++ b/tests/cl/store/store-kernels-global.inc
@@ -0,0 +1,21 @@
+typedef TYPE type_t;
+
+#if TYPE == double
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+
+kernel void store_global(global type_t *out, global type_t *in) {
+ out[0] = in[0];
+ out[1] = in[1];
+ out[2] = in[2];
+ out[3] = in[3];
+ out[4] = in[4];
+ out[5] = in[5];
+ out[6] = in[6];
+ out[7] = in[7];
+}
+
+kernel void store_global_wi(global type_t *out, global type_t *in) {
+ size_t id = get_global_id(0);
+ out[id] = in[id];
+}
diff --git a/tests/cl/store/store-kernels-local.inc b/tests/cl/store/store-kernels-local.inc
new file mode 100644
index 000000000..b3265f402
--- /dev/null
+++ b/tests/cl/store/store-kernels-local.inc
@@ -0,0 +1,14 @@
+typedef TYPE type_t;
+
+#if TYPE == double
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+
+kernel void store_local(global type_t *out, global type_t *in) {
+ local type_t local_data[8];
+ size_t id = get_local_id(0);
+ size_t store_index = (id + 1) % 8;
+ local_data[store_index] = in[store_index];
+ barrier(CLK_LOCAL_MEM_FENCE);
+ out[id] = local_data[id];
+}