user [email protected] usertags 1056840 python3.15 tags 1056840 patch thanks
Hi! While rebuilding the python related packages against the Python 3.15rc2 version I ran into this FTBFS [1]. I've managed to get it back into a buildable state, but it required quite a number of patches: - Migrate away from distutils - Add patch to fix hdf5 types - Avoid chambollepock.cu infinite loop, do not include complex.h unconditionally - Fix compiler-options args passing - Add cuda 12 texture reference compatibility patch (#1036680) This last patch was generated with LLM assistance, and would require review from someone that is more familiar with the package. However, while doing all of this work, I came to the conclusion that this package should likely be removed instead of fixed: - Upstream is inactive, no commits since 2023. - Popcon shows ~10 installations. - It already didn't make it to trixie - It has no strict reverse dependencies, only pan-tomography suggests it. I'm not filing the request for removal yet, though, as maybe someone else has a different point of view. Happy hacking, [1]: https://debusine.debian.net/debian/r-python-python3.15/artifact/4653058/ -- "Can you imagine what I would do if I could do all I can?" -- Sun Tzu Saludos /\/\ /\ >< `/
diff -Nru pyhst2-2020c/debian/changelog pyhst2-2020c/debian/changelog --- pyhst2-2020c/debian/changelog 2023-01-31 17:13:52.000000000 +0100 +++ pyhst2-2020c/debian/changelog 2026-09-11 17:35:16.000000000 +0200 @@ -1,3 +1,17 @@ +pyhst2 (2020c-8) UNRELEASED; urgency=medium + + * Team upload. + * Add patch to migrate away from distutils + * Add patch to fix hdf5 types + * Add patch to avoid chambollepock.cu infinite loop, do not include + complex.h unconditionally + * Add patch to fix compiler-options args passing + * Add patch to add cuda 12 texture reference compatibility patch + (Closes: #1036680) + * Fix autopkgtests + + -- Maximiliano Curia <[email protected]> Fri, 11 Sep 2026 17:35:16 +0200 + pyhst2 (2020c-7) unstable; urgency=medium * Team upload. diff -Nru pyhst2-2020c/debian/patches/0005-drop-numpy-distutils.patch pyhst2-2020c/debian/patches/0005-drop-numpy-distutils.patch --- pyhst2-2020c/debian/patches/0005-drop-numpy-distutils.patch 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/patches/0005-drop-numpy-distutils.patch 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,34 @@ +Description: Drop usage of deprecated numpy.distutils + Replace removed numpy.distutils get_numpy_include_dirs() with + numpy.get_include(). +Author: Maximiliano Curia <[email protected]> +Forwarded: no + +--- a/setup.py ++++ b/setup.py +@@ -56,7 +56,6 @@ import math + import glob + from distutils.core import Extension, setup + # from distutils.command.build_ext import build_ext +-from numpy.distutils.misc_util import get_numpy_include_dirs + from os.path import join as pjoin + import string + import stat +@@ -553,7 +552,7 @@ def do_pyhst(): + depends=depends, + library_dirs = mpilibs_dirs, + libraries=[ "mpi"], +- include_dirs=get_numpy_include_dirs()+ mpi_dirs+[], ++ include_dirs=[numpy.get_include()]+ mpi_dirs+[], + language="c++", # generate C++ code + extra_compile_args={'gcc':['-fopenmp',"-fPIC", "-O3"] }, + extra_link_args=['-fopenmp'] , +@@ -566,7 +565,7 @@ def do_pyhst(): + depends=depends, + library_dirs = mpilibs_dirs, + libraries=[ "mpi"], +- include_dirs=get_numpy_include_dirs()+ mpi_dirs+[], ++ include_dirs=[numpy.get_include()]+ mpi_dirs+[], + language="c++", # generate C++ code + extra_compile_args={'gcc':['-fopenmp',"-fPIC", "-O3"] }, + extra_link_args=['-fopenmp'] , diff -Nru pyhst2-2020c/debian/patches/0006-fix-hdf5-types-in-edftools.patch pyhst2-2020c/debian/patches/0006-fix-hdf5-types-in-edftools.patch --- pyhst2-2020c/debian/patches/0006-fix-hdf5-types-in-edftools.patch 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/patches/0006-fix-hdf5-types-in-edftools.patch 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,68 @@ +Description: Fix HDF5 types and buffer arguments in edftools + Use hsize_t for hyperslab offsets and counts, and pass the currents buffer + (instead of integer span) to H5Dread. Also fix dataset_current dereference. +Author: Maximiliano Curia <[email protected]> +Forwarded: no + +--- a/PyHST/Cspace/edftools.c ++++ b/PyHST/Cspace/edftools.c +@@ -505,14 +505,15 @@ void read_data_from_h5( char * filename, + hid_t dataspace_current = H5Dget_space(dataset_current); + hid_t memspace_current; + +- int myoffset = nproj ; +- int mycount = v_span; ++ hsize_t myoffset = nproj ; ++ hsize_t mycount = v_span; ++ hsize_t myoffset_mem = 0; + float currents[v_span] ; + + status = H5Sselect_hyperslab (dataspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); + memspace_current = H5Screate_simple (1, &mycount, NULL); +- status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); +- status = H5Dread (dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, v_span); ++ status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset_mem, NULL, &mycount, NULL); ++ status = H5Dread (dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, currents); + + status = H5Sclose (memspace_current); + status = H5Sclose (dataspace_current); +@@ -727,14 +728,15 @@ void read_data_from_h5_eli( char * filen + hid_t dataspace_current = H5Dget_space(*dataset_current); + hid_t memspace_current; + +- int myoffset = offset[0] ; +- int mycount = v_span; ++ hsize_t myoffset = offset[0] ; ++ hsize_t mycount = v_span; ++ hsize_t myoffset_mem = 0; + float currents[v_span] ; + + status = H5Sselect_hyperslab (dataspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); + memspace_current = H5Screate_simple (1, &mycount, NULL); +- status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); +- status = H5Dread (*dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, v_span); ++ status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset_mem, NULL, &mycount, NULL); ++ status = H5Dread (*dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, currents); + + status = H5Sclose (memspace_current); + status = H5Sclose (dataspace_current); +@@ -901,14 +903,15 @@ void read_projSequence_from_h5( char * f + if(*dataset_current>=0) { + dataspace_current = H5Dget_space(*dataset_current); + +- int myoffset = offset[0] ; +- int mycount = v_span; ++ hsize_t myoffset = offset[0] ; ++ hsize_t mycount = v_span; ++ hsize_t myoffset_mem = 0; + float currents[v_span] ; + + status = H5Sselect_hyperslab (dataspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); + memspace_current = H5Screate_simple (1, &mycount, NULL); +- status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset, NULL, &mycount, NULL); +- status = H5Dread (dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, v_span); ++ status = H5Sselect_hyperslab (memspace_current, H5S_SELECT_SET, &myoffset_mem, NULL, &mycount, NULL); ++ status = H5Dread (*dataset_current, H5T_NATIVE_FLOAT , memspace_current, dataspace_current, H5P_DEFAULT, currents); + + status = H5Sclose (memspace_current); + status = H5Sclose (dataspace_current); diff -Nru pyhst2-2020c/debian/patches/0007-fix-implicit-drand48.patch pyhst2-2020c/debian/patches/0007-fix-implicit-drand48.patch --- pyhst2-2020c/debian/patches/0007-fix-implicit-drand48.patch 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/patches/0007-fix-implicit-drand48.patch 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,16 @@ +Description: Define _DEFAULT_SOURCE in pyhst_nnfbp.c for drand48 +Author: Maximiliano Curia <[email protected]> +Forwarded: no + +Index: pyhst2/PyHST/Cspace/pyhst_nnfbp.c +=================================================================== +--- pyhst2.orig/PyHST/Cspace/pyhst_nnfbp.c ++++ pyhst2/PyHST/Cspace/pyhst_nnfbp.c +@@ -5,6 +5,7 @@ + //Contact: [email protected] + //----------------------------------------------------------------------- + ++#define _DEFAULT_SOURCE + + #include <stdio.h> + #include <stdlib.h> diff -Nru pyhst2-2020c/debian/patches/0008-fix-cuda-complex-hang.patch pyhst2-2020c/debian/patches/0008-fix-cuda-complex-hang.patch --- pyhst2-2020c/debian/patches/0008-fix-cuda-complex-hang.patch 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/patches/0008-fix-cuda-complex-hang.patch 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,45 @@ +Description: Prevent complex.h inclusion in CUDA and fix nvcc compiler options + In C++ / CUDA mode, including C standard <complex.h> inside extern "C" + causes cicc (CUDA device compiler) to enter an infinite loop. Wrap + <complex.h> in #ifndef FROMCU in CCspace.h. Also fix nvcc compiler options + in setup.py to properly pass comma-separated flags to the host compiler. +Author: Maximiliano Curia <[email protected]> +Forwarded: no + +Index: pyhst2/PyHST/Cspace/CCspace.h +=================================================================== +--- pyhst2.orig/PyHST/Cspace/CCspace.h ++++ pyhst2/PyHST/Cspace/CCspace.h +@@ -33,9 +33,9 @@ + + #ifndef FROMCU + #include <fftw3.h> ++#include <complex.h> + #endif + +-#include<complex.h> + #include<cufft.h> + #include"edftools.h" + +Index: pyhst2/setup.py +=================================================================== +--- pyhst2.orig/setup.py ++++ pyhst2/setup.py +@@ -464,7 +464,7 @@ def do_pyhst(): + # we're only going to use certain compiler args with nvcc and not with gcc + # the implementation of this trick is in customize_compiler() below + extra_compile_args={'gcc': ["-std=c99", "-g","-fPIC", "-O3","-Wall"], +- 'nvcc': CUDA["arch"] + [ "--compiler-options", "-fPIC", "-O3", "-g","-D_FORCE_INLINES" ]}, ++ 'nvcc': CUDA["arch"] + [ "--compiler-options", "-fPIC,-O3,-g", "-D_FORCE_INLINES" ]}, + include_dirs=[numpy.get_include(), CUDA['include'], 'PyHST/Cspace'] + hdf5_dirs) + return module + +@@ -486,7 +486,7 @@ def do_pyhst(): + libraries=['cudart','QtGui','QtCore','cudart','tiff'], + + extra_compile_args={'gcc': ["-std=c99","-g","-fPIC", "-O3","-Wall"], +- 'nvcc': CUDA["arch"] + [ "--compiler-options", "-fPIC", "-O3", "-g","-D_FORCE_INLINES" ]}, ++ 'nvcc': CUDA["arch"] + [ "--compiler-options", "-fPIC,-O3,-g", "-D_FORCE_INLINES" ]}, + include_dirs=[numpy.get_include(), CUDA['include'], 'PyHST/Cspace',"/usr/include/qt4"] ) + + diff -Nru pyhst2-2020c/debian/patches/0009-cuda12-texture-compat.patch pyhst2-2020c/debian/patches/0009-cuda12-texture-compat.patch --- pyhst2-2020c/debian/patches/0009-cuda12-texture-compat.patch 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/patches/0009-cuda12-texture-compat.patch 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,396 @@ +Description: CUDA 12 texture reference compatibility layer + CUDA 12 completely removed texture references (texture<...>, cudaBindTexture*, + cudaUnbindTexture). Provide a C++ template wrapper TextureRef with unified + memory (__managed__) and overloads for tex1D, tex2D, tex3D and cudaBindTexture* + so existing kernels compile and run without invasive changes. +Author: Maximiliano Curia <[email protected]> +Forwarded: no + +Index: pyhst2/PyHST/Cspace/c_hst_project_1over.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/c_hst_project_1over.cu ++++ pyhst2/PyHST/Cspace/c_hst_project_1over.cu +@@ -43,6 +43,8 @@ extern "C" { + #include<CCspace.h> + } + ++#include "cuda_texture_compat.h" ++ + + # define CUDA_SAFE_CALL_NO_SYNC( call) { \ + cudaError err = call; \ +Index: pyhst2/PyHST/Cspace/cudamedianfilter.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/cudamedianfilter.cu ++++ pyhst2/PyHST/Cspace/cudamedianfilter.cu +@@ -51,6 +51,7 @@ extern "C" { + #include"CCspace.h" + } + ++#include "cuda_texture_compat.h" + + texture<float, 2, cudaReadModeElementType> texData; + +@@ -668,12 +669,11 @@ void gpu_med( Gpu_med_Context * self, i + // semop(semidSend,&wait,1); + CUDA_SAFE_CALL_EXCEPTION( cudaMemcpyToArray(a_data, 0, 0, data , ny*nx*sizeof(float) , cudaMemcpyHostToDevice) ); + // semop(semidSend,&signal,1); ++ texData.filterMode = cudaFilterModePoint; + CUDA_SAFE_CALL_EXCEPTION( cudaBindTextureToArray(texData, a_data) ); + + + +- texData.filterMode = cudaFilterModePoint; +- + + size_t pitch ; + +Index: pyhst2/PyHST/Cspace/cudaunsharp.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/cudaunsharp.cu ++++ pyhst2/PyHST/Cspace/cudaunsharp.cu +@@ -16,6 +16,7 @@ + #include<iostream> + + #include"unsharp3D.h" ++#include "cuda_texture_compat.h" + + + #include<tiffio.h> +Index: pyhst2/PyHST/Cspace/gputomo.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/gputomo.cu ++++ pyhst2/PyHST/Cspace/gputomo.cu +@@ -54,6 +54,8 @@ using namespace std; + extern "C" { + #include<CCspace.h> + } ++ ++#include "cuda_texture_compat.h" + // #include <cutil.h> + #include<assert.h> + +@@ -2842,8 +2844,8 @@ int gpu_fbdl(Gpu_Context * self, float * + + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) self->a_Proje_voidptr, 0, 0,self->dev_Work_perproje , + self->nprojs_span*self->num_bins*sizeof(float) , cudaMemcpyDeviceToDevice) ); +- CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + texProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + + // da fare : testare prima la copia di memoria in blocco 2D ===> OK + // Allocare con pinned da CCspace.cx +@@ -7678,12 +7680,12 @@ int pro_gpu_main_conicity(Gpu_Cont + params.dstPos.z = nslices+1 ; + CUDA_SAFE_CALL( cudaMemcpy3D(¶ms);); + } +- CUDA_SAFE_CALL( cudaBindTextureToArray(texfor3D,a_SLICE_voidptr) ); + texfor3D.filterMode = cudaFilterModeLinear; + texfor3D.addressMode[0] = cudaAddressModeBorder; + texfor3D.addressMode[1] = cudaAddressModeBorder; + texfor3D.addressMode[2] = cudaAddressModeClamp; + texfor3D.normalized = false; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texfor3D,a_SLICE_voidptr) ); + + + int pezzo_z=8, pezzo_x=8, pezzo_y=8; +@@ -8445,8 +8447,8 @@ int gpu_main_2by2(Gpu_Context * self, fl + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) self->a_cProje_voidptr, 0, 0,self->dev_Work_perproje , // 2by2 cProj allocare, dichiarare anche texture + self->nprojs_span*self->num_bins*sizeof(float)*2 , cudaMemcpyDeviceToDevice) ); + +- CUDA_SAFE_CALL( cudaBindTextureToArray(texcProjes,(cudaArray*) self->a_cProje_voidptr) ); // !! unbind ?? // 2by2 texture texProjes_c + texcProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texcProjes,(cudaArray*) self->a_cProje_voidptr) ); // !! unbind ?? // 2by2 texture texProjes_c + + + +Index: pyhst2/PyHST/Cspace/gputomo16.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/gputomo16.cu ++++ pyhst2/PyHST/Cspace/gputomo16.cu +@@ -45,6 +45,8 @@ using namespace std; + extern "C" { + #include<CCspace.h> + } ++ ++#include "cuda_texture_compat.h" + // #include <cutil.h> + #include<assert.h> + # define CUDA_SAFE_CALL_NO_SYNC( call) { \ +@@ -613,8 +615,8 @@ int gpu_fbdl(Gpu_Context * self, float * + + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) self->a_Proje_voidptr, 0, 0,self->dev_Work_perproje , + self->num_projections*self->num_bins*sizeof(float) , cudaMemcpyDeviceToDevice) ); +- CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + texProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + + // da fare : testare prima la copia di memoria in blocco 2D ===> OK + // Allocare con pinned da CCspace.cx +@@ -663,8 +665,8 @@ int gpu_fbdl(Gpu_Context * self, float * + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) a_slice, 0, 0, self->d_SLICE, + dimrecy*dimrecx*sizeof(float) , + cudaMemcpyDeviceToDevice) ); +- CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) a_slice) ); // !! unbind ?? + texProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) a_slice) ); // !! unbind ?? + + printf("gputomo.cu: DOING slicerot_kernel with DETECTOR_DUTY_OVERSAMPLING=%d\n",DETECTOR_DUTY_OVERSAMPLING); + slicerot_kernel<<<dimGrid,dimBlock>>>( dimslice, +Index: pyhst2/PyHST/Cspace/gputomo32.cu +=================================================================== +--- pyhst2.orig/PyHST/Cspace/gputomo32.cu ++++ pyhst2/PyHST/Cspace/gputomo32.cu +@@ -45,6 +45,8 @@ using namespace std; + extern "C" { + #include<CCspace.h> + } ++ ++#include "cuda_texture_compat.h" + // #include <cutil.h> + #include<assert.h> + # define CUDA_SAFE_CALL_NO_SYNC( call) { \ +@@ -641,8 +643,8 @@ int gpu_fbdl(Gpu_Context * self, float * + + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) self->a_Proje_voidptr, 0, 0,self->dev_Work_perproje , + self->num_projections*self->num_bins*sizeof(float) , cudaMemcpyDeviceToDevice) ); +- CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + texProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) self->a_Proje_voidptr) ); // !! unbind ?? + + // da fare : testare prima la copia di memoria in blocco 2D ===> OK + // Allocare con pinned da CCspace.cx +@@ -691,8 +693,8 @@ int gpu_fbdl(Gpu_Context * self, float * + CUDA_SAFE_CALL( cudaMemcpyToArray((cudaArray*) a_slice, 0, 0, self->d_SLICE, + dimrecy*dimrecx*sizeof(float) , + cudaMemcpyDeviceToDevice) ); +- CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) a_slice) ); // !! unbind ?? + texProjes.filterMode = cudaFilterModeLinear; ++ CUDA_SAFE_CALL( cudaBindTextureToArray(texProjes,(cudaArray*) a_slice) ); // !! unbind ?? + + printf("gputomo.cu: DOING slicerot_kernel with DETECTOR_DUTY_OVERSAMPLING=%d\n",DETECTOR_DUTY_OVERSAMPLING); + slicerot_kernel<<<dimGrid,dimBlock>>>( dimslice, +Index: pyhst2/PyHST/Cspace/cuda_texture_compat.h +=================================================================== +--- /dev/null ++++ pyhst2/PyHST/Cspace/cuda_texture_compat.h +@@ -0,0 +1,219 @@ ++#ifndef CUDA_TEXTURE_COMPAT_H ++#define CUDA_TEXTURE_COMPAT_H ++ ++#if defined(__CUDACC__) ++ ++#include <cuda_runtime.h> ++#include <string.h> ++ ++template<typename T> ++struct ChannelDescHelper { ++ static __host__ __device__ cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{0, 0, 0, 0, cudaChannelFormatKindNone}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<float> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{32, 0, 0, 0, cudaChannelFormatKindFloat}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<float2> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{32, 32, 0, 0, cudaChannelFormatKindFloat}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<float4> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{32, 32, 32, 32, cudaChannelFormatKindFloat}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<int> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{32, 0, 0, 0, cudaChannelFormatKindSigned}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<unsigned int> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{32, 0, 0, 0, cudaChannelFormatKindUnsigned}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<short> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{16, 0, 0, 0, cudaChannelFormatKindSigned}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<unsigned short> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{16, 0, 0, 0, cudaChannelFormatKindUnsigned}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<char> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{8, 0, 0, 0, cudaChannelFormatKindSigned}; ++ } ++}; ++ ++template<> ++struct ChannelDescHelper<unsigned char> { ++ static __host__ __device__ constexpr cudaChannelFormatDesc desc() { ++ return cudaChannelFormatDesc{8, 0, 0, 0, cudaChannelFormatKindUnsigned}; ++ } ++}; ++ ++template<typename T, int dim = 2, enum cudaTextureReadMode readMode = cudaReadModeElementType> ++struct TextureRef { ++ cudaTextureObject_t obj; ++ cudaTextureFilterMode filterMode; ++ cudaTextureAddressMode addressMode[3]; ++ cudaTextureReadMode readModeVal; ++ int normalized; ++ ++ struct ChannelDescProxy { ++ __host__ __device__ operator cudaChannelFormatDesc() const { ++ return ChannelDescHelper<T>::desc(); ++ } ++ __host__ __device__ const cudaChannelFormatDesc* operator&() const { ++ static const cudaChannelFormatDesc d = ChannelDescHelper<T>::desc(); ++ return &d; ++ } ++ __host__ __device__ cudaChannelFormatDesc* operator&() { ++ static cudaChannelFormatDesc d = ChannelDescHelper<T>::desc(); ++ return &d; ++ } ++ } channelDesc; ++}; ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static __device__ __forceinline__ T tex1D(const TextureRef<T, dim, readMode>& tex, float x) { ++ return tex1D<T>(tex.obj, x); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static __device__ __forceinline__ T tex2D(const TextureRef<T, dim, readMode>& tex, float x, float y) { ++ return tex2D<T>(tex.obj, x, y); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static __device__ __forceinline__ T tex3D(const TextureRef<T, dim, readMode>& tex, float x, float y, float z) { ++ return tex3D<T>(tex.obj, x, y, z); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaBindTextureToArray(TextureRef<T, dim, readMode>& tex, const void* a_data, const cudaChannelFormatDesc* desc = NULL) { ++ if (tex.obj) { ++ cudaDestroyTextureObject(tex.obj); ++ tex.obj = 0; ++ } ++ struct cudaResourceDesc resDesc; ++ memset(&resDesc, 0, sizeof(resDesc)); ++ resDesc.resType = cudaResourceTypeArray; ++ resDesc.res.array.array = (cudaArray_t)a_data; ++ ++ struct cudaTextureDesc texDesc; ++ memset(&texDesc, 0, sizeof(texDesc)); ++ texDesc.addressMode[0] = tex.addressMode[0]; ++ texDesc.addressMode[1] = tex.addressMode[1]; ++ texDesc.addressMode[2] = tex.addressMode[2]; ++ texDesc.filterMode = tex.filterMode; ++ texDesc.readMode = tex.readModeVal ? tex.readModeVal : readMode; ++ texDesc.normalizedCoords = tex.normalized; ++ ++ return cudaCreateTextureObject(&tex.obj, &resDesc, &texDesc, NULL); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaBindTextureToArray(TextureRef<T, dim, readMode>& tex, const void* a_data, const cudaChannelFormatDesc& desc) { ++ return cudaBindTextureToArray(tex, a_data, &desc); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaBindTexture2D(size_t* offset, TextureRef<T, dim, readMode>& tex, const void* devPtr, const cudaChannelFormatDesc& desc, size_t width, size_t height, size_t pitch) { ++ if (tex.obj) { ++ cudaDestroyTextureObject(tex.obj); ++ tex.obj = 0; ++ } ++ struct cudaResourceDesc resDesc; ++ memset(&resDesc, 0, sizeof(resDesc)); ++ resDesc.resType = cudaResourceTypePitch2D; ++ resDesc.res.pitch2D.devPtr = const_cast<void*>(devPtr); ++ resDesc.res.pitch2D.desc = desc; ++ resDesc.res.pitch2D.width = width; ++ resDesc.res.pitch2D.height = height; ++ resDesc.res.pitch2D.pitchInBytes = pitch; ++ ++ struct cudaTextureDesc texDesc; ++ memset(&texDesc, 0, sizeof(texDesc)); ++ texDesc.addressMode[0] = tex.addressMode[0]; ++ texDesc.addressMode[1] = tex.addressMode[1]; ++ texDesc.addressMode[2] = tex.addressMode[2]; ++ texDesc.filterMode = tex.filterMode; ++ texDesc.readMode = tex.readModeVal ? tex.readModeVal : readMode; ++ texDesc.normalizedCoords = tex.normalized; ++ ++ if (offset) *offset = 0; ++ return cudaCreateTextureObject(&tex.obj, &resDesc, &texDesc, NULL); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaBindTexture2D(size_t* offset, TextureRef<T, dim, readMode>& tex, const void* devPtr, const cudaChannelFormatDesc* desc, size_t width, size_t height, size_t pitch) { ++ cudaChannelFormatDesc d = desc ? *desc : ChannelDescHelper<T>::desc(); ++ return cudaBindTexture2D(offset, tex, devPtr, d, width, height, pitch); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaBindTexture(size_t* offset, TextureRef<T, dim, readMode>& tex, const void* devPtr, const cudaChannelFormatDesc* desc = NULL, size_t size = 0) { ++ if (tex.obj) { ++ cudaDestroyTextureObject(tex.obj); ++ tex.obj = 0; ++ } ++ struct cudaResourceDesc resDesc; ++ memset(&resDesc, 0, sizeof(resDesc)); ++ resDesc.resType = cudaResourceTypeLinear; ++ resDesc.res.linear.devPtr = const_cast<void*>(devPtr); ++ resDesc.res.linear.desc = desc ? *desc : ChannelDescHelper<T>::desc(); ++ resDesc.res.linear.sizeInBytes = size; ++ ++ struct cudaTextureDesc texDesc; ++ memset(&texDesc, 0, sizeof(texDesc)); ++ texDesc.addressMode[0] = tex.addressMode[0]; ++ texDesc.addressMode[1] = tex.addressMode[1]; ++ texDesc.addressMode[2] = tex.addressMode[2]; ++ texDesc.filterMode = tex.filterMode; ++ texDesc.readMode = tex.readModeVal ? tex.readModeVal : readMode; ++ texDesc.normalizedCoords = tex.normalized; ++ ++ if (offset) *offset = 0; ++ return cudaCreateTextureObject(&tex.obj, &resDesc, &texDesc, NULL); ++} ++ ++template<typename T, int dim, enum cudaTextureReadMode readMode> ++static inline cudaError_t cudaUnbindTexture(TextureRef<T, dim, readMode>& tex) { ++ if (tex.obj) { ++ cudaError_t err = cudaDestroyTextureObject(tex.obj); ++ tex.obj = 0; ++ return err; ++ } ++ return cudaSuccess; ++} ++ ++#define texture __managed__ TextureRef ++ ++#endif /* __CUDACC__ */ ++ ++#endif /* CUDA_TEXTURE_COMPAT_H */ diff -Nru pyhst2-2020c/debian/patches/series pyhst2-2020c/debian/patches/series --- pyhst2-2020c/debian/patches/series 2023-01-31 17:13:52.000000000 +0100 +++ pyhst2-2020c/debian/patches/series 2026-09-11 17:35:16.000000000 +0200 @@ -3,3 +3,8 @@ 0003-set-debian-specific-pep-440-version.patch 0004-drop-custom-lib-search-path.patch typos.patch +0005-drop-numpy-distutils.patch +0006-fix-hdf5-types-in-edftools.patch +0007-fix-implicit-drand48.patch +0008-fix-cuda-complex-hang.patch +0009-cuda12-texture-compat.patch diff -Nru pyhst2-2020c/debian/tests/control pyhst2-2020c/debian/tests/control --- pyhst2-2020c/debian/tests/control 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/tests/control 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1,2 @@ +Test-Command: set -e ; for py in $(py3versions -i); do echo "=== $py ==="; $py -c "import PyHST2_2020c; print(PyHST2_2020c)"; done +Depends: @, python3-all diff -Nru pyhst2-2020c/debian/tests/pkg-python/import-name pyhst2-2020c/debian/tests/pkg-python/import-name --- pyhst2-2020c/debian/tests/pkg-python/import-name 1970-01-01 01:00:00.000000000 +0100 +++ pyhst2-2020c/debian/tests/pkg-python/import-name 2026-09-11 17:35:16.000000000 +0200 @@ -0,0 +1 @@ +PyHST2_2020c
