Using nabu from python to reconstruct a dataset with GPU¶

This notebook shows how to use the Nabu software for performing a basic reconstruction of a tomography dataset.
The computations are done on a local machine with a GPU and Cuda available.

This tutorial goes a bit further than nabu_basic_reconstruction.ipynb:

• GPU implementation of each component is used
• We see how to start from a configuration file and devise a simple processing chain accordingly

The same dataset is used (binned scan of a bamboo stick, thanks Ludovic Broche, ESRF ID19).

Note¶

In this notebook, several GPU operations are done manually: memory allocation, data transfers from/to GPU, ... It can be relevant if you have custom processing needs.

Nabu offers many facilities for dealing automatically with these operations. Please refer to the notebook nabu_pipeline_reconstruction.ipynb if you don't need finer control on the operations.

1 - Load the dataset informations¶

We must provide nabu with the the configuration file (nabu.conf), describing the path to the dataset and the processing steps. This is the equivalent of the .par file in PyHST2. In this file, no information is given on the detector size, energy, distance, etc: these informations are extracted from the dataset metadata.

In [1]:

import os
from nabu.testutils import utilstest, get_file
from nabu.pipeline.fullfield.processconfig import ProcessConfig

import os from nabu.testutils import utilstest, get_file from nabu.pipeline.fullfield.processconfig import ProcessConfig

In [2]:

print("Getting dataset (downloading if necessary) ...")
data_path = get_file("bamboo_reduced.nx")
print("... OK")

# Get the configuration file of this dataset
conf_fname = get_file("bamboo_reduced.conf")

# Change directory to the path where the data is located (only useful for this tutorial)
os.chdir(utilstest.data_home)

# Parse this configuration file
conf = ProcessConfig(conf_fname)

print("Getting dataset (downloading if necessary) ...") data_path = get_file("bamboo_reduced.nx") print("... OK") # Get the configuration file of this dataset conf_fname = get_file("bamboo_reduced.conf") # Change directory to the path where the data is located (only useful for this tutorial) os.chdir(utilstest.data_home) # Parse this configuration file conf = ProcessConfig(conf_fname)

Getting dataset (downloading if necessary) ...

ERROR:nabu:Cannot do SRCurrent normalization: missing flats and/or projections SRCurrent

... OK
Option 'double_flatfield_enabled' has been renamed 'double_flatfield' in [preproc]
This is deprecated since version 2025.1.0 and will result in an error in futures versions

Note that ProcessConfig will do quite a few things under the hood:

• Parse the configuration file and check parameters correctness
• Browse the dataset
• Get or compute the reduced flats/darks
• Estimate the center of rotation

The resulting object contains all necessary information to process the dataset.

In [3]:

# We can easily get information on the processing steps.
nabu_config = conf.nabu_config
from pprint import pprint
pprint(nabu_config)
# The same can be done with the dataset structure
dataset_info = conf.dataset_info
# print([getattr(dataset_info, attr) for attr in ["energy", "distance", "n_angles", "radio_dims"]])

# We can easily get information on the processing steps. nabu_config = conf.nabu_config from pprint import pprint pprint(nabu_config) # The same can be done with the dataset structure dataset_info = conf.dataset_info # print([getattr(dataset_info, attr) for attr in ["energy", "distance", "n_angles", "radio_dims"]])

