diff options
Diffstat (limited to 'tests/cl')
-rw-r--r-- | tests/cl/store/store-kernels-global.inc | 21 | ||||
-rw-r--r-- | tests/cl/store/store-kernels-local.inc | 14 |
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]; +} |