From 63d8bcafe26722a6104b44ad7706a619680e09fd Mon Sep 17 00:00:00 2001 From: Sangjoon Bob Lee Date: Thu, 8 Jan 2026 15:57:50 -0800 Subject: [PATCH 1/4] add CUDA-based disk to GPU load in 0.5s --- src/quantem/hpc/__init__.py | 23 + src/quantem/hpc/io.py | 578 ++++++++++++++++++ .../hpc/notebooks/performance_comparison.png | Bin 0 -> 21029 bytes 3 files changed, 601 insertions(+) create mode 100644 src/quantem/hpc/__init__.py create mode 100644 src/quantem/hpc/io.py create mode 100644 src/quantem/hpc/notebooks/performance_comparison.png diff --git a/src/quantem/hpc/__init__.py b/src/quantem/hpc/__init__.py new file mode 100644 index 000000000..c75b35fce --- /dev/null +++ b/src/quantem/hpc/__init__.py @@ -0,0 +1,23 @@ +""" +quantem.hpc: GPU-accelerated 4D-STEM loading and processing. + +Examples +-------- +>>> from quantem.hpc import load, bin +>>> data = load('lamella_data_000001.h5') +>>> binned = bin(data, factor=2) +""" + +from quantem.hpc.io import ( + GPUDecompressor, + load, + clear_memory, + bin, +) + +__all__ = [ + 'GPUDecompressor', + 'load', + 'clear_memory', + 'bin', +] diff --git a/src/quantem/hpc/io.py b/src/quantem/hpc/io.py new file mode 100644 index 000000000..afa92792b --- /dev/null +++ b/src/quantem/hpc/io.py @@ -0,0 +1,578 @@ +""" +GPU-accelerated HDF5 loading and binning for 4D-STEM diffraction data. + +Provides high-performance bitshuffle+LZ4 decompression directly to GPU memory. + +Examples +-------- +>>> from quantem.hpc import load, bin +>>> data = load('/path/to/file.h5') +>>> binned = bin(data, factor=2) +""" + +from __future__ import annotations + +import cupy as cp +import h5py +import hdf5plugin # noqa: F401 - registers bitshuffle filter +import numpy as np +from numba import njit, prange + +__all__ = ["GPUDecompressor", "load", "clear_memory", "bin"] + +# CUDA LZ4 decompression kernel (adapted from NVIDIA nvcomp, BSD-3-Clause) +_CUDA_LZ4_SOURCE = r''' +/* + * LZ4 decompression kernel extracted from NVIDIA nvcomp + * Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * BSD-3-Clause License + */ +typedef unsigned char uint8_t; +typedef unsigned short uint16_t; +typedef unsigned int uint32_t; +typedef unsigned long long uint64_t; +typedef long long int64_t; +typedef unsigned long size_t; + +using offset_type = uint16_t; +using word_type = uint32_t; +using position_type = uint32_t; +using double_word_type = uint64_t; +using item_type = uint32_t; + +constexpr const int DECOMP_THREADS_PER_CHUNK = 32; +constexpr const int DECOMP_CHUNKS_PER_BLOCK = 2; +constexpr const position_type DECOMP_INPUT_BUFFER_SIZE + = DECOMP_THREADS_PER_CHUNK * sizeof(double_word_type); +constexpr const position_type DECOMP_BUFFER_PREFETCH_DIST + = DECOMP_INPUT_BUFFER_SIZE / 2; + +inline __device__ void syncCTA() { + if (DECOMP_THREADS_PER_CHUNK > 32) __syncthreads(); + else __syncwarp(); +} + +inline __device__ int warpBallot(int vote) { + return __ballot_sync(0xffffffff, vote); +} + +inline __device__ offset_type readWord(const uint8_t* const address) { + offset_type word = 0; + for (size_t i = 0; i < sizeof(offset_type); ++i) + word |= address[i] << (8 * i); + return word; +} + +struct token_type { + position_type num_literals; + position_type num_matches; + + __device__ bool hasNumLiteralsOverflow() const { return num_literals >= 15; } + __device__ bool hasNumMatchesOverflow() const { return num_matches >= 19; } + + __device__ position_type numLiteralsOverflow() const { + return hasNumLiteralsOverflow() ? num_literals - 15 : 0; + } + __device__ uint8_t numLiteralsForHeader() const { + return hasNumLiteralsOverflow() ? 15 : num_literals; + } + __device__ position_type numMatchesOverflow() const { + return hasNumMatchesOverflow() ? num_matches - 19 : 0; + } + __device__ uint8_t numMatchesForHeader() const { + return hasNumMatchesOverflow() ? 15 : num_matches - 4; + } + __device__ position_type lengthOfLiteralEncoding() const { + if (hasNumLiteralsOverflow()) { + const position_type num = numLiteralsOverflow(); + return (num / 0xff) + 1; + } + return 0; + } + __device__ position_type lengthOfMatchEncoding() const { + if (hasNumMatchesOverflow()) { + const position_type num = numMatchesOverflow(); + return (num / 0xff) + 1; + } + return 0; + } +}; + +class BufferControl { +public: + __device__ BufferControl(uint8_t* const buffer, const uint8_t* const compData, + const position_type length) + : m_offset(0), m_length(length), m_buffer(buffer), m_compData(compData) {} + + inline __device__ position_type readLSIC(position_type& idx) const { + position_type num = 0; + uint8_t next = 0xff; + while (next == 0xff && idx < end()) { + next = rawAt(idx)[0]; + ++idx; + num += next; + } + while (next == 0xff) { + next = m_compData[idx]; + ++idx; + num += next; + } + return num; + } + + inline __device__ const uint8_t* raw() const { return m_buffer; } + inline __device__ const uint8_t* rawAt(const position_type i) const { + return raw() + (i - begin()); + } + + inline __device__ uint8_t operator[](const position_type i) const { + if (i >= m_offset && i - m_offset < DECOMP_INPUT_BUFFER_SIZE) + return m_buffer[i - m_offset]; + return m_compData[i]; + } + + inline __device__ void setAndAlignOffset(const position_type offset) { + const uint8_t* const alignedPtr = reinterpret_cast( + (reinterpret_cast(m_compData + offset) / sizeof(double_word_type)) + * sizeof(double_word_type)); + m_offset = alignedPtr - m_compData; + } + + inline __device__ void loadAt(const position_type offset) { + setAndAlignOffset(offset); + if (m_offset + DECOMP_INPUT_BUFFER_SIZE <= m_length) { + const double_word_type* const word_data + = reinterpret_cast(m_compData + m_offset); + double_word_type* const word_buffer + = reinterpret_cast(m_buffer); + word_buffer[threadIdx.x] = word_data[threadIdx.x]; + } else { + #pragma unroll + for (int i = threadIdx.x; i < DECOMP_INPUT_BUFFER_SIZE; + i += DECOMP_THREADS_PER_CHUNK) { + if (m_offset + i < m_length) + m_buffer[i] = m_compData[m_offset + i]; + } + } + syncCTA(); + } + + inline __device__ position_type begin() const { return m_offset; } + inline __device__ position_type end() const { return m_offset + DECOMP_INPUT_BUFFER_SIZE; } + +private: + int64_t m_offset; + const position_type m_length; + uint8_t* const m_buffer; + const uint8_t* const m_compData; +}; + +inline __device__ void coopCopyNoOverlap(uint8_t* const dest, const uint8_t* const source, + const position_type length) { + for (position_type i = threadIdx.x; i < length; i += blockDim.x) + dest[i] = source[i]; +} + +inline __device__ void coopCopyRepeat(uint8_t* const dest, const uint8_t* const source, + const position_type dist, const position_type length) { + for (position_type i = threadIdx.x; i < length; i += blockDim.x) + dest[i] = source[i % dist]; +} + +inline __device__ void coopCopyOverlap(uint8_t* const dest, const uint8_t* const source, + const position_type dist, const position_type length) { + if (dist < length) coopCopyRepeat(dest, source, dist, length); + else coopCopyNoOverlap(dest, source, length); +} + +inline __device__ token_type decodePair(const uint8_t num) { + return token_type{static_cast((num & 0xf0) >> 4), + static_cast(num & 0x0f)}; +} + +inline __device__ void decompressStream(uint8_t* buffer, uint8_t* decompData, + const uint8_t* compData, const position_type comp_end) { + BufferControl ctrl(buffer, compData, comp_end); + ctrl.loadAt(0); + position_type decomp_idx = 0; + position_type comp_idx = 0; + + while (comp_idx < comp_end) { + if (comp_idx + DECOMP_BUFFER_PREFETCH_DIST > ctrl.end()) + ctrl.loadAt(comp_idx); + + token_type tok = decodePair(*ctrl.rawAt(comp_idx)); + ++comp_idx; + + position_type num_literals = tok.num_literals; + if (tok.num_literals == 15) + num_literals += ctrl.readLSIC(comp_idx); + const position_type literalStart = comp_idx; + + if (num_literals + comp_idx > ctrl.end()) + coopCopyNoOverlap(decompData + decomp_idx, compData + comp_idx, num_literals); + else + coopCopyNoOverlap(decompData + decomp_idx, ctrl.rawAt(comp_idx), num_literals); + + comp_idx += num_literals; + decomp_idx += num_literals; + + if (comp_idx < comp_end) { + offset_type offset; + if (comp_idx + sizeof(offset_type) > ctrl.end()) + offset = readWord(compData + comp_idx); + else + offset = readWord(ctrl.rawAt(comp_idx)); + + comp_idx += sizeof(offset_type); + + position_type match = 4 + tok.num_matches; + if (tok.num_matches == 15) + match += ctrl.readLSIC(comp_idx); + + if (offset <= num_literals + && (ctrl.begin() <= literalStart && ctrl.end() >= literalStart + num_literals)) { + coopCopyOverlap(decompData + decomp_idx, + ctrl.rawAt(literalStart + (num_literals - offset)), offset, match); + syncCTA(); + } else { + syncCTA(); + coopCopyOverlap(decompData + decomp_idx, + decompData + decomp_idx - offset, offset, match); + } + decomp_idx += match; + } + } +} + +inline __device__ uint32_t read32be_batch(const uint8_t* address) { + return ((uint32_t)(255 & address[0]) << 24 | (uint32_t)(255 & address[1]) << 16 | + (uint32_t)(255 & address[2]) << 8 | (uint32_t)(255 & address[3])); +} + +extern "C" __global__ void h5lz4dc_batched( + const uint8_t* const compressed, const uint32_t* const chunk_offsets, + const uint32_t* const block_starts, const uint32_t* const block_counts, + const uint32_t* const block_offsets, const uint32_t blocksize, + const uint32_t frame_bytes, uint8_t* const decompressed +) { + const int frame_id = blockIdx.z; + const int block_in_frame = blockIdx.x * blockDim.y + threadIdx.y; + const uint32_t chunk_offset = chunk_offsets[frame_id]; + const uint32_t block_offset = block_offsets[frame_id]; + const uint32_t num_blocks = block_counts[frame_id]; + __shared__ uint8_t buffer[DECOMP_INPUT_BUFFER_SIZE * DECOMP_CHUNKS_PER_BLOCK]; + + if (block_in_frame < num_blocks) { + const uint32_t block_start = block_starts[block_offset + block_in_frame]; + const uint8_t* input = compressed + chunk_offset + block_start + 4; + const uint32_t comp_size = read32be_batch(compressed + chunk_offset + block_start); + uint8_t* output = decompressed + frame_id * frame_bytes + block_in_frame * blocksize; + decompressStream(buffer + threadIdx.y * DECOMP_INPUT_BUFFER_SIZE, output, input, comp_size); + } +} + +extern "C" __global__ void shuf_8192_32_batched( + const uint32_t* __restrict__ in, uint32_t* __restrict__ out, const uint32_t frame_u32s +) { + const int frame_id = blockIdx.z; + const uint32_t* frame_in = in + frame_id * frame_u32s; + uint32_t* frame_out = out + frame_id * frame_u32s; + __shared__ uint32_t smem[32][33]; + + smem[threadIdx.y][threadIdx.x] = frame_in[threadIdx.x + threadIdx.y * 64 + + blockIdx.x * 2048 + blockIdx.y * 32]; + __syncthreads(); + + uint32_t v = smem[threadIdx.x][threadIdx.y]; + #pragma unroll 32 + for (int i = 0; i < 32; i++) + smem[i][threadIdx.y] = __ballot_sync(0xFFFFFFFFU, v & (1U << i)); + __syncthreads(); + + frame_out[threadIdx.x + threadIdx.y * 32 + blockIdx.y * 1024 + blockIdx.x * 2048] = + smem[threadIdx.x][threadIdx.y]; +} +''' + +# Compile CUDA kernels at import time for fast first load() +_cuda_module = cp.RawModule(code=_CUDA_LZ4_SOURCE, options=("-std=c++11", "-w")) +_h5lz4dc_kernel = _cuda_module.get_function("h5lz4dc_batched") +_bitshuffle_kernel = _cuda_module.get_function("shuf_8192_32_batched") + +class GPUDecompressor: + """GPU-accelerated decompressor for bitshuffle+LZ4 HDF5 datasets. + + Uses pinned memory and CUDA kernels for maximum throughput. + CUDA kernels are compiled at module import time. + """ + + def __init__( + self, + max_compressed_bytes: int = 1024 * 1024 * 1024, + max_frames: int = 100000, + max_frame_bytes: int = 192 * 192 * 4, + n_blocks_per_frame: int = 18, + ): + """Initialize the decompressor with pre-allocated buffers. + + Parameters + ---------- + max_compressed_bytes : int, optional + Maximum size of compressed data, by default 1GB. + max_frames : int, optional + Maximum number of frames to support, by default 100000. + max_frame_bytes : int, optional + Maximum bytes per frame, by default 147456 (192x192 uint32). + n_blocks_per_frame : int, optional + LZ4 blocks per frame, by default 18 for 192x192 uint32. + """ + self.max_compressed_bytes = max_compressed_bytes + self.max_frames = max_frames + self.max_frame_bytes = max_frame_bytes + self.n_blocks_per_frame = n_blocks_per_frame + self._h5lz4dc = _h5lz4dc_kernel + self._shuf = _bitshuffle_kernel + # Pinned memory for fast CPU->GPU transfers + self._pinned_mem = cp.cuda.alloc_pinned_memory(max_compressed_bytes) + self._pinned_buffer = np.frombuffer( + self._pinned_mem, dtype=np.uint8, count=max_compressed_bytes + ) + # Pre-allocated metadata arrays + self._chunk_sizes = np.zeros(max_frames, dtype=np.uint32) + self._chunk_offsets = np.zeros(max_frames, dtype=np.uint32) + self._block_counts = np.zeros(max_frames, dtype=np.uint32) + self._block_starts_flat = np.zeros(max_frames * n_blocks_per_frame, dtype=np.uint32) + self._block_offsets = np.zeros(max_frames + 1, dtype=np.uint32) + # Pre-allocate all GPU buffers for fast first load() + self._concat_gpu = cp.empty(max_compressed_bytes, dtype=cp.uint8) + total_output_bytes = max_frames * max_frame_bytes + self._lz4_output = cp.empty(total_output_bytes, dtype=cp.uint8) + self._shuffled_output = cp.empty(total_output_bytes, dtype=cp.uint8) + + def load( + self, + filepath: str, + dataset_path: str = "entry/data/data", + ) -> cp.ndarray: + """Load and decompress a bitshuffle+LZ4 HDF5 dataset to GPU. + + Parameters + ---------- + filepath : str + Path to the HDF5 file. + dataset_path : str, optional + Path to the dataset within the HDF5 file, by default "entry/data/data". + + Returns + ------- + cp.ndarray + CuPy array on GPU with shape (n_frames, height, width). + """ + with h5py.File(filepath, "r") as f: + ds = f[dataset_path] + n_frames = ds.shape[0] + frame_shape = ds.shape[1:] + dtype = ds.dtype + frame_bytes = int(np.prod(frame_shape) * np.dtype(dtype).itemsize) + + # Reallocate GPU buffers only if dataset exceeds pre-allocated size + total_needed = n_frames * frame_bytes + if total_needed > len(self._lz4_output): + self._lz4_output = cp.empty(total_needed, dtype=cp.uint8) + self._shuffled_output = cp.empty(total_needed, dtype=cp.uint8) + # Read chunks into pinned memory + offset = 0 + for i in range(n_frames): + _, raw = ds.id.read_direct_chunk((i, 0, 0)) + chunk_len = len(raw) + self._chunk_offsets[i] = offset + self._chunk_sizes[i] = chunk_len + self._pinned_buffer[offset : offset + chunk_len] = np.frombuffer( + raw, dtype=np.uint8 + ) + offset += chunk_len + total_compressed = offset + # Parse headers + _parse_headers( + self._pinned_buffer, + self._chunk_sizes, + self._chunk_offsets, + self._block_starts_flat, + self._block_counts, + n_frames, + self.n_blocks_per_frame, + ) + # Compute block offsets + self._block_offsets[1 : n_frames + 1] = np.cumsum(self._block_counts[:n_frames]) + total_blocks = int(self._block_offsets[n_frames]) + # Transfer to GPU + self._concat_gpu[:total_compressed].set(self._pinned_buffer[:total_compressed]) + chunk_offsets_gpu = cp.asarray(self._chunk_offsets[:n_frames]) + block_starts_gpu = cp.asarray(self._block_starts_flat[:total_blocks]) + block_counts_gpu = cp.asarray(self._block_counts[:n_frames]) + block_offsets_gpu = cp.asarray(self._block_offsets[: n_frames + 1]) + # LZ4 decompress + block_size = 8192 + max_blocks = int(self._block_counts[:n_frames].max()) + max_batch = 10000 + for start in range(0, n_frames, max_batch): + end = min(start + max_batch, n_frames) + batch_n = end - start + byte_offset = start * frame_bytes + self._h5lz4dc( + ((max_blocks + 1) // 2, 1, batch_n), + (32, 2, 1), + ( + self._concat_gpu, + chunk_offsets_gpu[start:], + block_starts_gpu, + block_counts_gpu[start:], + block_offsets_gpu[start:], + np.uint32(block_size), + np.uint32(frame_bytes), + self._lz4_output[byte_offset:], + ), + ) + # Bitshuffle + n_8kb = frame_bytes // 8192 + frame_u32s = frame_bytes // 4 + for start in range(0, n_frames, max_batch): + end = min(start + max_batch, n_frames) + batch_n = end - start + byte_offset = start * frame_bytes + self._shuf( + (n_8kb, 2, batch_n), + (32, 32, 1), + ( + self._lz4_output[byte_offset:].view(cp.uint32), + self._shuffled_output[byte_offset:].view(cp.uint32), + np.uint32(frame_u32s), + ), + ) + cp.cuda.Device().synchronize() + total_bytes = n_frames * frame_bytes + return self._shuffled_output[:total_bytes].view(dtype).reshape( + (n_frames,) + frame_shape + ) + +@njit(cache=True, parallel=True) +def _parse_headers( + pinned_buffer, + chunk_sizes, + chunk_offsets, + block_starts_out, + block_counts_out, + n_frames, + n_blocks_per_frame, +): + """Parse bitshuffle+LZ4 chunk headers in parallel.""" + for i in prange(n_frames): + offset = chunk_offsets[i] + chunk = pinned_buffer[offset : offset + chunk_sizes[i]] + + # Parse header (first 12 bytes) + uncomp_size = ( + int(chunk[0]) << 56 + | int(chunk[1]) << 48 + | int(chunk[2]) << 40 + | int(chunk[3]) << 32 + | int(chunk[4]) << 24 + | int(chunk[5]) << 16 + | int(chunk[6]) << 8 + | int(chunk[7]) + ) + block_size = ( + int(chunk[8]) << 24 + | int(chunk[9]) << 16 + | int(chunk[10]) << 8 + | int(chunk[11]) + ) + n_blocks = (uncomp_size + block_size - 1) // block_size + block_counts_out[i] = n_blocks + pos = 12 + base_idx = i * n_blocks_per_frame + for b in range(n_blocks): + block_starts_out[base_idx + b] = pos + comp_size = ( + int(chunk[pos]) << 24 + | int(chunk[pos + 1]) << 16 + | int(chunk[pos + 2]) << 8 + | int(chunk[pos + 3]) + ) + pos += 4 + comp_size + +# Pre-allocate decompressor at import for fast first load() +_default_decompressor = GPUDecompressor() + +def load(filepath: str, dataset_path: str = "entry/data/data") -> cp.ndarray: + """Load a bitshuffle+LZ4 compressed HDF5 dataset directly to GPU. + + Parameters + ---------- + filepath : str + Path to the HDF5 file. + dataset_path : str, optional + Path to the dataset within the HDF5 file, by default "entry/data/data". + + Returns + ------- + cp.ndarray + CuPy array on GPU with shape (n_frames, height, width). + """ + return _default_decompressor.load(filepath, dataset_path) + +def bin( + data: cp.ndarray, + factor: int = 2, + dtype: type | np.dtype | None = None, + reduction: str = "sum", +) -> cp.ndarray: + """Apply spatial binning to a stack of 2D images on GPU. + + Parameters + ---------- + data : cp.ndarray + CuPy array with shape (n_frames, height, width) or (height, width). + factor : int, optional + Binning factor (2 for 2x2, 4 for 4x4, etc.), by default 2. + dtype : type or np.dtype, optional + Output dtype. If None, uses uint32 for int input (sum), float32 for mean. + reduction : str, optional + Reduction method - 'sum' (default) or 'mean'. + + Returns + ------- + cp.ndarray + Binned CuPy array with reduced spatial dimensions. + """ + if reduction not in ("sum", "mean"): + raise ValueError(f"reduction must be 'sum' or 'mean', got '{reduction}'") + if factor == 1: + return data + is_2d = data.ndim == 2 + if is_2d: + data = data[None, :, :] + if data.ndim != 3: + raise ValueError(f"Expected 2D or 3D array, got {data.ndim}D") + n, h, w = data.shape + if h % factor != 0 or w % factor != 0: + raise ValueError(f"Dimensions ({h}, {w}) not divisible by factor {factor}") + reshaped = data.reshape(n, h // factor, factor, w // factor, factor) + if dtype is None: + dtype = cp.float32 if reduction == "mean" else ( + cp.uint32 if cp.issubdtype(data.dtype, cp.integer) else cp.float32 + ) + if reduction == "mean": + result = reshaped.mean(axis=(2, 4), dtype=dtype) + else: + result = reshaped.sum(axis=(2, 4), dtype=dtype) + return result[0] if is_2d else result + +def clear_memory() -> None: + """Release GPU memory pools.""" + try: + cp.get_default_memory_pool().free_all_blocks() + cp.get_default_pinned_memory_pool().free_all_blocks() + except Exception: + pass \ No newline at end of file diff --git a/src/quantem/hpc/notebooks/performance_comparison.png b/src/quantem/hpc/notebooks/performance_comparison.png new file mode 100644 index 0000000000000000000000000000000000000000..5d88be09cd06f646e9e6609130464c91fe77434b GIT binary patch literal 21029 zcmeIacT`ka_brNct8GAA#Vjaop`eHX$$}VAiy&Dts7TIPK!OQdP{ER*A|Qz5AUO#b znk+d{wVftt0;} zG7781cPT!go;t2_<>^p|lXd`OVc~nWYLAz#yq7I^@_69l=8-iXPkr}DF3OSoW>eTvHRM>u z;Ap0A=i?14NS;jD>~--s+y#CPKG0Fz^z)aVG~eIJqvI!P*UiPuUwHQah7}+`7er5R z|NPaX`~RhvRfWf_WEp02cXyYTl~rksQ}xMZV}8-Fa7s#wt&>x8OPbM`Kxlx%>xH04 zk7CT4lQuYPVP4VxAW0`L>aV~4y6S7h#rAE%Ajz{e&&g`78uR8~hwM3c@LOha+PeQ2 zp6cIgWXCWudn?gVEs{&G1R7J(VDKPsrjT}c2tHj*rGhyZs(h6*`etd%`bKK z?4P_FeJ)+P~xhM9v*JZFbfhi|8Su`us&2_n~-qNad|m{>)1wLGp$)^%g3>g3pR zHTPwtr725FdY*eK{6N3((WA{8yc-j>vd+;|o}Xq7c1yDK_xF!nzJ7O%``omxgF`@2 zP|$GRh{S&F?1O3o2AX+}iFnLE{`kXwe%4uh{zv!d*x19oyu8*-i}-ungfwT|Ct8e5 zKfDar*3nUYA-w^Ac5cg-EiF|ulU*KW4Ka#wYRRD)1v4XsdNM(LmEYZN-@M6v=#b)p zSKqG3sl@)FrluxfQg=#GQL(9@%Y7zMGu@ce$;m0;fS&HV+iTQpZEZ6Jo(kKXIC*<@ z#{6J;z`p30FArL@We!CN{8dsyzqLtM;oeq(p08hH*6up}^|4XtgQ82dcen7{&hpZm z68DLTiLr~j94V1tF88yqY?{CI_V&Iw7a~ffH6>`+Rm|g?O`DTK{@ArkHT=YqTeoh- z+H@CB<9-=fhJm-YR%Siz&S=)HmpPCh86_8N`p%0rPU(g8eK+x=M~~o`kUOQZad236 zben#BeXYrT`SRttBLy=@90%1yj(O3xDXUBh$I&+2|ew(ANGvAfbZda}GE_~UJs|rB}^qPKr?-nq8|A^O){pb~i zSBE0c(4Ifnz&l-1=PY{6PqR<>gG>ZbM_I~?dp%Z}W#l9DsnoB&GYQxX0) z8!USLn;Xj*uYXw+V_5klXYK$7ox-Td5)Rf|)(KTmmSX-$Nb z4_UwZOp7+b03PMtB`G*SF4O(d0{X?PGOW6EkIjC&aa3H~%0)s?Pp>jWOpTM1Q`Bkr zX_V|^+R#8kv#zV0nB!pS;gK+p`Ph#iFUSQ8%B>L}_sg~Kx2>3b#B=zt(%dwoev{tS ziyhm{+Ouu89&>-?GC7=x6KPT%D&C~EadN0Fy7=BU`nxs9aAu$PeEP&mm6n%R!{0i5 z{%g71lOxmv1wSGN2b22l-d?lgys>c%i7-H=F-f+e=Zv6$je%~h9d#15_jU(3PAz$j$FieMqTamGrdDY7WGZJ4rUhZ z*@}5iBUP!<>W0FMzt+q5O!F^e9`#%CHRn%In{RORPnU;jP0^DJDbzI*KQ07jeyI%Z z+q_pN-$m7JdYptxL%yrie!YU(z(p#ST)TIh;uxf+rPW7=xtsD^A#==+g)N+=MV$+I z674J#u=j;YeN>Bi)%@&Ob;YF=-jK5nkY*3!~qGCkhk#3D^wx@z-IY%aURRdobd+EN9Y zO3EcE@re)mv`9UVxki_(L#-LDX-2+r=M&?Ex~|rLIbZzko3+D8r;b{Zb}hE{!I04* zq-)+q{uFUWgEB2O$;;cvryh&1nrW_q3&q<+7I+RGJlIz9VAmO20S@Le_2`5r-8{VN zBrSC?#Xp3^H!}1=Q^rUesaaVE{93t8O762u zpF4Lx&9IW~_uu_qPb+WBFn*|c| z3Uz4>*lqK|?=SI9nyHuXWOUuzs}gfzy?sZFTC%Ro{8XQS$NW{f(4!ibmX@^0qFN=izp>|>v9?=S__boue=Hlt{h9-A=>mrmglIi<|3 zekv_1D+|x+|Mcn8_(t8wdsXB7hrYjW8%#B;HkXl;%i!?2bLRnt?5kjKoRLI%;QoEo zmf>hs1%dVf{_>;l0}Zjf7OL^*D;@7ra-<8Q+G z?=kQFQM7{E37id?I!0TWgH!OEJ^xx#VwhYc&oQT2GWY`*I}K)F_Zb;ppfn z()lEC$A*39)0-ZIR0JKM*S(NIUaqut)wvj>os$|#kH0>j^Vs!bPNR{3j^@7T05u^z z8#_svL1w1E7$d0S}+Aqse(~*45fS*WRd3o`j|D#977GhVvJ?p$W zCP%IN_UWCSTfucynU3AF#n*?TOL5_P`c}+B|$D^=~1}!rCB%nNR9PXtIbRdF(%rq=s2><3JQLg^!2e| z`??#&DZNgKsy`@+fXI*)xZd2@YtZZqwX6JX9R_X73IAdOm)9wn_KO@xq1qJYW&G z3K4=kcX-Sy;)_frp=HM2!zv=s1^OayQHGIS?$4t z2jqK^#iB^^Foy(wW#y_>ynoV|u*b8E zxA@+q%%Y;#2o?foDj)NgZS;xudeaoOC&;)rfPCX?;7s2|+%+eFgzx6wIQ-OP96V6@(7xNE1+g^t% zFn`5zke%}t3|EzxP3Fb4gytcpjpI>(2I!^J|p{%a+fU2rV+9GH}=)2HY zl-S(4b7y+A;s9gBV?hh;L*R`YH!43odi3ao>zWp-ynHhO49}kJL}a%$5HgA7b)_(y z!y-xQk&ADCo%{A~jWdVS*BVaN?ZP!Ty;l^Iz6%KzJO?cjP~HTA*3^tjt;`r`^hw&u z`~W#T%kHu7-xamBwL?opgoR6silhJ}4u!|fJ{EIIV&otjAZZ_vm#>c4$E14{%&o;Z zE?n2mcQN_&_LivQ;3LQ3_Q!M^EN>&ZL)LAA+!7MHpFVwZ799!PuQiSv+jr>Dp%MuK zc+pb=5czlL2jay7fC@ziMVKx;-k+JLn2J9qCkJolA|J5V^QwCnWU zQbv1YyjpeYNd5ICD-&jp-@1A8EbUyFL}H45$pDJ?c^rC@WJE+o%iG$tkSXfPn%VVL z8MxjDgw_BYiPOx8dm0kb`{hd%ssoR(u==hu_ZtwzTsnI93{uq1EsGvTZWS<9r%(0u zxev7pxP33$gyN=}W7`X;5qa34{3}|GhGspF5R09g)~>yn)~D*?l8fpQb^6}cvc|>$ ztp&0bJAJd+*Kl*r&do8fU}hveq#54Jbs9-3Ej^2VR)!XMKraca-jAjiiLTCb(Jw1r zdAZF@n7qBYT;sl8!Bq`Xi|(_F9dVs7PXtJE7-&eztVWw~V<~I-VLm~@^Wefr#b==M z#ACmlllsxWkmUooHnf`^DO_m6+C}KOPf4LLH2^avhN(JEPmUxnS-G)iYL?unC;hi@ z&vkFBs?OoE7r3LiBA-oWNSP!J(YR#N-RE4;J@x~6k zlEP3208X$&@VAa9>*nVI9z~+4@7up$jxYvD{yp#sRlb4S@aNAF zsMBSYm2%PNLh1ls{r&td@H+ukPoQg=o0-(|bsA`(G7+C`Z#k8QT7$ooW7+Wg@6R+d zO#6|0M6Ex1QcPM>4G2bL^hbw%`LTsp{iP)Sk!fpsEDD0Hb}ci|OF1~%6s`*k3maoE z@N8-TPdLs|PXhZkNIm}a{$5rO4%i2C6+SAuBt>)c1hn-ugZs!@ zs%RaPC@9ye-oCzSu}Uu})IIW{=@Y@?u9+Ne;39(M(uf<8F)s#qub zjwihvbFfc5qN1(SE}#Fpg23Cm>A)*)Nf1tC9A5cwu;zB zR_?c&>~PX6c=xn@`m@i0BbQ74c*B()KX-R0R-Xgd&-^ zjmT2PskWF#zLv#z)S@L$z&*VnqWVuP%>5Lsq_6avZRX{evKyu1cRdn;_Mt!rL6k2PdkXi@8)ir5{o|Ei#T zJtp6EvNgeStDyNKa$tA3|2U46sC4DZ54jSsL0^Nd3Tp4H-(%lfzG0|6N11SDi9>+W zjkr27*{VHk?;aeXG-o%U<}F?iy7j>uh*TLEmFx&2cVR8|Lu-S=_1XEE5f`9OwoRM% zxVpN!0hVq(c)`qPn^sJb?KUkrUSbEAt51O-amd$P%y2hzdtt2eVV%F=r>Qt!4Kgm7b>Blp&)-QEy~wN zC##T@l!U4uNFnS5y*W7upjpte-B>&{1Of6jr!Q0)0k4{@D?at($5(8RpzBF$v(Uln z;T#I#xlnt+MUvK*ZSC4{r~z^Klh#xNA4)$#cqly$**5ylSFU^=Y)K=%H`>1y1Nvmm zS$aHwUvcmaa#6yYl4Iykk8k4BF3sF8ARtg;Ln8o^K8B3knywnFB(+~Ny(ZP5f@}a| zq^m2cqoP7)pITA_5NPs^WoryIxa{+ef=n}fzd3)R)vQVXF{cue1t9PNlhK|sp^$F$ zv30o)^i&IaeYEcrK?_s24sc(DdF!hRc}M{;JBQ**yWF8xektWExp{z$4E>VTRP0Vr zI(R7FF}pl=a^sJD3LO@4-9&lcxL-38H3of{aGL}!YpOoBkbpeFEi9b<+HA9csXzYR zi)jYE0(k`78sRZeFQU*n^a@*|Q6=FEeEb|)9JN^$h!otYGMFKrVH z+Lo7;>?skCK)3!3k!%9)8K3ZQc6JujCHJOHo3KqfpbBe&=8`{r_^^16m)tE&Ca;%< zUceGI{%hB+dHaaR3$;#G)Ms-(P?n2hq7qNBGZ!vAi=^wx2JHCXQ4QJdoy+p^+-K!R z`ICP{Z~R-@@)!Q!aVh?%i;e$(Q=@*K4conZs=k%2S`7vyo{EZRIj>3SEcU!~J>WlY zV#p`3{Fk*v|L4>G$6Cq%>Fd_wrUOYmWFv(OX;xK8W9z!4=T_!y4`eeHf{iT)6?{1#@FY*8!tNiynP{~ zL#y6#5B%ZB?^E;#b-K@&XVDVm1vzb0_4EPmAVVDe3zbS8ot(@P(N(W{ zD)JoMNIaS{TRXdb<70g!|DP**5jVE&7@GqRK3p>8kjL2zia{*T{C-`kyWrLZRllIv zPE7ZZfss+PQ0JAC;30Fbe7_X9@)k#!#5R*KA>^6|NdDD4J+t>P3NsM_@@J)(o_IoH*DP+ck|}21OpMceP93d z{q5wNVlBwGLG7ek-dHZC?T#+uPgi!8%M$O;K|>jERN`imVw@ zUGK<9m_a+*pBEs6bR6hS0Tcp|h;#!mkzrW5wF5*^P7>zv2#1eTml7q_2sIv>Yz*B)z>SdiKz$C3C(UtubDtix=DQk6pzLj#gmY@ z9#YW%O0*qfDodwnJ#>0}bf}9oZ5-e6rGGXjgboIDx;IF#XJ{x9?T*$j+|+@Z@EhQ- zLZZi{n}w7rhHZLE9~`az$@$pc4p1oi3*1FwwX(9Z1Ms&+ORALBtM6Xu*BfJ%qqVDh zp}Yd88;R(Gd8RfOw(yC8k#y#G_^`c-u!Gcr4=K-q?o#hM_dGar#AAN;XvmCV`?}9R z9GOv5y5!x#(N_X1hgUuJ(O$8QQ{~Vhr{Q+q633y|7m)I>d~zPM-)~ozM+AQ#8%xtr z^e@9I&{8{|YSX~?J?%RrQZ}Uh_?imfYHCO{4z#mqvj3PGbt#i!GY{<}J(p*m6vziY z-B&RDB<3<6LovwHQ=-;@7{*#JKL8;e?TX;g!rX|5PXrS(X45GY%2`@@&;k4; z|Bt}ScX5KYNh3km=g#L(Y*@c`?b@b3 zh)|6nTr-@<4ADXpyAjB7kJ=4Jsc4|$0eq)>9?GYqVJE^^w0u|-B7({Lk=7cTk$0{v z`6BKV%T;@_mn2`#`R-G(>fJ$y8@4B6E?w$;<=oTAT8n)mJfr%EN~sg zWTZ2HN_3MN>+luu>tGA@0=vzc5^5p06H*-7N%U`ff2smGQor32>qL{mws~_5xC4iw zR#li#l)$177cRJSQ7FbZ9wfOb1}$gaq70PS0XzmF-N^|N5fLFuJINwQha0wSGaC}K zY!^le3cigLY}%9%yC?L@324)qkHvQHE^Zw$1i5-SMg0h_;_tk4?Seda`*%ol%(e>s zbZfTFbFd7Mk`tjd_O}>?se)ee&OJrG-ZQequy;f-@vR_A$+gN1o9@%MZr`S(4@A{b zhNpnsk`Ac1woHq9u!2Z`l>IunDo~i8IcPsWeUFWejnWT70~Q1#m}+KQhZ9W>mu^zO z=d|Jk?gfSs0q4=42)>_7v<=E;(=Tg|KJ)iyBb%|6-|#N-a^euo0MqF01)z?H9CfZU zG>7_PaehC{EmB0lMrTP>b~CKgEr#F<)~xQI^{eN>#mz*>Mbk=_fC%fmJ!wfEd2aAE zkbWB7U5CBInl)?OU^Spcg6_DyJW+b_-%dTxqlAJ3HLwO`wFb^rfr%ahsxet_A9Z9q=dkUAvy5s;qQM_<2~edC{%1M?Z(MBgO2_W9)v@Jq0-XQ?5@3c-UE7->9^aXue=4GBok!B(&v8a++;>|x>*{vX&Wog8D zl9Q9uG4aZ2>-oRGoGq1T9hXv2LQ8Q}z~T~#8BBG|nq#vsV=DNU>si&i#@R(iaY;x-?$P+6*KhR1I@{L zp(RdG;Jb_O- zK)!wfRx=j)mM~_i>~;eUJo!I9`w;&DsxrJA;a;rUOy55`M5JTth_DlsDgon~orH`) zm)=yp6ItUdoot2YEU8b6fBD5>w1=9{XwniI+O}gyG8n5b^rI6K6Q=c1_n5qC%Y=x{ z=bMfro%>-aGV2ms%gJWFtPCn6DrzIrF|nl(kpVe<-MV%5J*8`jwv1AFI?x75K@ruN zfm3V-djYWwKz}7}8#oz=EI}%|_eXZ>zs)t|c9a)x?G(a8S+c$XHR_$t)NqS};BNw} zz8{l`$kduoD>4?J+Xwe4U!2CpvQJ7m( zlpA<#LA`g*Bl(l2(17Dq;{zG$wzjFz!p!)Gp`@hdbPPX~m#8Ixk4hxs;`{9TsG_z# ze@4lNB|(tnl(9&DHR??;Gi+yuL$cm6nK0atW}A+V z&OQB-`@N>a5Ja^~V7u9ZeAAR!A1$xdM%jWTEI;gC7fEfJ0LOSl#uZR0LlEk^K1@wH zab^&R1pBIn4;fW$e|K<|KFHx>EQ|#-(6=#4FT6i}`C=$6BC-W0{drl7h=n)})B@1Ke&lSOD=&|YG&*$aMx zi*6lJa@&eZO3cQ;))IXWmNN$&kvRD<2@Tp1l4ZqoJD?PLqNZkI>T#Py1S7i#i@1z` z4M&buAkhKWLC@e|0<;`0I+4OR@~ZEonzm&oUcY{Qbc(0P9B2Q@F%S2^ZX~cqL=ka_ z^n>?grp&?{VmO3;{Yq8uXbDp=h|*2RKEUwC109qAVasYva+I&AxV6NZnjsi;1#7GX-}{Z`A(|R4<=nGp4{`cz6}68+;on0gwn!u?o=uyAwDU($Dfb>d zs!bGbLiCr6z;zn1U+X;0ynwT&tc*&`wm3bzX(R6zHI76#RP0vSnH`4e@I7b&Qw8 z&{!3buPJ39c#(5x2r#0w5UUN6mdocmdw}_LkZDBQs&|exm7(&)VtI*&4+KGy1XF9` zwi(U$TETtseJ;BdY6U1+Ua zfVK{cIt=(5T(!5i&vY1WNh|rrUAW58i~+ zXrBYSiCzudXTsi=W+ZRc2CHgf3H*5mls$Xe?&e#6y8a!Vo?%X63O=XF2Eo_Tv>iUn zp>DVbXW<1E90QJNY@M$;xu)grTwmWD{DYgZ%xRg?k&!*)iUdbF@miMb%ULnzcvNic z^L`c@fzn!GZHTr!n<;E}+IBj6>~o< z^(Z=ZP$AOZ(+3BiT7=;EH6yNjy1So4E}H{}13K&OBCd~5pKRzt!SeY0W@I$tb0xKh zkB<-h3vuDV0+~3sn&YUkn|rs{(D(1(2^CNHe8}qD0s^Xptw*7?IIq=v8Z8+kUxw!6 z<3r3q{ge7WW^2j?6U2oGqBJ7|HZ>C? zFH&Xd_&yxBKz&GKim;ohz{-G{RT(u@Y*k7yOdy|*Y5tOt(;I!bDs1%hy0+6 z9dd%5xjM7pIt?8oq>+KLJz))KRUC*iG?!+{wqu8tTW2J7Jvf!J>1h-pB`CxVFE5^S zc5#V=HIP^&#hP>pH1>|g^TgXw z1MNle1Xef#!pf$cdq3*7x4ddqV!yjhI;uZnY}2%Q>M zL#zDR(F)?Ujim3hLwE8-KOacAYF+!PEr(A*pjE{&wf%m7Y#1$1MoSyu3l)H$SnlB& zC*~XUj%R6=eflZt<1onek5wOIO67Ka=jgjX$!Dm$ml z)EEztXJFd{E?;&XFj;cJboumjsOxa;jWrUFEcB2Uf-B#0wLK$}JH}tq3#!^b%$(-n z;2@Dp)F#64pqlu@lSuN!P(660?Lx&ILY7aN1|=n!TQ?ri(UR(sH$$HAy#?%ub4^qYze~ zNVw$+9N7Zh3|{(P7`M=Vj0hP>AXd~8d5TeGuuwAqouOL? ziWDm2c~XCeJMz-nqaO7m9lGFK1+MD_JO^4{VV4%EK!gAudc zLMBQ}y(t(#cja60=N)$|B= z(evfYZa74j{QB!x)RJZ-*3@?IJyb+#9O`rg;6PbNhYoUfoZ@q@rK`8c5ZM!QsmQQA znRb|ieOwbn90A`JVOUxdAb!$v19jOkbZskIS~O7e8xaEk&>vl0dJuwhaUSeeIRgnF zF>B(5=}Q5HM2*;vK7QPZ8zN>(#A6-OPvZ-jO~E$AYW4mxKPb{&a1{5btD*m&80kt@ ze119>EH&Y`LA#Q!l5~ymj$#Dm0dT(VvoL1s>aVhKSKf@aXP&~p$)N~7CP6yPjee>r zZw)x)ilA@)I{`Bd-NTz;P#$yOOJar#6z6|*vyXxl4EJCOh))U`4lt1yNsPn<$TRP| zcU6GSDj&T+a5ne5BG@(3ED-|+F0h#PMXv`NCN308-ip6eSc za7Q)V@Fwr@vcK(r+`p&gimNqn^_2US#x6NrT@@lmCOb%tB@mrV2E?euDhav~*hPH3 zWLA3MP!rMqJNPziYU+E0HYfkROB8VJdWa3zD|;= z1gJ0$`gJYa0aCgIGKi&}d8r>K;JroW9P2v;xSvReh$0Omg7Mdy z@FiOIYyTM-g(zfs>56cc-&y!YMEHMW=#;vHVOh*bz(*05O;~eRtwJYZmigEc6ik;`MLk@gOK;m z*4EaNi7^h>!IZ1j{08N0FdNJT&VehT02fn5CUR0LJ37R>1Rme|*LjLTcSVb|Dx4Fn zuy8UYcJsfPnlf`cJG+OzzP|key&L{@zl1{0?F)5f+qNyf9IsPeVd5<+Sx-Eh{Y&{dfDs)cb!~SErzfD3k1h>>adQJ~RPj2C=#jZVOzXJ&2V# zY%{S!BZm-=UYlqa3L^u}lPS_JaX!Liyboj_aX=AD0(h*J@ZKgwX{ zbwXhhJNAW(gDMKB3LVA6l<9=p? z|A&*u+1KP%3;!On^;~!7?}_Dhh<_;()upA}Z>1G#;OS4Swwe_QzV_e!^#dGd3eTZK zDSyTvKMv1d5Cw}7%yB~Gk*FA{DlFV(E)@5K9)RF-l8ls;d`}pKi z+Lyu!25qGRk`b&^QC;49{@y)LZOrhHszLDk7B9|s2fcGCJvWLj@%Y#Dhy8svB||r; z`d$0%!~PfZ=Klu zEJf$RLX}>&$r1+DL|K~~@mIHINykk(yVLq4r*^sFNO}jZhvGBEA^+QwSr6c}0J8O2 zf9_lKe$(d7+;`~gTsxXYhj(j_WCvU4cUvmalAI>~CU5_5p9&+z2*996_&%Mzr~^*kzb)wyM(9y)Ih~5zk)_oa!o>4R(2%CGt7#`Y z`G9Iu+_Z@877ikru|aA$k22nnq$5TODSCyyTA2xCga80_qr+p#P%YLc|S0 z3*VuPLu+1alm9%Mqz+KOs8TUVkW|{la7R4&uyPy`nLu7e`1;I?t@EE#{`~XLnjk%Q zo_7a(Mn;m+W*N2l;;B_(xg|Z?Pc(4Zgm9IOKxV=EMbj=qo_P%*=2@wPyt@m9ll~o3 z#npb7!PITiv5}M@ddg-(E^bBB*Q!0cmJoQa z^7HRyXJ?;?*!c-egq(Rt)dnKwxp#CnSZG9yKD|p^|!cRR#YXPAb@@g?JoNMoic6UvskIc{yfUE z3BPk+!A$iL$`+iJ)260z4B=)m>AfBv9w%wv(dcUUlEsJAaTp@_Jt!1gz|Xxh`K<_d z&YR4sf^TnGzryLo*RRGzAw$H&PHGc40fa0~1I${HcjQ?OI2tpcG1|OSW!Ilwzp$x+ zS%A$VDvBg98XhPWW(R2+BvA^!?eb;}tP-DSwZzPOmwQIHkr5u!tnv_)I9o zd|j3)KK6segt5S)TSEZ%SOHaH0mYHw9h&8MmMsbiIg;*i=zvacIt7cVO$smCg=dJ4 zcu4N#n`h^UlNQ7$+L)d6f<%!|E@!BR;H@Xai%?%NKX;#4sl9rJ;z^ z{D%$<-vga@?C1BNFI>7~F>h{mr;BXm3xoz$AGP1+l}XDJOX`uW{tsUn7gX!$`|ShRk5E*)99% zf_N3cnZ;Rw47Wi*p%E*TmIn4>X#5wSm&W$lSiNk4~}A<1Qm!zVx+F75S^Jf8FSBKRWn z#v)wEj4T;20d}j33c+sN16zXA5KR^dAgl4wA^mia&Q$FY?gBJOQJ5?oNUN5pU~IpxPy&9gD@-J7w z+(yzDKK{OL-5*qii1Y$*uG2 z6+V!f0gVGTaXrGq!o+*IzG`HV_YKh81pJbbO?bcwkEFpH$V-w84VHgt+Wa*A09lPr z^3bTj25o^f+9ivMRZ--UqM{<|beDuY*C&@g$D{25X!mvzC@ErliVSMk72tePiM?lK z&f#a@vU76>5N$zBT~woWO8j*QAAr!pbaNqjVSvNenm1qC-prFI(fFcQJ~ubVg>>)( zAiu2yWEMrIu$bSMScefM(n*jgWQY%(LflA7%X2zeKk^BNDMX^VDI$Nx}PS!>vtTEt?0n`F}S=6z4$Ic!KPMZbXpHQ3tp6W zU0ui891%=v=tzoTATx@DuWb7+)u<`g5MIUfK*_hZ=YP*o$~P98;)g+V{Rbe;$~Bfurfa2=r-io8PVzJ5saZa z!Mjj@a*cQ-?h^Eb>c30x%CY?{Ifs##?rlDFg4O_86AkFzHg|O2DV=d717CT9sW2Q- zq9F)aw5ot%u0@xZZdA=dNI)PO$f%+4QCaWHIOH?%XTNjI;4EU1i{mBm0VjPy# zyS>;iJ5X2TeHQ#UnEEg#y|K2pKLVe{z1OcLyI}$m02}1rtxgJUmge#RoV=n+h=_NQ z9$Squp{$5udZv&!I(k?$)7<10V8LnhRrmFM=!`1qf4%1DpGNXOe|gJ4bjg4ABBksK zb>6_E;*iVvs4JE)$6G)YWo2a_G==E_hmdjNs<2~W$=`8yX%NsOsSj6VAI2+$Ruju6 zepWQ2T&8VrdwKB&hBP8Dq6S^MG^(<>Im>S^KZb8z35Pmf)wkOujH7K zes#(g&X&Vnocbx#g)y6}OcRBtb}U*`^sL@GEaaUO#BakcYM-45 z^Corze`(t9bD)5c=JbZx)`1GEEfjBb$WB28K%7hk3 zIOMuyw5-Hn^vlv!HGsou;%bnz36}wlC>9hCc_9egN`$O{8MHTc7_Y-P0rqPExi)i% zSo8>u1K?W@%NmgeaPjH8o5P7H(ccAc6!Ackx$sa}$v#>#r%F8+*Ak5ZCy2~jlef;0 zjvM>c1Vl*+G4kJW1e zGG9d4KWwRS{RFUsX0_XM*gMKVm*AT!=9c}+!lJ$$Jf!FK@0jY;gqM>z2FTk*uneSx zkr!(aUUN^_RVAAF+ozC-&5$qf`WkQiPk;SlJQ{i12cUKx7!W(vdPKTE^eLQd!?t7f_yKP*fpT71k>lD zAj$3T_uz!-C0l^$$!ZrocI;S78?2_o`a*)6}e zNK4C`DzxZ~6IcG?)s1)VAe-09&4`MLb=ekdn|xSK{#m*U!xUzyeR$6ap{hYRK{ne< zICnB?PTneHplv7##Nt)~F8H+rGaj?`K{;s9$c!9$^G=qy9?&_ZKPo;7lpdLG2gsDp zbzj-_=h|KAP~`4$h-(siXWHGe_LT$0;xEbkvyB*eeaY}K}v{c+LJCp--F&P>C84r~pcQ zGGS;F%^)x4JgM&J=$K}8gIZl)H;)3*k5aC2e*wfREs}sHVmAr~hPAe~PBO(!HUE8Y z+*n>!ltBn=%&97nxIh`CVTIZsjgLnNP6xsDV<(^o!vHZhkjU{BpLitK2302t6-N^- zQaxUnv{c^uL}LOako1VV yi%1ss{dG&UIaye^cKk!4zqbGX-ni57!XjBG{??{dsx-1pRH<{Pl1^N>{@(!G@GKVq literal 0 HcmV?d00001 From 23f9126475473cfbf6eaa9a3fc940b4df9c71f99 Mon Sep 17 00:00:00 2001 From: Sangjoon Bob Lee Date: Thu, 8 Jan 2026 15:58:21 -0800 Subject: [PATCH 2/4] 'tests: mock h5 file and compare bin, load with numpy --- tests/hpc/__init__.py | 1 + tests/hpc/test_io.py | 53 +++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 54 insertions(+) create mode 100644 tests/hpc/__init__.py create mode 100644 tests/hpc/test_io.py diff --git a/tests/hpc/__init__.py b/tests/hpc/__init__.py new file mode 100644 index 000000000..cb91799a0 --- /dev/null +++ b/tests/hpc/__init__.py @@ -0,0 +1 @@ +# Tests for quantem.hpc module diff --git a/tests/hpc/test_io.py b/tests/hpc/test_io.py new file mode 100644 index 000000000..6aee15851 --- /dev/null +++ b/tests/hpc/test_io.py @@ -0,0 +1,53 @@ +"""Tests for quantem.hpc.io module.""" + +import numpy as np +import h5py +import hdf5plugin +import pytest + + +def _gpu_available(): + try: + import cupy as cp + cp.cuda.Device(0).compute_capability + return True + except Exception: + return False + +pytestmark = pytest.mark.skipif(not _gpu_available(), reason="CUDA GPU not available") + + +@pytest.fixture +def mock_h5_file(tmp_path): + """Create a temporary HDF5 file with bitshuffle+LZ4 compressed data.""" + filepath = tmp_path / "test_data.h5" + n_frames, height, width = 100, 192, 192 + rng = np.random.default_rng(42) + data = rng.integers(0, 1000, size=(n_frames, height, width), dtype=np.uint32) + with h5py.File(filepath, "w") as f: + f.create_dataset( + "entry/data/data", + data=data, + chunks=(1, height, width), + **hdf5plugin.Bitshuffle(nelems=0, cname="lz4"), + ) + return filepath, data + + +def test_load_and_bin(mock_h5_file): + """Test GPU load and binning with sum/mean reductions.""" + import cupy as cp + from quantem.hpc.io import load, bin + filepath, original = mock_h5_file + data = load(str(filepath)) + # Verify load + assert data.shape == original.shape + assert np.array_equal(cp.asnumpy(data), original) + # Verify bin sum (factor=2) + binned_sum = bin(data, factor=2, reduction="sum") + assert binned_sum.shape == (100, 96, 96) + assert np.isclose(float(data.sum()), float(binned_sum.sum()), rtol=1e-5) + # Verify bin mean (factor=2) + binned_mean = bin(data, factor=2, reduction="mean") + assert binned_mean.shape == (100, 96, 96) + assert np.isclose(float(data.mean()), float(binned_mean.mean()), rtol=1e-5) From 52f09d9463a019c0b3b7061c4cae007e7e47ccc9 Mon Sep 17 00:00:00 2001 From: Sangjoon Bob Lee Date: Thu, 8 Jan 2026 16:00:12 -0800 Subject: [PATCH 3/4] remove benchmark data --- .../hpc/notebooks/performance_comparison.png | Bin 21029 -> 0 bytes 1 file changed, 0 insertions(+), 0 deletions(-) delete mode 100644 src/quantem/hpc/notebooks/performance_comparison.png diff --git a/src/quantem/hpc/notebooks/performance_comparison.png b/src/quantem/hpc/notebooks/performance_comparison.png deleted file mode 100644 index 5d88be09cd06f646e9e6609130464c91fe77434b..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 21029 zcmeIacT`ka_brNct8GAA#Vjaop`eHX$$}VAiy&Dts7TIPK!OQdP{ER*A|Qz5AUO#b znk+d{wVftt0;} zG7781cPT!go;t2_<>^p|lXd`OVc~nWYLAz#yq7I^@_69l=8-iXPkr}DF3OSoW>eTvHRM>u z;Ap0A=i?14NS;jD>~--s+y#CPKG0Fz^z)aVG~eIJqvI!P*UiPuUwHQah7}+`7er5R z|NPaX`~RhvRfWf_WEp02cXyYTl~rksQ}xMZV}8-Fa7s#wt&>x8OPbM`Kxlx%>xH04 zk7CT4lQuYPVP4VxAW0`L>aV~4y6S7h#rAE%Ajz{e&&g`78uR8~hwM3c@LOha+PeQ2 zp6cIgWXCWudn?gVEs{&G1R7J(VDKPsrjT}c2tHj*rGhyZs(h6*`etd%`bKK z?4P_FeJ)+P~xhM9v*JZFbfhi|8Su`us&2_n~-qNad|m{>)1wLGp$)^%g3>g3pR zHTPwtr725FdY*eK{6N3((WA{8yc-j>vd+;|o}Xq7c1yDK_xF!nzJ7O%``omxgF`@2 zP|$GRh{S&F?1O3o2AX+}iFnLE{`kXwe%4uh{zv!d*x19oyu8*-i}-ungfwT|Ct8e5 zKfDar*3nUYA-w^Ac5cg-EiF|ulU*KW4Ka#wYRRD)1v4XsdNM(LmEYZN-@M6v=#b)p zSKqG3sl@)FrluxfQg=#GQL(9@%Y7zMGu@ce$;m0;fS&HV+iTQpZEZ6Jo(kKXIC*<@ z#{6J;z`p30FArL@We!CN{8dsyzqLtM;oeq(p08hH*6up}^|4XtgQ82dcen7{&hpZm z68DLTiLr~j94V1tF88yqY?{CI_V&Iw7a~ffH6>`+Rm|g?O`DTK{@ArkHT=YqTeoh- z+H@CB<9-=fhJm-YR%Siz&S=)HmpPCh86_8N`p%0rPU(g8eK+x=M~~o`kUOQZad236 zben#BeXYrT`SRttBLy=@90%1yj(O3xDXUBh$I&+2|ew(ANGvAfbZda}GE_~UJs|rB}^qPKr?-nq8|A^O){pb~i zSBE0c(4Ifnz&l-1=PY{6PqR<>gG>ZbM_I~?dp%Z}W#l9DsnoB&GYQxX0) z8!USLn;Xj*uYXw+V_5klXYK$7ox-Td5)Rf|)(KTmmSX-$Nb z4_UwZOp7+b03PMtB`G*SF4O(d0{X?PGOW6EkIjC&aa3H~%0)s?Pp>jWOpTM1Q`Bkr zX_V|^+R#8kv#zV0nB!pS;gK+p`Ph#iFUSQ8%B>L}_sg~Kx2>3b#B=zt(%dwoev{tS ziyhm{+Ouu89&>-?GC7=x6KPT%D&C~EadN0Fy7=BU`nxs9aAu$PeEP&mm6n%R!{0i5 z{%g71lOxmv1wSGN2b22l-d?lgys>c%i7-H=F-f+e=Zv6$je%~h9d#15_jU(3PAz$j$FieMqTamGrdDY7WGZJ4rUhZ z*@}5iBUP!<>W0FMzt+q5O!F^e9`#%CHRn%In{RORPnU;jP0^DJDbzI*KQ07jeyI%Z z+q_pN-$m7JdYptxL%yrie!YU(z(p#ST)TIh;uxf+rPW7=xtsD^A#==+g)N+=MV$+I z674J#u=j;YeN>Bi)%@&Ob;YF=-jK5nkY*3!~qGCkhk#3D^wx@z-IY%aURRdobd+EN9Y zO3EcE@re)mv`9UVxki_(L#-LDX-2+r=M&?Ex~|rLIbZzko3+D8r;b{Zb}hE{!I04* zq-)+q{uFUWgEB2O$;;cvryh&1nrW_q3&q<+7I+RGJlIz9VAmO20S@Le_2`5r-8{VN zBrSC?#Xp3^H!}1=Q^rUesaaVE{93t8O762u zpF4Lx&9IW~_uu_qPb+WBFn*|c| z3Uz4>*lqK|?=SI9nyHuXWOUuzs}gfzy?sZFTC%Ro{8XQS$NW{f(4!ibmX@^0qFN=izp>|>v9?=S__boue=Hlt{h9-A=>mrmglIi<|3 zekv_1D+|x+|Mcn8_(t8wdsXB7hrYjW8%#B;HkXl;%i!?2bLRnt?5kjKoRLI%;QoEo zmf>hs1%dVf{_>;l0}Zjf7OL^*D;@7ra-<8Q+G z?=kQFQM7{E37id?I!0TWgH!OEJ^xx#VwhYc&oQT2GWY`*I}K)F_Zb;ppfn z()lEC$A*39)0-ZIR0JKM*S(NIUaqut)wvj>os$|#kH0>j^Vs!bPNR{3j^@7T05u^z z8#_svL1w1E7$d0S}+Aqse(~*45fS*WRd3o`j|D#977GhVvJ?p$W zCP%IN_UWCSTfucynU3AF#n*?TOL5_P`c}+B|$D^=~1}!rCB%nNR9PXtIbRdF(%rq=s2><3JQLg^!2e| z`??#&DZNgKsy`@+fXI*)xZd2@YtZZqwX6JX9R_X73IAdOm)9wn_KO@xq1qJYW&G z3K4=kcX-Sy;)_frp=HM2!zv=s1^OayQHGIS?$4t z2jqK^#iB^^Foy(wW#y_>ynoV|u*b8E zxA@+q%%Y;#2o?foDj)NgZS;xudeaoOC&;)rfPCX?;7s2|+%+eFgzx6wIQ-OP96V6@(7xNE1+g^t% zFn`5zke%}t3|EzxP3Fb4gytcpjpI>(2I!^J|p{%a+fU2rV+9GH}=)2HY zl-S(4b7y+A;s9gBV?hh;L*R`YH!43odi3ao>zWp-ynHhO49}kJL}a%$5HgA7b)_(y z!y-xQk&ADCo%{A~jWdVS*BVaN?ZP!Ty;l^Iz6%KzJO?cjP~HTA*3^tjt;`r`^hw&u z`~W#T%kHu7-xamBwL?opgoR6silhJ}4u!|fJ{EIIV&otjAZZ_vm#>c4$E14{%&o;Z zE?n2mcQN_&_LivQ;3LQ3_Q!M^EN>&ZL)LAA+!7MHpFVwZ799!PuQiSv+jr>Dp%MuK zc+pb=5czlL2jay7fC@ziMVKx;-k+JLn2J9qCkJolA|J5V^QwCnWU zQbv1YyjpeYNd5ICD-&jp-@1A8EbUyFL}H45$pDJ?c^rC@WJE+o%iG$tkSXfPn%VVL z8MxjDgw_BYiPOx8dm0kb`{hd%ssoR(u==hu_ZtwzTsnI93{uq1EsGvTZWS<9r%(0u zxev7pxP33$gyN=}W7`X;5qa34{3}|GhGspF5R09g)~>yn)~D*?l8fpQb^6}cvc|>$ ztp&0bJAJd+*Kl*r&do8fU}hveq#54Jbs9-3Ej^2VR)!XMKraca-jAjiiLTCb(Jw1r zdAZF@n7qBYT;sl8!Bq`Xi|(_F9dVs7PXtJE7-&eztVWw~V<~I-VLm~@^Wefr#b==M z#ACmlllsxWkmUooHnf`^DO_m6+C}KOPf4LLH2^avhN(JEPmUxnS-G)iYL?unC;hi@ z&vkFBs?OoE7r3LiBA-oWNSP!J(YR#N-RE4;J@x~6k zlEP3208X$&@VAa9>*nVI9z~+4@7up$jxYvD{yp#sRlb4S@aNAF zsMBSYm2%PNLh1ls{r&td@H+ukPoQg=o0-(|bsA`(G7+C`Z#k8QT7$ooW7+Wg@6R+d zO#6|0M6Ex1QcPM>4G2bL^hbw%`LTsp{iP)Sk!fpsEDD0Hb}ci|OF1~%6s`*k3maoE z@N8-TPdLs|PXhZkNIm}a{$5rO4%i2C6+SAuBt>)c1hn-ugZs!@ zs%RaPC@9ye-oCzSu}Uu})IIW{=@Y@?u9+Ne;39(M(uf<8F)s#qub zjwihvbFfc5qN1(SE}#Fpg23Cm>A)*)Nf1tC9A5cwu;zB zR_?c&>~PX6c=xn@`m@i0BbQ74c*B()KX-R0R-Xgd&-^ zjmT2PskWF#zLv#z)S@L$z&*VnqWVuP%>5Lsq_6avZRX{evKyu1cRdn;_Mt!rL6k2PdkXi@8)ir5{o|Ei#T zJtp6EvNgeStDyNKa$tA3|2U46sC4DZ54jSsL0^Nd3Tp4H-(%lfzG0|6N11SDi9>+W zjkr27*{VHk?;aeXG-o%U<}F?iy7j>uh*TLEmFx&2cVR8|Lu-S=_1XEE5f`9OwoRM% zxVpN!0hVq(c)`qPn^sJb?KUkrUSbEAt51O-amd$P%y2hzdtt2eVV%F=r>Qt!4Kgm7b>Blp&)-QEy~wN zC##T@l!U4uNFnS5y*W7upjpte-B>&{1Of6jr!Q0)0k4{@D?at($5(8RpzBF$v(Uln z;T#I#xlnt+MUvK*ZSC4{r~z^Klh#xNA4)$#cqly$**5ylSFU^=Y)K=%H`>1y1Nvmm zS$aHwUvcmaa#6yYl4Iykk8k4BF3sF8ARtg;Ln8o^K8B3knywnFB(+~Ny(ZP5f@}a| zq^m2cqoP7)pITA_5NPs^WoryIxa{+ef=n}fzd3)R)vQVXF{cue1t9PNlhK|sp^$F$ zv30o)^i&IaeYEcrK?_s24sc(DdF!hRc}M{;JBQ**yWF8xektWExp{z$4E>VTRP0Vr zI(R7FF}pl=a^sJD3LO@4-9&lcxL-38H3of{aGL}!YpOoBkbpeFEi9b<+HA9csXzYR zi)jYE0(k`78sRZeFQU*n^a@*|Q6=FEeEb|)9JN^$h!otYGMFKrVH z+Lo7;>?skCK)3!3k!%9)8K3ZQc6JujCHJOHo3KqfpbBe&=8`{r_^^16m)tE&Ca;%< zUceGI{%hB+dHaaR3$;#G)Ms-(P?n2hq7qNBGZ!vAi=^wx2JHCXQ4QJdoy+p^+-K!R z`ICP{Z~R-@@)!Q!aVh?%i;e$(Q=@*K4conZs=k%2S`7vyo{EZRIj>3SEcU!~J>WlY zV#p`3{Fk*v|L4>G$6Cq%>Fd_wrUOYmWFv(OX;xK8W9z!4=T_!y4`eeHf{iT)6?{1#@FY*8!tNiynP{~ zL#y6#5B%ZB?^E;#b-K@&XVDVm1vzb0_4EPmAVVDe3zbS8ot(@P(N(W{ zD)JoMNIaS{TRXdb<70g!|DP**5jVE&7@GqRK3p>8kjL2zia{*T{C-`kyWrLZRllIv zPE7ZZfss+PQ0JAC;30Fbe7_X9@)k#!#5R*KA>^6|NdDD4J+t>P3NsM_@@J)(o_IoH*DP+ck|}21OpMceP93d z{q5wNVlBwGLG7ek-dHZC?T#+uPgi!8%M$O;K|>jERN`imVw@ zUGK<9m_a+*pBEs6bR6hS0Tcp|h;#!mkzrW5wF5*^P7>zv2#1eTml7q_2sIv>Yz*B)z>SdiKz$C3C(UtubDtix=DQk6pzLj#gmY@ z9#YW%O0*qfDodwnJ#>0}bf}9oZ5-e6rGGXjgboIDx;IF#XJ{x9?T*$j+|+@Z@EhQ- zLZZi{n}w7rhHZLE9~`az$@$pc4p1oi3*1FwwX(9Z1Ms&+ORALBtM6Xu*BfJ%qqVDh zp}Yd88;R(Gd8RfOw(yC8k#y#G_^`c-u!Gcr4=K-q?o#hM_dGar#AAN;XvmCV`?}9R z9GOv5y5!x#(N_X1hgUuJ(O$8QQ{~Vhr{Q+q633y|7m)I>d~zPM-)~ozM+AQ#8%xtr z^e@9I&{8{|YSX~?J?%RrQZ}Uh_?imfYHCO{4z#mqvj3PGbt#i!GY{<}J(p*m6vziY z-B&RDB<3<6LovwHQ=-;@7{*#JKL8;e?TX;g!rX|5PXrS(X45GY%2`@@&;k4; z|Bt}ScX5KYNh3km=g#L(Y*@c`?b@b3 zh)|6nTr-@<4ADXpyAjB7kJ=4Jsc4|$0eq)>9?GYqVJE^^w0u|-B7({Lk=7cTk$0{v z`6BKV%T;@_mn2`#`R-G(>fJ$y8@4B6E?w$;<=oTAT8n)mJfr%EN~sg zWTZ2HN_3MN>+luu>tGA@0=vzc5^5p06H*-7N%U`ff2smGQor32>qL{mws~_5xC4iw zR#li#l)$177cRJSQ7FbZ9wfOb1}$gaq70PS0XzmF-N^|N5fLFuJINwQha0wSGaC}K zY!^le3cigLY}%9%yC?L@324)qkHvQHE^Zw$1i5-SMg0h_;_tk4?Seda`*%ol%(e>s zbZfTFbFd7Mk`tjd_O}>?se)ee&OJrG-ZQequy;f-@vR_A$+gN1o9@%MZr`S(4@A{b zhNpnsk`Ac1woHq9u!2Z`l>IunDo~i8IcPsWeUFWejnWT70~Q1#m}+KQhZ9W>mu^zO z=d|Jk?gfSs0q4=42)>_7v<=E;(=Tg|KJ)iyBb%|6-|#N-a^euo0MqF01)z?H9CfZU zG>7_PaehC{EmB0lMrTP>b~CKgEr#F<)~xQI^{eN>#mz*>Mbk=_fC%fmJ!wfEd2aAE zkbWB7U5CBInl)?OU^Spcg6_DyJW+b_-%dTxqlAJ3HLwO`wFb^rfr%ahsxet_A9Z9q=dkUAvy5s;qQM_<2~edC{%1M?Z(MBgO2_W9)v@Jq0-XQ?5@3c-UE7->9^aXue=4GBok!B(&v8a++;>|x>*{vX&Wog8D zl9Q9uG4aZ2>-oRGoGq1T9hXv2LQ8Q}z~T~#8BBG|nq#vsV=DNU>si&i#@R(iaY;x-?$P+6*KhR1I@{L zp(RdG;Jb_O- zK)!wfRx=j)mM~_i>~;eUJo!I9`w;&DsxrJA;a;rUOy55`M5JTth_DlsDgon~orH`) zm)=yp6ItUdoot2YEU8b6fBD5>w1=9{XwniI+O}gyG8n5b^rI6K6Q=c1_n5qC%Y=x{ z=bMfro%>-aGV2ms%gJWFtPCn6DrzIrF|nl(kpVe<-MV%5J*8`jwv1AFI?x75K@ruN zfm3V-djYWwKz}7}8#oz=EI}%|_eXZ>zs)t|c9a)x?G(a8S+c$XHR_$t)NqS};BNw} zz8{l`$kduoD>4?J+Xwe4U!2CpvQJ7m( zlpA<#LA`g*Bl(l2(17Dq;{zG$wzjFz!p!)Gp`@hdbPPX~m#8Ixk4hxs;`{9TsG_z# ze@4lNB|(tnl(9&DHR??;Gi+yuL$cm6nK0atW}A+V z&OQB-`@N>a5Ja^~V7u9ZeAAR!A1$xdM%jWTEI;gC7fEfJ0LOSl#uZR0LlEk^K1@wH zab^&R1pBIn4;fW$e|K<|KFHx>EQ|#-(6=#4FT6i}`C=$6BC-W0{drl7h=n)})B@1Ke&lSOD=&|YG&*$aMx zi*6lJa@&eZO3cQ;))IXWmNN$&kvRD<2@Tp1l4ZqoJD?PLqNZkI>T#Py1S7i#i@1z` z4M&buAkhKWLC@e|0<;`0I+4OR@~ZEonzm&oUcY{Qbc(0P9B2Q@F%S2^ZX~cqL=ka_ z^n>?grp&?{VmO3;{Yq8uXbDp=h|*2RKEUwC109qAVasYva+I&AxV6NZnjsi;1#7GX-}{Z`A(|R4<=nGp4{`cz6}68+;on0gwn!u?o=uyAwDU($Dfb>d zs!bGbLiCr6z;zn1U+X;0ynwT&tc*&`wm3bzX(R6zHI76#RP0vSnH`4e@I7b&Qw8 z&{!3buPJ39c#(5x2r#0w5UUN6mdocmdw}_LkZDBQs&|exm7(&)VtI*&4+KGy1XF9` zwi(U$TETtseJ;BdY6U1+Ua zfVK{cIt=(5T(!5i&vY1WNh|rrUAW58i~+ zXrBYSiCzudXTsi=W+ZRc2CHgf3H*5mls$Xe?&e#6y8a!Vo?%X63O=XF2Eo_Tv>iUn zp>DVbXW<1E90QJNY@M$;xu)grTwmWD{DYgZ%xRg?k&!*)iUdbF@miMb%ULnzcvNic z^L`c@fzn!GZHTr!n<;E}+IBj6>~o< z^(Z=ZP$AOZ(+3BiT7=;EH6yNjy1So4E}H{}13K&OBCd~5pKRzt!SeY0W@I$tb0xKh zkB<-h3vuDV0+~3sn&YUkn|rs{(D(1(2^CNHe8}qD0s^Xptw*7?IIq=v8Z8+kUxw!6 z<3r3q{ge7WW^2j?6U2oGqBJ7|HZ>C? zFH&Xd_&yxBKz&GKim;ohz{-G{RT(u@Y*k7yOdy|*Y5tOt(;I!bDs1%hy0+6 z9dd%5xjM7pIt?8oq>+KLJz))KRUC*iG?!+{wqu8tTW2J7Jvf!J>1h-pB`CxVFE5^S zc5#V=HIP^&#hP>pH1>|g^TgXw z1MNle1Xef#!pf$cdq3*7x4ddqV!yjhI;uZnY}2%Q>M zL#zDR(F)?Ujim3hLwE8-KOacAYF+!PEr(A*pjE{&wf%m7Y#1$1MoSyu3l)H$SnlB& zC*~XUj%R6=eflZt<1onek5wOIO67Ka=jgjX$!Dm$ml z)EEztXJFd{E?;&XFj;cJboumjsOxa;jWrUFEcB2Uf-B#0wLK$}JH}tq3#!^b%$(-n z;2@Dp)F#64pqlu@lSuN!P(660?Lx&ILY7aN1|=n!TQ?ri(UR(sH$$HAy#?%ub4^qYze~ zNVw$+9N7Zh3|{(P7`M=Vj0hP>AXd~8d5TeGuuwAqouOL? ziWDm2c~XCeJMz-nqaO7m9lGFK1+MD_JO^4{VV4%EK!gAudc zLMBQ}y(t(#cja60=N)$|B= z(evfYZa74j{QB!x)RJZ-*3@?IJyb+#9O`rg;6PbNhYoUfoZ@q@rK`8c5ZM!QsmQQA znRb|ieOwbn90A`JVOUxdAb!$v19jOkbZskIS~O7e8xaEk&>vl0dJuwhaUSeeIRgnF zF>B(5=}Q5HM2*;vK7QPZ8zN>(#A6-OPvZ-jO~E$AYW4mxKPb{&a1{5btD*m&80kt@ ze119>EH&Y`LA#Q!l5~ymj$#Dm0dT(VvoL1s>aVhKSKf@aXP&~p$)N~7CP6yPjee>r zZw)x)ilA@)I{`Bd-NTz;P#$yOOJar#6z6|*vyXxl4EJCOh))U`4lt1yNsPn<$TRP| zcU6GSDj&T+a5ne5BG@(3ED-|+F0h#PMXv`NCN308-ip6eSc za7Q)V@Fwr@vcK(r+`p&gimNqn^_2US#x6NrT@@lmCOb%tB@mrV2E?euDhav~*hPH3 zWLA3MP!rMqJNPziYU+E0HYfkROB8VJdWa3zD|;= z1gJ0$`gJYa0aCgIGKi&}d8r>K;JroW9P2v;xSvReh$0Omg7Mdy z@FiOIYyTM-g(zfs>56cc-&y!YMEHMW=#;vHVOh*bz(*05O;~eRtwJYZmigEc6ik;`MLk@gOK;m z*4EaNi7^h>!IZ1j{08N0FdNJT&VehT02fn5CUR0LJ37R>1Rme|*LjLTcSVb|Dx4Fn zuy8UYcJsfPnlf`cJG+OzzP|key&L{@zl1{0?F)5f+qNyf9IsPeVd5<+Sx-Eh{Y&{dfDs)cb!~SErzfD3k1h>>adQJ~RPj2C=#jZVOzXJ&2V# zY%{S!BZm-=UYlqa3L^u}lPS_JaX!Liyboj_aX=AD0(h*J@ZKgwX{ zbwXhhJNAW(gDMKB3LVA6l<9=p? z|A&*u+1KP%3;!On^;~!7?}_Dhh<_;()upA}Z>1G#;OS4Swwe_QzV_e!^#dGd3eTZK zDSyTvKMv1d5Cw}7%yB~Gk*FA{DlFV(E)@5K9)RF-l8ls;d`}pKi z+Lyu!25qGRk`b&^QC;49{@y)LZOrhHszLDk7B9|s2fcGCJvWLj@%Y#Dhy8svB||r; z`d$0%!~PfZ=Klu zEJf$RLX}>&$r1+DL|K~~@mIHINykk(yVLq4r*^sFNO}jZhvGBEA^+QwSr6c}0J8O2 zf9_lKe$(d7+;`~gTsxXYhj(j_WCvU4cUvmalAI>~CU5_5p9&+z2*996_&%Mzr~^*kzb)wyM(9y)Ih~5zk)_oa!o>4R(2%CGt7#`Y z`G9Iu+_Z@877ikru|aA$k22nnq$5TODSCyyTA2xCga80_qr+p#P%YLc|S0 z3*VuPLu+1alm9%Mqz+KOs8TUVkW|{la7R4&uyPy`nLu7e`1;I?t@EE#{`~XLnjk%Q zo_7a(Mn;m+W*N2l;;B_(xg|Z?Pc(4Zgm9IOKxV=EMbj=qo_P%*=2@wPyt@m9ll~o3 z#npb7!PITiv5}M@ddg-(E^bBB*Q!0cmJoQa z^7HRyXJ?;?*!c-egq(Rt)dnKwxp#CnSZG9yKD|p^|!cRR#YXPAb@@g?JoNMoic6UvskIc{yfUE z3BPk+!A$iL$`+iJ)260z4B=)m>AfBv9w%wv(dcUUlEsJAaTp@_Jt!1gz|Xxh`K<_d z&YR4sf^TnGzryLo*RRGzAw$H&PHGc40fa0~1I${HcjQ?OI2tpcG1|OSW!Ilwzp$x+ zS%A$VDvBg98XhPWW(R2+BvA^!?eb;}tP-DSwZzPOmwQIHkr5u!tnv_)I9o zd|j3)KK6segt5S)TSEZ%SOHaH0mYHw9h&8MmMsbiIg;*i=zvacIt7cVO$smCg=dJ4 zcu4N#n`h^UlNQ7$+L)d6f<%!|E@!BR;H@Xai%?%NKX;#4sl9rJ;z^ z{D%$<-vga@?C1BNFI>7~F>h{mr;BXm3xoz$AGP1+l}XDJOX`uW{tsUn7gX!$`|ShRk5E*)99% zf_N3cnZ;Rw47Wi*p%E*TmIn4>X#5wSm&W$lSiNk4~}A<1Qm!zVx+F75S^Jf8FSBKRWn z#v)wEj4T;20d}j33c+sN16zXA5KR^dAgl4wA^mia&Q$FY?gBJOQJ5?oNUN5pU~IpxPy&9gD@-J7w z+(yzDKK{OL-5*qii1Y$*uG2 z6+V!f0gVGTaXrGq!o+*IzG`HV_YKh81pJbbO?bcwkEFpH$V-w84VHgt+Wa*A09lPr z^3bTj25o^f+9ivMRZ--UqM{<|beDuY*C&@g$D{25X!mvzC@ErliVSMk72tePiM?lK z&f#a@vU76>5N$zBT~woWO8j*QAAr!pbaNqjVSvNenm1qC-prFI(fFcQJ~ubVg>>)( zAiu2yWEMrIu$bSMScefM(n*jgWQY%(LflA7%X2zeKk^BNDMX^VDI$Nx}PS!>vtTEt?0n`F}S=6z4$Ic!KPMZbXpHQ3tp6W zU0ui891%=v=tzoTATx@DuWb7+)u<`g5MIUfK*_hZ=YP*o$~P98;)g+V{Rbe;$~Bfurfa2=r-io8PVzJ5saZa z!Mjj@a*cQ-?h^Eb>c30x%CY?{Ifs##?rlDFg4O_86AkFzHg|O2DV=d717CT9sW2Q- zq9F)aw5ot%u0@xZZdA=dNI)PO$f%+4QCaWHIOH?%XTNjI;4EU1i{mBm0VjPy# zyS>;iJ5X2TeHQ#UnEEg#y|K2pKLVe{z1OcLyI}$m02}1rtxgJUmge#RoV=n+h=_NQ z9$Squp{$5udZv&!I(k?$)7<10V8LnhRrmFM=!`1qf4%1DpGNXOe|gJ4bjg4ABBksK zb>6_E;*iVvs4JE)$6G)YWo2a_G==E_hmdjNs<2~W$=`8yX%NsOsSj6VAI2+$Ruju6 zepWQ2T&8VrdwKB&hBP8Dq6S^MG^(<>Im>S^KZb8z35Pmf)wkOujH7K zes#(g&X&Vnocbx#g)y6}OcRBtb}U*`^sL@GEaaUO#BakcYM-45 z^Corze`(t9bD)5c=JbZx)`1GEEfjBb$WB28K%7hk3 zIOMuyw5-Hn^vlv!HGsou;%bnz36}wlC>9hCc_9egN`$O{8MHTc7_Y-P0rqPExi)i% zSo8>u1K?W@%NmgeaPjH8o5P7H(ccAc6!Ackx$sa}$v#>#r%F8+*Ak5ZCy2~jlef;0 zjvM>c1Vl*+G4kJW1e zGG9d4KWwRS{RFUsX0_XM*gMKVm*AT!=9c}+!lJ$$Jf!FK@0jY;gqM>z2FTk*uneSx zkr!(aUUN^_RVAAF+ozC-&5$qf`WkQiPk;SlJQ{i12cUKx7!W(vdPKTE^eLQd!?t7f_yKP*fpT71k>lD zAj$3T_uz!-C0l^$$!ZrocI;S78?2_o`a*)6}e zNK4C`DzxZ~6IcG?)s1)VAe-09&4`MLb=ekdn|xSK{#m*U!xUzyeR$6ap{hYRK{ne< zICnB?PTneHplv7##Nt)~F8H+rGaj?`K{;s9$c!9$^G=qy9?&_ZKPo;7lpdLG2gsDp zbzj-_=h|KAP~`4$h-(siXWHGe_LT$0;xEbkvyB*eeaY}K}v{c+LJCp--F&P>C84r~pcQ zGGS;F%^)x4JgM&J=$K}8gIZl)H;)3*k5aC2e*wfREs}sHVmAr~hPAe~PBO(!HUE8Y z+*n>!ltBn=%&97nxIh`CVTIZsjgLnNP6xsDV<(^o!vHZhkjU{BpLitK2302t6-N^- zQaxUnv{c^uL}LOako1VV yi%1ss{dG&UIaye^cKk!4zqbGX-ni57!XjBG{??{dsx-1pRH<{Pl1^N>{@(!G@GKVq From ed9a5e8742582c46d7e84f7e73136bd7f19a5254 Mon Sep 17 00:00:00 2001 From: Sangjoon Bob Lee Date: Fri, 9 Jan 2026 03:29:45 -0800 Subject: [PATCH 4/4] Delete tests/hpc/__init__.py no __init__ needed for test file. --- tests/hpc/__init__.py | 1 - 1 file changed, 1 deletion(-) delete mode 100644 tests/hpc/__init__.py diff --git a/tests/hpc/__init__.py b/tests/hpc/__init__.py deleted file mode 100644 index cb91799a0..000000000 --- a/tests/hpc/__init__.py +++ /dev/null @@ -1 +0,0 @@ -# Tests for quantem.hpc module