{'about': {},
 'dataset': {'binning': 1,
             'binning_z': 1,
             'darks_flats_dir': None,
             'exclude_projections': None,
             'flip_lr': 'auto',
             'flip_ud': 'auto',
             'hdf5_entry': None,
             'location': '/tmp/nabu_testdata_pierre/bamboo_reduced.nx',
             'nexus_version': None,
             'overwrite_metadata': '',
             'projections_subsampling': (1, 0)},
 'output': {'file_format': 'hdf5',
            'file_prefix': 'bamboo_reduced_rec',
            'float_clip_values': None,
            'jpeg2000_compression_ratio': None,
            'keep_existing_files': False,
            'location': '/tmp/nabu_testdata_pierre',
            'overwrite_results': True,
            'tiff_single_file': False,
            'zarr_options': ''},
 'phase': {'ctf_advanced_params': 'length_scale=1e-5; lim1=1e-5; lim2=0.2; '
                                  'normalize_by_mean=True',
           'ctf_geometry': 'z1_v=None; z1_h=None; detec_pixel_size=None; '
                           'magnification=True',
           'delta_beta': 100.0,
           'material_density': None,
           'material_formula': None,
           'method': 'paganin',
           'padding_type': 'edge',
           'unsharp_coeff': 0.0,
           'unsharp_method': 'gaussian',
           'unsharp_sigma': 0.0},
 'pipeline': {'ignore_checkpoint_config': False,
              'processing_margin': None,
              'resume_from_step': None,
              'save_steps': None,
              'steps_file': None,
              'verbosity': 'info'},
 'postproc': {'histogram_bins': 1000000, 'output_histogram': False},
 'preproc': {'autotilt_options': None,
             'ccd_filter_enabled': False,
             'ccd_filter_threshold': 0.04,
             'detector_distortion_correction': None,
             'detector_distortion_correction_options': None,
             'dff_sigma': None,
             'double_flatfield': False,
             'double_flatfield_enabled': '0',
             'flat_distortion_correction_enabled': False,
             'flat_distortion_params': 'tile_size=100; '
                                       "interpolation_kind='linear'; "
                                       "padding_mode='edge'; "
                                       'correction_spike_threshold=None',
             'flatfield': True,
             'flatfield_loading_mode': 'load_if_present',
             'log_max_clip': 10.0,
             'log_min_clip': 1e-06,
             'normalize_srcurrent': True,
             'processes_file': None,
             'rotate_projections_center': None,
             'sino_normalization': None,
             'sino_normalization_file': '',
             'sino_rings_correction': 'munch',
             'sino_rings_options': None,
             'take_logarithm': True,
             'target_dark_mean': None,
             'tilt_correction': None},
 'reconstruction': {'angle_offset': 0.0,
                    'angles_file': None,
                    'axis_correction_file': None,
                    'centered_axis': False,
                    'clip_outer_circle': False,
                    'cor_options': "side='from_file'",
                    'cor_slice': None,
                    'crop_filtered_data': True,
                    'enable_halftomo': 'auto',
                    'end_x': -1,
                    'end_y': -1,
                    'end_z': -1,
                    'expand_support_factor': 1.0,
                    'fbp_filter_cutoff': 1.0,
                    'fbp_filter_type': 'ramlak',
                    'hbp_legs': 4,
                    'hbp_reduction_steps': 2,
                    'implementation': None,
                    'iterations': 200,
                    'method': 'FBP',
                    'optim_algorithm': 'chambolle-pock',
                    'outer_circle_value': 0.0,
                    'padding_type': 'edge',
                    'positivity_constraint': True,
                    'preconditioning_filter': True,
                    'regularization_weight': 0.0,
                    'rotation_axis_position': 'sliding-window',
                    'sample_detector_dist': None,
                    'source_sample_dist': None,
                    'start_x': 0,
                    'start_y': 0,
                    'start_z': 0,
                    'translation_movements_file': None},
 'resources': {'gpu_id': [],
               'gpus': 1,
               'memory_fraction': (90.0, True),
               'method': 'local',
               'num_threads': 24,
               'workers': 1}}

2 - Chunk processing¶

Nabu processes data by chunks of radios (see the documentation for more explanations).
In a first step, we define how to read chunks of radios.

In [4]:

from nabu.io.reader import NXTomoReader

from nabu.io.reader import NXTomoReader

What is the largest chunk size we can process ?
The answer is given by inspecting the current GPU memory, and the processing steps.

In [5]:

from nabu.cuda.utils import collect_cuda_gpus
from nabu.pipeline.fullfield.computations import estimate_max_chunk_size

from nabu.cuda.utils import collect_cuda_gpus from nabu.pipeline.fullfield.computations import estimate_max_chunk_size

In [6]:

# Pick the fist GPU
gpu0 = collect_cuda_gpus()[0]
chunk_size = estimate_max_chunk_size(
    gpu0["memory_avail_GB"], 
    conf
)
print("Chunk_size = %d" % chunk_size)

# Pick the fist GPU gpu0 = collect_cuda_gpus()[0] chunk_size = estimate_max_chunk_size( gpu0["memory_avail_GB"], conf ) print("Chunk_size = %d" % chunk_size)

Chunk_size = 540

In [7]:

# Load the first 'chunk_size' lines of all the radios
# i.e do projections_data[:, 0:chunk_size, :]
sub_region = (
    slice(None), 
    slice(0, chunk_size),
    slice(None)
)
projections_reader = NXTomoReader(
    data_path,
    sub_region=sub_region, 
)

# Load the first 'chunk_size' lines of all the radios # i.e do projections_data[:, 0:chunk_size, :] sub_region = ( slice(None), slice(0, chunk_size), slice(None) ) projections_reader = NXTomoReader( data_path, sub_region=sub_region, )

In [8]:

# Load the current chunk
print("Loading data", end="")
projections = projections_reader.load_data() # takes some time
print("... OK")

# Load the current chunk print("Loading data", end="") projections = projections_reader.load_data() # takes some time print("... OK")

Loading data

... OK

3 - Initialize the GPU¶

Most of the processing can be done on GPU (or many-core CPU if using OpenCL).

• For Cuda, nabu uses the cupy Cuda wrapper.
• For OpenCL, nabu uses pyopencl.array.

In both cases, we manipulate array objects with memory residing on device. This allows to avoid extraneous host <-> device copies.

In [9]:

