"""Independent coordinate oracle, exact kernel host shim and closed pair budgets.""" import ctypes from contextlib import nullcontext from copy import deepcopy from dataclasses import replace import hashlib import json from pathlib import Path import subprocess import tempfile import unittest from unittest.mock import patch import numpy as np import run_case as rc import shear_axis_square as axis import shear_centered as candidate import test_shear_wave as legacy from test_axis_square import masks from CelerisLab.lbm.hierarchy import build_hierarchy from CelerisLab.lbm.blocks import BlockLayout2D from CelerisLab.lbm.refinement_tables import projected_explosion_tables from CelerisLab.lbm.refinement import ACTIVE, OUTER_GHOST from tests.unit import test_projected_explosion as old REPO = Path(__file__).resolve().parents[3] SOURCE = REPO/'src/CelerisLab/lbm/kernels/blocks.cu' def oracle(c, pair, centered=True): """Uses only dense coordinates/legal donors, never production neighbor/mask arrays.""" out = np.empty((9, len(pair.ghost_fine))) legal = set(map(int, pair.ghost_coarse)) for r, (donor, target) in enumerate(zip(pair.ghost_coarse, pair.ghost_fine)): y, x = divmod(int(donor), 128); fy, fx = divmod(int(target), 256) delta = np.array([(fx+.5)/2-x-.5, (fy+.5)/2-y-.5]) usable = [(128*(y-dy)+x-dx in legal and 128*(y+dy)+x+dx in legal) for dx,dy in [(1,0),(0,1)]] straight = ((x in (30,31,96,97) and 36 <= y < 92 and usable == [False,True]) or (y in (30,31,96,97) and 36 <= x < 92 and usable == [True,False])) for q, e in enumerate(old.E): grad = np.array([(c[q,y+dy,x+dx]-c[q,y-dy,x-dx])/2 if ok else 0. for (dx,dy),ok in zip([(1,0),(0,1)],usable)]) if q: grad -= e*(grad@e)/(e@e) out[q,r] = c[q,y,x] + (2 if centered and straight and np.all(e) else 1)*(grad@delta) return out class CenteredTests(unittest.TestCase): @classmethod def setUpClass(cls): cls.h = build_hierarchy(rc.config(128,100), axis.REFINEMENTS) cls.p = cls.h.pairs[0] cls.mask = candidate.routing(cls.h) cls.layout = BlockLayout2D.build(128,128,8,8,periodic=(True,True)) cls.ids, cls.offsets = projected_explosion_tables(cls.p,cls.layout,cls.h.kinds[0]) cls.neighbors = old.decode(cls.layout,cls.ids) def test_mask_constants_quartets_and_transpose(self): self.assertEqual(candidate.metadata(self.mask), dict(variant=candidate.VARIANT, mask_sha256='112de63ce9bafd2e205a9e7029ef6ce5683fe14407b6ac04444579c84960b7ea', ghost_rows=2112, selected_rows=1792, rule='diagonal projected correction x2; audited straight faces only')) c=np.random.default_rng(47).normal(0,.001,(9,128,128)) actual=oracle(c,self.p); baseline=oracle(c,self.p,False) base,g=old.reconstruct(c,self.p,self.neighbors,self.offsets) g[np.ix_(np.flatnonzero(np.all(old.E,axis=1)),np.flatnonzero(self.mask),[0,1])]*=2 expected=c.reshape(9,-1)[:,self.p.ghost_coarse]+np.einsum('qra,ra->qr',g,self.offsets) np.testing.assert_allclose(actual,expected,atol=1e-18,rtol=0) np.testing.assert_array_equal(actual[:,~self.mask.astype(bool)],baseline[:,~self.mask.astype(bool)]) for d in np.unique(self.p.ghost_coarse): rows=self.p.ghost_coarse==d np.testing.assert_allclose(actual[:,rows].mean(axis=1),c.reshape(9,-1)[:,d],atol=1e-18) const=np.broadcast_to(old.W[:,None,None],c.shape) np.testing.assert_array_equal(oracle(const,self.p),np.broadcast_to(old.W[:,None],actual.shape)) qswap=[list(map(tuple,old.E)).index(tuple(e[::-1])) for e in old.E] index={int(t):r for r,t in enumerate(self.p.ghost_fine)} swap=[index[int(t%256)*256+int(t//256)] for t in self.p.ghost_fine] np.testing.assert_array_equal(self.mask,self.mask[swap]) np.testing.assert_allclose(oracle(c[qswap].transpose(0,2,1),self.p),actual[qswap][:,swap],atol=1e-18) qreflect=[list(map(tuple,old.E)).index(tuple((-e[0],e[1]))) for e in old.E] reflect=[index[int(t//256)*256+255-int(t%256)] for t in self.p.ghost_fine] np.testing.assert_array_equal(self.mask,self.mask[reflect]) np.testing.assert_allclose(oracle(c[qreflect,:,::-1],self.p),actual[qreflect][:,reflect],atol=1e-18) # Central dependency corridors remain in a constant mask regime beyond radius8*2.5. for name,(x,y) in axis.GROUPS.items(): tangent = y if name.endswith(('left','right')) else x self.assertGreater(min(tangent+.25-36,92-(tangent+.75)),20) def test_actual_generated_kernel_shifted_unshifted(self): child=self.p.fine_layout((8,4)) # unequal owner strides, unused slots poisoned patched=candidate.patch_source(SOURCE.read_text(),self.mask) kernel=patched.split('__device__ __constant__ unsigned char centered_face_mask',1)[1] kernel='unsigned char centered_face_mask'+kernel.split('// Targets are receiver-local SOURCE slots',1)[0] kernel=kernel.replace('extern "C" __global__','extern "C"') code='#include \nusing fpxx=float;\n' code+=f'constexpr unsigned NQ=9,TOTAL_CELLS={np.prod(child.local_shape)};\n' code+='struct Id{unsigned x;}; Id blockIdx{0},blockDim{1},threadIdx{0};\n' code+='unsigned index_f(unsigned k,unsigned q){return q*TOTAL_CELLS+k;}\n' code+='int d_cx[]={0,1,-1,0,0,1,-1,1,-1},d_cy[]={0,0,0,1,-1,1,-1,-1,1};\n'+kernel original=SOURCE.read_text().split('extern "C" __global__ void BlockExplodeProjected(',1)[1].split('// Targets are receiver-local SOURCE slots',1)[0] code+='void Original('+original code+='extern "C" void invoke(float*c,float*f,uint64_t*t,uint64_t*d,unsigned n,unsigned long l,uint64_t*b,float*o){for(unsigned r=0;r