Repository: silx-kit/silx
Issue: [OpenCL] Avoid the usage of pyopencl.array.zeros

This command triggers the usage of enqueue_fill_buffer in opencl which is not (yet) properly implemented in the cuda backend of PortableCL:
https://github.com/pocl/pocl/issues/725

Changed files:
silx/opencl/reconstruction.py, silx/opencl/common.py, silx/opencl/sift/test/test_keypoints.py, silx/opencl/convolution.py, silx/opencl/sift/test/test_image.py, silx/opencl/sift/test/test_algebra.py, silx/opencl/sift/test/test_reductions.py, silx/opencl/test/test_addition.py, silx/opencl/test/test_convolution.py, silx/math/fft/clfft.py, silx/opencl/sift/test/test_matching.py, silx/opencl/test/test_kahan.py, silx/opencl/test/test_sparse.py, silx/opencl/projection.py, silx/opencl/test/test_linalg.py, silx/opencl/linalg.py

diff --git a/silx/math/fft/clfft.py b/silx/math/fft/clfft.py
index af75aa6ad..dad8ec15a 100644
--- a/silx/math/fft/clfft.py
+++ b/silx/math/fft/clfft.py
@@ -113,7 +113,9 @@ class CLFFT(BaseFFT):
         self.axes = self.axes[::-1]
 
     def _allocate(self, shape, dtype):
-        return parray.zeros(self.queue, shape, dtype=dtype)
+        ary = parray.empty(self.queue, shape, dtype=dtype)
+        ary.fill(0)
+        return ary
 
 
     def check_array(self, array, shape, dtype, copy=True):
@@ -232,8 +234,8 @@ class CLFFT(BaseFFT):
             Whether to perform operation in asynchronous mode.
             Default is False, meaning that we wait for transform to complete.
         """
-        data_in = self.set_input_data(array, copy=False)
-        data_out = self.set_output_data(output, copy=False)
+        self.set_input_data(array, copy=False)
+        self.set_output_data(output, copy=False)
         self.update_forward_plan_arrays()
         event, = self.plan_forward.enqueue()
         if not(do_async):
@@ -259,8 +261,8 @@ class CLFFT(BaseFFT):
             Whether to perform operation in asynchronous mode.
             Default is False, meaning that we wait for transform to complete.
         """
-        data_in = self.set_output_data(array, copy=False)
-        data_out = self.set_input_data(output, copy=False)
+        self.set_output_data(array, copy=False)
+        self.set_input_data(output, copy=False)
         self.update_inverse_plan_arrays()
         event, = self.plan_inverse.enqueue(forward=False)
         if not(do_async):
diff --git a/silx/opencl/common.py b/silx/opencl/common.py
index d9f1a716f..b2f77e2c3 100644
--- a/silx/opencl/common.py
+++ b/silx/opencl/common.py
@@ -34,7 +34,7 @@ __author__ = "Jerome Kieffer"
 __contact__ = "Jerome.Kieffer@ESRF.eu"
 __license__ = "MIT"
 __copyright__ = "2012-2017 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "07/06/2019"
+__date__ = "01/08/2019"
 __status__ = "stable"
 __all__ = ["ocl", "pyopencl", "mf", "release_cl_buffers", "allocate_cl_buffers",
            "measure_workgroup_size", "kernel_workgroup_size"]
@@ -260,7 +260,8 @@ def _measure_workgroup_size(device_or_context, fast=False):
     max_valid_wg = 1
     data = numpy.random.random(shape).astype(numpy.float32)
     d_data = pyopencl.array.to_device(queue, data)
-    d_data_1 = pyopencl.array.zeros_like(d_data) + 1
+    d_data_1 = pyopencl.array.empty_like(d_data)
+    d_data_1.fill(numpy.float32(1.0))
 
     program = pyopencl.Program(ctx, get_opencl_code("addition")).build()
     if fast:
diff --git a/silx/opencl/convolution.py b/silx/opencl/convolution.py
index 18232f4b1..9ee07aab7 100644
--- a/silx/opencl/convolution.py
+++ b/silx/opencl/convolution.py
@@ -29,7 +29,7 @@ from __future__ import absolute_import, print_function, with_statement, division
 
 __authors__ = ["P. Paleo"]
 __license__ = "MIT"