import cupy
import numpy as np

import cupy import numpy as np

In [10]:

# Create a Cuda context on current GPU
# By default, all following GPU processings will be bound on this context
cupy.cuda.Device(gpu0["device_id"]).use()

# Create a Cuda context on current GPU # By default, all following GPU processings will be bound on this context cupy.cuda.Device(gpu0["device_id"]).use()

Out[10]:

<CUDA Device 0>

In [11]:

n_angles, n_z, n_x = projections.shape
# transfer the chunk on GPU
d_radios = cupy.array(projections)

n_angles, n_z, n_x = projections.shape # transfer the chunk on GPU d_radios = cupy.array(projections)

4 - Pre-processing¶

Pre-processing utilities are available in the nabu.preproc module.
Utilities available with the cuda backend are implemented in a module with a _cuda suffix.

4.1 - Flat-field¶

In [12]:

from nabu.preproc.flatfield_cuda import CudaFlatField

from nabu.preproc.flatfield_cuda import CudaFlatField

In [13]:

radios_indices = sorted(conf.dataset_info.projections.keys())
# Configure the `FlatField` processor
cuda_flatfield = CudaFlatField(
    d_radios.shape, 
    dataset_info.get_reduced_flats(sub_region=sub_region), 
    dataset_info.get_reduced_darks(sub_region=sub_region), 
    radios_indices=radios_indices,
)

radios_indices = sorted(conf.dataset_info.projections.keys()) # Configure the `FlatField` processor cuda_flatfield = CudaFlatField( d_radios.shape, dataset_info.get_reduced_flats(sub_region=sub_region), dataset_info.get_reduced_darks(sub_region=sub_region), radios_indices=radios_indices, )

---------------------------------------------------------------------------
AttributeError                            Traceback (most recent call last)
Cell In[13], line 3
      1 radios_indices = sorted(conf.dataset_info.projections.keys())
      2 # Configure the `FlatField` processor
----> 3 cuda_flatfield = CudaFlatField(
      4     d_radios.shape, 
      5     dataset_info.get_reduced_flats(sub_region=sub_region), 
      6     dataset_info.get_reduced_darks(sub_region=sub_region), 
      7     radios_indices=radios_indices,
      8 )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/flatfield_cuda.py:44, in CudaFlatFieldArrays.__init__(self, radios_shape, flats, darks, radios_indices, interpolation, distortion_correction, nan_value, radios_srcurrent, flats_srcurrent, cuda_options)
     32 super().__init__(
     33     radios_shape,
     34     flats,
   (...)     41     nan_value=nan_value,
     42 )
     43 self.cuda_processing = CudaProcessing(**(cuda_options or {}))
---> 44 self._init_cuda_kernels()
     45 self._load_flats_and_darks_on_gpu()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/flatfield_cuda.py:59, in CudaFlatFieldArrays._init_cuda_kernels(self)
     57 if self.nan_value is not None:
     58     options.append("-DNAN_VALUE=%f" % self.nan_value)
---> 59 self.cuda_kernel = self.cuda_processing.kernel(
     60     "flatfield_normalization", self._cuda_fname, options=tuple(options)
     61 )
     62 self._nx = np.int32(self.shape[1])
     63 self._ny = np.int32(self.shape[0])

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/processing.py:66, in CudaProcessing.kernel(self, kernel_name, filename, src, automation_params, **build_kwargs)
     65 def kernel(self, kernel_name, filename=None, src=None, automation_params=None, **build_kwargs):
---> 66     return CudaKernel(  # pylint: disable=E0606
     67         kernel_name,
     68         filename=filename,
     69         src=src,
     70         automation_params=automation_params,
     71         **build_kwargs,
     72     )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:48, in CudaKernel.__init__(self, kernel_name, filename, src, automation_params, silent_compilation_warnings, extern_c, **sourcemodule_kwargs)
     45 if extern_c:
     46     # pycuda/pyopencl do that automatically, not cupy
     47     self.src = patch_sourcecode_add_externC(self.src, filename=filename)
---> 48 self.compile_kernel_source(kernel_name, sourcemodule_kwargs)

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:59, in CudaKernel.compile_kernel_source(self, kernel_name, sourcemodule_kwargs)
     57 with catch_warnings(action=("ignore" if self.silent_compilation_warnings else None)):  # pylint: disable=E1123
     58     self.module = RawModule(code=self.src, **self.sourcemodule_kwargs)
---> 59     self.module.compile()
     60 self.func = self.module.get_function(kernel_name)

File cupy/_core/raw.pyx:437, in cupy._core.raw.RawModule.compile()

File cupy/_core/raw.pyx:415, in cupy._core.raw.RawModule._module()

File cupy/_util.pyx:68, in cupy._util.memoize.decorator.ret()

File cupy/_core/raw.pyx:549, in cupy._core.raw._get_raw_module()

