Conversation
Port HEonGPU (Fully Homomorphic Encryption on GPU) to AMD GPUs via HIP.
Build configuration:
- Add USE_HIP CMake option to enable HIP compilation
- Link hip::host (not hip::device) to avoid propagating HIP compile flags
to downstream pure C++ consumers -- the INTERFACE_COMPILE_OPTIONS on
hip::device would cause g++ to receive -x hip --offload-arch=gfx90a
- Compile test .cpp files as HIP sources since they transitively include
rocThrust headers via heongpu.hpp
CUDA-to-HIP adaptations:
- cuda_to_hip.h compatibility header maps CUDA runtime symbols to HIP
- hip_compat/ shim headers redirect cuda_runtime.h and curand_kernel.h
includes from thirdparty code to HIP equivalents
- Replace PTX inline assembly with portable __umul64hi intrinsic
- Use hipPointerAttribute_t.type (HIP differs from CUDA's .memoryType)
- hiprand replaces curand; rocThrust replaces Thrust
RMM HIP stub (thirdparty/rmm_hip_stub/):
- Minimal RMM implementation for HIP (real RMM requires CUDA)
- Implements device_uvector, device_buffer, pool_memory_resource,
statistics_resource_adaptor, pinned_memory_resource
- Checks hipError_t return values; throws std::runtime_error on failure
Submodule updates (same author, separate repos):
- GPU-NTT: PTX->__umul64hi, clang syntax fixes
- GPU-FFT: CMake HIP support
- RNGonGPU: CUDA->HIP compat header, hiprand linkage
Test Plan:
```bash
cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DHEonGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)
cd build && ctest --output-on-failure
```
This port was authored with AI assistance (Claude).
The submodule gitlinks pointed at three commits that exist in no
repository, so the branch could not be cloned and built at all. This
resets thirdparty/GPU-FFT, thirdparty/GPU-NTT and thirdparty/RNGonGPU to
the upstream commits the parent branch was based on, and carries their
AMD GPU support as patches applied on top of those pinned commits.
GPU-NTT, GPU-FFT and RNGonGPU are separate upstream repositories, so
their changes cannot live in the checked-out submodule tree. thirdparty/
build.sh applies thirdparty/patches/<name>.patch after initialising the
submodules, and only when configured for AMD GPUs; each apply is skipped
when the patch is already in place, so repeated configures stay
idempotent and an NVIDIA build sees the submodules untouched.
The patches themselves cover three things. The build files gain a
USE_HIP option that selects the HIP language, compiles the .cu sources
as HIP and links hip::host and hipRAND in place of the CUDA runtime and
cuRAND. hip::host rather than hip::device is deliberate: hip::device
exports -x hip --offload-arch on its INTERFACE, which reaches pure C++
consumers compiled by g++ and breaks them. In GPU-NTT's modular
arithmetic the PTX carry-chain sequences are replaced, on AMD only, by
the equivalent C++: mul.lo/mul.hi becomes a 64-bit product plus
__umul64hi, and sub.cc/subc becomes an explicit borrow subtraction where
the borrow out of the low limb is exactly (x < other.x). The NVIDIA path
keeps the original assembly. Finally clang rejects the chained
comparisons in nttparameters.cu (0 < logn <= 25 parses as (0 < logn) <=
25), so those are spelled out as explicit conjunctions.
The header shims grow curand_mtgp32_host.h and curand_mtgp32dc_p_11213.h
so RNGonGPU's cuRAND path still resolves its includes; hipRAND has the
former and has no equivalent of the latter, whose MTGP32 parameter set
that path never actually uses. cuda_runtime.h now also defines
CUDART_VERSION, because hipPointerAttribute_t exposes .type and has no
.memoryType, which is the shape version-guarded code expects from CUDA
10.0 onward.
Review the build.sh and thirdparty/CMakeLists.txt wiring first, then the
three patches.
Test Plan:
```bash
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
```
Builds clean on gfx1100: libheongpu.a, libntt-1.0.a, libfft-1.0.a,
librngongpu-1.0.a and all 15 test executables. A fresh clone of this
branch now configures and builds without manual submodule repair.
```bash
ctest --test-dir build --output-on-failure
```
2 of 20 tests pass (the BFV and CKKS encoding cases). The remaining 18
fail on incorrect results, identically to what a gfx90a run produces, so
the outstanding fault is not wavefront-width dependent. That defect is
tracked separately and is not addressed here.
This work was authored with an AI assistant (Claude).
The README and the getting-started guide each carry a CUDA build block and an architecture table, so each gains the parallel AMD GPU block in the same place and the same style: the USE_HIP configure line, a CMAKE_HIP_ARCHITECTURES table covering CDNA2, CDNA3, RDNA3 and RDNA4, and a note that hipRAND and rocThrust stand in for cuRAND and Thrust while RMM is replaced by a bundled minimal implementation. ROCm joins the CUDA Toolkit in both requirement lists, each now marked with the vendor it applies to. This work was authored with an AI assistant (Claude).
Configuring for AMD GPUs applies thirdparty/patches/<name>.patch into the submodule working trees, which then read as modified in every git status for the rest of the checkout. The patches are what is tracked here, so that content difference is expected rather than a local edit someone forgot; ignore = dirty says so and keeps a real submodule bump still visible. This work was authored with an AI assistant (Claude).
The Box-Muller sampler in the bundled random number generator wrote its
result with `static_cast<T>(z)`, where `T` is the unsigned residue type
and `z` is a signed sample, relying on the following `+ (flag & modulus)`
to turn a negative sample into its representative. Converting an
out-of-range floating-point value to an unsigned integer is undefined,
and the two GPU back ends resolve it differently: NVIDIA's `cvt.rzi.u64.f64`
saturates a negative source to zero, while the AMD lowering of a 64-bit
`fptoui` keeps only the low 32 bits of the two's-complement value. On an
MI210 (gfx90a) that turned roughly half of every error polynomial into
values near 2^32 instead of single digits, so every ciphertext decrypted
to noise while plaintext encoding, which never samples an error, was
unaffected.
Routing the truncation through the signed type first makes it well
defined on both back ends and produces the value the modulus addition is
written to consume. `truncate_signed` in the shared header is the only
new code; the four Box-Muller kernels and the curand-backed normal
kernels each swap one cast for a call to it. Note that this is not
behaviour-neutral on NVIDIA: it replaces the saturated zero with the
correct negative representative, which is the value the surrounding
arithmetic already assumes.
Measured on gfx90a with a 4096-coefficient ring and standard deviation
3.2, the largest centred error goes from 4294967295 to 13 and the mean
absolute error from 1.6e9 to 2.07. The CKKS suites go from failing to
passing; the BFV and TFHE suites still fail for reasons unrelated to
this change.
Prepared with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON .
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
Before: 2 of 20 suites pass. After: 11 of 20 pass, including every CKKS
suite (addition, encoding, encryption, multiplication, relinearization
and both rotation methods).
The device Barrett reduction in GPU-NTT builds its 128-bit intermediates out
of a pair of 64-bit limbs and shifts them by modulus.bit + 3. For a modulus
of 61 bits or more that count reaches 64, and the limb shifts it expands to
are then shifts of a 64-bit value by 64 or more, which C++ leaves undefined
and the two back ends resolve differently: PTX produces zero, while AMDGPU
keeps only the low 6 bits of the shift count, so a shift of 64 acts as a
shift of 0 and the high limb is OR-ed into the low one instead of replacing
it. Every modular multiplication at such a modulus then returned an
unreduced, unrelated value.
BFV decryption is the visible casualty: its gamma correction term is a
61-bit prime, so ciphertexts decrypted to noise on AMD GPUs while CKKS,
whose moduli here are all 37 bits or fewer, was unaffected. TFHE uses a
61-bit prime as well.
The shift operators now handle a count of 0, a count below 64, and a count
of 64 or more explicitly. NVIDIA behaviour is unchanged, since PTX already
yielded the values these branches now compute.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
Direct comparison of the device operators against a __uint128_t host
reference over 20000 random pairs plus the modulus boundary values, at
20-, 36-, 37-, 60-, 61- and 62-bit moduli. Before: every multiplication at
61 and 62 bits wrong. After: no mismatches at any width.
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
On gfx90a with ROCm 7.2.1 this takes the suite from 11/20 to 19/20; all
seven BFV suites and BFV encryption/decryption now pass.
The TFHE encryptor allocates its per-thread random states with room for n_
of them, which is 512, then initializes and uses 512 * 32 of them: both the
initialization kernel and encrypt_lwe_kernel index the buffer by a global
thread id over a 32-block, 512-thread launch. Every state past the first
512 was written and read out of bounds, 32 times more memory than was
allocated.
On AMD GPUs this corrupts whatever follows the allocation and the LWE
ciphertexts come back as noise, so no gate test could pass; NVIDIA appears
to tolerate it. Allocating total_state entries is what the two kernels
already assume. The allocation is now checked as well, since a failure here
was previously silent.
This is a memory-safety defect rather than a portability one, so it is a
fix on both platforms.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
On gfx90a with ROCm 7.2.1, HEonGPU.TFHE_Gate_Boots goes from failing every
gate to passing. A direct check of encrypt followed by decrypt with no
bootstrapping went from 35 of 64 bits wrong to 0 of 64.
test_bfv_addition computes its expected value with a conditional subtraction
guarded by a strict inequality, so a coefficient pair summing to exactly the
plain modulus was expected to decode as the modulus rather than as zero. The
library is right and the reference is wrong, and with 4096 coefficients per
parameter set the case comes up in roughly one run in eighty, which reads as
a flaky GPU failure.
This is platform independent and equally wrong on NVIDIA.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
Forty encrypt-add-decrypt rounds over 4096 coefficients were compared
against (m1 + m2) mod t: no mismatches. The same data compared against the
reference as it stood mismatched once, on the pair 957528 + 74665 = 1032193,
which the library decoded as 0.
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
find_package(HEonGPU) against an installation built with USE_HIP=ON could
not configure. The exported package configs asked for the CUDA Toolkit
unconditionally; the bundled memory-manager stub installed its export set
without a config file to find it by; the exported link interface named
hip::hiprand without the find_package that defines it; and the shim headers
that resolve the cuda_runtime.h and curand_kernel.h includes in the
installed headers were neither installed nor on the interface include path.
Each of these only bites a consumer of an installed library, which is why
the in-tree tests never saw them.
The configs now select their GPU runtime dependency from the option the
library was built with, so the CUDA path is untouched. Review the config
templates first, then the two install rules that back them.
The downstream integration section of the documentation gains the HIP
counterpart of the CMake snippet it already carries for CUDA.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
Installed the library and built a separate CMake project against it using
exactly the documented snippet, then ran it on gfx90a with ROCm 7.2.1; it
generated a BFV context and printed its parameters.
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=$PREFIX
cmake --build build -j$(nproc)
cmake --install build
cmake -S consumer -B consumer/build -DCMAKE_PREFIX_PATH=$PREFIX \
-DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_BUILD_TYPE=Release
cmake --build consumer/build
./consumer/build/consumer
```
The full suite still passes, 20 of 20.
```
ctest --test-dir build --output-on-failure
```
The guarded shift operators landed earlier in this branch with the claim
that NVIDIA behaviour was unchanged. That is true only for shift counts up
to and including 64, and the code comment did not say so. This corrects
both the comment and the record.
PTX clamps a shift count to the register width, so the original unguarded
form is already right at a count of exactly 64: the low limb reads
value.y and the high limb reads zero, which is what the explicit branch
computes. From 65 upward it is not right. There 64 - shift underflows, PTX
clamps every term to the width, and the whole 128-bit result comes back as
zero, while the correct value, and what the explicit branch now computes,
is value.y >> (shift - 64) in the low limb.
Barrett reduction shifts by modulus.bit + 3, so a count of 65 is reached at
a 62-bit modulus, which BarrettOperations documents as supported for
Data64. A CUDA build therefore also returns garbage from every modular
multiplication at a 62-bit modulus, and this change fixes that too. The
measured sweep quoted with the original fix already showed 62-bit moduli
failing before it and clean after; only the description was wrong.
No code behaviour changes here beyond the comment, which is carried in the
submodule patch.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
20/20 suites pass on gfx90a with ROCm 7.2.1. The regenerated patch applies
to a pristine submodule checkout and reverse-applies cleanly, so repeated
configures stay idempotent:
```
git -C thirdparty/GPU-NTT checkout -- .
bash thirdparty/build.sh ON && bash thirdparty/build.sh ON
git -C thirdparty/GPU-NTT diff | diff - thirdparty/patches/GPU-NTT.patch
```
HEonGPU_BUILD_EXAMPLES=ON with USE_HIP=ON did not build: the example
targets linked the HIP runtime but were still compiled as plain C++, so
the first HEonGPU header pulled in cuda_runtime.h and then rocThrust,
which needs a HIP compilation context. HEonGPU_BUILD_BENCHMARKS=ON did not
even configure, because the benchmark targets linked CUDA::cudart
unconditionally.
Both now do what the tests already do: the source is compiled as HIP, the
USE_HIP definition reaches it so the compatibility header maps the CUDA
runtime calls the examples make directly, and the shim include directory
is on the path for the third-party headers that include CUDA headers by
name. OpenMP needs one extra step here that it does not need in the tests:
OpenMP::OpenMP_CXX only decorates the CXX language, so a source compiled
as HIP gets the flag explicitly on both the compile and the link line,
otherwise the OpenMP example fails to link. The CUDA branches are
unchanged.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
./build/bin/examples/basic/1_basic_bfv
./build/bin/examples/basic/9_multi_stream_usage_way1
./build/bin/examples/basic/15_basic_tfhe
./build/bin/examples/mpc/1_multiparty_computation_bfv
./build/bin/benchmark/tfhe_benchmark
```
All 24 example targets and all 3 benchmark targets build on gfx90a with
ROCm 7.2.1. The runs above produce correct output, including the OpenMP
multi-stream example and the bootstrapped TFHE gates.
RMM has no HIP build, so the AMD path substitutes a small stand-in for the
handful of RMM types HEonGPU uses. Its device_uvector deleted the copy
constructor and stopped there, but RMM also offers an explicit copy onto a
given stream and memory resource, and DeviceVector's own copy constructor
in the installed public headers forwards straight to it. Nothing in the
library instantiates that constructor, so the build stayed green while any
consumer copy-constructing a DeviceVector would compile against CUDA and
fail against the AMD build.
The stub now provides the same constructor RMM does, with the stream
required rather than defaulted so it cannot become ambiguous with the
deleted copy constructor.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
A consumer translation unit that fills a heongpu::DeviceVector<Data64>,
copy-constructs a second one from it and reads the copy back reports
`copy-construct OK` with distinct device pointers and matching contents.
Against the previous stub the same file fails to compile with `no matching
constructor for initialization of 'rmm::device_uvector<unsigned long>'` at
devicevector.cuh:35.
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
20/20 suites pass on gfx90a with ROCm 7.2.1.
Housekeeping over the earlier commits in this branch, no behaviour change.
The compatibility header defined kWarpSize and FULL_WARP_MASK on both
paths. Neither is referenced anywhere; the wavefront-width fixes that
landed all read the runtime warpSize instead, which is what a header of
compile-time constants cannot express when one binary targets several
architectures. The root CMakeLists set two cache variables nothing reads.
Two comments described something other than the code under them. The warp
reduction said HIP needs a 64-bit shuffle mask where the call passes no
mask at all, and small_ntt.cu said it was kept for explicit instantiation
after the instantiations had been removed from it, leaving an empty
translation unit still in the kernel source list. Its device functions are
defined inline in small_ntt.cuh for both back ends, so the file goes.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
20/20 suites pass on gfx90a with ROCm 7.2.1, and the examples and
benchmarks build.
The verification section spells out the configure commands for tests,
examples and benchmarks with the CUDA architecture flag, so say once that
the AMD options substitute for it there as they do in the build steps
above, now that all three groups build and run on AMD GPUs.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
Documentation only, no build change. The commands it describes were run:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
The examples' AMD branch passed ${OpenMP_CXX_FLAGS} to the HIP compile and link
and also kept OpenMP::OpenMP_CXX on the target. FindOpenMP guards that target's
compile options with $<COMPILE_LANGUAGE:CXX>, so on a source compiled as HIP it
contributed nothing but its interface link libraries, which are unguarded and are
the C++ compiler's OpenMP runtime. Every example binary therefore listed both
libgomp.so.1 and libomp.so in DT_NEEDED.
That resolves to a single runtime only while the loader finds ROCm's
lib/llvm/lib first, where libgomp.so.1 is a symlink to libomp.so. With GNU's
libgomp ahead of it both load: the clang-compiled object's __kmpc_ calls still
fork a real team out of LLVM's runtime, while omp_get_thread_num answers from
libgomp, which knows nothing about that team and returns 0 for every thread.
Checked with a minimal HIP plus OpenMP binary linked the same way -- four
distinct OS thread ids, omp_get_thread_num reporting 0 0 0 0. In
9_multi_stream_usage_way1 the thread id selects the stream, so all workers would
share stream 0 and the example would stop being multi-stream without failing.
Dropping the imported target from the AMD branch leaves the explicit flags, which
are what a HIP-compiled source needs; libheongpu.a references no OpenMP symbol,
and the NVIDIA branch, where libgomp is the right runtime for an nvcc and g++
build, is untouched. The flags now use the SHELL: form FindOpenMP itself uses, so
a compiler whose OpenMP flag is more than one token is not passed as one
argument.
Authored with the assistance of Claude (Anthropic).
Test Plan:
```
cmake -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_BUILD_TYPE=Release \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON ..
cmake --build . -j$(nproc)
ctest --test-dir build # 20/20, four consecutive runs
readelf -d bin/examples/basic/9_multi_stream_usage_way1 | grep -E 'gomp|omp'
./bin/examples/basic/9_multi_stream_usage_way1
./bin/examples/basic/1_basic_bfv
./bin/examples/basic/15_basic_tfhe
./bin/examples/bootstrapping/3_ckks_bit_bootstrapping
./bin/examples/mpc/1_multiparty_computation_bfv
```
Run on gfx90a with ROCm 7.2.1. readelf now reports libomp.so alone, and no
executable under bin/ links libgomp.
The NVIDIA branch of the compatibility header included <curand_kernel.h>.
util.cuh includes that header and is reached from heongpu.hpp, so every host
.cpp translation unit ended up with cuRAND. curand_mtgp32_kernel.h declares
threadIdx and blockDim with C++ linkage and device_launch_parameters.h then
declares them with C linkage; g++ rejects the conflict, so lib/heongpu.cpp
failed to compile with nvcc 12.8 and gcc 13. The project otherwise reaches
cuRAND only from kernel/*.cuh, which no host translation unit includes.
The include was also redundant: the kernel headers that use cuRAND include it
themselves, and the AMD branch of this header keeps its own hipRAND include.
Only the NVIDIA branch changes, so no AMD object code moves.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
NVIDIA, compile only (no NVIDIA GPU on this host):
```
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
```
AMD:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
Both build clean, and 20/20 suites pass on an MI210 (gfx90a) with ROCm 7.2.1.
SmallForwardNTT and SmallInverseNTT were moved into small_ntt.cuh earlier in
this branch, because the AMD build compiles without relocatable device code and
a __device__ function defined in another translation unit does not link there:
lld reports "undefined hidden symbol". That move was correct for AMD but broke
the NVIDIA build, since keygeneration.cuh and bootstrapping.cuh reach this
header from heongpu.hpp, so g++ compiling lib/heongpu.cpp had to parse device
bodies and failed on __syncthreads.
The definitions now sit behind __CUDACC__ || __HIPCC__ and the declarations
stay unconditional, so a host compiler sees exactly what it saw before this
branch while every calling translation unit -- all of them .cu, compiled by
nvcc or as HIP -- still gets the bodies inline. The alternative, restoring
lib/kernel/small_ntt.cu for NVIDIA and keeping the header definitions for AMD,
would put 145 lines of NTT butterfly code in two places that must not drift;
the guard keeps one copy. The condition tests the compiler rather than the
target because that is the actual requirement: only a device compiler can parse
a device body, and both back ends compile these callers with one.
No AMD code changes: the callers are device translation units where the guard
is true. Rebuilt for gfx90a and compared against the binaries from the previous
branch head with llvm-objdump: exported symbols and device ISA are identical on
all 42 binaries.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
NVIDIA, compile only (no NVIDIA GPU on this host):
```
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
```
AMD:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
Both build clean, and 20/20 suites pass on an MI210 (gfx90a) with ROCm 7.2.1.
The comment added with that change named the wrong condition. It said the
C++-linkage threadIdx and blockDim in curand_mtgp32_kernel.h conflict with the
C-linkage ones in device_launch_parameters.h whenever a host compiler sees
both, which the tree disproves: every host translation unit still sees both
declarations today -- the C-linkage one arrives from GPU-NTT's
common/modular_arith.cuh and the C++-linkage one from kernel/encryption.cuh
via heongpu.hpp -- and g++ 13 accepts it.
The condition is the ORDER, and only one direction fails. Against CUDA 12.8
headers, curand_kernel.h followed by cuda_runtime.h and
device_launch_parameters.h gives the two conflicting-declaration errors, while
the reverse order is clean. That is why this header specifically was the wrong
home for the include and why the kernel headers are the right one: util.cuh
pulls this header in ahead of everything else, so a curand include here reached
the compiler before anything had declared the builtins, and the
<cuda_runtime.h> on the next line does not help because cuda_runtime.h includes
device_launch_parameters.h only under __CUDACC__. The kernel headers that use
cuRAND are reached later, in the order that compiles.
Comment only; no compiled output changes on either back end.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
g++-13 -fsyntax-only -I$CUDA/include probe.cpp
```
with probe.cpp including curand_kernel.h first (two errors) and last (clean).
Both project builds were rebuilt and re-checked:
```
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
Showing the SmallForwardNTT and SmallInverseNTT bodies to device compilers
only, rather than keeping them in a translation unit of their own, is the more
intrusive of two ways to get this header to link on AMD, and the branch never
recorded why the other one was turned down.
The AMD build could enable relocatable device code, which is what the NVIDIA
side already does through CUDA_SEPARABLE_COMPILATION in
heongpu_set_gpu_properties. That is a build-flag change and it would have left
this header alone. It was rejected on performance: -fgpu-rdc keeps each
butterfly a call across translation units that the compiler cannot inline into
the kernel running it, and these two functions run inside a shared-memory NTT
where the call overhead lands on every stage of every butterfly. The
header-definition form measured 20 of 20 test suites passing on an MI210, while
an -fgpu-rdc build was never measured, so the guard also keeps the AMD
configuration the same one the results come from.
Comment only; no compiled output changes on either back end.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
```
SmallForwardNTT and SmallInverseNTT are __device__ functions defined in
lib/kernel/small_ntt.cu and called from kernels in other translation units. A
HIP compile resolves device symbols one translation unit at a time, so it stops
with "lld: error: undefined hidden symbol: void SmallForwardNTT<unsigned long
long>(...)", referenced by the calling kernel. This branch had answered that by
moving both bodies into kernel/small_ntt.cuh behind a device-compiler guard.
Turning on relocatable device code instead is the counterpart of the
CUDA_SEPARABLE_COMPILATION this build already sets for NVIDIA, and it costs
three lines of CMake, so small_ntt.cuh and small_ntt.cu go back to exactly what
they were before this branch and the NVIDIA build returns to compiling them the
way it always did.
The form being replaced was justified by a claim that is false on AMD: that
-fgpu-rdc leaves each butterfly a cross-unit call the compiler cannot inline.
HIP's device link is a bitcode link followed by optimization of the whole
image, so such a call is inlined and the callee's standalone body does not
survive it. A reduction of exactly the shape used here -- a __device__ function
template with its explicit instantiation in one translation unit, called from a
__global__ in another -- links to a device image containing only the kernel,
with no call instructions in it and the callee's barriers inline.
What -fgpu-rdc does cost was measured rather than assumed, on gfx90a with ROCm
7.2.1, each configuration built clean:
20 of 20 test suites pass either way.
benchmark/tfhe_benchmark, the binary whose gates run these two functions:
NAND 16.52-16.68 ms and MUX 29.67-29.99 ms, two runs of each build, the
ranges overlapping.
Default build (library only): 53.7 s before, 52.1 s after.
Tests, examples and benchmarks all enabled (34 executables): 193.6 s before,
326.6 s after.
Device code in test/bfv_multiplication_testcases: 1,575,096 bytes before,
1,684,264 after, of which the part belonging to this library grows from
864,312 to 973,480.
The build time is spent in the device link, which every executable performs
over the whole library, so it lands on a build that enables the tests, examples
and benchmarks and not on the default one. The larger image is the whole-image
optimization inlining more than a per-unit compile could.
Carrying the flag on the interface, and into the installed export, keeps the
device link a detail of this library: the downstream CMake in
docs/advanced_topics.rst needs no change. Verified by installing this build and
building and running an encrypt-decrypt program against the installed package
with that snippet unmodified.
Review src/CMakeLists.txt first; the other two files are reverts to the
pre-branch text.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
./build/bin/benchmark/tfhe_benchmark
```
20/20 suites pass on an MI210 (gfx90a) with ROCm 7.2.1.
NVIDIA, compile only (no NVIDIA GPU on this host):
```
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
```
Installed-package consumer, the snippet from docs/advanced_topics.rst verbatim:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=<prefix>
cmake --build build -j$(nproc)
cmake --install build
cmake -S consumer -B consumer/build -DCMAKE_PREFIX_PATH="<prefix>;/opt/rocm" \
-DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_BUILD_TYPE=Release
cmake --build consumer/build
./consumer/build/consumer
```
The comment stated two things the build disproves. It said host translation
units reach this header first of all, and that a curand include here broke
every host .cpp. Preprocessing the library's four host translation units with
their own compile flags puts the numbers on it: in lib/heongpu.cpp this header
arrives at line 18 and device_launch_parameters.h only at 58820, the order that
fails, but lib/kernel/contextpool.cpp reaches the launch parameters at 58806
and this header at 59557, because kernel/contextpool.hpp includes GPU-NTT's
nttparameters.cuh ahead of util.cuh; lib/util/serializer.cpp reaches neither
header and lib/util/defaultmodulus.cpp does not reach this one. The failing
build agreed and named one object, lib/heongpu.cpp.o.
The order rule itself is unchanged and is what the comment now gives, together
with the translation unit it was observed on. Reaching this header before the
launch parameters is enough to condemn it as the home for the include; a
universal claim only invites a counterexample.
Comment only; no compiled output changes on either back end.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
g++-13 -E <the target's own defines, includes and flags> src/lib/heongpu.cpp \
| grep -n 'device_launch_parameters.h"\|cuda_to_hip.h"'
```
repeated for lib/kernel/contextpool.cpp, lib/util/serializer.cpp and
lib/util/defaultmodulus.cpp, giving the four results above. Both builds were
rebuilt and re-checked:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
```
The comment on the -fgpu-rdc interface options claimed the flag "keeps that a
detail of this library rather than something every consumer has to know". That
is not true, and the AMD half of the downstream snippet has no counterpart to
the line that tells a CUDA consumer to set CUDA_SEPARABLE_COMPILATION, so a
reader comparing the two concludes the AMD side needs nothing.
An archive of relocatable device code carries no complete device image, so the
device link happens at the consumer's final link and only the HIP compiler
driver in HIP link mode performs it. Measured against an installed export with
ROCm 7.2.1 on gfx90a: g++ -no-pie over the installed archive stops with
undefined reference to __hip_gpubin_handle_<hash> and __hip_fatbin_<hash>, and
so does hipcc without -fgpu-rdc. The exported INTERFACE_LINK_OPTIONS hides this
from a consumer whose own sources are HIP -- the documented snippet builds and
runs unchanged -- but not from one whose target has no HIP source of its own:
CMake adds --hip-link only for targets that have HIP sources, so an application
.cpp linking a HIP library that links this archive gets the same undefined
__hip_fatbin_* references. Setting LINKER_LANGUAGE HIP does not add the flag;
target_link_options(<target> PRIVATE --hip-link) does, and that program then
links and runs.
The comment now states the requirement instead of denying it, and
docs/advanced_topics.rst says what a consumer needs, next to the snippet, as the
CUDA paragraph above it already does. Comment and documentation only: all 42
binaries of the AMD build are byte-identical across this change, and the CUDA
configure and build have nothing to redo.
Two corrections to the record while I am here. Earlier messages on this branch
report results on an "MI210"; the machine that produced every AMD measurement
reports AMD Instinct MI250X / MI250, SKU D65209, GFX version gfx90a, so read
those results as gfx90a on an MI250X/MI250. And the build-cost figure that came
with the -fgpu-rdc change said 34 test, example and benchmark executables where
the configuration builds 42 (15 + 15 + 5 + 4 + 3): the 193.6 s to 326.6 s clean
build is 133 s of extra device link spread over 42 executables, about 3.2 s
each. The runtime comparison behind that change covered the BFV and CKKS
benchmarks as well as the TFHE gates, and none of them moved outside
run-to-run noise.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j64
ctest --test-dir build --output-on-failure
```
20/20 suites pass on gfx90a (AMD Instinct MI250X / MI250) with ROCm 7.2.1, and
the 42 binaries carry the same SHA-256 as before the change. The consumer
claims were checked against an installed tree:
```
cmake -S . -B build-inst -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=$PREFIX
cmake --build build-inst -j64
cmake --install build-inst
```
then building three consumers against it with CMAKE_PREFIX_PATH=$PREFIX: the
documented snippet, which links and runs a BFV encrypt/decrypt roundtrip; the
same program split so the executable has no HIP source, which fails with
undefined __hip_fatbin_*; and that one with --hip-link added, which links and
runs. The CUDA path was reconfigured and rebuilt with nvcc against CUDA 12.8,
with nothing to redo and no errors.
The note added with -fgpu-rdc told a consumer whose target has no HIP source of its own to add --hip-link, which is right for a consumer that reaches the library through find_package but not for one that imports the archive by hand. Rebuilding both shapes against the installed package shows why: install(EXPORT) records HIP as the link interface language of HEonGPU::heongpu, so any CMake target linking it, directly or through a plain C++ library of its own, is already driven by the HIP compiler and already receives -fgpu-rdc, and only HIP link mode is missing. A target that names the archive by path gets none of the three, so there --hip-link alone is simply a second option the C++ driver does not recognize, and the link needs LINKER_LANGUAGE HIP plus --hip-link plus -fgpu-rdc.
Documentation and comments only: advanced_topics.rst now separates the two shapes and gives the remedy for each, and the comment beside the interface options names both failure modes instead of one. No build file or source behaviour changes.
Assistance from an AI coding agent was used to prepare this change.
Test Plan:
Consumer shapes rebuilt against the installed package (ROCm 7.2.1, gfx90a, CMake 4.0.3), a target per row, each configured with the install prefix on CMAKE_PREFIX_PATH:
```
cmake -S consumers -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_PREFIX_PATH=<install-prefix>
cmake --build build --target <target> -j8
head -1 build/CMakeFiles/<target>.dir/link.txt
```
Results: a target with one HIP source of its own links and runs; a plain C++ target linking HEonGPU::heongpu is driven by the ROCm clang++ but fails with undefined __hip_fatbin_* and __hip_gpubin_handle_* until --hip-link is added, and the same holds when it reaches the library through a plain C++ library or through a local HIP-source wrapper; the archive named by path is driven by /usr/bin/c++ and needs LINKER_LANGUAGE HIP, --hip-link and -fgpu-rdc together. LINKER_LANGUAGE HIP alone never adds --hip-link and never fixes a HIP-driven link.
The consumer notes claimed the installed package hands the HIP link driver to every CMake target that links HEonGPU::heongpu. That holds only when the consuming project has HIP among its enabled languages: CMake picks a linker driver from the languages the project enabled, so a project declaring only CXX links with the C++ compiler, which rejects the -fgpu-rdc the export carries on its link interface. That is exactly the reader the paragraph addresses, because a target with no HIP source of its own also has no reason to have listed HIP in project().
Measured on gfx90a with ROCm 7.2.1 and CMake 4.0.3 against the installed package, one source tree with only the project() line differing. With LANGUAGES CXX HIP, a plain C++ executable linking HEonGPU::heongpu is linked by /opt/rocm/lib/llvm/bin/clang++ and needs only --hip-link. With LANGUAGES CXX, the same target is linked by /usr/bin/c++ and fails on -fgpu-rdc before --hip-link would matter; adding --hip-link there just earns a second rejected option; find_package(hip) and hip::host do not change it; and LINKER_LANGUAGE HIP is not available either, since without the language CMake has no HIP link rule and stops at generate time with CMAKE_HIP_LINK_EXECUTABLE unset.
The library cannot rescue such a consumer from its own side, so the documentation now says so plainly: enable the HIP language first, then add --hip-link on a target that compiles none of its own sources as HIP. The by-path case keeps its own sentence, since it needs the driver and both flags by hand. The comment beside the interface options carries the same condition. Documentation and comments only; no build setting changes.
An AI coding agent assisted with this change.
Test Plan:
Reproduce both project shapes against an installed HEonGPU built with USE_HIP=ON, with one main.cpp holding int main(){return 0;} and a CMakeLists.txt that differs only in the project() line:
```
cmake_minimum_required(VERSION 3.26)
project(consumer LANGUAGES CXX HIP) # or: LANGUAGES CXX
find_package(HEonGPU REQUIRED)
add_executable(app main.cpp)
target_link_options(app PRIVATE -Wl,-u,_ZN7heongpu10modInverseEmm)
target_link_libraries(app PRIVATE HEonGPU::heongpu)
```
```
cmake -S consumer -B build -DCMAKE_BUILD_TYPE=Release \
-DCMAKE_PREFIX_PATH=<install-prefix> \
-DCMAKE_CXX_COMPILER=/usr/bin/c++ -DCMAKE_HIP_ARCHITECTURES=gfx90a
cmake --build build
head -1 build/CMakeFiles/app.dir/link.txt
```
With HIP enabled the link driver is /opt/rocm/lib/llvm/bin/clang++ and the failure is undefined __hip_gpubin_handle_*, which target_link_options(app PRIVATE --hip-link) fixes. Without it the driver is /usr/bin/c++ and the failure is unrecognized command-line option '-fgpu-rdc', which no target-level setting fixes.
The library build and its tests are unaffected, since this commit changes only documentation and comments:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
```
The note for a consumer whose project does not enable the HIP language said the library cannot rescue the link from its side. That is false, and an earlier commit message in this series repeated it, so here is the corrected statement: a package configuration file is included at the caller's file scope, so enable_language(HIP) as the first line of cmake/Config.cmake.in enables the language in the consumer's scope. Measured against the installed package with the configuration file as the only difference (ROCm 7.2.1, gfx90a, CMake 4.0.3), a project(app LANGUAGES CXX) consumer with a plain C++ executable and only --hip-link on its target goes from `unrecognized command-line option '-fgpu-rdc'` to a clean configure and link driven by /opt/rocm/lib/llvm/bin/clang++, and the same holds when find_package is called from an add_subdirectory. The installed package deliberately does not do this, so the requirement on the consumer stands and the advice around it is unchanged. enable_language must be called in file scope and in the highest directory common to all targets using the language, so a consumer that wraps find_package in a function silently gets no language and the same failing C++ link, and every consumer would pay HIP compiler detection at configure time, 0.9s to 2.9s on the tree measured here. What the documentation now says is what the installed package does, not what the library is unable to do. Documentation only. No source, build system or behaviour change, so the library and its tests are unaffected. This change was prepared with the assistance of an AI coding agent. Test Plan: Both consumers are the same source tree, a C++-only project with a plain C++ executable that links HEonGPU::heongpu and sets --hip-link on its own target, and both use the same installed prefix. The alternate prefix differs from it in one file: enable_language(HIP) prepended to lib/cmake/HEonGPU-1.1/HEonGPUConfig.cmake. ``` cmake -S consumer -B consumer/b1 -DCMAKE_PREFIX_PATH="$PREFIX" -DCMAKE_CXX_COMPILER=/usr/bin/c++ cmake --build consumer/b1 # fails: c++: error: unrecognized command-line option '-fgpu-rdc' cmake -S consumer -B consumer/b2 -DCMAKE_PREFIX_PATH="$ALT_PREFIX" -DCMAKE_CXX_COMPILER=/usr/bin/c++ cmake --build consumer/b2 # links; head -c40 consumer/b2/CMakeFiles/app.dir/link.txt is /opt/rocm/lib/llvm/bin/clang++ ```
SmallForwardNTT and SmallInverseNTT are __device__ functions called from kernels in other translation units, and the AMD build can resolve that in two ways: relocatable device code (-fgpu-rdc, the counterpart of the CUDA_SEPARABLE_COMPILATION the NVIDIA build sets), or keeping the definitions in the header where every device compiler that needs them sees them. This build used the first; it now uses the second again.
Both work. Measured on gfx90a with ROCm 7.2.1, 20 of 20 test suites pass either way and the BFV, CKKS and TFHE benchmarks are equal within run-to-run noise, because the device link inlines the cross-unit calls just as the header does. The performance objection recorded when the header form was first chosen was wrong and was retracted; it is not the reason here.
The reason is what each form asks of a project that links the installed library. Relocatable device code leaves no complete device image in the archive, so the device link happens wherever the library is linked, which puts -fgpu-rdc on the library's link interface and requires the consuming project to enable the HIP language. A project that does not gets its link driven by the C++ compiler, which rejects the flag, and nothing set on the consuming target makes up for it. Keeping the definitions in the header puts nothing on the link interface, exactly as the NVIDIA path puts nothing there.
Measured against a fresh install of this change: a consumer that is project(app LANGUAGES CXX) with a plain C++ executable linking HEonGPU::heongpu and no further options configures, links with /usr/bin/c++, runs, and carries a complete gfx90a code object.
The cost is that the definitions no longer live in a translation unit of their own: lib/kernel/small_ntt.cu is removed and its contents sit in the header behind a guard that admits only a device compiler, since heongpu.hpp reaches this header from host translation units too. The NVIDIA path compiles the same definitions from the header, and nothing else about it changes.
Authored with the assistance of Claude, an AI assistant by Anthropic.
Test Plan:
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \
-DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build -j$(nproc)
ctest --test-dir build --output-on-failure
# 20/20 suites pass on gfx90a (AMD Instinct MI250X / MI250), ROCm 7.2.1.
cmake -S . -B build-install -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_BUILD_TYPE=Release -DCMAKE_INSTALL_PREFIX=/tmp/heongpu
cmake --build build-install -j$(nproc)
cmake --install build-install
cat > /tmp/consumer/CMakeLists.txt <<'CM'
cmake_minimum_required(VERSION 3.26)
project(app LANGUAGES CXX)
find_package(HEonGPU REQUIRED)
add_executable(app main.cpp)
target_link_libraries(app PRIVATE HEonGPU::heongpu)
CM
cmake -S /tmp/consumer -B /tmp/consumer/build -DCMAKE_PREFIX_PATH=/tmp/heongpu
cmake --build /tmp/consumer/build && /tmp/consumer/build/app
cmake -S . -B build-cuda -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=OFF \
-DCMAKE_CUDA_ARCHITECTURES=80 -DCMAKE_CXX_COMPILER=g++-13 \
-DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON \
-DHEonGPU_BUILD_BENCHMARKS=ON
cmake --build build-cuda -j$(nproc)
# CUDA 12.8, nvcc: 81 objects and 42 executables, no NVIDIA GPU on this host.
```
The host side of the library assumed a glibc host in three places, so a
Windows build stopped before any device code was reached.
The memory pool sized its host arena with sysinfo(), which exists only on
Linux. Query the closest Windows analogue through GlobalMemoryStatusEx():
MEMORYSTATUSEX::ullAvailPhys is the physical memory still available to the
process. It is not the same quantity as freeram * mem_unit, which excludes
the page cache while ullAvailPhys counts reclaimable standby pages, but it
is the figure a Windows caller would size an arena from. The POSIX branch
is untouched, so behaviour there is unchanged. random.cuh also included
<sys/sysinfo.h> but uses nothing from it, so that include is removed
rather than guarded.
Two more spellings kept the build from finishing under clang-cl.
u_int32_t is a BSD typedef that the Microsoft headers do not provide;
uint32_t is the same type and is what the surrounding code already uses.
The GNU-style -g and -O3 additions are unused arguments for clang-cl,
which already gets /Zi and /O2 /Ob2 from CMake, and an unused argument is
fatal in a subproject that compiles with warnings as errors, so they are
now added only for compilers that are not MSVC-like.
This change was prepared with the assistance of an AI coding agent.
Test Plan:
Linux (code path unchanged by this commit):
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
-DHEonGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)
cd build && ctest --output-on-failure
```
Windows 11, clang-cl host compiler, Radeon 8060S (gfx1151):
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1151 \
-DHEonGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j6
```
The library and all 15 test executables compile and link.
The CRT constants that RNS composition depends on are built with GMP, and
the moduli reach it through the *_ui entry points. Those take an unsigned
long, which is 64 bits under the LP64 data model but only 32 under LLP64,
so on Windows every prime wider than 32 bits was truncated to its low half
before GMP ever saw it. calculate_Mi(), calculate_M() and
calculate_upper_half_threshold() then described a different modulus chain
than the one actually in use, as did the BFV coefficient-divided-by-plain-
modulus table.
The symptom was narrow enough to be misleading. Encoding produced correct
residues, and BFV encoding round-tripped correctly because it works modulo
the plain modulus alone, which is small enough to survive the truncation.
Everything that composes an RNS value back into an integer -- BFV and CKKS
decryption, and CKKS decoding -- instead returned values spread across the
whole modulus product. Importing the value as a single 64-bit word with
mpz_import is exact under both data models and leaves LP64 results bit for
bit unchanged.
The same width assumption appears in the two CKKS helpers that apply a
gaussian-integer constant, add_constant_plain_ckks_v2() and
multiply_const_plain_ckks_v2(), which move a modulus and a residue through
NTL's long conversions. They now go through the byte representation, again
unchanged under LP64. No test executable reaches either helper.
add_constant_plain_ckks_v2() is called only from add_plain_v2(), which in
turn is used by eval_mod(), gen_power() and
evaluate_poly_from_polynomial_basis(), and multiply_const_plain_ckks_v2()
only from multiply_plain_v2() and scale_up_ckks(); the test suite calls
none of these, and the runs below configure tests only, with
HEonGPU_BUILD_EXAMPLES left at its default of OFF. That hunk is therefore
a fix by inspection, and the results below do not exercise it.
This change was prepared with the assistance of an AI coding agent.
Test Plan:
Linux (results unchanged by this commit):
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
-DHEonGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)
cd build && ctest --output-on-failure
```
Windows 11, clang-cl host compiler, Radeon 8060S (gfx1151):
```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1151 \
-DHEonGPU_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release -G Ninja
cmake --build build -j6
cd build && ctest --output-on-failure
```
All 15 test executables pass, 20 of 20 cases. Before this commit only
bfv_encoding passed; the other 14 failed on mismatched values.
Collaborator
Author
|
To approve this port, leave a comment containing this line by itself: To send it back to the porter instead: Both are commands because GitHub greys out the Approve and Request Changes buttons for a pull request's author, and this PR was opened on your credentials. Your latest command is the one that stands, and a command quoted in a code fence -- like the two above -- is ignored. Either box works: a review comment (Review changes -> Comment, or The title and body above are what gets opened upstream, verbatim, so approving here approves all three: the code, the title and the body. Anything pushed afterwards, or any edit to the title or body, voids it and needs a fresh one. |
Collaborator
Author
|
Review finding: the rocThrust rationale on the
Suggested resolution, cheap while Repro commands are in the MOAT session log for 2026-08-17; the two key ones: Note: this finding was investigated and drafted by an AI assistant; posted at Jeff Daily's request. |
jeffdaily
commented
Aug 17, 2026
jeffdaily
left a comment
jeffdaily
left a comment
Collaborator
Author
There was a problem hiding this comment.
Help me understand why we have src/include/heongpu/cuda_to_hip.h but also thirdparty/hip_compat/cuda_runtime.h et al that seem to repeat the same hipify stuff.
A big part of this port seems to be the rmm_hip_stub . Doesn't the ROCm-DS project have an implementation of rmm? See https://github.com/AMD-Ecosystem/hipMM.
/moat changes-requested
The comments in the test, example, and benchmark CMake files and the sentence in docs/advanced_topics.rst said sources including the HEonGPU headers must be compiled as HIP because the headers pull in rocThrust. That attribution is wrong: rocThrust's thrust/host_vector.h -- the only Thrust header these headers reach -- compiles cleanly under a plain host compiler, matching NVIDIA Thrust, where host containers never needed the GPU compiler. The real constraint is device code in the public header chain: __umul64hi in GPU-NTT's modular_arith.cuh and warpSize in util.cuh's warp_reduce are device intrinsics no host compiler defines. The same headers require nvcc in a CUDA build for the same reason, so the build arrangement is unchanged; only the stated reason was off. No functional change. Prepared with the help of an AI coding agent. Test Plan: Verified the attribution before rewording, on gfx90a with ROCm 7.2.1: ``` # rocThrust alone under a plain host compiler: compiles g++ -std=c++17 -c thrust_host.cpp -isystem <rocm>/include -o /dev/null # the public header chain under a plain host compiler with all include # paths present: fails on __umul64hi and warpSize, not on thrust g++ -std=c++17 -DUSE_HIP -D__HIP_PLATFORM_AMD__=1 -c heongpu_hdr.cpp \ -I src/include -I thirdparty/rmm_hip_stub/include \ -I thirdparty/GPU-NTT/src/include -I thirdparty/GPU-FFT/src/include \ -I thirdparty/RNGonGPU/src/include \ -isystem thirdparty/hip_compat -isystem <rocm>/include -o /dev/null ``` Full rebuild and suite at this commit: ``` cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \ -DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON cmake --build build -j64 ctest --test-dir build # 100% passed, 20/20 ```
The USE_HIP branch of target_link_libraries(heongpu ...) repeated the whole upstream list to change one entry. A generator expression keeps the single upstream block and selects hip::hiprand or CUDA::curand in the same position, so the link line is unchanged on both platforms. The duplicated target_compile_options block is dropped entirely: the only difference was the HIP branch omitting the two --generate-line-info entries, and those are already guarded by $<COMPILE_LANGUAGE:CUDA>, which is false in a HIP build. Upstream's block is restored verbatim. No functional change on either platform. Prepared with the help of an AI coding agent. Test Plan: ``` cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_BUILD_TYPE=Release -DHEonGPU_BUILD_TESTS=ON \ -DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON cmake --build build -j64 # clean, no warnings or errors ctest --test-dir build # 100% passed, 20/20 HIP_VISIBLE_DEVICES=0 ./build/bin/examples/basic/1_basic_bfv # runs, exit 0 ```
cuda_to_hip.h and thirdparty/hip_compat/cuda_runtime.h carried two copies of the same CUDA-to-HIP alias table. Every other header in the tree already includes cuda_runtime.h by name and picks up the shim on HIP builds, and util.cuh was the alias header's only consumer, so this keeps the shim as the single mechanism: its table gains the six aliases only the alias header had (cudaMemGetInfo, the function-attribute and occupancy trio, the device-limit pair), util.cuh includes cuda_runtime.h like its siblings, and cuda_to_hip.h is deleted. The deliberate omission of curand_kernel.h from util.cuh is preserved and its reason documented at the include site: host translation units reach util.cuh through the umbrella header before device_launch_parameters.h, and that inclusion order does not compile under nvcc. Prepared with the help of an AI coding agent. Test Plan: ``` # HIP: clean build, tests, on MI250X (gfx90a) cmake -S . -B build -DUSE_HIP=ON -DCMAKE_BUILD_TYPE=Release \ -DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON cmake --build build -j64 # clean ctest --test-dir build # 100% passed, 20/20 # CUDA: compile-only no-regression (nvcc 12.8, arch 80) -- exercises the # host-TU inclusion order through lib/heongpu.cpp cmake -S . -B build-cuda -DUSE_HIP=OFF -DCMAKE_CUDA_ARCHITECTURES=80 ... cmake --build build-cuda -j64 # rc=0, zero error lines ```
The USE_HIP block forced CMAKE_HIP_ARCHITECTURES to gfx90a whenever the user did not set it, which silently built for the wrong card on every other machine and defeated CMake's own detection. CMake auto-detects the architecture from the GPU in the build machine when the variable is unset, so the forced default is removed; setting the variable remains the way to cross-compile. Docs updated in the next commit alongside the version floor. Prepared with the help of an AI coding agent. Test Plan: ``` # no -DCMAKE_HIP_ARCHITECTURES: configure reports the detected gfx90a # and the full build + suite pass on MI250X cmake -S . -B build -DUSE_HIP=ON -DCMAKE_BUILD_TYPE=Release \ -DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON cmake --build build -j64 # clean ctest --test-dir build # 100% passed, 20/20 ```
The docs claimed ROCm 6.0 or higher, which nothing in the port's testing backs: every validated configuration ran ROCm 7.2 or newer. State 7.2 as the floor. The architecture instructions now say CMAKE_HIP_ARCHITECTURES is auto-detected from the GPU in the build machine and is set explicitly to cross-compile, matching the previous commit's behavior. Prepared with the help of an AI coding agent. Test Plan: Documentation-only; rendered locally and covered by the builds in the previous two commits. ``` cmake --build build -j64 && ctest --test-dir build # 20/20 ```
Collaborator
Author
|
Both review-body questions, then the round summary. (Drafted by an AI assistant.) Why both cuda_to_hip.h and hip_compat/cuda_runtime.h? There was no good reason left. The two mechanisms started distinct — hip_compat shims resolve rmm_hip_stub vs ROCm-DS hipMM: hipMM (AMD-Ecosystem/hipMM) is a real RMM port — derived from RMM 25.10, MIT, deliberately mirroring RMM's structure and API naming — so replacing the bundled stub with it is credible and would shrink the port. It is also a scope change: it swaps a ~5-type internal stub for an external dependency the upstream proposal would then carry (fetch/pin story included), and costs a build+revalidation round on every platform. I've registered it as a deferral ( Round summary — two rounds since the
Each line thread has a reply naming its fix commit. Reviewed and revalidated on linux-gfx90a at |
The HIP build carried a minimal reimplementation of the RMM types the library uses (device_uvector, device_buffer, pool/statistics/pinned memory resources). hipMM -- the ROCm-DS port of RMM, derived from RMM 25.10 -- keeps RMM's target name, rmm:: namespace, and header layout, so the sources compile unchanged against the real library and the stub is deleted. The HIP arm of the FetchContent block now mirrors the CUDA arm: rapidsai/rmm there, AMD-Ecosystem/hipMM here. hipMM's own sources expect a HIP-capable compiler (its hip dependency injects -x hip into their compile flags), so they are marked LANGUAGE HIP after the fetch, the same way this port compiles its own GPU sources; nothing propagates to plain C++ consumers of the library. hipMM installs its own rmm CMake package, so the installed HEonGPU config's find_dependency(rmm) resolves exactly as it does on CUDA. Prepared with the help of an AI coding agent. Test Plan: ``` # HIP (MI250X gfx90a, ROCm 7.2.1): clean fetch+build, full suite, example cmake -S . -B build -DUSE_HIP=ON -DCMAKE_BUILD_TYPE=Release \ -DHEonGPU_BUILD_TESTS=ON -DHEonGPU_BUILD_EXAMPLES=ON -DHEonGPU_BUILD_BENCHMARKS=ON cmake --build build -j64 # clean ctest --test-dir build # 100% passed, 20/20 ./build/bin/examples/basic/1_basic_bfv # exit 0 # installed package: find_package(HEonGPU) consumer builds and runs # context generation on GPU against the hipMM-backed install # CUDA: compile-only no-regression (nvcc 12.8, arch 80): rc=0, real RMM # fetched exactly as before ```
hipMM publishes no release tags and release/rocmds-26.03 is its moving development branch, so the fetch is pinned to the exact commit the port was validated against; the comment names the branch line for whoever bumps it. Prepared with the help of an AI coding agent. Test Plan: ``` cmake -S . -B build -DUSE_HIP=ON ... # fresh tree, fetch resolves the pin cmake --build build -j64 # clean ctest --test-dir build # 100% passed, 20/20 ```
Collaborator
Author
|
Following up on the rmm question: done, per your direction, rather than deferred. (Drafted by an AI assistant.)
Validated on gfx90a from a fresh tree: clean build, 20/20 suite, example run, CUDA compile gate clean (real RMM untouched), and an installed-package consumer builds and runs GPU context generation against the hipMM-backed install. One heads-up: hipMM's README says Linux-only. The Windows revalidation (gfx1151) will answer whether that holds in practice; if it does, that comes back here as a finding with options rather than a quiet revert. |
The HIP build fetches hipMM in place of RMM, and hipMM in turn fetches rapids_logger. Both are built shared by default, and neither one works that way with an MSVC-style toolchain. Their export macros expand to the ELF visibility attribute under __GNUC__ and to nothing otherwise, so the DLLs export no symbols and every executable link ends in undefined rmm:: and rapids_logger:: symbols. rapids_logger also hides its spdlog symbols with the GNU-only linker option --exclude-libs, whose operand lld-link reads as an input file name, so its own link is the first thing to fail. Building both static for a Windows HIP build sidesteps both problems, and keeps their code in the executables rather than in DLLs that Windows, having no RPATH, would need copied next to every binary before it could start. The setting is a directory-scoped normal variable so that it covers the fetch and leaves the rest of the project alone. Under CMP0077 that also keeps RMM's own option(BUILD_SHARED_LIBS ... ON) from writing the cache, which is what otherwise flips later fetches, GoogleTest included, to shared. A cache entry would work here as well, but it would override a -DBUILD_SHARED_LIBS the user set for their own build and would outlive a later reconfigure of the same tree with USE_HIP=OFF. Linux and the CUDA path are untouched: the block is guarded on both the HIP build and an MSVC-style compiler. A fix for the --exclude-libs call is proposed upstream at AMD-Ecosystem/rocmds-logger#2. The missing exports are a separate defect, so the static build stays useful after it lands. This change was prepared with the assistance of an AI coding agent. Test Plan: Windows 11, Radeon 8060S (gfx1151), ROCm 7.14 with clang-cl as the C, C++ and HIP compiler. Configure and build with no extra options beyond the toolchain settings: ``` cmake -S . -B build -G Ninja -D USE_HIP=ON -D CMAKE_HIP_ARCHITECTURES=gfx1151 \ -D HEonGPU_BUILD_TESTS=ON -D HEonGPU_BUILD_EXAMPLES=ON -D HEonGPU_BUILD_BENCHMARKS=ON \ -D CMAKE_BUILD_TYPE=Release cmake --build build -j16 ``` All 15 test executables, all 3 benchmarks and every example link, and no DLL is produced or has to be staged. Then, on the GPU: ``` ctest --test-dir build --output-on-failure ```
The memory pool defaults in defines.h and the matching bullet in the advanced topics guide quote percentages that no longer agree with the constants they describe. The device pool is initialized at 0.9 and capped at 0.95, and the host pool at 0.3 and 0.4, but the surrounding text still says 50, 80, 10 and 20 percent, which sends a reader looking for a knob that is not there. Bring the comment text and the prose in line with the values that are actually compiled in. Only comments and documentation change; every constant keeps its current value, so there is no behavioral difference. This change was prepared with assistance from an AI coding agent. Test Plan: Comment and documentation text only. Confirm that no value moved: ``` git diff -U0 -- src/include/heongpu/kernel/defines.h docs/advanced_topics.rst ``` Every changed line is a comment or a sentence; the four constants read 0.9f, 0.95f, 0.3f and 0.4f before and after.
The device memory pool is sized as a fraction of the memory the runtime reports for the device. On a discrete card that is dedicated VRAM, so reserving 90% of it up front costs nothing that anything else wants. On a GPU that shares its memory with the system, the same query returns the whole machine's memory, so the default reserves 90% of everything the machine has. That reservation is expensive rather than merely large: context creation grows from about a second to about ten, and memory-heavy workloads can take roughly twice as long. Describe the effect and the remedy where the pool defaults are already documented, in the advanced topics guide, and leave a short pointer in the AMD GPU section of the README. The remedy is a per-application setting and needs no rebuild: a MemoryPoolConfig with a smaller initial_device_fraction, or an explicit initial_device_bytes, which the existing memory pool example already demonstrates. An initial pool of 10% or less, or a fixed gigabyte or two, performs as well as no pool at all, while 50% is already about twice as slow as 10% on heavy work and so is not a safe middle ground. The compiled defaults are deliberately left alone. They remain the right choice for a discrete GPU, and this is guidance for parts that share system memory rather than a reason to change what everyone else gets. This change was prepared with assistance from an AI coding agent. Test Plan: Documentation only. No header, source, or build file changes, so no build artifact can differ: ``` git show --stat HEAD ``` The figures quoted were measured on an AMD Radeon 8060S (gfx1151, unified memory) under ROCm 7.14, comparing the default pool against `initial_device_fraction` values of 0.5 and 0.1 and against `use_memory_pool = false`.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This adds a HIP build so HEonGPU runs on AMD GPUs through ROCm, alongside the existing CUDA build. The CUDA path is unchanged by default: everything new is behind
-D USE_HIP=ON, andUSE_CUDA=ONremains the default.The bulk of the change is build system and a compatibility header.
src/include/heongpu/cuda_to_hip.hmaps the CUDA runtime and library calls the project uses onto their HIP equivalents, so the kernel and host sources keep their existing CUDA spelling.src/CMakeLists.txtand the example, benchmark and test lists grow a HIP branch next to the CUDA one, and the pinned GPU-NTT, GPU-FFT and RNGonGPU submodules are built for AMD through patches underthirdparty/patches/.cmake/Config.cmake.incarries the same choice into the installed package. Hardcoded warp sizes are replaced withwarpSizewhere the value has to follow the hardware, since AMD parts are 64 lanes on CDNA and 32 on RDNA.Defects found while porting
Four of these are ordinary porting fixes. The first is not, and it needs your attention before anything else here.
The Box-Muller error sampler was undefined behaviour, and fixing it changes NVIDIA results. The sampler wrote
static_cast<T>(z)whereTis the unsigned residue type andzis a signed sample, relying on the following+ (flag & modulus)to turn a negative sample into its representative. Converting an out-of-range floating-point value to an unsigned integer is undefined, and the two back ends resolve it differently: NVIDIA'scvt.rzi.u64.f64saturates a negative source to zero, while the AMD lowering of a 64-bitfptouikeeps the low 32 bits of the two's-complement value. On AMD that turned roughly half of every error polynomial into values near 2^32, so every ciphertext decrypted to noise. Routing the truncation through the signed type first makes it well defined on both back ends. This is deliberately not behaviour-neutral on NVIDIA: it replaces the saturated zero with the correct negative representative, which is the value the surrounding arithmetic already assumes. Since a saturated zero makes the sample congruent to zero rather thanq - 2, the CUDA build has been drawing from a thinner error distribution than intended, which is a security-relevant property of an RLWE scheme rather than a cosmetic difference. Please review this one on its own terms; we would rather you decide how it should behave than have us pick for you.The remaining fixes are AMD-only or test-only:
thirdparty/patches/GPU-NTT.patch.test/test_bfv_addition.cppcomputed its expected value without reducing at the plaintext modulus, so the reference itself was out of range. This corrects the test, not the library.MemoryPool::get_host_avaliable_memory()and two headers used<sys/sysinfo.h>, which is Linux-only. Guarded, withGlobalMemoryStatusExon Windows.mpz_*_uientry points takeunsigned long, which is 32 bits under LLP64. Every coefficient modulus wider than 32 bits was silently truncated before GMP saw it, so on Windows the CRT constants described a different modulus chain than the one in use. Values now go throughmpz_import, which is exact under both data models and leaves LP64 results bit for bit unchanged.Testing
ctestpasses 20 of 20 on each of the following, run twice from a clean build:Both wavefront widths are covered, which is what the
warpSizechanges above need. Examples and benchmarks were also built and run:1_basic_bfv,2_basic_ckks,15_basic_tfhe,bootstrapping/3_ckks_bit_bootstrapping,bootstrapping/5_ckks_regular_bootstrapping_v2,mpc/1_multiparty_computation_bfvandbenchmark/tfhe_benchmark, with decrypted values checked rather than exit codes alone. The eight TFHE gate outputs were verified bit for bit against their truth tables.The CUDA build was reconfigured and rebuilt with nvcc at this revision to confirm it still compiles and links.
One known gap: on Windows the
9_multi_stream_usage_way1example does not link, because the OpenMP flags are not propagated to the linker by that generator. The library, all tests and the other examples build and run there.Note on authorship
This port was prepared with the assistance of an AI coding agent, with the results above measured on the listed hardware. Please review it as you would any other outside contribution.
This pull request was prepared with the help of an AI assistant acting as a coding agent and was read and approved by a person before it was opened. It comes from an ongoing effort to add AMD GPU support to widely used CUDA projects, one repository at a time: https://github.com/AMD-Ecosystem/moat -- that repository describes how the work is done and what a person checks before anything is submitted.
If you would rather not receive pull requests from this effort, say so here or open an issue at https://github.com/AMD-Ecosystem/moat/issues/new/choose and we will close this and stop. That can cover this repository alone or everything you own, whichever you prefer.