diff options
Diffstat (limited to 'silx/opencl/backprojection.py')
| -rw-r--r-- | silx/opencl/backprojection.py | 33 |
1 files changed, 15 insertions, 18 deletions
diff --git a/silx/opencl/backprojection.py b/silx/opencl/backprojection.py index 5a4087b..65a9836 100644 --- a/silx/opencl/backprojection.py +++ b/silx/opencl/backprojection.py @@ -164,9 +164,7 @@ class Backprojection(OpenclProcessing): def _allocate_memory(self): # Host memory self.slice = np.zeros(self.dimrec_shape, dtype=np.float32) - self.is_cpu = False - if self.device.type == "CPU": - self.is_cpu = True + self._use_textures = self.check_textures_availability() # Device memory self.buffers = [ @@ -180,7 +178,7 @@ class Backprojection(OpenclProcessing): self.d_sino = self.cl_mem["d_sino"] # shorthand # Texture memory (if relevant) - if not(self.is_cpu): + if self._use_textures: self._allocate_textures() # Local memory @@ -199,7 +197,14 @@ class Backprojection(OpenclProcessing): self.cl_mem["d_axes"][:] = np.ones(self.num_projs, dtype="f") * self.axis_pos def _init_kernels(self): - OpenclProcessing.compile_kernels(self, self.kernel_files) + compile_options = None + if not(self._use_textures): + compile_options = "-DDONT_USE_TEXTURES" + OpenclProcessing.compile_kernels( + self, + self.kernel_files, + compile_options=compile_options + ) # check that workgroup can actually be (16, 16) self.compiletime_workgroup_size = self.kernels.max_workgroup_size("backproj_cpu_kernel") # Workgroup and ndrange sizes are always the same @@ -209,7 +214,7 @@ class Backprojection(OpenclProcessing): _idivup(int(self.dimrec_shape[0]), 32) * self.wg[1] ) # Prepare arguments for the kernel call - if self.is_cpu: + if not(self._use_textures): d_sino_ref = self.d_sino.data else: d_sino_ref = self.d_sino_tex @@ -242,15 +247,7 @@ class Backprojection(OpenclProcessing): """ Allocate the texture for the sinogram. """ - self.d_sino_tex = pyopencl.Image( - self.ctx, - mf.READ_ONLY | mf.USE_HOST_PTR, - pyopencl.ImageFormat( - pyopencl.channel_order.INTENSITY, - pyopencl.channel_type.FLOAT - ), - hostbuf=np.zeros(self.shape[::-1], dtype=np.float32) - ) + self.d_sino_tex = self.allocate_texture(self.shape) def _init_filter(self, filter_name): """Filter initialization @@ -289,7 +286,7 @@ class Backprojection(OpenclProcessing): sino2 = sino if not(sino.flags["C_CONTIGUOUS"] and sino.dtype == np.float32): sino2 = np.ascontiguousarray(sino, dtype=np.float32) - if self.is_cpu: + if not(self._use_textures): ev = pyopencl.enqueue_copy( self.queue, self.d_sino.data, @@ -309,7 +306,7 @@ class Backprojection(OpenclProcessing): return EventDescription(what, ev) def _transfer_device_to_texture(self, d_sino): - if self.is_cpu: + if not(self._use_textures): if id(self.d_sino) == id(d_sino): return ev = pyopencl.enqueue_copy( @@ -343,7 +340,7 @@ class Backprojection(OpenclProcessing): with self.sem: events.append(self._transfer_to_texture(sino)) # Call the backprojection kernel - if self.is_cpu: + if not(self._use_textures): kernel_to_call = self.kernels.backproj_cpu_kernel else: kernel_to_call = self.kernels.backproj_kernel |