File cupy/_core/core.pyx:2534, in cupy._core.core.compile_with_cache()

File cupy/_core/core.pyx:2552, in cupy._core.core.compile_with_cache()

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:566, in _compile_module_with_cache(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, jitify, to_ltoir)
    562     return _compile_with_cache_hip(
    563         source, options, arch, cache_dir, extra_source, backend,
    564         name_expressions, log_stream, cache_in_memory)
    565 else:
--> 566     return _compile_with_cache_cuda(
    567         source, options, arch, cache_dir, extra_source, backend,
    568         enable_cooperative_groups, name_expressions, log_stream,
    569         cache_in_memory, jitify, to_ltoir)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:618, in _compile_with_cache_cuda(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, cache_in_memory, jitify, to_ltoir)
    615 base = _empty_file_preprocess_cache.get(env, None)
    616 if base is None:
    617     # This is for checking NVRTC/NVCC compiler internal version
--> 618     base = _preprocess('', options, arch, backend)
    619     _empty_file_preprocess_cache[env] = base
    621 key_src = '%s %s %s %s %s' % (
    622     env, base, source, extra_source, _get_cupy_cache_key())

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:513, in _preprocess(source, options, arch, backend)
    511 try:
    512     options = options + ('-o', 'preprocess.ptx')
--> 513     result = compile_using_nvcc(source, options, arch, 'preprocess.cu',
    514                                 code_type='ptx')
    515 except CompileException as e:
    516     dump = _get_bool_env_variable(
    517         'CUPY_DUMP_CUDA_SOURCE_ON_ERROR', False)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:425, in compile_using_nvcc(source, options, arch, filename, code_type, separate_compilation, log_stream)
    423 _nvcc = get_nvcc_path()
    424 # split() is needed because _nvcc could come from the env var NVCC
--> 425 cmd = _nvcc.split()
    426 cmd.append(arch_str)
    428 with tempfile.TemporaryDirectory() as root_dir:

AttributeError: 'NoneType' object has no attribute 'split'

In [14]:

# Perform the normalization on GPU
if nabu_config["preproc"]["flatfield"]:
    print("Doing flat-field", end="")
    cuda_flatfield.normalize_radios(d_radios)
    print("... OK")

# Perform the normalization on GPU if nabu_config["preproc"]["flatfield"]: print("Doing flat-field", end="") cuda_flatfield.normalize_radios(d_radios) print("... OK")

Doing flat-field

---------------------------------------------------------------------------
NameError                                 Traceback (most recent call last)
Cell In[14], line 4
      2 if nabu_config["preproc"]["flatfield"]:
      3     print("Doing flat-field", end="")
----> 4     cuda_flatfield.normalize_radios(d_radios)
      5     print("... OK")

NameError: name 'cuda_flatfield' is not defined

4.2 - Phase retrieval¶

In [15]:

from nabu.preproc.phase_cuda import CudaPaganinPhaseRetrieval

from nabu.preproc.phase_cuda import CudaPaganinPhaseRetrieval

In [16]:

energy = dataset_info.energy
# Phase retrieval is done on each radio individually, with the sub-region specified above
if (nabu_config["phase"]["method"] or "").lower() == "paganin":
    print("Doing phase retrieval", end="")
    cudapaganin = CudaPaganinPhaseRetrieval(
        (n_z, n_x),
        distance=dataset_info.distance,
        energy=energy,
        delta_beta=nabu_config["phase"]["delta_beta"],
        pixel_size=dataset_info.pixel_size * 1e6,
    )
    for i in range(n_angles):
        cudapaganin.apply_filter(d_radios[i], output=d_radios[i])
    print("... OK")

energy = dataset_info.energy # Phase retrieval is done on each radio individually, with the sub-region specified above if (nabu_config["phase"]["method"] or "").lower() == "paganin": print("Doing phase retrieval", end="") cudapaganin = CudaPaganinPhaseRetrieval( (n_z, n_x), distance=dataset_info.distance, energy=energy, delta_beta=nabu_config["phase"]["delta_beta"], pixel_size=dataset_info.pixel_size * 1e6, ) for i in range(n_angles): cudapaganin.apply_filter(d_radios[i], output=d_radios[i]) print("... OK")

Doing phase retrieval

---------------------------------------------------------------------------
AttributeError                            Traceback (most recent call last)
Cell In[16], line 5
      3 if (nabu_config["phase"]["method"] or "").lower() == "paganin":
      4     print("Doing phase retrieval", end="")
----> 5     cudapaganin = CudaPaganinPhaseRetrieval(
      6         (n_z, n_x),
      7         distance=dataset_info.distance,
      8         energy=energy,
      9         delta_beta=nabu_config["phase"]["delta_beta"],
     10         pixel_size=dataset_info.pixel_size * 1e6,
     11     )
     12     for i in range(n_angles):
     13         cudapaganin.apply_filter(d_radios[i], output=d_radios[i])

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/phase_cuda.py:45, in CudaPaganinPhaseRetrieval.__init__(self, shape, distance, energy, delta_beta, pixel_size, padding, cuda_options, fftw_num_threads, fft_num_threads, fft_backend)
     43 self._init_gpu_arrays()
     44 self._init_fft(fft_backend)
