diff --git a/tests/validation/multigrid/spacetime_gpu.cu b/tests/validation/multigrid/spacetime_gpu.cu index 6b6de1c..b10f433 100644 --- a/tests/validation/multigrid/spacetime_gpu.cu +++ b/tests/validation/multigrid/spacetime_gpu.cu @@ -9,18 +9,36 @@ #include "operators/collision_srt.cuh" #include "streaming/pull_double_buffer.cuh" -#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0 -#error "Space-time experiment requires unshifted FP32 D2Q9 SRT double buffers without LES" +#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0 +#error "Space-time experiment requires FP32 D2Q9 SRT double buffers without LES" #endif __device__ __forceinline__ void collide(float* f, float omega) { float rho, ux, uy, feq[NQ]; compute_rho_u(f, rho, ux, uy); +#if USE_DDF_SHIFTING + // Same perturbation-preserving SRT branch as production BlockStep. + const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8]; + const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8]; + const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8]; + compute_feq_from_momentum(drho, jx, jy, ux, uy, feq); +#else compute_feq(rho, ux, uy, feq); +#endif collide_srt_no_force(f, feq, omega); } -__device__ __forceinline__ void health(float value, unsigned int* counts) { +__device__ __forceinline__ double physical(double value, unsigned int q) { +#if USE_DDF_SHIFTING + // Match population_storage.py's FP64 weights, not promoted FP32 d_w. + return value + (q == 0 ? 4.0/9.0 : (q < 5 ? 1.0/9.0 : 1.0/36.0)); +#else + return value; +#endif +} + +__device__ __forceinline__ void health(float stored, unsigned int q, unsigned int* counts) { + const double value = physical((double)stored, q); if (!isfinite(value)) atomicAdd(counts + 4, 1u); if (value <= 0.0f) atomicAdd(counts + 5, 1u); } @@ -41,7 +59,7 @@ extern "C" __global__ void CoarseStart( : coarse[index_f(k, q)]; } collide(f, omega); - for (int q = 0; q < NQ; ++q) if (!covered) health(f[q], counts); + for (int q = 0; q < NQ; ++q) if (!covered) health(f[q], q, counts); stream_pull_store(k, f, out); } @@ -56,7 +74,7 @@ extern "C" __global__ void Collide( if (active == nullptr || active[k]) { for (int q = 0; q < NQ; ++q) f[q] = in[index_f(k, q)]; collide(f, omega); - if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts); + if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], q, counts); } else { for (int q = 0; q < NQ; ++q) f[q] = NAN; } @@ -74,7 +92,7 @@ extern "C" __global__ void Pull( unsigned long j[NQ]; compute_neighbors(x, y, j); stream_pull_load(k, f, in, j); - if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts); + if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], q, counts); } else { for (int q = 0; q < NQ; ++q) f[q] = NAN; } @@ -90,6 +108,7 @@ extern "C" __global__ void Interpolate( unsigned int group = blockIdx.x * blockDim.x + threadIdx.x; if (group >= groups) return; const int size = sizes[group]; + const unsigned int q = origins[group] / TOTAL_CELLS; for (int child = 0; child < 4; ++child) { double value = frozen ? (double)g0[origins[group]] : 0.0; if (!frozen) for (int j = 0; j < size; ++j) { @@ -99,10 +118,10 @@ extern "C" __global__ void Interpolate( const float stored = __double2float_rn(value); pre[4u*group+child] = value; h[4u*group+child] = stored; - if (!isfinite(value)) atomicAdd(counts, 1u); - if (value <= 0.0) atomicAdd(counts+1, 1u); - if (!isfinite(stored)) atomicAdd(counts+2, 1u); - if (stored <= 0.0f) atomicAdd(counts+3, 1u); + if (!isfinite(physical(value, q))) atomicAdd(counts, 1u); + if (physical(value, q) <= 0.0) atomicAdd(counts+1, 1u); + if (!isfinite(physical((double)stored, q))) atomicAdd(counts+2, 1u); + if (physical((double)stored, q) <= 0.0) atomicAdd(counts+3, 1u); } } @@ -142,5 +161,5 @@ extern "C" __global__ void ApplyRegisters( if (row >= count) return; const unsigned int slot = targets[row]; predicted[slot] = __double2float_rn((double)predicted[slot] + registers[row]); - health(predicted[slot], counts); + health(predicted[slot], slot / TOTAL_CELLS, counts); } diff --git a/tests/validation/multigrid/spacetime_gpu.py b/tests/validation/multigrid/spacetime_gpu.py index 053ba12..1e8269a 100644 --- a/tests/validation/multigrid/spacetime_gpu.py +++ b/tests/validation/multigrid/spacetime_gpu.py @@ -6,6 +6,7 @@ from pathlib import Path import numpy as np import spacetime_reference as ref +from CelerisLab.lbm.population_storage import _encode_physical, _decode_stored, _WEIGHTS PINS = {"spacetime_reference.py": "1fe13242b9a41062a1227db4a6f9b695165aefab7f4c0b3b2adff0c9023552bf", "cpu_reference.py": "bba2a06a9f6517b32079ffff15e4060d8c3b7ed35f16ca134e1877226c9e4e1e", @@ -24,9 +25,9 @@ def validate_config(config): if (cfg.nx != cfg.ny or not 8 <= cfg.nx <= 64 or cfg.nz != 1 or cfg.collision != "SRT" or not cfg.is_d2q9 or cfg.streaming != "double_buffer" or cfg.store_precision != "FP32" - or cfg.data_type != "FP32" or cfg.ddf_shifting or cfg.les_enabled + or cfg.data_type != "FP32" or type(cfg.ddf_shifting) is not bool or cfg.les_enabled or cfg.velocity != 0 or cfg.wall_slot_width or cfg.wall_normal_velocity): - raise ValueError("experiment requires periodic unforced square N8..64, unshifted FP32 D2Q9/SRT") + raise ValueError("experiment requires periodic unforced square N8..64, FP32 D2Q9/SRT") fine = replace(cfg, nx=2*cfg.nx, ny=2*cfg.ny, viscosity=2*cfg.viscosity) fine.validate() return cfg, fine @@ -79,7 +80,7 @@ def table_hash(data): return digest.hexdigest() -def rounding_ledger(plan, tables, before, after, g0, predicted, fine_phases, pre, h, registers): +def rounding_ledger(plan, tables, before, after, g0, predicted, fine_phases, pre, h, registers, *, shifted=False): """Audit actual stored FP32 phases; never call a missing packet 'rounding'.""" maps, n = plan.maps, plan.n f64 = lambda a: np.asarray(a, dtype=np.float64) @@ -118,8 +119,9 @@ def rounding_ledger(plan, tables, before, after, g0, predicted, fine_phases, pre net = total(after)-total(before) raw = net-collision unexplained = raw-group-store_error-register_error - limit = float(64*np.finfo(float).eps*total(before).sum()) + limit = float(64*np.finfo(float).eps*(total(before).sum()+int(shifted)*n*n)) np.testing.assert_allclose(unexplained, 0, rtol=0, atol=limit) + offset = _WEIGHTS.ravel()[tables['origins']//(n*n),None] if shifted else 0. return dict(raw=raw.tolist(), net_change=net.tolist(), physical_collision=collision.tolist(), coarse_collision=coarse_collision.tolist(), fine0_collision=fine_collision[0].tolist(), fine1_collision=fine_collision[1].tolist(), child_interpolation=interpolation.tolist(), @@ -128,7 +130,8 @@ def rounding_ledger(plan, tables, before, after, g0, predicted, fine_phases, pre for a in cell_mass], child_rounding=group.tolist(), receiver_rounding=store_error.tolist(), register_rounding=register_error.tolist(), unexplained=unexplained.tolist(), - unexplained_bound=limit, minimum_pre=float(pre.min()), minimum_stored=float(h.min())) + unexplained_bound=limit, minimum_pre=float((pre+offset).min()), + minimum_stored=float((f64(h)+offset).min())) class SpaceTimeGPU: @@ -210,11 +213,18 @@ class SpaceTimeGPU: if (self.status not in ('new', 'ready') or state.phase != 'arrival_f' or isinstance(state.tick, bool) or not isinstance(state.tick, Integral) or state.tick != 0): raise ValueError("initial synchronized arrival state with tick=0 required") - arrays = [np.array(a, dtype=np.float32, copy=True, order='C') for a in (state.coarse_f,state.fine_f)] - for a, mask in zip(arrays, (self.plan.maps.coarse_active,self.plan.maps.fine_active)): - if a.shape != (9,*mask.shape) or not np.isfinite(a[:,mask]).all() or (a[:,mask] <= 0).any(): + shifted = self.configs[0].ddf_shifting + arrays = [] + for values, mask in zip((state.coarse_f,state.fine_f), + (self.plan.maps.coarse_active,self.plan.maps.fine_active)): + if np.shape(values) != (9,*mask.shape): + raise ValueError("positive finite leaf populations and exact shapes required") + a = _encode_physical(values,shifted) + physical = _decode_stored(a,shifted) + if not np.isfinite(physical[:,mask]).all() or (physical[:,mask] <= 0).any(): raise ValueError("positive finite leaf populations and exact shapes required") a[:,~mask] = np.nan + arrays.append(a) with self.ctx.activate(): try: self.cuda.memcpy_htod(self.coarse[0], arrays[0]) @@ -225,7 +235,8 @@ class SpaceTimeGPU: self.status, self.step_count = 'ready', 0 self.health_counts.fill(0) - def download(self): + def _download_stored(self): + """Private audit/reset snapshot; values remain in their storage representation.""" if self.status != 'ready': raise RuntimeError(f"GPU owner is {self.status}") try: @@ -239,12 +250,18 @@ class SpaceTimeGPU: a[:,~mask] = np.nan return ref.cpu.State(*arrays, tick=2*self.step_count) + def download(self): + """Physical arrival populations; shifted values are reconstructed in FP64.""" + stored = self._download_stored() + return ref.cpu.State(*(_decode_stored(a,self.configs[0].ddf_shifting) + for a in (stored.coarse_f,stored.fine_f)),tick=stored.tick) + def step(self, *, mutation=0, audit=True): """One interval; audit=False skips only host ledger snapshots, never health checks.""" if (self.status != 'ready' or type(mutation) is not int or mutation not in range(5) or type(audit) is not bool or (mutation and not audit)): raise RuntimeError("ready owner and known test mutation required") - before = self.download() if audit else None + before = self._download_stored() if audit else None n, nf, t = self.plan.n, self.plan.nf, self.device_tables ng, nr = len(self.plan.groups), len(self.plan.receivers) u32, null = np.uint32, np.uintp(0) @@ -288,7 +305,8 @@ class SpaceTimeGPU: after = ref.cpu.State(coarse_after,self.read(self.fine[0],(9,nf,nf)),tick=2*(self.step_count+1)) ledger = rounding_ledger(self.plan,self.tables,before,after,coarse_g,prediction,fine_phases, self.read(self.pre,(ng,4),np.float64),self.read(self.h,(ng,4)), - self.read(self.registers,(nr,),np.float64)) + self.read(self.registers,(nr,),np.float64), + shifted=self.configs[0].ddf_shifting) ledger['boundary_budget_status'] = 'measured' ledger['health_counts'] = counts.tolist() self.coarse[0], self.coarse[2] = predicted, fc diff --git a/tests/validation/multigrid/test_spacetime_gpu.py b/tests/validation/multigrid/test_spacetime_gpu.py index bcd1e02..1f9c225 100644 --- a/tests/validation/multigrid/test_spacetime_gpu.py +++ b/tests/validation/multigrid/test_spacetime_gpu.py @@ -34,8 +34,8 @@ def record(value): ref.write_receipt(OUTPUT,REPORT) -def config(n): - return LBMConfig(nx=n,ny=n,viscosity=.0192,velocity=0.,ddf_shifting=False) +def config(n, shifted=False): + return LBMConfig(nx=n,ny=n,viscosity=.0192,velocity=0.,ddf_shifting=shifted) def initial(plan, tagged=True): @@ -60,11 +60,63 @@ class Host(unittest.TestCase): c,f = gpu.validate_config(config(n)) self.assertAlmostEqual(c.omega,1/.5576) self.assertAlmostEqual(f.omega,1/.6152) - for change in (dict(ddf_shifting=True),dict(collision='TRT'),dict(nx=128,ny=128), + for change in (dict(ddf_shifting=1),dict(collision='TRT'),dict(nx=128,ny=128), dict(ny=16),dict(les_enabled=True),dict(velocity=.03)): with self.subTest(change=change), self.assertRaises(ValueError): gpu.validate_config(replace(config(12),**change)) + def test_shifted_codec_and_uploaded_physical_validation(self): + for n in (12,64): + self.assertTrue(gpu.validate_config(config(n,True))[1].ddf_shifting) + physical=gpu._WEIGHTS+2.**-35 + stored=gpu._encode_physical(physical,True) + np.testing.assert_array_equal(stored,np.full((9,1,1),2.**-35,np.float32)) + np.testing.assert_array_equal(gpu._decode_stored(stored,True),physical) + run=gpu.SpaceTimeGPU.__new__(gpu.SpaceTimeGPU) + run.status='new'; run.configs=gpu.validate_config(config(12,True)); run.plan=self.plan + run.ctx=SimpleNamespace(activate=nullcontext); run.cuda=Mock() + run.coarse=['coarse']; run.fine=['fine']; run.health_counts=np.zeros(6,np.uint64) + state=initial(self.plan,False) + run.initialize(state) + for call in run.cuda.memcpy_htod.call_args_list: + self.assertTrue((call.args[1][np.isfinite(call.args[1])]==0).all()) + state.coarse_f[:]-=.01; state.fine_f[:]-=.01 + run.initialize(state) # Valid negative stored values. + with self.assertRaisesRegex(ValueError,'exact shapes'): + run.initialize(ref.cpu.State(state.coarse_f[:1],state.fine_f)) + for invalid in (-.01,np.nan): + bad=ref.cpu.State(state.coarse_f.copy(),state.fine_f.copy()) + bad.coarse_f[0,0,0]=invalid + with self.assertRaisesRegex(ValueError,'positive finite'): + run.initialize(bad) + + def test_shifted_transfer_translation_and_raw_ledger(self): + p=self.plan; t=gpu.pack_plan(p) + rng=np.random.default_rng(19) + for g in p.groups: + values=rng.uniform(.02,.45,len(g.samples)); weight=gpu._WEIGHTS[g.q,0,0] + np.testing.assert_allclose(g.weights@(values-weight),g.weights@values-weight,rtol=0,atol=2e-14) + # Uniform nonzero stored packets expose register omissions; exact rest would not. + arrays=[np.full((9,n,n),-.01,np.float32) for n in (p.n,p.nf)] + for a,m in zip(arrays,(p.maps.coarse_active,p.maps.fine_active)): + a[:,~m]=np.nan + state=ref.cpu.State(*arrays); g0=np.full((9,p.n,p.n),-.01,np.float32) + pre=np.repeat(g0.ravel()[t['origins'],None].astype(float),4,axis=1) + h=pre.astype(np.float32); registers=np.zeros(len(p.receivers)) + args=(p,t,state,state,g0,g0,[(state.fine_f,state.fine_f)]*2,pre,h,registers) + ledger=gpu.rounding_ledger(*args,shifted=True) + self.assertEqual(max(map(abs,ledger['unexplained'])),0.) + self.assertGreater(ledger['unexplained_bound'],0.) + self.assertGreater(ledger['minimum_stored'],0.) + expected_mass=p.n*p.n*(1.+9*float(np.float32(-.01))) + self.assertEqual(ledger['unexplained_bound'],64*np.finfo(float).eps*expected_mass) + decoded=ref.cpu.State(*(gpu._decode_stored(a,True) for a in arrays)) + with self.assertRaises(AssertionError): + gpu.rounding_ledger(p,t,decoded,*args[3:],shifted=True) + registers[0]+=.25*float(np.float32(-.01)) + with self.assertRaisesRegex(AssertionError,'register mismatch'): + gpu.rounding_ledger(*args,shifted=True) + def test_packed_weights_and_packet_ownership(self): for plan in (self.plan,ref.Plan(16,(4,5,11,12))): t = gpu.pack_plan(plan) @@ -261,6 +313,11 @@ class Host(unittest.TestCase): class Device(unittest.TestCase): + shifted = False + + def config(self,n): + return config(n,self.shifted) + @classmethod def setUpClass(cls): if not GPU_ADMITTED: @@ -275,12 +332,14 @@ class Device(unittest.TestCase): def test_tiny_equivalence_budgets_and_phase(self): for mode,identity in (('spacetime',False),('spacetime',True),('frozen',False),('frozen',True)): with self.subTest(mode=mode,identity=identity), gpu.SpaceTimeGPU( - config(12),(3,3,9,8),context=self.context,mode=mode,identity=identity) as run: + self.config(12),(3,3,9,8),context=self.context,mode=mode,identity=identity) as run: spec = ref.Spec(n_coarse=12,patch_lo=1,halo=2,tau_coarse=.5576,tau_fine=.6152, collision='identity' if identity else 'srt') for tagged,steps in ((False,2),(True,8)): expected = initial(run.plan,tagged); run.initialize(expected) before = run.download() + if self.shifted: + expected = before for k in range(1,steps+1): expected,_ = ref.advance(spec,run.plan,expected,mode) ledger = run.step(); actual = run.download() @@ -298,14 +357,19 @@ class Device(unittest.TestCase): wrong = ref.cpu._collide(spec,b,np.isfinite(b),tau,mask,np.zeros(mask.shape,int)) self.assertGreater(float(np.nanmax(abs(a-wrong))),3e-6) if k in (1,2,8): - record(dict(case='tiny',mode=mode,identity=identity,tagged=tagged, + record(dict(case='tiny',shifted=self.shifted,mode=mode,identity=identity,tagged=tagged, step=k,device_bytes=run.device_bytes,ledger=ledger)) def test_omitted_packets_and_ghosts_fail(self): for mutation in (1,2,3,4): with self.subTest(mutation=mutation), gpu.SpaceTimeGPU( - config(12),(3,3,9,8),context=self.context) as run: - run.initialize(initial(run.plan)) + self.config(12),(3,3,9,8),context=self.context,identity=self.shifted) as run: + state=initial(run.plan,not self.shifted) + if self.shifted: + state.coarse_f[:]+=1e-4; state.fine_f[:]+=1e-4 + for a,mask in zip((state.coarse_f,state.fine_f),(run.plan.maps.coarse_active,run.plan.maps.fine_active)): + self.assertTrue((gpu._encode_physical(a,True)[:,mask]>0).all()) + run.initialize(state) with self.assertRaises(AssertionError): run.step(mutation=mutation) self.assertEqual(run.step_count,0) @@ -314,33 +378,37 @@ class Device(unittest.TestCase): run.download() def test_frozen_donor_trap(self): - with gpu.SpaceTimeGPU(config(16),(4,4,12,12),context=self.context,mode='frozen',identity=True) as run: + with gpu.SpaceTimeGPU(self.config(16),(4,4,12,12),context=self.context,mode='frozen',identity=True) as run: state = initial(run.plan,False) state.coarse_f[5,4,3]=1.; state.coarse_f[5,5,3]=2. - run.initialize(state); run.step() - self.assertEqual(run.download().fine_f[5,11,8],1.) - # Deliberately negative boundary data are reported and remain unclipped. + run.initialize(state); donor=run.download().coarse_f[5,4,3]; run.step() + self.assertEqual(run.download().fine_f[5,11,8],donor) + # Negative stored values can be physical in shifted mode; never clip either case. with self.context.activate(): - data=np.full((9,16,16),-.01,np.float32) - run.cuda.memcpy_htod(run.coarse[1],data); run.cuda.memcpy_htod(run.coarse[3],data) - run.cuda.memset_d32_async(run.counts,0,6,run.stream) - t=run.device_tables; ng=len(run.plan.groups) - run.launch(0,'Interpolate',run.coarse[1],run.coarse[3],t['samples'],t['origins'], - t['sizes'],t['weights'],run.pre,run.h,np.uint32(ng),np.uint32(0),run.counts,count=ng) - counts=run.read(run.counts,(6,),np.uint32) - self.assertEqual(counts[1],4*ng); self.assertEqual(counts[3],4*ng) - self.assertTrue((run.read(run.h,(ng,4))<0).all()) + for physical_negative in ((False,True) if self.shifted else (True,)): + data=np.full((9,16,16),-.01,np.float32) + if self.shifted and physical_negative: + data=np.broadcast_to(-gpu._WEIGHTS-.01,(9,16,16)).astype(np.float32).copy() + run.cuda.memcpy_htod(run.coarse[1],data); run.cuda.memcpy_htod(run.coarse[3],data) + run.cuda.memset_d32_async(run.counts,0,6,run.stream) + t=run.device_tables; ng=len(run.plan.groups) + run.launch(0,'Interpolate',run.coarse[1],run.coarse[3],t['samples'],t['origins'], + t['sizes'],t['weights'],run.pre,run.h,np.uint32(ng),np.uint32(0),run.counts,count=ng) + counts=run.read(run.counts,(6,),np.uint32) + self.assertEqual(counts[1],4*ng*physical_negative) + self.assertEqual(counts[3],4*ng*physical_negative) + self.assertTrue((run.read(run.h,(ng,4))<0).all()) def test_n64_one_interval_only(self): sys.path.insert(0,str(Path(__file__).parents[1]/'taylor_green')) import oracle as tg - with gpu.SpaceTimeGPU(config(64),(8,8,56,48),context=self.context) as run: + with gpu.SpaceTimeGPU(self.config(64),(8,8,56,48),context=self.context) as run: arrays = [tg.initial((64,64),0,64,.0192),tg.initial((128,128),1,64,.0192)] for a,mask in zip(arrays,(run.plan.maps.coarse_active,run.plan.maps.fine_active)): a[:,~mask]=np.nan state = ref.cpu.State(*arrays); run.initialize(state) spec = ref.Spec(n_coarse=64,patch_lo=8,halo=16,tau_coarse=.5576,tau_fine=.6152,collision='srt') - expected,_ = ref.advance(spec,run.plan,state) + expected,_ = ref.advance(spec,run.plan,run.download() if self.shifted else state) started=time.monotonic(); ledger=run.step(); seconds=time.monotonic()-started actual=run.download() errors=[] @@ -348,28 +416,34 @@ class Device(unittest.TestCase): np.testing.assert_array_equal(np.isnan(a),np.isnan(b)) np.testing.assert_allclose(a,b,rtol=0,atol=3e-6,equal_nan=True) errors.append(float(np.nanmax(abs(a-b)))) - record(dict(case='N64_one_interval',errors=errors,ledger=ledger,seconds=seconds, + record(dict(case='N64_one_interval',shifted=self.shifted,errors=errors,ledger=ledger,seconds=seconds, device_bytes=run.device_bytes,table_sha256=gpu.table_hash(run.tables), build_metadata=run.build_metadata,device=self.context.name)) def test_unmeasured_path_and_warm_reset(self): - with gpu.SpaceTimeGPU(config(12),(3,3,9,8),context=self.context) as run: + with gpu.SpaceTimeGPU(self.config(12),(3,3,9,8),context=self.context) as run: state=initial(run.plan) run.initialize(state) + raw_initial=run._download_stored() for _ in range(2): run.step() - expected=run.download() + expected=run._download_stored() run.initialize(state) for _ in range(2): result=run.step(audit=False) self.assertEqual(result['boundary_budget_status'],'not_measured') - actual=run.download() + actual=run._download_stored() for a,b in zip((actual.coarse_f,actual.fine_f),(expected.coarse_f,expected.fine_f)): - np.testing.assert_array_equal(a,b) + np.testing.assert_array_equal(a.view(np.uint32),b.view(np.uint32)) reset,warm=trajectory.warm_and_reset(run,state,lambda:None) - for a,b in zip((reset.coarse_f,reset.fine_f),(state.coarse_f,state.fine_f)): - np.testing.assert_array_equal(a,b.astype(np.float32)) - record(dict(case='unmeasured_bitwise_and_reset',warmup=warm,health_counts=run.health_counts.tolist())) + raw_reset=run._download_stored() + for a,b in zip((raw_reset.coarse_f,raw_reset.fine_f),(raw_initial.coarse_f,raw_initial.fine_f)): + np.testing.assert_array_equal(a.view(np.uint32),b.view(np.uint32)) + record(dict(case='unmeasured_bitwise_and_reset',shifted=self.shifted,warmup=warm,health_counts=run.health_counts.tolist())) + + +class ShiftedDevice(Device): + shifted = True def main(): @@ -399,6 +473,7 @@ def main(): suite=unittest.TestLoader().loadTestsFromTestCase(Host) if args.gpu_admitted: suite.addTests(unittest.TestLoader().loadTestsFromTestCase(Device)) + suite.addTests(unittest.TestLoader().loadTestsFromTestCase(ShiftedDevice)) result=unittest.TextTestRunner(verbosity=2).run(suite) report.update(completed=True,passed=result.wasSuccessful(),tests=result.testsRun, failures=[dict(test=str(test),error=error) for test,error in result.failures+result.errors],