summaryrefslogtreecommitdiff
path: root/tests/cl
diff options
context:
space:
mode:
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];
+}