---> 45 self._init_padding_kernel()
     46 self._init_mult_kernel()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/phase_cuda.py:70, in CudaPaganinPhaseRetrieval._init_padding_kernel(self)
     69 def _init_padding_kernel(self):
---> 70     self.padding_kernel = CudaPadding(
     71         shape=self.shape,
     72         pad_width=(
     73             (self.pad_top_len, self.pad_bottom_len),
     74             (self.pad_left_len, self.pad_right_len),
     75         ),
     76         mode=self.padding,
     77     )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/processing/padding_cuda.py:17, in CudaPadding.__init__(self, shape, pad_width, mode, cuda_options, **kwargs)
     15 super().__init__(shape, pad_width, mode=mode, **kwargs)
     16 self.cuda_processing = self.processing = CudaProcessing(**(cuda_options or {}))
---> 17 self._init_cuda_coordinate_transform()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/processing/padding_cuda.py:25, in CudaPadding._init_cuda_coordinate_transform(self)
     21     self.d_padded_array_constant = self.processing.to_device(
     22         "d_padded_array_constant", self.padded_array_constant
     23     )
     24     return
---> 25 self._coords_transform_kernel = self.processing.kernel(
     26     "coordinate_transform",
     27     filename=get_cuda_srcfile("padding.cu"),
     28 )
     29 self._coords_transform_block = (32, 32, 1)
     30 self._coords_transform_grid = [
     31     updiv(a, b) for a, b in zip(self.padded_shape[::-1], self._coords_transform_block)
     32 ]

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/processing.py:66, in CudaProcessing.kernel(self, kernel_name, filename, src, automation_params, **build_kwargs)
     65 def kernel(self, kernel_name, filename=None, src=None, automation_params=None, **build_kwargs):
---> 66     return CudaKernel(  # pylint: disable=E0606
     67         kernel_name,
     68         filename=filename,
     69         src=src,
     70         automation_params=automation_params,
     71         **build_kwargs,
     72     )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:48, in CudaKernel.__init__(self, kernel_name, filename, src, automation_params, silent_compilation_warnings, extern_c, **sourcemodule_kwargs)
     45 if extern_c:
     46     # pycuda/pyopencl do that automatically, not cupy
     47     self.src = patch_sourcecode_add_externC(self.src, filename=filename)
---> 48 self.compile_kernel_source(kernel_name, sourcemodule_kwargs)

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:59, in CudaKernel.compile_kernel_source(self, kernel_name, sourcemodule_kwargs)
     57 with catch_warnings(action=("ignore" if self.silent_compilation_warnings else None)):  # pylint: disable=E1123
     58     self.module = RawModule(code=self.src, **self.sourcemodule_kwargs)
---> 59     self.module.compile()
     60 self.func = self.module.get_function(kernel_name)

File cupy/_core/raw.pyx:437, in cupy._core.raw.RawModule.compile()

File cupy/_core/raw.pyx:415, in cupy._core.raw.RawModule._module()

File cupy/_util.pyx:68, in cupy._util.memoize.decorator.ret()

File cupy/_core/raw.pyx:549, in cupy._core.raw._get_raw_module()

File cupy/_core/core.pyx:2534, in cupy._core.core.compile_with_cache()

File cupy/_core/core.pyx:2552, in cupy._core.core.compile_with_cache()

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:566, in _compile_module_with_cache(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, jitify, to_ltoir)
    562     return _compile_with_cache_hip(
    563         source, options, arch, cache_dir, extra_source, backend,
    564         name_expressions, log_stream, cache_in_memory)
    565 else:
--> 566     return _compile_with_cache_cuda(
    567         source, options, arch, cache_dir, extra_source, backend,
    568         enable_cooperative_groups, name_expressions, log_stream,
    569         cache_in_memory, jitify, to_ltoir)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:618, in _compile_with_cache_cuda(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, cache_in_memory, jitify, to_ltoir)
    615 base = _empty_file_preprocess_cache.get(env, None)
    616 if base is None:
    617     # This is for checking NVRTC/NVCC compiler internal version
--> 618     base = _preprocess('', options, arch, backend)
    619     _empty_file_preprocess_cache[env] = base
    621 key_src = '%s %s %s %s %s' % (
    622     env, base, source, extra_source, _get_cupy_cache_key())

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:513, in _preprocess(source, options, arch, backend)
    511 try:
    512     options = options + ('-o', 'preprocess.ptx')
