Fix for issue in the opencl backend where array offsets for the value…

arrayfire · christophe-murphy · 3c5a4f478312cdf32f991a8edce50626ede5af3c
commit 3c5a4f478312cdf32f991a8edce50626ede5af3c
diff --git a/src/backend/opencl/kernel/csr2dense.cl b/src/backend/opencl/kernel/csr2dense.cl
@@ -9,13 +9,16 @@
 
 kernel void csr2Dense(global T *output, global const T *values,
                       global const int *rowidx, global const int *colidx,
-                      const int M) {
+                      const int M, const int v_off, const int r_off, const int c_off) {
+    T *v = values + v_off;
+    int *r = rowidx + r_off;
+    int *c = colidx + c_off;
     int lid = get_local_id(0);
     for (int rowId = get_group_id(0); rowId < M; rowId += get_num_groups(0)) {
-        int colStart = rowidx[rowId];
-        int colEnd   = rowidx[rowId + 1];
+        int colStart = r[rowId];
+        int colEnd   = r[rowId + 1];
         for (int colId = colStart + lid; colId < colEnd; colId += THREADS) {
-            output[rowId + colidx[colId] * M] = values[colId];
+            output[rowId + c[colId] * M] = v[colId];
         }
     }
 }
diff --git a/src/backend/opencl/kernel/sparse.hpp b/src/backend/opencl/kernel/sparse.hpp
@@ -84,7 +84,10 @@ void csr2dense(Param output, const Param values, const Param rowIdx,
     cl::NDRange global(local[0] * groups_x, 1);
 
     csr2dense(cl::EnqueueArgs(getQueue(), global, local), *output.data,
-              *values.data, *rowIdx.data, *colIdx.data, M);
+              *values.data, *rowIdx.data, *colIdx.data, M,
+	      static_cast<int>(values.info.offset),
+	      static_cast<int>(rowIdx.info.offset),
+	      static_cast<int>(colIdx.info.offset));
     CL_DEBUG_FINISH(getQueue());
 }