diff options
| author | Tom Stellard <thomas.stellard@amd.com> | 2013-07-31 19:10:58 -0700 |
|---|---|---|
| committer | Tom Stellard <thomas.stellard@amd.com> | 2013-08-09 20:40:19 -0700 |
| commit | 2a43c94060e3683682b3ebe41318aa479bd58d9b (patch) | |
| tree | 1db42570ecfda417ea6a26137165561932a68b93 /generated_tests/cl | |
| parent | 4869588d96758ace8a84599e36a12b5e0a72561d (diff) | |
cl: Add generated tests for global and local stores v2
v2:
- Fix coding style
Acked-by: Dylan Baker <baker.dylan.c@gmail.com>
Diffstat (limited to 'generated_tests/cl')
| -rw-r--r-- | generated_tests/cl/store/store-kernels-global.inc | 21 | ||||
| -rw-r--r-- | generated_tests/cl/store/store-kernels-local.inc | 14 |
2 files changed, 35 insertions, 0 deletions
diff --git a/generated_tests/cl/store/store-kernels-global.inc b/generated_tests/cl/store/store-kernels-global.inc new file mode 100644 index 000000000..b6220d0e3 --- /dev/null +++ b/generated_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/generated_tests/cl/store/store-kernels-local.inc b/generated_tests/cl/store/store-kernels-local.inc new file mode 100644 index 000000000..7d70d13bb --- /dev/null +++ b/generated_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] = store_index; + barrier(CLK_LOCAL_MEM_FENCE); + out[id] = local_data[id]; +} |