--> 513     result = compile_using_nvcc(source, options, arch, 'preprocess.cu',
    514                                 code_type='ptx')
    515 except CompileException as e:
    516     dump = _get_bool_env_variable(
    517         'CUPY_DUMP_CUDA_SOURCE_ON_ERROR', False)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:425, in compile_using_nvcc(source, options, arch, filename, code_type, separate_compilation, log_stream)
    423 _nvcc = get_nvcc_path()
    424 # split() is needed because _nvcc could come from the env var NVCC
--> 425 cmd = _nvcc.split()
    426 cmd.append(arch_str)
    428 with tempfile.TemporaryDirectory() as root_dir:

AttributeError: 'NoneType' object has no attribute 'split'

4.3 - Logarithm¶

In [17]:

from nabu.preproc.ccd_cuda import CudaLog

from nabu.preproc.ccd_cuda import CudaLog

In [18]:

if nabu_config["preproc"]["take_logarithm"]:
    print("Taking logarithm", end="")
    cuda_log = CudaLog(d_radios.shape, clip_min=0.01)
    cuda_log.take_logarithm(d_radios)
    print("... OK")

if nabu_config["preproc"]["take_logarithm"]: print("Taking logarithm", end="") cuda_log = CudaLog(d_radios.shape, clip_min=0.01) cuda_log.take_logarithm(d_radios) print("... OK")

Taking logarithm

---------------------------------------------------------------------------
AttributeError                            Traceback (most recent call last)
Cell In[18], line 3
      1 if nabu_config["preproc"]["take_logarithm"]:
      2     print("Taking logarithm", end="")
----> 3     cuda_log = CudaLog(d_radios.shape, clip_min=0.01)
      4     cuda_log.take_logarithm(d_radios)
      5     print("... OK")

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/ccd_cuda.py:106, in CudaLog.__init__(self, radios_shape, clip_min, clip_max)
     93 """
     94 Initialize a Log processing.
     95 
   (...)    103     Data bigger than this value is replaced by this value.
    104 """
    105 super().__init__(radios_shape, clip_min=clip_min, clip_max=clip_max)
--> 106 self._init_kernels()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/preproc/ccd_cuda.py:121, in CudaLog._init_kernels(self)
    118 self._nthreadsperblock = (16, 16, 4)  # TODO tune ?
    119 self._nblocks = tuple([updiv(n, p) for n, p in zip([nx, ny, nz], self._nthreadsperblock)])
--> 121 self.nlog_kernel = CudaKernel(  # pylint: disable=E0606
    122     "nlog",
    123     filename=self._nlog_srcfile,
    124     options=(
    125         "-DDO_CLIP_MIN=%d" % self._do_clip_min,
    126         "-DDO_CLIP_MAX=%d" % self._do_clip_max,
    127     ),
    128 )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:48, in CudaKernel.__init__(self, kernel_name, filename, src, automation_params, silent_compilation_warnings, extern_c, **sourcemodule_kwargs)
     45 if extern_c:
     46     # pycuda/pyopencl do that automatically, not cupy
     47     self.src = patch_sourcecode_add_externC(self.src, filename=filename)
---> 48 self.compile_kernel_source(kernel_name, sourcemodule_kwargs)

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:59, in CudaKernel.compile_kernel_source(self, kernel_name, sourcemodule_kwargs)
     57 with catch_warnings(action=("ignore" if self.silent_compilation_warnings else None)):  # pylint: disable=E1123
     58     self.module = RawModule(code=self.src, **self.sourcemodule_kwargs)
---> 59     self.module.compile()
     60 self.func = self.module.get_function(kernel_name)

File cupy/_core/raw.pyx:437, in cupy._core.raw.RawModule.compile()

File cupy/_core/raw.pyx:415, in cupy._core.raw.RawModule._module()

File cupy/_util.pyx:68, in cupy._util.memoize.decorator.ret()

File cupy/_core/raw.pyx:549, in cupy._core.raw._get_raw_module()

File cupy/_core/core.pyx:2534, in cupy._core.core.compile_with_cache()

File cupy/_core/core.pyx:2552, in cupy._core.core.compile_with_cache()

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:566, in _compile_module_with_cache(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, jitify, to_ltoir)
    562     return _compile_with_cache_hip(
    563         source, options, arch, cache_dir, extra_source, backend,
    564         name_expressions, log_stream, cache_in_memory)
    565 else:
--> 566     return _compile_with_cache_cuda(
    567         source, options, arch, cache_dir, extra_source, backend,
    568         enable_cooperative_groups, name_expressions, log_stream,
    569         cache_in_memory, jitify, to_ltoir)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:618, in _compile_with_cache_cuda(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, cache_in_memory, jitify, to_ltoir)
    615 base = _empty_file_preprocess_cache.get(env, None)
    616 if base is None:
    617     # This is for checking NVRTC/NVCC compiler internal version
