aboutsummaryrefslogtreecommitdiff
path: root/handcoded-opencl/sql7.cl
diff options
context:
space:
mode:
authorTom Gall <tom.gall@linaro.org>2014-04-02 16:33:12 +0000
committerTom Gall <tom.gall@linaro.org>2014-04-02 16:33:12 +0000
commit53a45813e8c4e4918c160c14f0373e07166afd2b (patch)
treef27663d07798ed06c3baf34b479572bead1c7fe4 /handcoded-opencl/sql7.cl
Initial v.01 commit
Represents and early early but functional prototype. See README for various notes. See LICENSE for license info.
Diffstat (limited to 'handcoded-opencl/sql7.cl')
-rw-r--r--handcoded-opencl/sql7.cl38
1 files changed, 38 insertions, 0 deletions
diff --git a/handcoded-opencl/sql7.cl b/handcoded-opencl/sql7.cl
new file mode 100644
index 0000000..ef90654
--- /dev/null
+++ b/handcoded-opencl/sql7.cl
@@ -0,0 +1,38 @@
+
+#define workUnits 64
+#define workUnitsM1 63
+
+typedef struct tag_my_struct {
+ int id;
+ int w;
+ int y;
+ } Row;
+
+__kernel void x1_search_kernel(int totalRows,
+ __global Row *data,
+ __global Row *resultArray,
+ __global int *roffsetResult) {
+
+ int i = get_global_id(0);
+ size_t offset = i * (totalRows/workUnits);
+ size_t endRow = (totalRows/workUnits);
+ size_t roffset = offset;
+ int tmp;
+
+ if (i == workUnitsM1) {
+ endRow = (totalRows/workUnits) + (totalRows % workUnits);
+ }
+ do {
+ tmp = data[offset].w * data[offset].y;
+ if (tmp >= -5 && tmp <= 5) {
+ resultArray[roffset].id = data[offset].id;
+ resultArray[roffset].w = data[offset].w;
+ resultArray[roffset].y = data[offset].y;
+ roffset++;
+ }
+ offset++;
+ endRow--;
+ } while (endRow);
+
+ roffsetResult[i] = roffset- (i * (totalRows/workUnits));
+}