-__date__ = "11/02/2019"
+__date__ = "01/08/2019"
 
 import numpy as np
 from math import ceil
@@ -262,7 +262,8 @@ class Convolution(OpenclProcessing):
         # Allocate arrays
         for option_name, array_name in option_array_names.items():
             if self.extra_options[option_name]:
-                value = parray.zeros(self.queue, self.shape, "f")
+                value = parray.empty(self.queue, self.shape, np.float32)
+                value.fill(np.float32(0.0))
             else:
                 value = None
             setattr(self, array_name, value)
diff --git a/silx/opencl/linalg.py b/silx/opencl/linalg.py
index fba3e9c62..a64122aa1 100644
--- a/silx/opencl/linalg.py
+++ b/silx/opencl/linalg.py
@@ -29,7 +29,7 @@ from __future__ import absolute_import, print_function, with_statement, division
 
 __authors__ = ["P. Paleo"]
 __license__ = "MIT"
-__date__ = "10/08/2017"
+__date__ = "01/08/2019"
 
 import numpy as np
 
@@ -63,8 +63,10 @@ class LinAlg(OpenclProcessing):
                                   platformid=platformid, deviceid=deviceid,
                                   profile=profile)
 
-        self.d_gradient = parray.zeros(self.queue, shape, np.complex64)
-        self.d_image = parray.zeros(self.queue, shape, np.float32)
+        self.d_gradient = parray.empty(self.queue, shape, np.complex64)
+        self.d_gradient.fill(np.complex64(0.0))
+        self.d_image = parray.empty(self.queue, shape, np.float32)
+        self.d_image.fill(np.float32(0.0))
         self.add_to_cl_mem({
             "d_gradient": self.d_gradient,
             "d_image": self.d_image
diff --git a/silx/opencl/projection.py b/silx/opencl/projection.py
index 0505d80b6..da8752f77 100644
--- a/silx/opencl/projection.py
+++ b/silx/opencl/projection.py
@@ -29,7 +29,7 @@ from __future__ import absolute_import, print_function, with_statement, division
 
 __authors__ = ["A. Mirone, P. Paleo"]
 __license__ = "MIT"
-__date__ = "28/02/2018"
+__date__ = "01/08/2019"
 
 import logging
 import numpy as np
@@ -141,10 +141,11 @@ class Projection(OpenclProcessing):
             BufferDescription("d_strideJoseph", self._dimrecy * 2, np.int32, mf.READ_ONLY),
             BufferDescription("d_strideLine", self._dimrecy * 2, np.int32, mf.READ_ONLY),
         ]
+        d_axis_corrections = parray.empty(self.queue, self.nprojs, np.float32)
+        d_axis_corrections.fill(np.float32(0.0))
         self.add_to_cl_mem(
             {
-                "d_axis_corrections": parray.zeros(self.queue,
-                                                   self.nprojs, np.float32)
+                "d_axis_corrections": d_axis_corrections
             }
         )
         self._tmp_extended_img = np.zeros((self.shape[0] + 2, self.shape[1] + 2),
@@ -193,7 +194,9 @@ class Projection(OpenclProcessing):
         pyopencl.enqueue_copy(self.queue, self.cl_mem["d_angles"], angles2)
 
     def allocate_slice(self):
-            self.add_to_cl_mem({"d_slice": parray.zeros(self.queue, (self.shape[1] + 2, self.shape[1] + 2), np.float32)})
+        ary = parray.zeros(self.queue, (self.shape[1] + 2, self.shape[1] + 2), np.float32)
+        ary.fill(0)
+        self.add_to_cl_mem({"d_slice": ary})
 
     def allocate_textures(self):
         self.d_image_tex = pyopencl.Image(
diff --git a/silx/opencl/reconstruction.py b/silx/opencl/reconstruction.py
index cb68680ae..2c84aee81 100644
--- a/silx/opencl/reconstruction.py
+++ b/silx/opencl/reconstruction.py
@@ -29,7 +29,7 @@ from __future__ import absolute_import, print_function, with_statement, division
 
 __authors__ = ["P. Paleo"]
 __license__ = "MIT"
-__date__ = "19/09/2017"
+__date__ = "01/08/2019"
 
 import logging
 import numpy as np
@@ -98,19 +98,23 @@ class ReconstructionAlgorithm(OpenclProcessing):
         self.sino_shape = sino_shape
         self.is_cpu = self.backprojector.is_cpu
         # Arrays
-        self.d_data = parray.zeros(self.queue, sino_shape, dtype=np.float32)
-        self.d_sino = parray.zeros_like(self.d_data)
-        self.d_x = parray.zeros(self.queue,
+        self.d_data = parray.empty(self.queue, sino_shape, dtype=np.float32)
+        self.d_data.fill(0.0)
+        self.d_sino = parray.empty_like(self.d_data)
+        self.d_sino.fill(0.0)
+        self.d_x = parray.empty(self.queue,
                                 self.backprojector.slice_shape,
                                 dtype=np.float32)
-        self.d_x_old = parray.zeros_like(self.d_x)
+        self.d_x.fill(0.0)
+        self.d_x_old = parray.empty_like(self.d_x)
+        self.d_x_old.fill(0.0)
 
         self.add_to_cl_mem({
-            "d_data": self.d_data,
-            "d_sino": self.d_sino,
-            "d_x": self.d_x,
-            "d_x_old": self.d_x_old,
-        })
+                            "d_data": self.d_data,
+                            "d_sino": self.d_sino,
+                            "d_x": self.d_x,
+                            "d_x_old": self.d_x_old,
+                            })
 
     def proj(self, d_slice, d_sino):
         """
@@ -276,10 +280,13 @@ class TV(ReconstructionAlgorithm):
         )
         # Additional arrays
         self.linalg.gradient(self.d_x)
-        self.d_p = parray.zeros_like(self.linalg.cl_mem["d_gradient"])
-        self.d_q = parray.zeros_like(self.d_data)
+        self.d_p = parray.empty_like(self.linalg.cl_mem["d_gradient"])
+        self.d_q = parray.empty_like(self.d_data)
         self.d_g = self.linalg.d_image
-        self.d_tmp = parray.zeros_like(self.d_x)
+        self.d_tmp = parray.empty_like(self.d_x)
+        self.d_p.fill(0)
+        self.d_q.fill(0)
+        self.d_tmp.fill(0)
         self.add_to_cl_mem({
             "d_p": self.d_p,
             "d_q": self.d_q,
diff --git a/silx/opencl/sift/test/test_algebra.py b/silx/opencl/sift/test/test_algebra.py
index 614208465..b5ed1fde1 100644
--- a/silx/opencl/sift/test/test_algebra.py
+++ b/silx/opencl/sift/test/test_algebra.py
@@ -36,7 +36,7 @@ __authors__ = ["Jérôme Kieffer"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/07/2018"
+__date__ = "01/08/2019"
 
 import os
 import time
@@ -158,7 +158,8 @@ class TestAlgebra(unittest.TestCase):
         gpu_keypoints = pyopencl.array.to_device(self.queue, keypoints)
         output = pyopencl.array.empty(self.queue, (nbkeypoints, 4), dtype=numpy.float32, order="C")
         output.fill(-1.0, self.queue)
-        counter = pyopencl.array.zeros(self.queue, (1,), dtype=numpy.int32, order="C")
+        counter = pyopencl.array.empty(self.queue, (1,), dtype=numpy.int32)
+        counter.fill(0)
         wg = self.wg_compact,
         shape = calc_size((keypoints.shape[0],), wg)
         nbkeypoints = numpy.int32(nbkeypoints)
diff --git a/silx/opencl/sift/test/test_image.py b/silx/opencl/sift/test/test_image.py
index 2e1d51d78..4cefc2b38 100755
--- a/silx/opencl/sift/test/test_image.py
+++ b/silx/opencl/sift/test/test_image.py
@@ -37,7 +37,7 @@ __authors__ = ["Jérôme Kieffer", "Pierre Paleo"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/07/2018"
+__date__ = "01/08/2019"
 
 import time
 import logging
@@ -141,8 +141,9 @@ class TestImage(unittest.TestCase):
         self.s = numpy.int32(s)  # 1, 2, 3 ... not 4 nor 0.
         self.gpu_dogs = pyopencl.array.to_device(self.queue, DOGS)
         self.output = pyopencl.array.empty(self.queue, (nb_keypoints, 4), dtype=numpy.float32, order="C")
-        self.output.fill(-1.0, self.queue)  # memset for invalid keypoints
-        self.counter = pyopencl.array.zeros(self.queue, (1,), dtype=numpy.int32, order="C")
+        self.output.fill(-1.0)  # memset for invalid keypoints
+        self.counter = pyopencl.array.empty(self.queue, (1,), dtype=numpy.int32)
+        self.counter.fill(0)
         nb_keypoints = numpy.int32(nb_keypoints)
         self.shape = calc_size((DOGS.shape[1], DOGS.shape[0] * DOGS.shape[2]), self.wg)  # it's a 3D vector !!
 
diff --git a/silx/opencl/sift/test/test_keypoints.py b/silx/opencl/sift/test/test_keypoints.py
index 967ab79cc..d5d6f9f9b 100755
--- a/silx/opencl/sift/test/test_keypoints.py
+++ b/silx/opencl/sift/test/test_keypoints.py
@@ -40,7 +40,7 @@ __authors__ = ["Jérôme Kieffer", "Pierre Paleo"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/07/2018"
+__date__ = "01/08/2019"
 
 import os
 import unittest
@@ -265,23 +265,23 @@ class TestKeypoints(unittest.TestCase):
             logger.warning("test_descriptor: Skipping test of WG=%s when maximum for this kernel is %s ", wg, max_wg)
             return
         gpu_keypoints = pyopencl.array.to_device(self.queue, keypoints)
-        gpu_descriptors = pyopencl.array.zeros(self.queue, (keypoints_end - keypoints_start, 128), dtype=numpy.uint8, order="C")
+        gpu_descriptors = pyopencl.array.empty(self.queue, (keypoints_end - keypoints_start, 128), dtype=numpy.uint8, order="C")
+        gpu_descriptors.fill(0)
         gpu_grad = pyopencl.array.to_device(self.queue, grad)
         gpu_ori = pyopencl.array.to_device(self.queue, ori)
         keypoints_start, keypoints_end = numpy.int32(keypoints_start), numpy.int32(keypoints_end)
         grad_height, grad_width = numpy.int32(grad.shape)
         counter = pyopencl.array.to_device(self.queue, keypoints_end)
-        kargs = [
-            gpu_keypoints.data,
-            gpu_descriptors.data,
-            gpu_grad.data,
-            gpu_ori.data,
-            numpy.int32(octsize),
-            keypoints_start,
-            counter.data,
-            grad_width,
-            grad_height
-        ]
+        kargs = [gpu_keypoints.data,
+                 gpu_descriptors.data,
+                 gpu_grad.data,
+                 gpu_ori.data,
+                 numpy.int32(octsize),
+                 keypoints_start,
+                 counter.data,
+                 grad_width,
+                 grad_height
+                 ]
 
         # Call the kernel
         t0 = time.time()
diff --git a/silx/opencl/sift/test/test_matching.py b/silx/opencl/sift/test/test_matching.py
index d11c9d3ae..427e56c85 100755
--- a/silx/opencl/sift/test/test_matching.py
+++ b/silx/opencl/sift/test/test_matching.py
@@ -36,7 +36,7 @@ __authors__ = ["Jérôme Kieffer", "Pierre Paleo"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013-2017 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "25/06/2018"
+__date__ = "01/08/2019"
 
 import os
 import unittest
@@ -139,10 +139,12 @@ class TestMatching(unittest.TestCase):
 
         gpu_keypoints1 = pyopencl.array.to_device(self.queue, ref_sift)
         gpu_keypoints2 = pyopencl.array.to_device(self.queue, ref_sift_2)
-        gpu_matchings = pyopencl.array.zeros(self.queue, (keypoints_end - keypoints_start, 2), dtype=numpy.int32, order="C")
+        gpu_matchings = pyopencl.array.empty(self.queue, (keypoints_end - keypoints_start, 2), dtype=numpy.int32, order="C")
+        gpu_matchings.fill(0)
         keypoints_start, keypoints_end = numpy.int32(keypoints_start), numpy.int32(keypoints_end)
         nb_keypoints = numpy.int32(10000)
-        counter = pyopencl.array.zeros(self.queue, (1, 1), dtype=numpy.int32, order="C")
+        counter = pyopencl.array.empty(self.queue, (1, 1), dtype=numpy.int32, order="C")
+        counter.fill(0)
 
         t0 = time.time()
         k1 = self.program.matching(self.queue, shape, wg,
diff --git a/silx/opencl/sift/test/test_reductions.py b/silx/opencl/sift/test/test_reductions.py
index 730be5a6e..de192058d 100755
--- a/silx/opencl/sift/test/test_reductions.py
+++ b/silx/opencl/sift/test/test_reductions.py
@@ -37,7 +37,7 @@ __authors__ = ["Jérôme Kieffer", "Pierre Paleo"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/07/2018"
+__date__ = "01/08/2019"
 
 import os
 import unittest
@@ -114,7 +114,8 @@ class TestReduction(unittest.TestCase):
             return
 
         size = wg * wg
-        max_min_gpu = pyopencl.array.zeros(self.queue, (wg, 2), dtype=numpy.float32, order="C")
+        max_min_gpu = pyopencl.array.empty(self.queue, (wg, 2), dtype=numpy.float32, order="C")
+        max_min_gpu.fill(0)
 #        max_min_gpu = pyopencl.array.empty(self.queue, (wg, 2), dtype=numpy.float32, order="C")
         max_gpu = pyopencl.array.empty(self.queue, (1,), dtype=numpy.float32, order="C")
         min_gpu = pyopencl.array.empty(self.queue, (1,), dtype=numpy.float32, order="C")
diff --git a/silx/opencl/test/test_addition.py b/silx/opencl/test/test_addition.py
index 3283a4deb..49cc0b469 100644
--- a/silx/opencl/test/test_addition.py
+++ b/silx/opencl/test/test_addition.py
@@ -35,7 +35,7 @@ __authors__ = ["Henri Payno, Jérôme Kieffer"]
 __contact__ = "jerome.kieffer@esrf.eu"
 __license__ = "MIT"
 __copyright__ = "2013 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "05/07/2018"
+__date__ = "01/08/2019"
 
 import logging
 import numpy
@@ -80,7 +80,8 @@ class TestAddition(unittest.TestCase):
         self.shape = 4096
         self.data = numpy.random.random(self.shape).astype(numpy.float32)
         self.d_array_img = pyopencl.array.to_device(self.queue, self.data)
-        self.d_array_5 = pyopencl.array.zeros_like(self.d_array_img) - 5
+        self.d_array_5 = pyopencl.array.empty_like(self.d_array_img)
+        self.d_array_5.fill(-5)
         self.program = pyopencl.Program(self.ctx, get_opencl_code("addition")).build()
 
     def tearDown(self):
diff --git a/silx/opencl/test/test_convolution.py b/silx/opencl/test/test_convolution.py
index 8acd385b3..9be1df679 100644
--- a/silx/opencl/test/test_convolution.py
+++ b/silx/opencl/test/test_convolution.py
@@ -34,7 +34,7 @@ __authors__ = ["Pierre Paleo"]
 __contact__ = "pierre.paleo@esrf.fr"
 __license__ = "MIT"
 __copyright__ = "2019 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "15/02/2019"
+__date__ = "01/08/2019"
 
 import logging
 from itertools import product
@@ -178,7 +178,8 @@ class TestConvolution(unittest.TestCase):
         else:
             data_ref = data
         if self.param["output_on_device"]:
-            d_res = parray.zeros_like(conv.data_out)
+            d_res = parray.empty_like(conv.data_out)
+            d_res.fill(0)
             res = d_res
         else:
             res = None
diff --git a/silx/opencl/test/test_kahan.py b/silx/opencl/test/test_kahan.py
index bb9ea3fc9..167640c9f 100644
--- a/silx/opencl/test/test_kahan.py
+++ b/silx/opencl/test/test_kahan.py
@@ -34,7 +34,7 @@ __author__ = "Jérôme Kieffer"
 __contact__ = "Jerome.Kieffer@ESRF.eu"
 __license__ = "MIT"
 __copyright__ = "European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "30/01/2019"
+__date__ = "01/08/2019"
 
 
 import unittest
@@ -120,7 +120,8 @@ class TestKahan(unittest.TestCase):
         """
         prg = pyopencl.Program(self.ctx, read_cl_file("kahan.cl") + src).build(self.args)
         ones_d = pyopencl.array.to_device(self.queue, data)
-        res_d = pyopencl.array.zeros(self.queue, 2, numpy.float32)
+        res_d = pyopencl.array.empty(self.queue, 2, numpy.float32)
+        res_d.fill(0)
         evt = prg.summation(self.queue, (1,), (1,), ones_d.data, numpy.int32(N), res_d.data)
         evt.wait()
         res = res_d.get().sum(dtype=numpy.float64)
@@ -202,7 +203,8 @@ class TestKahan(unittest.TestCase):
 
         prg = pyopencl.Program(self.ctx, read_cl_file("kahan.cl") + src).build(self.args)
         ones_d = pyopencl.array.to_device(self.queue, data)
-        res_d = pyopencl.array.zeros(self.queue, 2, numpy.float32)
+        res_d = pyopencl.array.empty(self.queue, 2, numpy.float32)
+        res_d.fill(0)
         evt = prg.test_dot16(self.queue, (1,), (1,), ones_d.data, numpy.int32(N), res_d.data)
         evt.wait()
         res = res_d.get().sum(dtype="float64")
diff --git a/silx/opencl/test/test_linalg.py b/silx/opencl/test/test_linalg.py
index 7d0398314..b72fa20c6 100644
--- a/silx/opencl/test/test_linalg.py
+++ b/silx/opencl/test/test_linalg.py
@@ -30,7 +30,7 @@ from __future__ import division, print_function
 __authors__ = ["Pierre paleo"]
 __license__ = "MIT"
 __copyright__ = "2013-2017 European Synchrotron Radiation Facility, Grenoble, France"
-__date__ = "14/06/2017"
+__date__ = "01/08/2019"
 
 
 import time
@@ -109,7 +109,8 @@ class TestLinAlg(unittest.TestCase):
         self.div_ref = divergence(self.grad_ref)
         self.image2 = np.zeros_like(self.image)
         # Device images
-        self.gradient_parray = parray.zeros(self.la.queue, self.image.shape, np.complex64)
+        self.gradient_parray = parray.empty(self.la.queue, self.image.shape, np.complex64)
+        self.gradient_parray.fill(0)
         # we should be using cl.Buffer(self.la.ctx, cl.mem_flags.READ_WRITE, size=self.image.nbytes*2),
         # but platforms not suporting openCL 1.2 have a problem with enqueue_fill_buffer,
         # so we use the parray "fill" utility
diff --git a/silx/opencl/test/test_sparse.py b/silx/opencl/test/test_sparse.py
index 41650691b..76a6a0a67 100644
--- a/silx/opencl/test/test_sparse.py
+++ b/silx/opencl/test/test_sparse.py
@@ -115,9 +115,12 @@ class TestCSR(unittest.TestCase):
         else:
             arr = array
         if output_on_device:
-            d_data = parray.zeros_like(csr.data)
-            d_indices = parray.zeros_like(csr.indices)
-            d_indptr = parray.zeros_like(csr.indptr)
+            d_data = parray.empty_like(csr.data)
+            d_indices = parray.empty_like(csr.indices)
+            d_indptr = parray.empty_like(csr.indptr)
+            d_data.fill(0)
+            d_indices.fill(0)
+            d_indptr.fill(0)
             output = (d_data, d_indices, d_indptr)
         else:
             output = None
@@ -169,7 +172,8 @@ class TestCSR(unittest.TestCase):
             indices = ref_sparse.indices
             indptr = ref_sparse.indptr
         if output_on_device:
-            d_arr = parray.zeros_like(csr.array)
+            d_arr = parray.empty_like(csr.array)
+            d_arr.fill(0)
             output = d_arr
         else:
             output = None