--> 618     base = _preprocess('', options, arch, backend)
    619     _empty_file_preprocess_cache[env] = base
    621 key_src = '%s %s %s %s %s' % (
    622     env, base, source, extra_source, _get_cupy_cache_key())

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:513, in _preprocess(source, options, arch, backend)
    511 try:
    512     options = options + ('-o', 'preprocess.ptx')
--> 513     result = compile_using_nvcc(source, options, arch, 'preprocess.cu',
    514                                 code_type='ptx')
    515 except CompileException as e:
    516     dump = _get_bool_env_variable(
    517         'CUPY_DUMP_CUDA_SOURCE_ON_ERROR', False)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:425, in compile_using_nvcc(source, options, arch, filename, code_type, separate_compilation, log_stream)
    423 _nvcc = get_nvcc_path()
    424 # split() is needed because _nvcc could come from the env var NVCC
--> 425 cmd = _nvcc.split()
    426 cmd.append(arch_str)
    428 with tempfile.TemporaryDirectory() as root_dir:

AttributeError: 'NoneType' object has no attribute 'split'

5 - Reconstruction¶

We use the filtered backprojection with nabu.reconstruction.fbp

In [19]:

from nabu.reconstruction.fbp import Backprojector

from nabu.reconstruction.fbp import Backprojector

In [20]:

rec_options = conf.processing_options["reconstruction"]
B = Backprojector(
    (n_angles, n_x), 
    angles=rec_options["angles"], 
    rot_center=rec_options["rotation_axis_position"],
    padding_mode="edges",
    # extra_options={"use_textures": False}
)
d_recs = cupy.zeros((n_z, n_x, n_x), "f")

