ModelDensityGPU's gather had arithmetic of its own: per-thread offsets in Cartesian, the atom's
nearest image to the brick, r2 = x x + y y + z z, the radius squared in float. It now does what
UseDensityBoxesByPlane + AtomDensity (ModelGrid.cpp) and ModelDensityMetal.metal do, in their order:
- each brick's atoms carry their box (c - d, wrapped start, 2 d); a point is in the box by a
subtraction, and f = centre - (u, v, w) in the box's own unwrapped index;
- x = m00 fx + (m01 fy + m02 fz), y = m11 fy + m12 fz, z = m22 fz with gemmi's orth_n
(orth * (1/n), as Grid sets it) in float; r2 = (y y + z z) + x x against max_r2 = float(r^2)
rounded from the double radius (new ModelDensityAtom field);
- gemmi's expapprox, the five Gaussians in order, atoms in model order per point. Device code is
already compiled with --fmad=false (top-level CMakeLists), so nothing is contracted.
- The centres are no longer wrapped into the cell (ModelStructureFactorsGPU, RigidBodyGPU's
place_kernel): float(fractional * n), as MakeDensityBox holds them; the mask keeps wrapped
fractional coordinates. The 'narrow' nearest-image branch is gone: Supports() already leaves
only boxes narrower than the cell.
The target is PutModelDensityOnGrid() with no orbits to symmetrize (one copy of the content, which
is what the gather grids; symmetry stays composed in reciprocal space). A rounding tie in the box
centre c (float on the GPU, double on the CPU) changes no point: either box covers the sphere with
half a spacing to spare. New ModelDensityGPU_BitIdenticalToCPU ([ModelValidation][gpu]) memcmps the
grids: triclinic P1 and P 21 21 21 cells, 3000 atoms over -0.5..1.5 of the cell, a quarter
anisotropic - 0 points differ.
RTX 5080, alternating base (8d199cceb) / this, under the GPU lock:
- gather kernel (nsys): 4nwv --model median 2.06 -> 2.55 ms, total 0.79 -> 0.97 s;
9s02 --mr median 0.30 -> 0.37 ms, total 0.94 -> 1.17 s (+~25%, the box test and per-point
index arithmetic).
- wall, two each: 4nwv --model 76.0/68.9 -> 74.4/71.0 s, --mr 71.3/71.1 -> 70.9/71.6 s;
9s02 --model 25.1/22.6 -> 23.3/23.5 s, --mr 23.3/23.6 -> 23.3/23.2 s; 5reo --model
6.2/5.5 -> 5.5/5.7 s. Model validation 4nwv 12.4/12.3 -> 12.3/11.7 s; MR 4nwv 15.4/15.2 ->
15.4/15.1 s, 9s02 2.3/2.4 -> 2.4/2.4 s: no change beyond noise.
- outputs (each build run-to-run identical): p.hkl identical on all; R-work/R-free and every
decision unchanged (4nwv validation 0.2982/0.2969, 9s02 MR AGREES P 21 21 2 0.2503/0.2517,
4nwv MR UNDECIDED I 2 2 2 / I 21 21 21). Moved: MODEL_FIT_SIGMA 9s02 99.21 -> 101.68, 5reo
103.96 -> 104.98; MR placement Z 9s02 142.1 -> 145.6, 4nwv 46.7 -> 46.2; 4nwv anomalous site
4.01 -> 4.02 sigma; 4nwv placed model 1322/125462 lines in the last digit (2fo-fc max |d| 0.04
rms, CC 0.99999994); 9s02/5reo models identical, maps to 2e-5 rms.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PXc1zD1yxHqaR1DGdCTHL4
123 lines
5.4 KiB
C++
123 lines
5.4 KiB
C++
// SPDX-FileCopyrightText: 2026 Filip Leonarski, Paul Scherrer Institute <filip.leonarski@psi.ch>
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#include <catch2/catch_all.hpp>
|
|
#include "../common/CUDAWrapper.h"
|
|
|
|
#ifdef JFJOCH_USE_CUDA
|
|
|
|
#include <array>
|
|
#include <cstring>
|
|
#include <random>
|
|
#include <vector>
|
|
|
|
#include <gemmi/dencalc.hpp>
|
|
#include <gemmi/it92.hpp>
|
|
|
|
#include "../image_analysis/structure_refinement/ModelDensityGPU.h"
|
|
#include "../image_analysis/structure_refinement/ModelGrid.h"
|
|
#include "../image_analysis/structure_refinement/ModelStructureFactorsGPU.h"
|
|
|
|
// The CUDA gather is the CPU's density walk, bit for bit (ModelDensityGPU.h): PutModelDensityOnGrid() with no
|
|
// orbits to symmetrize - one copy of the content, which is what the gather puts on the grid - against
|
|
// ModelDensityGPU fed as ModelStructureFactorsGPU feeds it. The grids are compared as bit patterns, as in
|
|
// ModelDensityMetalTest.cpp.
|
|
|
|
namespace {
|
|
// Atoms scattered over the whole cell and a half cell beyond it on every side - so that boxes wrap at every
|
|
// face and centres lie outside the cell - of a few elements, a quarter of them anisotropic, with occupancies
|
|
// below one.
|
|
gemmi::Structure random_crystal(const gemmi::UnitCell &cell, const char *group) {
|
|
std::mt19937 rng(11);
|
|
std::uniform_real_distribution<double> frac(-0.5, 1.5), b(8.0, 60.0), occ(0.3, 1.0), u(0.05, 0.4);
|
|
const gemmi::El elements[] = {gemmi::El::C, gemmi::El::N, gemmi::El::O, gemmi::El::S};
|
|
gemmi::Structure st;
|
|
st.cell = cell;
|
|
st.spacegroup_hm = group;
|
|
st.setup_cell_images();
|
|
gemmi::Model model(1);
|
|
gemmi::Chain chain("A");
|
|
gemmi::Residue res;
|
|
res.name = "UNK";
|
|
res.seqid = gemmi::SeqId(1, ' ');
|
|
for (int i = 0; i < 3000; i++) {
|
|
gemmi::Atom a;
|
|
a.name = "X";
|
|
a.element = elements[i % 4];
|
|
a.pos = cell.orthogonalize(gemmi::Fractional(frac(rng), frac(rng), frac(rng)));
|
|
a.b_iso = static_cast<float>(b(rng));
|
|
a.occ = static_cast<float>(occ(rng));
|
|
if (i % 4 == 3) {
|
|
const float d1 = static_cast<float>(u(rng)), d2 = static_cast<float>(u(rng)),
|
|
d3 = static_cast<float>(u(rng));
|
|
a.aniso = {d1, d2, d3, 0.3f * std::min(d1, d2), -0.2f * std::min(d1, d3), 0.1f * std::min(d2, d3)};
|
|
}
|
|
res.atoms.push_back(a);
|
|
}
|
|
chain.residues.push_back(res);
|
|
model.chains.push_back(chain);
|
|
st.models.push_back(model);
|
|
return st;
|
|
}
|
|
|
|
void check_same_grid(const gemmi::UnitCell &cell, const char *group, double d_min) {
|
|
const gemmi::Structure st = random_crystal(cell, group);
|
|
gemmi::DensityCalculator<gemmi::IT92<float>, float> dc;
|
|
dc.d_min = d_min;
|
|
dc.rate = 1.5;
|
|
dc.set_grid_cell_and_spacegroup(st);
|
|
dc.set_refmac_compatible_blur(st.models[0]);
|
|
PutModelDensityOnGrid(dc, st.models[0], {}, 4, false);
|
|
const gemmi::Grid<float> &cpu = dc.grid;
|
|
|
|
ModelDensityGrid grid;
|
|
grid.nu = cpu.nu;
|
|
grid.nv = cpu.nv;
|
|
grid.nw = cpu.nw;
|
|
for (int i = 0; i < 3; i++)
|
|
for (int j = 0; j < 3; j++) {
|
|
grid.orth[3 * i + j] = cell.orth.mat[i][j];
|
|
grid.frac[3 * i + j] = cell.frac.mat[i][j];
|
|
}
|
|
std::vector<ModelDensityAtom> atoms;
|
|
std::vector<int> mask_atom;
|
|
std::vector<float> mask_radius;
|
|
ModelDensityAtoms(st.models[0], dc, atoms, mask_atom, mask_radius);
|
|
REQUIRE(ModelDensityGPU::Supports(grid, atoms));
|
|
// ModelStructureFactorsGPU::Compute()'s centres: the fractional position times the grid size.
|
|
std::vector<std::array<float, 4>> pos;
|
|
for (const gemmi::Atom &a : st.models[0].chains[0].residues[0].atoms) {
|
|
const gemmi::Fractional f = cell.fractionalize(a.pos);
|
|
pos.push_back({static_cast<float>(f.x * grid.nu), static_cast<float>(f.y * grid.nv),
|
|
static_cast<float>(f.z * grid.nw), 0.0f});
|
|
}
|
|
|
|
CudaStream stream;
|
|
ModelDensityGPU density(stream, atoms.size(), ModelDensityGPU::PairBound(grid, atoms),
|
|
ModelDensityGPU::Bricks(grid.nu, grid.nv, grid.nw));
|
|
density.SetAtoms(grid, atoms);
|
|
const size_t n = cpu.data.size();
|
|
CudaDevicePtr<float4> pos_d(pos.size());
|
|
CudaDevicePtr<float> grid_d(n);
|
|
REQUIRE(cudaMemcpyAsync(pos_d, pos.data(), pos.size() * sizeof(float4), cudaMemcpyHostToDevice, stream) == cudaSuccess);
|
|
density.Compute(pos_d, grid_d);
|
|
std::vector<float> gpu(n);
|
|
REQUIRE(cudaMemcpyAsync(gpu.data(), grid_d, n * sizeof(float), cudaMemcpyDeviceToHost, stream) == cudaSuccess);
|
|
REQUIRE(cudaStreamSynchronize(stream) == cudaSuccess);
|
|
size_t different = 0;
|
|
for (size_t i = 0; i < n; i++)
|
|
different += std::memcmp(&cpu.data[i], &gpu[i], sizeof(float)) != 0;
|
|
CHECK(different == 0);
|
|
}
|
|
}
|
|
|
|
TEST_CASE("ModelDensityGPU_BitIdenticalToCPU", "[ModelValidation][gpu]") {
|
|
if (get_gpu_count() == 0)
|
|
SKIP("No GPU");
|
|
// Triclinic, so that every term of the orthogonalization is there; then an orthorhombic group's grid.
|
|
check_same_grid(gemmi::UnitCell(61, 68, 77, 81, 96, 107), "P 1", 2.0);
|
|
check_same_grid(gemmi::UnitCell(70, 80, 90, 90, 90, 90), "P 21 21 21", 2.0);
|
|
}
|
|
|
|
#endif
|