rec_options = conf.processing_options["reconstruction"] B = Backprojector( (n_angles, n_x), angles=rec_options["angles"], rot_center=rec_options["rotation_axis_position"], padding_mode="edges", # extra_options={"use_textures": False} ) d_recs = cupy.zeros((n_z, n_x, n_x), "f")

---------------------------------------------------------------------------
AttributeError                            Traceback (most recent call last)
Cell In[20], line 2
      1 rec_options = conf.processing_options["reconstruction"]
----> 2 B = Backprojector(
      3     (n_angles, n_x), 
      4     angles=rec_options["angles"], 
      5     rot_center=rec_options["rotation_axis_position"],
      6     padding_mode="edges",
      7     # extra_options={"use_textures": False}
      8 )
      9 d_recs = cupy.zeros((n_z, n_x, n_x), "f")

File ~/.venv/py313/lib/python3.13/site-packages/nabu/reconstruction/fbp_base.py:122, in BackprojectorBase.__init__(self, sino_shape, slice_shape, angles, rot_center, padding_mode, halftomo, filter_name, slice_roi, extra_options, backend_options)
    120 self._check_textures_availability()
    121 self._init_geometry(sino_shape, slice_shape, angles, rot_center, halftomo, slice_roi)
--> 122 self._init_filter(filter_name)
    123 self._allocate_memory()
    124 self._compute_angles()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/reconstruction/fbp_base.py:267, in BackprojectorBase._init_filter(self, filter_name)
    265 #
    266 sinofilter_other_kwargs = self._get_filter_init_extra_options()
--> 267 self.sino_filter = self.SinoFilterClass(
    268     self.sino_shape,
    269     filter_name=self.filter_name,
    270     padding_mode=self.padding_mode,
    271     extra_options={
    272         "cutoff": self.extra_options.get("filter_cutoff", 1.0),
    273         "pixel_size_cm": self.extra_options.get("pixel_size_cm"),
    274     },
    275     **sinofilter_other_kwargs,
    276 )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/reconstruction/filtering_cuda.py:31, in CudaSinoFilter.__init__(self, sino_shape, filter_name, padding_mode, crop_filtered_data, extra_options, cuda_options)
     23 self.cuda = CudaProcessing(**self._cuda_options)
     24 super().__init__(
     25     sino_shape,
     26     filter_name=filter_name,
   (...)     29     extra_options=extra_options,
     30 )
---> 31 self._init_kernels()

File ~/.venv/py313/lib/python3.13/site-packages/nabu/reconstruction/filtering_cuda.py:57, in CudaSinoFilter._init_kernels(self)
     55 else:
     56     kernel_name = "inplace_complex_mul_3Dby1D"
---> 57 self.mult_kernel = self.cuda.kernel(kernel_name, filename=fname)
     58 self.kern_args = (self.d_sino_f, self.d_filter_f)
     59 self.kern_args += self.d_sino_f.shape[::-1]

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/processing.py:66, in CudaProcessing.kernel(self, kernel_name, filename, src, automation_params, **build_kwargs)
     65 def kernel(self, kernel_name, filename=None, src=None, automation_params=None, **build_kwargs):
---> 66     return CudaKernel(  # pylint: disable=E0606
     67         kernel_name,
     68         filename=filename,
     69         src=src,
     70         automation_params=automation_params,
     71         **build_kwargs,
     72     )

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:48, in CudaKernel.__init__(self, kernel_name, filename, src, automation_params, silent_compilation_warnings, extern_c, **sourcemodule_kwargs)
     45 if extern_c:
     46     # pycuda/pyopencl do that automatically, not cupy
     47     self.src = patch_sourcecode_add_externC(self.src, filename=filename)
---> 48 self.compile_kernel_source(kernel_name, sourcemodule_kwargs)

File ~/.venv/py313/lib/python3.13/site-packages/nabu/cuda/kernel.py:59, in CudaKernel.compile_kernel_source(self, kernel_name, sourcemodule_kwargs)
     57 with catch_warnings(action=("ignore" if self.silent_compilation_warnings else None)):  # pylint: disable=E1123
     58     self.module = RawModule(code=self.src, **self.sourcemodule_kwargs)
---> 59     self.module.compile()
     60 self.func = self.module.get_function(kernel_name)

File cupy/_core/raw.pyx:437, in cupy._core.raw.RawModule.compile()

File cupy/_core/raw.pyx:415, in cupy._core.raw.RawModule._module()

File cupy/_util.pyx:68, in cupy._util.memoize.decorator.ret()

File cupy/_core/raw.pyx:549, in cupy._core.raw._get_raw_module()

File cupy/_core/core.pyx:2534, in cupy._core.core.compile_with_cache()

File cupy/_core/core.pyx:2552, in cupy._core.core.compile_with_cache()

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:566, in _compile_module_with_cache(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, jitify, to_ltoir)
    562     return _compile_with_cache_hip(
    563         source, options, arch, cache_dir, extra_source, backend,
    564         name_expressions, log_stream, cache_in_memory)
    565 else:
--> 566     return _compile_with_cache_cuda(
    567         source, options, arch, cache_dir, extra_source, backend,
    568         enable_cooperative_groups, name_expressions, log_stream,
    569         cache_in_memory, jitify, to_ltoir)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:618, in _compile_with_cache_cuda(source, options, arch, cache_dir, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, cache_in_memory, jitify, to_ltoir)
    615 base = _empty_file_preprocess_cache.get(env, None)
    616 if base is None:
    617     # This is for checking NVRTC/NVCC compiler internal version
--> 618     base = _preprocess('', options, arch, backend)
    619     _empty_file_preprocess_cache[env] = base
    621 key_src = '%s %s %s %s %s' % (
    622     env, base, source, extra_source, _get_cupy_cache_key())

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:513, in _preprocess(source, options, arch, backend)
    511 try:
    512     options = options + ('-o', 'preprocess.ptx')
--> 513     result = compile_using_nvcc(source, options, arch, 'preprocess.cu',
    514                                 code_type='ptx')
    515 except CompileException as e:
    516     dump = _get_bool_env_variable(
    517         'CUPY_DUMP_CUDA_SOURCE_ON_ERROR', False)

File ~/.venv/py313/lib/python3.13/site-packages/cupy/cuda/compiler.py:425, in compile_using_nvcc(source, options, arch, filename, code_type, separate_compilation, log_stream)
    423 _nvcc = get_nvcc_path()
    424 # split() is needed because _nvcc could come from the env var NVCC
--> 425 cmd = _nvcc.split()
    426 cmd.append(arch_str)
    428 with tempfile.TemporaryDirectory() as root_dir:

AttributeError: 'NoneType' object has no attribute 'split'

In [21]:

print("Reconstructing...", end="")
for i in range(n_z):
    B.fbp(d_radios[:, i, :], output=d_recs[i])
recs = d_recs.get()
print(" ... OK")

print("Reconstructing...", end="") for i in range(n_z): B.fbp(d_radios[:, i, :], output=d_recs[i]) recs = d_recs.get() print(" ... OK")

Reconstructing...

---------------------------------------------------------------------------
NameError                                 Traceback (most recent call last)
Cell In[21], line 3
      1 print("Reconstructing...", end="")
      2 for i in range(n_z):
----> 3     B.fbp(d_radios[:, i, :], output=d_recs[i])
      4 recs = d_recs.get()
      5 print(" ... OK")

NameError: name 'B' is not defined

6 - Visualize¶

In [22]:

import matplotlib.pyplot as plt

import matplotlib.pyplot as plt

In [23]:

plt.figure()
plt.imshow(recs[0], cmap="gray")
plt.show()

plt.figure() plt.imshow(recs[0], cmap="gray") plt.show()

---------------------------------------------------------------------------
NameError                                 Traceback (most recent call last)
Cell In[23], line 2
      1 plt.figure()
----> 2 plt.imshow(recs[0], cmap="gray")
      3 plt.show()

NameError: name 'recs' is not defined

<Figure size 640x480 with 0 Axes>

In [ ]: