From 6b6acc4e415a3af00dab587c2648ead3a072ae6f Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 14 May 2026 10:19:07 -0700 Subject: [PATCH 01/34] Fixed issue where GPU code was called before it was built --- mcdc/code_factory/gpu/program_builder.py | 27 ++++++ mcdc/code_factory/gpu/transport/simulation.py | 93 +++++++++++++++++++ mcdc/code_factory/numba_layers_generator.py | 61 +++++++++--- mcdc/main.py | 10 +- mcdc/object_/simulation.py | 10 ++ 5 files changed, 188 insertions(+), 13 deletions(-) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index 3e5570c75..f62a6d538 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -296,6 +296,33 @@ def step(program: nb.uintp, particle_input: particle_gpu): complete = src_fns["complete"] clear_flags = src_fns["clear_flags"] set_device = src_fns["set_device"] + + src_alloc_program = src_fns["alloc_program"] + src_free_program = src_fns["free_program"] + src_load_global = src_fns["load_state_device_global"] + src_store_global = src_fns["store_state_device_global"] + src_store_pointer_global = src_fns["store_pointer_state_device_global"] + src_load_data = src_fns["load_state_device_data"] + src_store_data = src_fns["store_state_device_data"] + src_store_pointer_data = src_fns["store_pointer_state_device_data"] + src_init_program = src_fns["init_program"] + src_exec_program = src_fns["exec_program"] + src_complete = src_fns["complete"] + src_clear_flags = src_fns["clear_flags"] + src_set_device = src_fns["set_device"] + + # ================================================================================== + # + # ================================================================================== + + """ + global loop_source + loop_source = gpu_loop_source + # + # Overwrite function + for impl in target_rosters["cpu"].values(): + overwrite_func(impl, impl) + """ # ====================================================================================== diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 591ec618a..50dc60f96 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -68,6 +68,99 @@ def source_loop(seed, simulation, data): gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) # Recover the original program state + src_load_constant(mcdc, mcdc["gpu_state_pointer"]) + src_load_data(data, mcdc["gpu_state_pointer"]) + src_clear_flags(mcdc["source_program_pointer"]) + + mcdc["mpi_work_size"] = full_work_size + + particle_bank_module.set_bank_size(mcdc["bank_active"], 0) + + # ===================================================================== + # Closeout (Moved out of the typical particle loop) + # ===================================================================== + + source_closeout(mcdc, 1, 1, data) + + if mcdc["technique"]["domain_decomposition"]: + source_dd_resolution(data, mcdc) + + +def build_gpu_progs(input_deck, args): + + STRAT = args.gpu_strategy + + src_spec = gpu_sources_spec() + + adapt.harm.RuntimeSpec.bind_specs() + + rank = MPI.COMM_WORLD.Get_rank() + device_id = rank % args.gpu_share_stride + + if MPI.COMM_WORLD.Get_size() > 1: + MPI.COMM_WORLD.Barrier() + + adapt.harm.RuntimeSpec.load_specs() + + if STRAT == "async": + args.gpu_arena_size = args.gpu_arena_size // 32 + src_fns = src_spec.async_functions() + pre_fns = pre_spec.async_functions() + else: + src_fns = src_spec.event_functions() + pre_fns = pre_spec.event_functions() + + ARENA_SIZE = args.gpu_arena_size + BLOCK_COUNT = args.gpu_block_count + + global alloc_state, free_state + alloc_state = src_fns["alloc_state"] + free_state = src_fns["free_state"] + + global src_alloc_program, src_free_program + global src_load_global, src_store_global, src_load_data, src_store_data, src_store_pointer_data + global src_init_program, src_exec_program, src_complete, src_clear_flags + src_alloc_program = src_fns["alloc_program"] + src_free_program = src_fns["free_program"] + src_load_global = src_fns["load_state_device_global"] + src_store_global = src_fns["store_state_device_global"] + src_store_pointer_global = src_fns["store_pointer_state_device_global"] + src_load_data = src_fns["load_state_device_data"] + src_store_data = src_fns["store_state_device_data"] + src_store_pointer_data = src_fns["store_pointer_state_device_data"] + src_init_program = src_fns["init_program"] + src_exec_program = src_fns["exec_program"] + src_complete = src_fns["complete"] + src_clear_flags = src_fns["clear_flags"] + src_set_device = src_fns["set_device"] + + global pre_alloc_program, pre_free_program + global pre_load_global, pre_store_global, pre_load_data, pre_store_data + global pre_init_program, pre_exec_program, pre_complete, pre_clear_flags + pre_alloc_state = pre_fns["alloc_state"] + pre_free_state = pre_fns["free_state"] + pre_alloc_program = pre_fns["alloc_program"] + pre_free_program = pre_fns["free_program"] + pre_load_global = pre_fns["load_state_device_global"] + pre_store_global = pre_fns["store_state_device_global"] + pre_load_data = pre_fns["load_state_device_data"] + pre_store_data = pre_fns["store_state_device_data"] + pre_init_program = pre_fns["init_program"] + pre_exec_program = pre_fns["exec_program"] + pre_complete = pre_fns["complete"] + pre_clear_flags = pre_fns["clear_flags"] + + @njit + def real_setup_gpu(mcdc_array, data_tally): + mcdc = mcdc_array[0] + + print("STATE POINTER {mcdc['gpu_meta']['state_pointer']}") + print("GLOBAL POINTER {mcdc['gpu_meta']['global_pointer']}") + print("TALLY POINTER {mcdc['gpu_meta']['tally_pointer']}") + src_set_device(device_id) + arena_size = ARENA_SIZE + mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) + # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) if config.gpu_state_storage == "separate": harmonize.memcpy_device_to_host( simulation, simulation["gpu_meta"]["state_pointer"] diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 4dd538c10..534596ab7 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -42,6 +42,18 @@ np.uintp: np.uintp, } +size_map = { + bool: 1, + float: 8, + int: 8, + str: 32, + np.bool_: 1, + np.float64: 8, + np.int64: 8, + np.uint64:8, + np.str_: 32, +} + bank_names = ["bank_active", "bank_census", "bank_source", "bank_future"] @@ -285,6 +297,24 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data) set_object(simulation, annotations, structures, records, data) + print("\n\n\nA\n\n\n",flush=True) + # Allocate the flattened data and re-set the objects + data["array"], data["pointer"] = create_data_array(data["size"], type_map[float],size_map[float]) + print("\n\n\nB\n\n\n",flush=True) + + data["size"] = 0 + records = {} + for mcdc_class in mcdc_classes: + if issubclass(mcdc_class, ObjectNonSingleton): + records[mcdc_class.label] = [] + else: + records[mcdc_class.label] = {} + records["simulation"] = records.pop("simulation") + + for object_ in objects: + set_object(object_, annotations, structures, records, data, set_data=True) + set_object(simulation, annotations, structures, records, data, set_data=True) + # ================================================================================== # Finalize the simulation object structure and set record # ================================================================================== @@ -375,6 +405,8 @@ def generate_numba_layers(simulation): simulation_dtype ) mcdc_simulation = mcdc_simulation_container[0] + mcdc_simulation["gpu_meta"]["global_pointer"] = mcdc_simulation_pointer + mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] record = records["simulation"] structure = structures["simulation"] @@ -756,18 +788,23 @@ def set_object( # ============================================================================= -def create_data_array(size): - if not config.target == "gpu": - data = np.zeros(size, dtype=np.float64) - return data, 0 - else: - return create_data_array_on_gpu(size * 8) - - -@njit -def create_data_array_on_gpu(size): - if config.gpu_state_storage == "managed": - data_ptr = gpu_builder.alloc_managed_bytes(size) +def create_data_array(size, dtype, itemsize): + if config.target == "gpu": + import mcdc.code_factory.gpu.adapt as adapt + import harmonize, numba + + print("\n\n\nW\n\n\n",flush=True) + print(f"Tally size is {size} with itemsize {itemsize}",flush=True) + if config.gpu_state_storage == "managed": + print("\n\n\nX\n\n\n",flush=True) + data_tally_ptr = harmonize.alloc_managed_bytes(size*itemsize) + else: + print("\n\n\nY\n\n\n",flush=True) + data_tally_ptr = harmonize.alloc_device_bytes(size*itemsize) + print("\n\n\nZ\n\n\n",flush=True) + data_tally_uint = adapt.voidptr_to_uintp(data_tally_ptr) + data_tally = numba.carray(data_tally_ptr, (size,), dtype) + return data_tally, data_tally_uint else: data_ptr = gpu_builder.alloc_device_bytes(size) data_uint = voidptr_to_uintp(data_ptr) diff --git a/mcdc/main.py b/mcdc/main.py index 2497571d1..d26946c78 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -136,7 +136,6 @@ def prepare(simulationPy: Simulation): simulation = simulation_container[0] # Pick Python-version RNG if needed - import mcdc.config as config import mcdc.transport.rng as rng if config.mode == "python": @@ -170,6 +169,15 @@ def prepare(simulationPy: Simulation): # simulation["bank_source"]["size"] = N_local # MPI.COMM_WORLD.Barrier() + # ================================================================================== + # Setup GPU-Related Data Structures, if Necessary + # ================================================================================== + + if config.target == "gpu": + from mcdc.code_factory.gpu.program_builder import build_gpu_program + setup_gpu_program(mcdc_container, data) + + # ================================================================================== # Finalize # ================================================================================== diff --git a/mcdc/object_/simulation.py b/mcdc/object_/simulation.py index bd9cbd06d..5de5de093 100644 --- a/mcdc/object_/simulation.py +++ b/mcdc/object_/simulation.py @@ -435,6 +435,16 @@ def _finalize_compilation(self) -> None: self.bank_census.size[0] = int(settings.census_bank_buffer_ratio * N_work) self.bank_source.size[0] = int(settings.source_bank_buffer_ratio * N_work) self.bank_future.size[0] = int(settings.future_bank_buffer_ratio * N_work) + + # ================================================================================== + # Platform targeting, adapters, and toggles for portability + # ================================================================================== + + import mcdc.config as config + # Build GPU program if desired + if config.target == "gpu": + from mcdc.code_factory.gpu.program_builder import build_gpu_program + # Initialize run state derived from the compiled settings self.k_eff = settings.k_init From 6f8623daae626f1213b67ff3b7f4eb8ac69a85cd Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 28 May 2026 00:24:36 -0700 Subject: [PATCH 02/34] Have proof of concept for on-gpu array return helper --- mcdc/code_factory/array_return.py | 445 +++++++++++++++++++ mcdc/code_factory/jit.py | 62 +++ mcdc/code_factory/numba_layers_generator.py | 1 + mcdc/transport/geometry/interface.py | 2 +- mcdc/transport/geometry/root_solve.py | 172 +++++++ mcdc/transport/geometry/surface/interface.py | 1 + mcdc/transport/geometry/surface/torus.py | 17 +- mcdc/transport/geometry/surface/torus_x.py | 18 +- mcdc/transport/geometry/surface/torus_y.py | 18 +- mcdc/transport/geometry/surface/torus_z.py | 18 +- mcdc/transport/physics/util.py | 5 +- mcdc/transport/rng.py | 1 + 12 files changed, 737 insertions(+), 23 deletions(-) create mode 100644 mcdc/code_factory/array_return.py create mode 100644 mcdc/code_factory/jit.py create mode 100644 mcdc/transport/geometry/root_solve.py diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py new file mode 100644 index 000000000..6c26ffcf7 --- /dev/null +++ b/mcdc/code_factory/array_return.py @@ -0,0 +1,445 @@ +from numba import njit, jit, objmode, literal_unroll, types +from numba.extending import intrinsic +import numba as nb +import numpy as np + +import cffi +ffi = cffi.FFI() + + +# ============================================================================= +# uintp/voidptr casters +# ============================================================================= + + +@intrinsic +def cast_any_to_voidptr(typingctx, src): + # create the expected type signature + result_type = types.voidptr + sig = result_type(src) + + # defines the custom code generation + def codegen(context, builder, signature, args): + # llvm IRBuilder code here + [src] = args + rtype = signature.return_type + llrtype = context.get_value_type(rtype) + return builder.bitcast(src, llrtype) + + return sig, codegen + + +@intrinsic +def cast_uintp_to_voidptr(typingctx, src): + # check for accepted types + if isinstance(src, types.Integer): + # create the expected type signature + result_type = types.voidptr + sig = result_type(types.uintp) + + # defines the custom code generation + def codegen(context, builder, signature, args): + # llvm IRBuilder code here + [src] = args + rtype = signature.return_type + llrtype = context.get_value_type(rtype) + return builder.inttoptr(src, llrtype) + + return sig, codegen + + +@intrinsic +def cast_voidptr_to_uintp(typingctx, src): + # check for accepted types + if isinstance(src, types.RawPointer): + # create the expected type signature + result_type = types.uintp + sig = result_type(types.voidptr) + + # defines the custom code generation + def codegen(context, builder, signature, args): + # llvm IRBuilder code here + [src] = args + rtype = signature.return_type + llrtype = context.get_value_type(rtype) + return builder.ptrtoint(src, llrtype) + + return sig, codegen + +@njit() +def uintp_to_voidptr(value): + val = nb.uintp(value) + return cast_uintp_to_voidptr(val) + +@njit() +def voidptr_to_uintp(value): + return cast_voidptr_to_uintp(value) + + +@njit() +def into_voidptr(value): + return into_voidptr_python(value) + +def into_voidptr_python(value): + raise RuntimeError("`into_voidptr` is only supported in nopython mode.") + +@nb.extending.overload(into_voidptr_python) +def into_voidptr_overload(value): + + if isinstance(value,nb.types.Array) : + def impl(value): + ptr = (ffi.from_buffer(value)) + vptr = cast_any_to_voidptr(ptr) + return vptr + return impl + elif isinstance(value,nb.types.CPointer) : + def impl(value): + return cast_any_to_voidptr(value) + return impl + elif isinstance(value,nb.types.Integer) : + def impl(value): + return cast_uintp_to_voidptr(value) + return impl + else : + raise RuntimeError(f"`into_voidptr` is not supported for type '{value}'") + + + + + +############################################################################### +# New Code +############################################################################### + + + +def array_result(array): + return array + +@nb.extending.overload(array_result) +def array_result_overload(array): + + if not isinstance(array,types.Array): + raise nb.core.errors.TypingError(f"Expected array type argument for array_result, got {array}.") + + print("OVER LOADED") + def impl(array): + return (into_voidptr(array),len(array)) + return impl + + + +def context_guard(context): + if isinstance(context, nb.core.typing.context.Context): + pass + elif isinstance(context, nb.cuda.target.CUDATypingContext): + pass + elif isinstance(context, nb.hip.target.HIPTypingContext): + pass + else: + raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") + + + + +def array_return_typing(fn,elem_type): + + + from inspect import signature + param_count = len(signature(fn).parameters) + template = "def typer({arg_list}):\n return nb.types.Array(dtype=elem_type,ndim=1,layout='C')({arg_list})" + + gns = globals() | {"elem_type":elem_type} + lns = {} + arg_list = ",".join([f"param_{i}" for i in range(param_count)]) + exec(template.format(arg_list=arg_list),gns,lns) + typer = lns["typer"] + + def typer_factory(context): + from numba.np.numpy_support import as_dtype + + context_guard(context) + + return typer + + nb.extending.type_callable(fn)(typer_factory) + + +def array_return_lowering(fn,elem_type): + + + from inspect import signature + param_count = len(signature(fn).parameters) + retty = nb.types.Array(dtype=elem_type,ndim=1,layout="C") + sig = retty(*([nb.types.Any]*param_count)) + + jit_fn = nb.njit(fn) + + def builtin(context, builder, sig, args): + + thing, data = args + thing_type, data_type = sig.args + + + import llvmlite.binding as ll + from llvmlite import ir + + try: + import numba.hip as hip + ROCM_AVAILABLE = True + except: + ROCM_AVAILABLE = False + + if not ROCM_AVAILABLE: + try: + import numba.cuda as cuda + CUDA_AVAILABLE = True + except: + CUDA_AVAILABLE = False + else: + CUDA_AVAILABLE = False + + lmod = builder.module + retty = nb.types.Tuple([nb.types.voidptr,nb.types.uintp]) + ptr_sig = retty(*sig.args) + + res = context.compile_internal(builder,jit_fn.py_func,ptr_sig,args) + ptr_res = builder.extract_value(res,0) + size_res = builder.extract_value(res,1) + shape = [size_res] + dtype = data_type.dtype + + if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + targetdata = ll.create_target_data( + nb.hip.amdgcn.DATA_LAYOUT + ) + elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) + lldtype = context.get_data_type(dtype) + if isinstance(context,nb.core.cpu.CPUContext): + itemsize = context.get_abi_sizeof(lldtype) + elif ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + itemsize = lldtype.get_abi_size(targetdata) + elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + itemsize = lldtype.get_abi_size(targetdata) + else: + raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") + + kstrides = [context.get_constant(types.intp,itemsize)] + + aryty = types.Array(dtype=elem_type, ndim=1, layout="C") + ary = context.make_array(aryty)(context, builder) + + dataptr = builder.addrspacecast( + ptr_res, ir.PointerType(ir.IntType(8)), "generic" + ) + + kshape = [size_res] + context.populate_array( + ary, + data=builder.bitcast(dataptr, ary.data.type), + shape=kshape, + strides=kstrides, + itemsize=context.get_constant(types.intp, itemsize), + meminfo=None, + ) + print(ary._getvalue()) + return ary._getvalue() + + nb.extending.lower_builtin(fn,*sig.args)(builtin) + + + + +def array_return(sig): + def array_return_true_decorator(fn): + array_return_typing(fn,sig) + array_return_lowering(fn,sig) + return fn + return array_return_true_decorator + + + + + + +#@nb.extending.type_callable(attribute_all) +def type_attribute_all(context): + from numba.np.numpy_support import as_dtype + + context_guard(context) + + def typer(thing, data): + thing_dtype = as_dtype(thing) + if not thing_dtype == thing_type: + raise nb.core.errors.TypingError(f"First argument is type {thing_dtype}, but must be type {thing_type}.") + if not isinstance(data,nb.types.Array): + raise nb.core.errors.TypingError(f"Second argument is type {data_dtype}, but must be an Array of type {data_type}.") + data_dtype = as_dtype(data.dtype) + if not data_dtype == data_type: + raise nb.core.errors.TypingError(f"Second argument has element dtype {data_dtype}, but must be type {data_type}.") + elif not data.ndim == 1: + raise nb.core.errors.TypingError(f"Second argument should be one-dimensional.") + elif not data.layout == "C": + raise nb.core.errors.TypingError(f"Second argument should have C layout.") + + + retty = types.Array(dtype=data.dtype,ndim=1,layout="C") + sig = retty(thing,data) + return sig + + return typer + + +#@nb.extending.lower_builtin(attribute_all,nb.from_dtype(thing_type),nb.types.Array(dtype=nb.from_dtype(data_type),ndim=1,layout="C")) +def builtin_attribute_all(context, builder, sig, args): + + thing, data = args + thing_type, data_type = sig.args + + + import llvmlite.binding as ll + from llvmlite import ir + + try: + import numba.hip as hip + ROCM_AVAILABLE = True + except: + raise nb.core.errors.UnsupportedError(f"ROCM ain't available.") + ROCM_AVAILABLE = False + + if not ROCM_AVAILABLE: + try: + import numba.cuda as cuda + CUDA_AVAILABLE = True + except: + CUDA_AVAILABLE = False + else: + CUDA_AVAILABLE = False + + lmod = builder.module + retty = nb.types.Tuple([nb.types.voidptr,nb.types.uintp]) + ptr_sig = retty(*sig.args) + + res = context.compile_internal(builder,attribute_all_ptr.py_func,ptr_sig,args) + ptr_res = builder.extract_value(res,0) + size_res = builder.extract_value(res,1) + shape = [size_res] + dtype = data_type.dtype + + if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + targetdata = ll.create_target_data( + nb.hip.amdgcn.DATA_LAYOUT + ) + elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) + lldtype = context.get_data_type(dtype) + if isinstance(context,nb.core.cpu.CPUContext): + itemsize = context.get_abi_sizeof(lldtype) + elif ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + itemsize = lldtype.get_abi_size(targetdata) + elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + itemsize = lldtype.get_abi_size(targetdata) + else: + raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") + + kstrides = [context.get_constant(types.intp,itemsize)] + + aryty = types.Array(dtype=dtype, ndim=1, layout="C") + ary = context.make_array(aryty)(context, builder) + + dataptr = builder.addrspacecast( + ptr_res, ir.PointerType(ir.IntType(8)), "generic" + ) + + kshape = [size_res] + context.populate_array( + ary, + data=builder.bitcast(dataptr, ary.data.type), + shape=kshape, + strides=kstrides, + itemsize=context.get_constant(types.intp, itemsize), + meminfo=None, + ) + print(ary._getvalue()) + return ary._getvalue() + + + + + + +from numba import hip + + +data_type = np.float64 + +thing_a_type = np.dtype([ + ("attribute_a_offset",np.intp), + ("attribute_a_size", np.intp), +]) + +thing_b_type = np.dtype([ + ("attribute_b_offset",np.intp), + ("attribute_b_size", np.intp), +]) + + +# We're cooking with gas now! >B) + +@array_return(nb.float64) +def attribute_a_all(thing, data): + start = thing["attribute_a_offset"] + size = thing["attribute_a_size"] + end = start + size + return array_result(data[start:end]) + + +@array_return(nb.float64) +def attribute_b_all(thing, data): + start = thing["attribute_b_offset"] + size = thing["attribute_b_size"] + end = start + size + return array_result(data[start:end]) + + + + +@njit() +def check(thing_a,thing_b,data): + a = attribute_a_all(thing_a,data) + b = attribute_b_all(thing_b,data) + print(a) + print(b) + + +@nb.hip.jit() +def hip_check(thing_a,thing_b,data): + a = attribute_a_all(thing_a,data) + b = attribute_b_all(thing_b,data) + for i in range(len(a)): + a[i] = -1 + for i in range(len(b)): + b[i] = -2 + + + +data = np.zeros((10,),dtype=data_type) +for i in range(10): + data[i] = i + +thing_a_arr = np.zeros((1,),thing_a_type) +thing_a = thing_a_arr[0] +thing_a["attribute_a_offset"] = 4 +thing_a["attribute_a_size"] = 3 +thing_b_arr = np.zeros((1,),thing_b_type) +thing_b = thing_b_arr[0] +thing_b["attribute_b_offset"] = 8 +thing_b["attribute_b_size"] = 1 + +check(thing_a,thing_b,data) +hip_check[1,1](thing_a,thing_b,data) + +print(data) + + diff --git a/mcdc/code_factory/jit.py b/mcdc/code_factory/jit.py new file mode 100644 index 000000000..224f58d18 --- /dev/null +++ b/mcdc/code_factory/jit.py @@ -0,0 +1,62 @@ +import numba as nb +from inspect import getfullargspec + +ARG_CHECK = True + +njit_trace_template = """ +def arg_check_{name}({args}): + return fn({args}) + +@nb.extending.overload(arg_check_{name}) +def arg_check_{name}_overload({args}): + arg_list = [{args}] + arg_name_list = [{arg_names}] + for idx in range(len(arg_list)): + arg = arg_list[idx] + arg_name = arg_name_list[idx] + if isinstance(arg,nb.types.Record): + raise RuntimeError(f"Argument {{arg_name}} has a Record type. Records should be passed in an array.") + elif isinstance(arg,nb.types.Optional): + raise RuntimeError(f"Argument {{arg_name}} has a Record type. Records should be passed in an array.") + return fn +""" + + +def wrap_with_check(fn,njit_fn): + nb.extending.register_jitable(fn) + arg_names = getfullargspec(fn).args + args_str = ",".join(arg_names) + arg_names_str = ",".join(f'"{n}"' for n in arg_names) + name = fn.__name__ + print(f"wrapping {name}") + gns = {"nb":nb,"fn":fn,"njit_fn":njit_fn} + lns = {} + code = njit_trace_template.format( + args=args_str, + arg_names=arg_names_str, + name=name + ) + exec(code,gns,lns) + return lns[f"arg_check_{name}"] + + +def njit(*args,**kwargs): + + if (len(args) == 1) and (len(kwargs) == 0): + if not ARG_CHECK: + return njit(args[0]) + fn = args[0] + njit_fn = nb.njit(args[0]) + return wrap_with_check(fn,njit_fn) + + else: + if not ARG_CHECK: + return nb.njit(*args,**kwargs) + + def wrapper(fn): + njit_fn = nb.njit(*args,**kwargs)(fn) + return wrap_with_check(fn,njit_fn) + + + + diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 534596ab7..a660927d1 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -1213,6 +1213,7 @@ def _accessor_1d_element(object_name, attribute_name, setter=False, cast_to_int= return text + def _accessor_1d_all(object_name, attribute_name, size, setter=False): text = f"@njit\n" if setter: diff --git a/mcdc/transport/geometry/interface.py b/mcdc/transport/geometry/interface.py index b904dbe3f..0050383af 100644 --- a/mcdc/transport/geometry/interface.py +++ b/mcdc/transport/geometry/interface.py @@ -369,7 +369,7 @@ def _check_cell(particle_container, speed, cell, simulation, data): return True # Create local value array - value = util.local_array(literals.rpn_evaluation_buffer_size(), np.bool_) + value = util.local_array(100, np.bool_) N_value = 0 # March forward through RPN tokens diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py new file mode 100644 index 000000000..b6ba4990d --- /dev/null +++ b/mcdc/transport/geometry/root_solve.py @@ -0,0 +1,172 @@ +import math +import numpy as np +import numba as nb + +import mcdc.transport.util as util + +from numba import njit + + +@njit() +def modulus(x): + return math.sqrt(x.real**2+x.imag**2) + +@njit() +def sqrt(x): + return math.sqrt(x.real) * (x+x.real) / modulus(x+x.real) + + + +@njit() +def nth_root(x,n,index): + # First, convert to polar form + r = modulus(x) + a = math.atan2(x.imag,x.real) + + # Apply de Moivre's Formula + root_modulus = math.pow(r,1./n) + root_argument = (a+2*math.pi*index)/n + + # ...then convert back to rectangular form + real = root_modulus * math.cos(root_argument) + imag = root_modulus * math.sin(root_argument) + return complex(real,imag) + + +@njit() +def principal_nth_root(x,n): + return nth_root(x,n,0) + + +@njit() +def solve_quadratic(coeff,roots): + a = coeff[2] + b = coeff[1] + c = coeff[0] + # standard quadratic formula, but with discriminant + # calculated separately for re-use + discriminant = sqrt(b**2-4*a*c) + roots[0] = ((-b) + discriminant) / (2*a) + roots[1] = ((-b) - discriminant) / (2*a) + + +@njit() +def solve_biquadratic(coeff,roots): + # Move each coefficient down to one-half it's power + coeff[1] = coeff[2] + coeff[2] = coeff[4] + + # Solve as quadratic equation, where the variable is + # actually x^2 + solve_quadratic(coeff,roots) + + # Yield roots for x by taking square roots of the x^2 + # solution. + roots[4] = sqrt(roots[1]) + roots[3] = -roots[2] + roots[0] = sqrt(roots[0]) + roots[1] = -roots[0] + + # Restore the original positions of the coefficients + coeff[4] = coeff[2] + coeff[2] = coeff[1] + coeff[1] = 0.0j + + + +@njit() +def solve_cubic(coeff,roots): + # TODO + # General soluton not needed for quartic solve + pass + + +@njit() +def solve_depressed_quartic(coeff,roots): + a = coeff[2] + b = coeff[1] + c = coeff[0] + + # To solve the depressed quartic, one must first find one + # root of a cubic polynomial. + + p = (-(a**2)/12) - c + q = (-(a**3)/108) + (a*c/3) - ((b**2)/8) + + cube_const = (-q/2) + sqrt_body = ((q**2)/4) + ((p**3)/27) + w_pos = principal_nth_root( cube_const + sqrt(sqrt_body), 3) + w_neg = principal_nth_root( cube_const - sqrt(sqrt_body), 3) + + # It's reccomended to opt for the larger w when + # calculating the root. + if abs(w_pos) > abs(w_neg): + w = w_pos + else: + w = w_neg + + # A root of the cubic + y = (a/6) + w - (p/(3*w)) + + # The different roots are found by flipping the signs + # of some terms in a formula. There are three sections + # unaffected by these flips, represented below by + # alpha, beta, and gamma + + alpha = sqrt(2*y-a) + beta = -2*y-a + gamma = (2*b) / sqrt(2*y-a) + + roots[0] = ( (-alpha) + sqrt(beta+gamma) ) # - + + + roots[1] = ( (-alpha) - sqrt(beta+gamma) ) # - - + + roots[2] = ( ( alpha) + sqrt(beta-gamma) ) # + + - + roots[3] = ( ( alpha) - sqrt(beta-gamma) ) # + - - + + +@njit() +def solve_quartic(coeff,roots): + # Algorithm logic derived from Wikipedia's quartic + # equation article. (^-^)=b + + # Coefficients for the general solve + a = coeff[4] + b = coeff[3] + c = coeff[2] + d = coeff[1] + e = coeff[0] + + # Coefficients for the sub-solve + # The sub-solve turns the equation into a depressed + # quartic by making u the new variable, with: + # + # x = u - (bg/(4*ag)) + # + # Once the roots for the depressed quartic are found, + # they can be plugged into this equation to yeild the + # roots for x. + + sub_coeff = util.local_array(4,np.complex128) + sub_coeff[4] = 1.0 + 0.0j + sub_coeff[3] = 0.0j + sub_coeff[2] = (-3*(b**2)) / (8*(a**2)) + c/a + sub_coeff[1] = (b**3) / (8*(a**3)) - (b*c) / (2*(a**2)) + d/a + sub_coeff[0] = (-3*(b**4)) / (256*(a**4)) + (c*(b**2)) / (16*(a**3)) - (b*d) / (4*(a**2)) + e/a + + # Get roots of sub-solve + sub_roots = util.local_array(4,np.complex128) + if sub_coeff[1] == 0 : + # If the linear term coefficient is zero, the + # normal depressed quartic solver won't work. + # Instead, it is a biquadratic, and can be + # solved as such. + + solve_biquadratic(sub_coeff,sub_roots) + else: + solve_depressed_quartic(sub_coeff,sub_roots) + + for idx in range(4): + roots[idx] = sub_roots[idx] - b/(4*a) + + + + diff --git a/mcdc/transport/geometry/surface/interface.py b/mcdc/transport/geometry/surface/interface.py index 62e31604b..04e00e47d 100644 --- a/mcdc/transport/geometry/surface/interface.py +++ b/mcdc/transport/geometry/surface/interface.py @@ -21,6 +21,7 @@ import mcdc.transport.geometry.surface.torus_y as torus_y import mcdc.transport.geometry.surface.torus_z as torus_z import mcdc.transport.geometry.surface.torus as torus +import mcdc.transport.util as util from mcdc.constant import ( COINCIDENCE_TOLERANCE, diff --git a/mcdc/transport/geometry/surface/torus.py b/mcdc/transport/geometry/surface/torus.py index 2dcf4e584..1e81cd6c9 100644 --- a/mcdc/transport/geometry/surface/torus.py +++ b/mcdc/transport/geometry/surface/torus.py @@ -29,6 +29,9 @@ from numba import njit +import mcdc.transport.util as util +import mcdc.transport.geometry.root_solve as root_solve + from mcdc.constant import ( COINCIDENCE_TOLERANCE, INF, @@ -164,12 +167,14 @@ def get_distance(particle_container, surface): a0 = (I + R * R - r * r) ** 2 - 4.0 * R * R * L # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; - # np.roots is sufficient for now. - coefficients = np.array( - [a4 + 0.0j, a3 + 0.0j, a2 + 0.0j, a1 + 0.0j, a0 + 0.0j], - dtype=np.complex128, - ) - roots = np.roots(coefficients) + coefficients = util.local_array(5,np.complex128) + coefficients[0] = a4 + 0.0j + coefficients[1] = a3 + 0.0j + coefficients[2] = a2 + 0.0j + coefficients[3] = a1 + 0.0j + coefficients[4] = a0 + 0.0j + roots = util.local_array(4,np.complex128) + root_solve.solve_quartic(coefficients,roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_x.py b/mcdc/transport/geometry/surface/torus_x.py index 79134d192..43723e7c6 100644 --- a/mcdc/transport/geometry/surface/torus_x.py +++ b/mcdc/transport/geometry/surface/torus_x.py @@ -15,6 +15,9 @@ from numba import njit +import mcdc.transport.util as util +import mcdc.transport.geometry.root_solve as root_solve + from mcdc.constant import ( COINCIDENCE_TOLERANCE, INF, @@ -189,11 +192,16 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = np.array( - [a4 + 0.0j, a3 + 0.0j, a2 + 0.0j, a1 + 0.0j, a0 + 0.0j], - dtype=np.complex128, - ) - roots = np.roots(coefficients) + coefficients = util.local_array(5,np.complex128) + coefficients[0] = a4 + 0.0j + coefficients[1] = a3 + 0.0j + coefficients[2] = a2 + 0.0j + coefficients[3] = a1 + 0.0j + coefficients[4] = a0 + 0.0j + roots = util.local_array(4,np.complex128) + #root_solve.solve_quartic(coefficients,roots) + for idx in range(5): + roots[idx] = 0 min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_y.py b/mcdc/transport/geometry/surface/torus_y.py index 571771b6e..19d3f1ebf 100644 --- a/mcdc/transport/geometry/surface/torus_y.py +++ b/mcdc/transport/geometry/surface/torus_y.py @@ -15,6 +15,9 @@ from numba import njit +import mcdc.transport.util as util +import mcdc.transport.geometry.root_solve as root_solve + from mcdc.constant import ( COINCIDENCE_TOLERANCE, INF, @@ -189,11 +192,16 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = np.array( - [a4 + 0.0j, a3 + 0.0j, a2 + 0.0j, a1 + 0.0j, a0 + 0.0j], - dtype=np.complex128, - ) - roots = np.roots(coefficients) + coefficients = util.local_array(5,np.complex128) + coefficients[0] = a4 + 0.0j + coefficients[1] = a3 + 0.0j + coefficients[2] = a2 + 0.0j + coefficients[3] = a1 + 0.0j + coefficients[4] = a0 + 0.0j + roots = util.local_array(4,np.complex128) + #root_solve.solve_quartic(coefficients,roots) + for idx in range(5): + roots[idx] = 0 min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_z.py b/mcdc/transport/geometry/surface/torus_z.py index a4c906cfa..10ed14e0a 100644 --- a/mcdc/transport/geometry/surface/torus_z.py +++ b/mcdc/transport/geometry/surface/torus_z.py @@ -15,6 +15,9 @@ from numba import njit +import mcdc.transport.util as util +import mcdc.transport.geometry.root_solve as root_solve + from mcdc.constant import ( COINCIDENCE_TOLERANCE, INF, @@ -189,11 +192,16 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = np.array( - [a4 + 0.0j, a3 + 0.0j, a2 + 0.0j, a1 + 0.0j, a0 + 0.0j], - dtype=np.complex128, - ) - roots = np.roots(coefficients) + coefficients = util.local_array(5,np.complex128) + coefficients[0] = a4 + 0.0j + coefficients[1] = a3 + 0.0j + coefficients[2] = a2 + 0.0j + coefficients[3] = a1 + 0.0j + coefficients[4] = a0 + 0.0j + roots = util.local_array(4,np.complex128) + #root_solve.solve_quartic(coefficients,roots) + for idx in range(5): + roots[idx] = 0 min_t = INF diff --git a/mcdc/transport/physics/util.py b/mcdc/transport/physics/util.py index 3475a1510..12db31f87 100644 --- a/mcdc/transport/physics/util.py +++ b/mcdc/transport/physics/util.py @@ -24,7 +24,10 @@ def evaluate_neutron_xs_energy_grid(e, nuclide, data): @njit def evaluate_electron_xs_energy_grid(e, element, data): - energy_grid = mcdc_get.element.electron_xs_energy_grid_all(element, data) + offset = element["electron_xs_energy_grid_offset"] + length = element["electron_xs_energy_grid_length"] + energy_grid = data[offset : offset + length] + # Above is equivalent to: energy_grid = mcdc_get.element.electron_xs_energy_grid_all(element, data) idx = find_bin(e, energy_grid) e0 = energy_grid[idx] e1 = energy_grid[idx + 1] diff --git a/mcdc/transport/rng.py b/mcdc/transport/rng.py index 4a8415e1e..3f16ea603 100644 --- a/mcdc/transport/rng.py +++ b/mcdc/transport/rng.py @@ -2,6 +2,7 @@ import numpy as np from numba import uint64, njit +#from mcdc.code_factory.jit import njit # ====================================================================================== # Random number generator From eefc6ef59bf755d0d49eea6f5ec9f13f16363f11 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 02:42:04 -0700 Subject: [PATCH 03/34] Added preliminary array-returning functionality --- mcdc/code_factory/array_return.py | 186 +----------------------------- 1 file changed, 1 insertion(+), 185 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 6c26ffcf7..87edb0fca 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -146,12 +146,11 @@ def array_return_typing(fn,elem_type): from inspect import signature - param_count = len(signature(fn).parameters) + arg_list = ",".join([param for param in signature(fn).parameters]) template = "def typer({arg_list}):\n return nb.types.Array(dtype=elem_type,ndim=1,layout='C')({arg_list})" gns = globals() | {"elem_type":elem_type} lns = {} - arg_list = ",".join([f"param_{i}" for i in range(param_count)]) exec(template.format(arg_list=arg_list),gns,lns) typer = lns["typer"] @@ -260,186 +259,3 @@ def array_return_true_decorator(fn): - - - -#@nb.extending.type_callable(attribute_all) -def type_attribute_all(context): - from numba.np.numpy_support import as_dtype - - context_guard(context) - - def typer(thing, data): - thing_dtype = as_dtype(thing) - if not thing_dtype == thing_type: - raise nb.core.errors.TypingError(f"First argument is type {thing_dtype}, but must be type {thing_type}.") - if not isinstance(data,nb.types.Array): - raise nb.core.errors.TypingError(f"Second argument is type {data_dtype}, but must be an Array of type {data_type}.") - data_dtype = as_dtype(data.dtype) - if not data_dtype == data_type: - raise nb.core.errors.TypingError(f"Second argument has element dtype {data_dtype}, but must be type {data_type}.") - elif not data.ndim == 1: - raise nb.core.errors.TypingError(f"Second argument should be one-dimensional.") - elif not data.layout == "C": - raise nb.core.errors.TypingError(f"Second argument should have C layout.") - - - retty = types.Array(dtype=data.dtype,ndim=1,layout="C") - sig = retty(thing,data) - return sig - - return typer - - -#@nb.extending.lower_builtin(attribute_all,nb.from_dtype(thing_type),nb.types.Array(dtype=nb.from_dtype(data_type),ndim=1,layout="C")) -def builtin_attribute_all(context, builder, sig, args): - - thing, data = args - thing_type, data_type = sig.args - - - import llvmlite.binding as ll - from llvmlite import ir - - try: - import numba.hip as hip - ROCM_AVAILABLE = True - except: - raise nb.core.errors.UnsupportedError(f"ROCM ain't available.") - ROCM_AVAILABLE = False - - if not ROCM_AVAILABLE: - try: - import numba.cuda as cuda - CUDA_AVAILABLE = True - except: - CUDA_AVAILABLE = False - else: - CUDA_AVAILABLE = False - - lmod = builder.module - retty = nb.types.Tuple([nb.types.voidptr,nb.types.uintp]) - ptr_sig = retty(*sig.args) - - res = context.compile_internal(builder,attribute_all_ptr.py_func,ptr_sig,args) - ptr_res = builder.extract_value(res,0) - size_res = builder.extract_value(res,1) - shape = [size_res] - dtype = data_type.dtype - - if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): - targetdata = ll.create_target_data( - nb.hip.amdgcn.DATA_LAYOUT - ) - elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): - targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) - lldtype = context.get_data_type(dtype) - if isinstance(context,nb.core.cpu.CPUContext): - itemsize = context.get_abi_sizeof(lldtype) - elif ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): - itemsize = lldtype.get_abi_size(targetdata) - elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): - itemsize = lldtype.get_abi_size(targetdata) - else: - raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") - - kstrides = [context.get_constant(types.intp,itemsize)] - - aryty = types.Array(dtype=dtype, ndim=1, layout="C") - ary = context.make_array(aryty)(context, builder) - - dataptr = builder.addrspacecast( - ptr_res, ir.PointerType(ir.IntType(8)), "generic" - ) - - kshape = [size_res] - context.populate_array( - ary, - data=builder.bitcast(dataptr, ary.data.type), - shape=kshape, - strides=kstrides, - itemsize=context.get_constant(types.intp, itemsize), - meminfo=None, - ) - print(ary._getvalue()) - return ary._getvalue() - - - - - - -from numba import hip - - -data_type = np.float64 - -thing_a_type = np.dtype([ - ("attribute_a_offset",np.intp), - ("attribute_a_size", np.intp), -]) - -thing_b_type = np.dtype([ - ("attribute_b_offset",np.intp), - ("attribute_b_size", np.intp), -]) - - -# We're cooking with gas now! >B) - -@array_return(nb.float64) -def attribute_a_all(thing, data): - start = thing["attribute_a_offset"] - size = thing["attribute_a_size"] - end = start + size - return array_result(data[start:end]) - - -@array_return(nb.float64) -def attribute_b_all(thing, data): - start = thing["attribute_b_offset"] - size = thing["attribute_b_size"] - end = start + size - return array_result(data[start:end]) - - - - -@njit() -def check(thing_a,thing_b,data): - a = attribute_a_all(thing_a,data) - b = attribute_b_all(thing_b,data) - print(a) - print(b) - - -@nb.hip.jit() -def hip_check(thing_a,thing_b,data): - a = attribute_a_all(thing_a,data) - b = attribute_b_all(thing_b,data) - for i in range(len(a)): - a[i] = -1 - for i in range(len(b)): - b[i] = -2 - - - -data = np.zeros((10,),dtype=data_type) -for i in range(10): - data[i] = i - -thing_a_arr = np.zeros((1,),thing_a_type) -thing_a = thing_a_arr[0] -thing_a["attribute_a_offset"] = 4 -thing_a["attribute_a_size"] = 3 -thing_b_arr = np.zeros((1,),thing_b_type) -thing_b = thing_b_arr[0] -thing_b["attribute_b_offset"] = 8 -thing_b["attribute_b_size"] = 1 - -check(thing_a,thing_b,data) -hip_check[1,1](thing_a,thing_b,data) - -print(data) - - From 023e834751f064c7736a1453c626cae962e35c6b Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 02:43:02 -0700 Subject: [PATCH 04/34] Updated root solvers to use in-house power function for complex numbers. It appears the builtin implementation is not provided on AMD. This, or something like it, would need to be included until there is a real fix. --- mcdc/transport/geometry/root_solve.py | 21 +++++++++++++-------- 1 file changed, 13 insertions(+), 8 deletions(-) diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py index b6ba4990d..948a0aa2e 100644 --- a/mcdc/transport/geometry/root_solve.py +++ b/mcdc/transport/geometry/root_solve.py @@ -15,7 +15,12 @@ def modulus(x): def sqrt(x): return math.sqrt(x.real) * (x+x.real) / modulus(x+x.real) - +@njit() +def power(x,n): + result = 1 + for i in range(n): + result = result * x + return result @njit() def nth_root(x,n,index): @@ -45,7 +50,7 @@ def solve_quadratic(coeff,roots): c = coeff[0] # standard quadratic formula, but with discriminant # calculated separately for re-use - discriminant = sqrt(b**2-4*a*c) + discriminant = sqrt(power(b,2)-4*a*c) roots[0] = ((-b) + discriminant) / (2*a) roots[1] = ((-b) - discriminant) / (2*a) @@ -90,11 +95,11 @@ def solve_depressed_quartic(coeff,roots): # To solve the depressed quartic, one must first find one # root of a cubic polynomial. - p = (-(a**2)/12) - c - q = (-(a**3)/108) + (a*c/3) - ((b**2)/8) + p = (-power(a,2)/12) - c + q = (-power(a,3)/108) + (a*c/3) - (power(b,2)/8) cube_const = (-q/2) - sqrt_body = ((q**2)/4) + ((p**3)/27) + sqrt_body = (power(q,2)/4) + (power(p,3)/27) w_pos = principal_nth_root( cube_const + sqrt(sqrt_body), 3) w_neg = principal_nth_root( cube_const - sqrt(sqrt_body), 3) @@ -148,9 +153,9 @@ def solve_quartic(coeff,roots): sub_coeff = util.local_array(4,np.complex128) sub_coeff[4] = 1.0 + 0.0j sub_coeff[3] = 0.0j - sub_coeff[2] = (-3*(b**2)) / (8*(a**2)) + c/a - sub_coeff[1] = (b**3) / (8*(a**3)) - (b*c) / (2*(a**2)) + d/a - sub_coeff[0] = (-3*(b**4)) / (256*(a**4)) + (c*(b**2)) / (16*(a**3)) - (b*d) / (4*(a**2)) + e/a + sub_coeff[2] = (-3*power(b,2)) / (8*power(a,2)) + c/a + sub_coeff[1] = power(b,3) / (8*power(a,3)) - (b*c) / (2*power(a,2)) + d/a + sub_coeff[0] = (-3*power(b,4)) / (256*power(a,4)) + (c*power(b,2)) / (16*power(a,3)) - (b*d) / (4*power(a,2)) + e/a # Get roots of sub-solve sub_roots = util.local_array(4,np.complex128) From 8a2f0ee21d59cdebde632119f331894435610a83 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 02:48:25 -0700 Subject: [PATCH 05/34] Fixed missing allocator function references in gpu program builder --- mcdc/code_factory/gpu/program_builder.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index f62a6d538..e48235d46 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -324,6 +324,9 @@ def step(program: nb.uintp, particle_input: particle_gpu): overwrite_func(impl, impl) """ + alloc_managed_bytes = harmonize.alloc_managed_bytes + alloc_device_bytes = harmonize.alloc_device_bytes + # ====================================================================================== # Setup GPU From e0b30c7439cc26962f0329db2117e69d65b84e6d Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 02:49:03 -0700 Subject: [PATCH 06/34] Added preliminary array returning functionality to numba objects generator --- mcdc/code_factory/numba_layers_generator.py | 20 +++++++++++++++----- 1 file changed, 15 insertions(+), 5 deletions(-) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index a660927d1..9eb9606a7 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -1090,7 +1090,13 @@ def generate_mcdc_access(targets): text_getter += "from numpy import int64\n" text_getter += "from numba import njit\n\n\n" text_setter += "from numba import njit\n\n\n" + text_getter += ( + "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" + ) + text_getter += ( + "import numba as nb\n\n\n" + ) for attribute in targets[object_name]: attribute_name = attribute.name shape = attribute.shape @@ -1100,7 +1106,7 @@ def generate_mcdc_access(targets): text_getter += _accessor_1d_element( object_name, attribute_name, cast_to_int=cast_to_int ) - text_getter += _accessor_1d_all(object_name, attribute_name, shape[0]) + text_getter += _accessor_1d_all(object_name, attribute_name, shape[0], nb.types.float64) text_getter += _accessor_1d_last( object_name, attribute_name, @@ -1110,7 +1116,7 @@ def generate_mcdc_access(targets): text_setter += _accessor_1d_element(object_name, attribute_name, True) text_setter += _accessor_1d_all( - object_name, attribute_name, shape[0], True + object_name, attribute_name, shape[0], nb.types.float64, True ) text_setter += _accessor_1d_last( object_name, attribute_name, shape[0], True @@ -1174,6 +1180,9 @@ def generate_mcdc_access(targets): text = "# The following is automatically generated by code_factory.py\n\n" for i, object_name in enumerate(targets.keys()): text += f"import mcdc.mcdc_{key}.{object_name} as {object_name}\n" + text += ( + "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" + ) if i < len(targets.keys()) - 1: text += "\n" f.write(text) @@ -1214,11 +1223,12 @@ def _accessor_1d_element(object_name, attribute_name, setter=False, cast_to_int= -def _accessor_1d_all(object_name, attribute_name, size, setter=False): - text = f"@njit\n" +def _accessor_1d_all(object_name, attribute_name, size, dtype, setter=False): if setter: + text = f"@njit\n" text += f"def {attribute_name}_all({object_name}, data, value):\n" else: + text = f"@array_return(nb.types.float64)\n" text += f"def {attribute_name}_all({object_name}, data):\n" text += f' start = {object_name}["{attribute_name}_offset"]\n' text += accessor_dimension("size", size, object_name) @@ -1226,7 +1236,7 @@ def _accessor_1d_all(object_name, attribute_name, size, setter=False): if setter: text += f" data[start:end] = value\n\n\n" else: - text += f" return data[start:end]\n\n\n" + text += f" return array_result(data[start:end])\n\n\n" return text From b77ebe480f3a6c6e349f14ac97770180d616b92c Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 02:50:18 -0700 Subject: [PATCH 07/34] Updated main to use gpu setup function --- mcdc/main.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mcdc/main.py b/mcdc/main.py index d26946c78..742540f13 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -174,8 +174,8 @@ def prepare(simulationPy: Simulation): # ================================================================================== if config.target == "gpu": - from mcdc.code_factory.gpu.program_builder import build_gpu_program - setup_gpu_program(mcdc_container, data) + from mcdc.code_factory.gpu.program_builder import setup_gpu_program + setup_gpu_program(simulation_container, data) # ================================================================================== From 37f590af53b534f481b7eee94448040a1631b78f Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 3 Jun 2026 15:15:42 -0700 Subject: [PATCH 08/34] Back in black --- mcdc/code_factory/array_return.py | 102 +++++++++--------- mcdc/code_factory/gpu/program_builder.py | 2 +- mcdc/code_factory/jit.py | 30 ++---- mcdc/code_factory/numba_layers_generator.py | 17 ++- mcdc/main.py | 1 + mcdc/transport/geometry/root_solve.py | 111 ++++++++++---------- mcdc/transport/geometry/surface/torus.py | 6 +- mcdc/transport/geometry/surface/torus_x.py | 6 +- mcdc/transport/geometry/surface/torus_y.py | 6 +- mcdc/transport/geometry/surface/torus_z.py | 6 +- mcdc/transport/rng.py | 3 +- 11 files changed, 144 insertions(+), 146 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 87edb0fca..c976ceae1 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -4,6 +4,7 @@ import numpy as np import cffi + ffi = cffi.FFI() @@ -66,11 +67,13 @@ def codegen(context, builder, signature, args): return sig, codegen + @njit() def uintp_to_voidptr(value): val = nb.uintp(value) return cast_uintp_to_voidptr(val) + @njit() def voidptr_to_uintp(value): return cast_voidptr_to_uintp(value) @@ -80,57 +83,65 @@ def voidptr_to_uintp(value): def into_voidptr(value): return into_voidptr_python(value) + def into_voidptr_python(value): raise RuntimeError("`into_voidptr` is only supported in nopython mode.") + @nb.extending.overload(into_voidptr_python) def into_voidptr_overload(value): - if isinstance(value,nb.types.Array) : + if isinstance(value, nb.types.Array): + def impl(value): - ptr = (ffi.from_buffer(value)) + ptr = ffi.from_buffer(value) vptr = cast_any_to_voidptr(ptr) return vptr + return impl - elif isinstance(value,nb.types.CPointer) : + elif isinstance(value, nb.types.CPointer): + def impl(value): return cast_any_to_voidptr(value) + return impl - elif isinstance(value,nb.types.Integer) : + elif isinstance(value, nb.types.Integer): + def impl(value): return cast_uintp_to_voidptr(value) + return impl - else : + else: raise RuntimeError(f"`into_voidptr` is not supported for type '{value}'") - - - ############################################################################### # New Code ############################################################################### - def array_result(array): return array + @nb.extending.overload(array_result) def array_result_overload(array): - if not isinstance(array,types.Array): - raise nb.core.errors.TypingError(f"Expected array type argument for array_result, got {array}.") + if not isinstance(array, types.Array): + raise nb.core.errors.TypingError( + f"Expected array type argument for array_result, got {array}." + ) print("OVER LOADED") + def impl(array): - return (into_voidptr(array),len(array)) - return impl + return (into_voidptr(array), len(array)) + return impl def context_guard(context): - if isinstance(context, nb.core.typing.context.Context): + if isinstance(context, nb.core.typing.context.Context): pass elif isinstance(context, nb.cuda.target.CUDATypingContext): pass @@ -140,37 +151,35 @@ def context_guard(context): raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") +def array_return_typing(fn, elem_type): - -def array_return_typing(fn,elem_type): - - from inspect import signature + arg_list = ",".join([param for param in signature(fn).parameters]) template = "def typer({arg_list}):\n return nb.types.Array(dtype=elem_type,ndim=1,layout='C')({arg_list})" - - gns = globals() | {"elem_type":elem_type} + + gns = globals() | {"elem_type": elem_type} lns = {} - exec(template.format(arg_list=arg_list),gns,lns) + exec(template.format(arg_list=arg_list), gns, lns) typer = lns["typer"] def typer_factory(context): from numba.np.numpy_support import as_dtype context_guard(context) - + return typer nb.extending.type_callable(fn)(typer_factory) -def array_return_lowering(fn,elem_type): - +def array_return_lowering(fn, elem_type): from inspect import signature + param_count = len(signature(fn).parameters) - retty = nb.types.Array(dtype=elem_type,ndim=1,layout="C") - sig = retty(*([nb.types.Any]*param_count)) + retty = nb.types.Array(dtype=elem_type, ndim=1, layout="C") + sig = retty(*([nb.types.Any] * param_count)) jit_fn = nb.njit(fn) @@ -179,12 +188,12 @@ def builtin(context, builder, sig, args): thing, data = args thing_type, data_type = sig.args - import llvmlite.binding as ll from llvmlite import ir try: import numba.hip as hip + ROCM_AVAILABLE = True except: ROCM_AVAILABLE = False @@ -192,39 +201,40 @@ def builtin(context, builder, sig, args): if not ROCM_AVAILABLE: try: import numba.cuda as cuda + CUDA_AVAILABLE = True except: CUDA_AVAILABLE = False else: - CUDA_AVAILABLE = False + CUDA_AVAILABLE = False - lmod = builder.module - retty = nb.types.Tuple([nb.types.voidptr,nb.types.uintp]) - ptr_sig = retty(*sig.args) + lmod = builder.module + retty = nb.types.Tuple([nb.types.voidptr, nb.types.uintp]) + ptr_sig = retty(*sig.args) - res = context.compile_internal(builder,jit_fn.py_func,ptr_sig,args) - ptr_res = builder.extract_value(res,0) - size_res = builder.extract_value(res,1) + res = context.compile_internal(builder, jit_fn.py_func, ptr_sig, args) + ptr_res = builder.extract_value(res, 0) + size_res = builder.extract_value(res, 1) shape = [size_res] dtype = data_type.dtype if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): - targetdata = ll.create_target_data( - nb.hip.amdgcn.DATA_LAYOUT - ) + targetdata = ll.create_target_data(nb.hip.amdgcn.DATA_LAYOUT) elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) lldtype = context.get_data_type(dtype) - if isinstance(context,nb.core.cpu.CPUContext): + if isinstance(context, nb.core.cpu.CPUContext): itemsize = context.get_abi_sizeof(lldtype) elif ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): itemsize = lldtype.get_abi_size(targetdata) elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): itemsize = lldtype.get_abi_size(targetdata) else: - raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") + raise nb.core.errors.UnsupportedError( + f"Unsupported target context {context}." + ) - kstrides = [context.get_constant(types.intp,itemsize)] + kstrides = [context.get_constant(types.intp, itemsize)] aryty = types.Array(dtype=elem_type, ndim=1, layout="C") ary = context.make_array(aryty)(context, builder) @@ -245,17 +255,13 @@ def builtin(context, builder, sig, args): print(ary._getvalue()) return ary._getvalue() - nb.extending.lower_builtin(fn,*sig.args)(builtin) - - + nb.extending.lower_builtin(fn, *sig.args)(builtin) def array_return(sig): def array_return_true_decorator(fn): - array_return_typing(fn,sig) - array_return_lowering(fn,sig) + array_return_typing(fn, sig) + array_return_lowering(fn, sig) return fn - return array_return_true_decorator - - + return array_return_true_decorator diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index e48235d46..f4e0fed28 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -325,7 +325,7 @@ def step(program: nb.uintp, particle_input: particle_gpu): """ alloc_managed_bytes = harmonize.alloc_managed_bytes - alloc_device_bytes = harmonize.alloc_device_bytes + alloc_device_bytes = harmonize.alloc_device_bytes # ====================================================================================== diff --git a/mcdc/code_factory/jit.py b/mcdc/code_factory/jit.py index 224f58d18..bf36e963d 100644 --- a/mcdc/code_factory/jit.py +++ b/mcdc/code_factory/jit.py @@ -22,41 +22,33 @@ def arg_check_{name}_overload({args}): """ -def wrap_with_check(fn,njit_fn): +def wrap_with_check(fn, njit_fn): nb.extending.register_jitable(fn) - arg_names = getfullargspec(fn).args - args_str = ",".join(arg_names) + arg_names = getfullargspec(fn).args + args_str = ",".join(arg_names) arg_names_str = ",".join(f'"{n}"' for n in arg_names) name = fn.__name__ print(f"wrapping {name}") - gns = {"nb":nb,"fn":fn,"njit_fn":njit_fn} + gns = {"nb": nb, "fn": fn, "njit_fn": njit_fn} lns = {} - code = njit_trace_template.format( - args=args_str, - arg_names=arg_names_str, - name=name - ) - exec(code,gns,lns) + code = njit_trace_template.format(args=args_str, arg_names=arg_names_str, name=name) + exec(code, gns, lns) return lns[f"arg_check_{name}"] -def njit(*args,**kwargs): +def njit(*args, **kwargs): if (len(args) == 1) and (len(kwargs) == 0): if not ARG_CHECK: return njit(args[0]) fn = args[0] njit_fn = nb.njit(args[0]) - return wrap_with_check(fn,njit_fn) + return wrap_with_check(fn, njit_fn) else: if not ARG_CHECK: - return nb.njit(*args,**kwargs) + return nb.njit(*args, **kwargs) def wrapper(fn): - njit_fn = nb.njit(*args,**kwargs)(fn) - return wrap_with_check(fn,njit_fn) - - - - + njit_fn = nb.njit(*args, **kwargs)(fn) + return wrap_with_check(fn, njit_fn) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 9eb9606a7..6f17abd79 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -1090,13 +1090,9 @@ def generate_mcdc_access(targets): text_getter += "from numpy import int64\n" text_getter += "from numba import njit\n\n\n" text_setter += "from numba import njit\n\n\n" - text_getter += ( - "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" - ) + text_getter += "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" - text_getter += ( - "import numba as nb\n\n\n" - ) + text_getter += "import numba as nb\n\n\n" for attribute in targets[object_name]: attribute_name = attribute.name shape = attribute.shape @@ -1106,7 +1102,9 @@ def generate_mcdc_access(targets): text_getter += _accessor_1d_element( object_name, attribute_name, cast_to_int=cast_to_int ) - text_getter += _accessor_1d_all(object_name, attribute_name, shape[0], nb.types.float64) + text_getter += _accessor_1d_all( + object_name, attribute_name, shape[0], nb.types.float64 + ) text_getter += _accessor_1d_last( object_name, attribute_name, @@ -1180,9 +1178,7 @@ def generate_mcdc_access(targets): text = "# The following is automatically generated by code_factory.py\n\n" for i, object_name in enumerate(targets.keys()): text += f"import mcdc.mcdc_{key}.{object_name} as {object_name}\n" - text += ( - "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" - ) + text += "from mcdc.code_factory.array_return import array_return, array_result\n\n\n" if i < len(targets.keys()) - 1: text += "\n" f.write(text) @@ -1222,7 +1218,6 @@ def _accessor_1d_element(object_name, attribute_name, setter=False, cast_to_int= return text - def _accessor_1d_all(object_name, attribute_name, size, dtype, setter=False): if setter: text = f"@njit\n" diff --git a/mcdc/main.py b/mcdc/main.py index 742540f13..7874c2eec 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -175,6 +175,7 @@ def prepare(simulationPy: Simulation): if config.target == "gpu": from mcdc.code_factory.gpu.program_builder import setup_gpu_program + setup_gpu_program(simulation_container, data) diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py index 948a0aa2e..fa7aa0546 100644 --- a/mcdc/transport/geometry/root_solve.py +++ b/mcdc/transport/geometry/root_solve.py @@ -9,62 +9,65 @@ @njit() def modulus(x): - return math.sqrt(x.real**2+x.imag**2) + return math.sqrt(x.real**2 + x.imag**2) + @njit() def sqrt(x): - return math.sqrt(x.real) * (x+x.real) / modulus(x+x.real) + return math.sqrt(x.real) * (x + x.real) / modulus(x + x.real) + @njit() -def power(x,n): +def power(x, n): result = 1 for i in range(n): result = result * x return result + @njit() -def nth_root(x,n,index): +def nth_root(x, n, index): # First, convert to polar form r = modulus(x) - a = math.atan2(x.imag,x.real) - + a = math.atan2(x.imag, x.real) + # Apply de Moivre's Formula - root_modulus = math.pow(r,1./n) - root_argument = (a+2*math.pi*index)/n - + root_modulus = math.pow(r, 1.0 / n) + root_argument = (a + 2 * math.pi * index) / n + # ...then convert back to rectangular form real = root_modulus * math.cos(root_argument) imag = root_modulus * math.sin(root_argument) - return complex(real,imag) - + return complex(real, imag) + @njit() -def principal_nth_root(x,n): - return nth_root(x,n,0) +def principal_nth_root(x, n): + return nth_root(x, n, 0) @njit() -def solve_quadratic(coeff,roots): +def solve_quadratic(coeff, roots): a = coeff[2] b = coeff[1] c = coeff[0] # standard quadratic formula, but with discriminant # calculated separately for re-use - discriminant = sqrt(power(b,2)-4*a*c) - roots[0] = ((-b) + discriminant) / (2*a) - roots[1] = ((-b) - discriminant) / (2*a) + discriminant = sqrt(power(b, 2) - 4 * a * c) + roots[0] = ((-b) + discriminant) / (2 * a) + roots[1] = ((-b) - discriminant) / (2 * a) @njit() -def solve_biquadratic(coeff,roots): +def solve_biquadratic(coeff, roots): # Move each coefficient down to one-half it's power coeff[1] = coeff[2] coeff[2] = coeff[4] # Solve as quadratic equation, where the variable is # actually x^2 - solve_quadratic(coeff,roots) - + solve_quadratic(coeff, roots) + # Yield roots for x by taking square roots of the x^2 # solution. roots[4] = sqrt(roots[1]) @@ -78,30 +81,29 @@ def solve_biquadratic(coeff,roots): coeff[1] = 0.0j - @njit() -def solve_cubic(coeff,roots): +def solve_cubic(coeff, roots): # TODO # General soluton not needed for quartic solve pass @njit() -def solve_depressed_quartic(coeff,roots): +def solve_depressed_quartic(coeff, roots): a = coeff[2] b = coeff[1] c = coeff[0] - + # To solve the depressed quartic, one must first find one # root of a cubic polynomial. - - p = (-power(a,2)/12) - c - q = (-power(a,3)/108) + (a*c/3) - (power(b,2)/8) - - cube_const = (-q/2) - sqrt_body = (power(q,2)/4) + (power(p,3)/27) - w_pos = principal_nth_root( cube_const + sqrt(sqrt_body), 3) - w_neg = principal_nth_root( cube_const - sqrt(sqrt_body), 3) + + p = (-power(a, 2) / 12) - c + q = (-power(a, 3) / 108) + (a * c / 3) - (power(b, 2) / 8) + + cube_const = -q / 2 + sqrt_body = (power(q, 2) / 4) + (power(p, 3) / 27) + w_pos = principal_nth_root(cube_const + sqrt(sqrt_body), 3) + w_neg = principal_nth_root(cube_const - sqrt(sqrt_body), 3) # It's reccomended to opt for the larger w when # calculating the root. @@ -111,25 +113,25 @@ def solve_depressed_quartic(coeff,roots): w = w_neg # A root of the cubic - y = (a/6) + w - (p/(3*w)) + y = (a / 6) + w - (p / (3 * w)) # The different roots are found by flipping the signs # of some terms in a formula. There are three sections # unaffected by these flips, represented below by # alpha, beta, and gamma - alpha = sqrt(2*y-a) - beta = -2*y-a - gamma = (2*b) / sqrt(2*y-a) + alpha = sqrt(2 * y - a) + beta = -2 * y - a + gamma = (2 * b) / sqrt(2 * y - a) - roots[0] = ( (-alpha) + sqrt(beta+gamma) ) # - + + - roots[1] = ( (-alpha) - sqrt(beta+gamma) ) # - - + - roots[2] = ( ( alpha) + sqrt(beta-gamma) ) # + + - - roots[3] = ( ( alpha) - sqrt(beta-gamma) ) # + - - + roots[0] = (-alpha) + sqrt(beta + gamma) # - + + + roots[1] = (-alpha) - sqrt(beta + gamma) # - - + + roots[2] = (alpha) + sqrt(beta - gamma) # + + - + roots[3] = (alpha) - sqrt(beta - gamma) # + - - @njit() -def solve_quartic(coeff,roots): +def solve_quartic(coeff, roots): # Algorithm logic derived from Wikipedia's quartic # equation article. (^-^)=b @@ -150,28 +152,29 @@ def solve_quartic(coeff,roots): # they can be plugged into this equation to yeild the # roots for x. - sub_coeff = util.local_array(4,np.complex128) + sub_coeff = util.local_array(4, np.complex128) sub_coeff[4] = 1.0 + 0.0j sub_coeff[3] = 0.0j - sub_coeff[2] = (-3*power(b,2)) / (8*power(a,2)) + c/a - sub_coeff[1] = power(b,3) / (8*power(a,3)) - (b*c) / (2*power(a,2)) + d/a - sub_coeff[0] = (-3*power(b,4)) / (256*power(a,4)) + (c*power(b,2)) / (16*power(a,3)) - (b*d) / (4*power(a,2)) + e/a + sub_coeff[2] = (-3 * power(b, 2)) / (8 * power(a, 2)) + c / a + sub_coeff[1] = power(b, 3) / (8 * power(a, 3)) - (b * c) / (2 * power(a, 2)) + d / a + sub_coeff[0] = ( + (-3 * power(b, 4)) / (256 * power(a, 4)) + + (c * power(b, 2)) / (16 * power(a, 3)) + - (b * d) / (4 * power(a, 2)) + + e / a + ) # Get roots of sub-solve - sub_roots = util.local_array(4,np.complex128) - if sub_coeff[1] == 0 : + sub_roots = util.local_array(4, np.complex128) + if sub_coeff[1] == 0: # If the linear term coefficient is zero, the # normal depressed quartic solver won't work. # Instead, it is a biquadratic, and can be # solved as such. - solve_biquadratic(sub_coeff,sub_roots) + solve_biquadratic(sub_coeff, sub_roots) else: - solve_depressed_quartic(sub_coeff,sub_roots) + solve_depressed_quartic(sub_coeff, sub_roots) for idx in range(4): - roots[idx] = sub_roots[idx] - b/(4*a) - - - - + roots[idx] = sub_roots[idx] - b / (4 * a) diff --git a/mcdc/transport/geometry/surface/torus.py b/mcdc/transport/geometry/surface/torus.py index 1e81cd6c9..6c3e5c096 100644 --- a/mcdc/transport/geometry/surface/torus.py +++ b/mcdc/transport/geometry/surface/torus.py @@ -167,14 +167,14 @@ def get_distance(particle_container, surface): a0 = (I + R * R - r * r) ** 2 - 4.0 * R * R * L # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; - coefficients = util.local_array(5,np.complex128) + coefficients = util.local_array(5, np.complex128) coefficients[0] = a4 + 0.0j coefficients[1] = a3 + 0.0j coefficients[2] = a2 + 0.0j coefficients[3] = a1 + 0.0j coefficients[4] = a0 + 0.0j - roots = util.local_array(4,np.complex128) - root_solve.solve_quartic(coefficients,roots) + roots = util.local_array(4, np.complex128) + root_solve.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_x.py b/mcdc/transport/geometry/surface/torus_x.py index 43723e7c6..178e11e3e 100644 --- a/mcdc/transport/geometry/surface/torus_x.py +++ b/mcdc/transport/geometry/surface/torus_x.py @@ -192,14 +192,14 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = util.local_array(5,np.complex128) + coefficients = util.local_array(5, np.complex128) coefficients[0] = a4 + 0.0j coefficients[1] = a3 + 0.0j coefficients[2] = a2 + 0.0j coefficients[3] = a1 + 0.0j coefficients[4] = a0 + 0.0j - roots = util.local_array(4,np.complex128) - #root_solve.solve_quartic(coefficients,roots) + roots = util.local_array(4, np.complex128) + # root_solve.solve_quartic(coefficients,roots) for idx in range(5): roots[idx] = 0 diff --git a/mcdc/transport/geometry/surface/torus_y.py b/mcdc/transport/geometry/surface/torus_y.py index 19d3f1ebf..75a9c18a2 100644 --- a/mcdc/transport/geometry/surface/torus_y.py +++ b/mcdc/transport/geometry/surface/torus_y.py @@ -192,14 +192,14 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = util.local_array(5,np.complex128) + coefficients = util.local_array(5, np.complex128) coefficients[0] = a4 + 0.0j coefficients[1] = a3 + 0.0j coefficients[2] = a2 + 0.0j coefficients[3] = a1 + 0.0j coefficients[4] = a0 + 0.0j - roots = util.local_array(4,np.complex128) - #root_solve.solve_quartic(coefficients,roots) + roots = util.local_array(4, np.complex128) + # root_solve.solve_quartic(coefficients,roots) for idx in range(5): roots[idx] = 0 diff --git a/mcdc/transport/geometry/surface/torus_z.py b/mcdc/transport/geometry/surface/torus_z.py index 10ed14e0a..283592eeb 100644 --- a/mcdc/transport/geometry/surface/torus_z.py +++ b/mcdc/transport/geometry/surface/torus_z.py @@ -192,14 +192,14 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. - coefficients = util.local_array(5,np.complex128) + coefficients = util.local_array(5, np.complex128) coefficients[0] = a4 + 0.0j coefficients[1] = a3 + 0.0j coefficients[2] = a2 + 0.0j coefficients[3] = a1 + 0.0j coefficients[4] = a0 + 0.0j - roots = util.local_array(4,np.complex128) - #root_solve.solve_quartic(coefficients,roots) + roots = util.local_array(4, np.complex128) + # root_solve.solve_quartic(coefficients,roots) for idx in range(5): roots[idx] = 0 diff --git a/mcdc/transport/rng.py b/mcdc/transport/rng.py index 3f16ea603..afb799d16 100644 --- a/mcdc/transport/rng.py +++ b/mcdc/transport/rng.py @@ -2,7 +2,8 @@ import numpy as np from numba import uint64, njit -#from mcdc.code_factory.jit import njit + +# from mcdc.code_factory.jit import njit # ====================================================================================== # Random number generator From 330fe58133bd5750ac9cde682b3e136bbb857973 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 11 Jun 2026 11:52:12 -0700 Subject: [PATCH 09/34] Fixing separate array usage --- mcdc/code_factory/array_return.py | 3 - mcdc/code_factory/gpu/program_builder.py | 46 +++++++++------ mcdc/code_factory/gpu/transport/simulation.py | 8 +-- mcdc/code_factory/numba_layers_generator.py | 58 ++++++++++--------- 4 files changed, 63 insertions(+), 52 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index c976ceae1..4cfe6c308 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -132,8 +132,6 @@ def array_result_overload(array): f"Expected array type argument for array_result, got {array}." ) - print("OVER LOADED") - def impl(array): return (into_voidptr(array), len(array)) @@ -252,7 +250,6 @@ def builtin(context, builder, sig, args): itemsize=context.get_constant(types.intp, itemsize), meminfo=None, ) - print(ary._getvalue()) return ary._getvalue() nb.extending.lower_builtin(fn, *sig.args)(builtin) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index f4e0fed28..e20aba343 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -344,7 +344,6 @@ def setup_gpu_program(simulation_container, data): set_device(device_id) simulation["gpu_meta"]["state_pointer"] = cast_voidptr_to_uintp(alloc_state()) - if config.gpu_state_storage == "separate": store_pointer_state_device_simulation( simulation["gpu_meta"]["state_pointer"], @@ -377,24 +376,33 @@ def teardown_gpu_program(simulation): # ====================================================================================== -def create_data_array(size, dtype): - if config.gpu_state_storage == "managed": - data_tally_ptr = harmonize.alloc_managed_bytes(size) - else: - data_tally_ptr = harmonize.alloc_device_bytes(size) - data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) - data_tally = nb.carray(data_tally_ptr, (size,), dtype) - return data_tally, data_tally_uint - - -def create_mcdc_container(dtype): - if config.gpu_state_storage == "managed": - mcdc_ptr = harmonize.alloc_managed_bytes(dtype.itemsize) - else: - mcdc_ptr = harmonize.alloc_device_bytes(dtype.itemsize) - mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) - mcdc_container = nb.carray(mcdc_ptr, (1,), dtype) - return mcdc_container, mcdc_uint +#def create_data_array(size, dtype): +# if config.gpu_state_storage == "managed": +# data_tally_ptr = harmonize.alloc_managed_bytes(size) +# else: +# data_tally_ptr = harmonize.alloc_device_bytes(size) +# data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) +# +# if config.gpu_state_storage == "separate": +# data_tally = nb.zeros( (size,),dtype=dtype) +# else: +# data_tally = nb.carray(data_tally_ptr, (size,), dtype) +# return data_tally, data_tally_uint + + +#def create_mcdc_container(dtype): +# if config.gpu_state_storage == "managed": +# mcdc_ptr = harmonize.alloc_managed_bytes(dtype.itemsize) +# else: +# mcdc_ptr = harmonize.alloc_device_bytes(dtype.itemsize) +# +# mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) +# if config.gpu_state_storage == "separate": +# mcdc_tally = nb.zeros((size,),dtype=dtype) +# else: +# mcdc_tally = nb.carray(mcdc_ptr, (size,), dtype) +# mcdc_container = nb.carray(mcdc_ptr, (1,), dtype) +# return mcdc_container, mcdc_uint # ====================================================================================== diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 50dc60f96..f0692eddb 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -40,10 +40,10 @@ def source_loop(seed, simulation, data): # Store the global state to the GPU if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: harmonize.memcpy_host_to_device( - simulation["gpu_meta"]["state_pointer"], simulation + simulation["gpu_meta"]["simulation_pointer"], simulation ) harmonize.memcpy_host_to_device( - simulation["gpu_meta"]["state_pointer"], data + simulation["gpu_meta"]["data_pointer"], data ) # Execute the program, and continue to do so until it is done @@ -163,10 +163,10 @@ def real_setup_gpu(mcdc_array, data_tally): # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) if config.gpu_state_storage == "separate": harmonize.memcpy_device_to_host( - simulation, simulation["gpu_meta"]["state_pointer"] + simulation, simulation["gpu_meta"]["simulation_pointer"] ) harmonize.memcpy_device_to_host( - data, simulation["gpu_meta"]["state_pointer"] + data, simulation["gpu_meta"]["data_pointer"] ) gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 6f17abd79..ce24e5b1b 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -443,6 +443,9 @@ def generate_numba_layers(simulation): # Manually set particle bank attributes for name in bank_names: mcdc_simulation[name]["tag"] = getattr(simulation, name).tag + + mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer + mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] # GPU program setup if config.target == "gpu": @@ -788,28 +791,27 @@ def set_object( # ============================================================================= -def create_data_array(size, dtype, itemsize): - if config.target == "gpu": - import mcdc.code_factory.gpu.adapt as adapt - import harmonize, numba - - print("\n\n\nW\n\n\n",flush=True) - print(f"Tally size is {size} with itemsize {itemsize}",flush=True) - if config.gpu_state_storage == "managed": - print("\n\n\nX\n\n\n",flush=True) - data_tally_ptr = harmonize.alloc_managed_bytes(size*itemsize) - else: - print("\n\n\nY\n\n\n",flush=True) - data_tally_ptr = harmonize.alloc_device_bytes(size*itemsize) - print("\n\n\nZ\n\n\n",flush=True) - data_tally_uint = adapt.voidptr_to_uintp(data_tally_ptr) - data_tally = numba.carray(data_tally_ptr, (size,), dtype) - return data_tally, data_tally_uint +def create_data_array(size): + if not config.target == "gpu": + data = np.zeros(size, dtype=np.float64) + return data, 0 + else: + return create_data_array_on_gpu(nb.types.float64,size,size*8) + + +@njit +def create_data_array_on_gpu(dtype,size,byte_size): + if config.gpu_state_storage == "managed": + data_tally_ptr = gpu_builder.alloc_managed_bytes(byte_size) else: - data_ptr = gpu_builder.alloc_device_bytes(size) - data_uint = voidptr_to_uintp(data_ptr) - data = nb.carray(data_ptr, (size,), dtype=np.float64) - return data, data_uint + data_tally_ptr = gpu_builder.alloc_device_bytes(byte_size) + data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) + + if config.gpu_state_storage == "separate": + data_tally = np.zeros( (size,),dtype=dtype) + else: + data_tally = nb.carray(data_tally_ptr, (size,), dtype) + return data_tally, data_tally_uint def create_simulation_container(dtype): @@ -823,12 +825,16 @@ def create_simulation_container(dtype): @njit def create_simulation_container_on_gpu(dtype, size): if config.gpu_state_storage == "managed": - simulation_ptr = gpu_builder.alloc_managed_bytes(size) + mcdc_ptr = gpu_builder.alloc_managed_bytes(size) + else: + mcdc_ptr = gpu_builder.alloc_device_bytes(size) + mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) + + if config.gpu_state_storage == "separate": + mcdc_container = np.zeros((1,),dtype=dtype) else: - simulation_ptr = gpu_builder.alloc_device_bytes(size) - simulation_uint = voidptr_to_uintp(simulation_ptr) - simulation = nb.carray(simulation_ptr, (1,), dtype) - return simulation, simulation_uint + mcdc_container = nb.carray(mcdc_ptr, (1,), dtype) + return mcdc_container, mcdc_uint # ============================================================================= From d9fd753e8fa042c20a42a00c995685dc2a808a4f Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 16 Jun 2026 00:26:01 -0700 Subject: [PATCH 10/34] Reached close output between GPU and CPU in Kobayashi-TD --- mcdc/code_factory/gpu/program_builder.py | 8 +++-- mcdc/code_factory/gpu/transport/simulation.py | 5 +++ mcdc/code_factory/gpu/transport/util.py | 23 +++++++++--- mcdc/main.py | 2 ++ mcdc/numba_types.py | 35 ++++++++++--------- mcdc/object_/particle.py | 2 ++ mcdc/object_/simulation.py | 3 ++ mcdc/transport/particle_bank.py | 31 +++++++++------- mcdc/transport/simulation.py | 16 +++++++++ mcdc/transport/util.py | 3 +- 10 files changed, 93 insertions(+), 35 deletions(-) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index e20aba343..6e3d37868 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -217,8 +217,7 @@ def make_work(program: nb.uintp) -> nb.boolean: data_ptr = access_data_ptr(program) data = harmonize.array_from_ptr(data_ptr, shape, nb.float64) - util.atomic_add(simulation["mpi_work_iter"], 0, 1) - idx_work = simulation["mpi_work_iter"][0] + idx_work = util.atomic_add(simulation["mpi_work_iter"], 0, 1) if idx_work >= simulation["mpi_work_size"]: return False @@ -252,6 +251,11 @@ def step(program: nb.uintp, particle_input: particle_gpu): particle_container = util.local_array(1, type_.particle) particle_container[0] = particle_input particle = particle_container[0] + particle["alive"] = True + particle["material_ID"] = -1 + particle["cell_ID"] = -1 + particle["surface_ID"] = -1 + particle["event"] = -1 particle["fresh"] = False step_particle(particle_container, program, data) if particle["alive"]: diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index f0692eddb..fedb03820 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -25,6 +25,8 @@ def source_loop(seed, simulation, data): full_work_size = simulation["mpi_work_size"] + print("Gen count before: ",simulation["gen_count"][0]) + simulation["gen_count"][0] = 0 if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: phase_size = 1000000000 else: @@ -39,6 +41,7 @@ def source_loop(seed, simulation, data): # Store the global state to the GPU if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: + print("STORING!") harmonize.memcpy_host_to_device( simulation["gpu_meta"]["simulation_pointer"], simulation ) @@ -162,6 +165,7 @@ def real_setup_gpu(mcdc_array, data_tally): mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) if config.gpu_state_storage == "separate": + print("LOADING!") harmonize.memcpy_device_to_host( simulation, simulation["gpu_meta"]["simulation_pointer"] ) @@ -176,3 +180,4 @@ def real_setup_gpu(mcdc_array, data_tally): particle_bank_module.set_bank_size(simulation["bank_active"], 0) source_closeout(simulation, 1, 1, data) + print("\nGen count after: ",simulation["gen_count"][0]) diff --git a/mcdc/code_factory/gpu/transport/util.py b/mcdc/code_factory/gpu/transport/util.py index 2acb013c8..321307832 100644 --- a/mcdc/code_factory/gpu/transport/util.py +++ b/mcdc/code_factory/gpu/transport/util.py @@ -5,10 +5,25 @@ from numba import njit, types -@njit -def atomic_add(array, idx, value): - harmonize.array_atomic_add(array, idx, value) - +def atomic_add(array,idx,value): + result = array[idx] + array[idx] += value + return result + + +@nb.extending.overload(atomic_add,target="gpu") +def overload_atomic_add_gpu(array,idx,value): + def impl(array, idx, value): + return harmonize.array_atomic_add(array, idx, value) + return impl + +@nb.extending.overload(atomic_add,target="cpu") +def overload_atomic_add_cpu(array,idx,value): + def impl(array, idx, value): + result = array[idx] + array[idx] += value + return result + return impl # ============================================================================= # Generic GPU/CPU local array variable constructors diff --git a/mcdc/main.py b/mcdc/main.py index 7874c2eec..0d2e915cb 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -65,8 +65,10 @@ def run_simulation(simulationPy: Simulation): import mcdc.transport.simulation as simulation_module if settings.neutron_eigenvalue_mode: + print("Eigenvalue!!!!") simulation_module.eigenvalue_simulation(simulation_container, data) else: + print("Fixed source!!!!") simulation_module.fixed_source_simulation(simulation_container, data) # TIMER: simulation diff --git a/mcdc/numba_types.py b/mcdc/numba_types.py index 28559d832..4b4ed807a 100644 --- a/mcdc/numba_types.py +++ b/mcdc/numba_types.py @@ -22,15 +22,17 @@ ('w', float64), ('particle_type', int64), ('rng_seed', uint64), + ('step_count', uint64), ]) particle = into_dtype([ ('cell_ID', int64), ('material_ID', int64), ('surface_ID', int64), - ('alive', bool), - ('fresh', bool), + ('alive', bool_), + ('fresh', bool_), ('event', int64), + ('step_count', uint64), ('x', float64), ('y', float64), ('z', float64), @@ -384,7 +386,7 @@ ('atomic_number', int64), ('mass_number', int64), ('atomic_weight_ratio', float64), - ('fissionable', bool), + ('fissionable', bool_), ('excitation_level', int64), ('neutron_xs_energy_grid_offset', int64), ('neutron_xs_energy_grid_length', int64), @@ -518,19 +520,19 @@ ('N_active', int64), ('N_cycle', int64), ('k_init', float64), - ('use_gyration_radius', bool), + ('use_gyration_radius', bool_), ('gyration_radius_type', int64), - ('use_source_file', bool), + ('use_source_file', bool_), ('source_file_name', 'U32'), ('time_boundary', float64), ('output_name', 'U32'), - ('use_progress_bar', bool), + ('use_progress_bar', bool_), ('N_census', int64), ('census_time_offset', int64), ('census_time_length', int64), - ('use_census_based_tally', bool), + ('use_census_based_tally', bool_), ('census_tally_frequency', int64), - ('save_particle', bool), + ('save_particle', bool_), ('active_bank_buffer', int64), ('census_bank_buffer_ratio', float64), ('source_bank_buffer_ratio', float64), @@ -549,7 +551,7 @@ ]) implicit_capture = into_dtype([ - ('active', bool), + ('active', bool_), ]) weighted_emission = into_dtype([ @@ -564,7 +566,7 @@ ]) weight_windows = into_dtype([ - ('active', bool), + ('active', bool_), ('energy_bounds_offset', int64), ('energy_bounds_length', int64), ('Ne', int64), @@ -622,7 +624,7 @@ ('time_range', float64, (2,)), ('particle_type', int64), ('probability', float64), - ('moving', bool), + ('moving', bool_), ('N_move', int64), ('N_move_grid', int64), ('move_velocities_offset', int64), @@ -652,13 +654,13 @@ ('J', float64), ('R', float64), ('r', float64), - ('linear', bool), - ('quadric', bool), - ('quartic', bool), + ('linear', bool_), + ('quadric', bool_), + ('quartic', bool_), ('nx', float64), ('ny', float64), ('nz', float64), - ('moving', bool), + ('moving', bool_), ('N_move', int64), ('N_move_grid', int64), ('move_velocities_offset', int64), @@ -852,7 +854,7 @@ def make_simulation_type(N: dict): ('cycle_active', bool), ('mpi_size', int64), ('mpi_rank', int64), - ('mpi_master', bool), + ('mpi_master', bool_), ('mpi_work_start', int64), ('mpi_work_size', int64), ('mpi_work_size_total', int64), @@ -863,5 +865,6 @@ def make_simulation_type(N: dict): ('runtime_output', float64), ('runtime_bank_management', float64), ('source_seed', int64), + ('gen_count', int64, (1,)), ]) diff --git a/mcdc/object_/particle.py b/mcdc/object_/particle.py index c28b925f0..50d65fe36 100644 --- a/mcdc/object_/particle.py +++ b/mcdc/object_/particle.py @@ -29,6 +29,7 @@ class ParticleData(MCDCBase): w: float = 0.0 particle_type: int = PARTICLE_NEUTRON rng_seed: uint64 = uint64(1) + step_count: uint64 = uint64(0) @dataclass @@ -54,6 +55,7 @@ class Particle(ParticleData): alive: bool = False fresh: bool = False event: int = -1 + step_count: uint64 = uint64(0) class ParticleBank(MCDCBase): diff --git a/mcdc/object_/simulation.py b/mcdc/object_/simulation.py index 5de5de093..8a46d278d 100644 --- a/mcdc/object_/simulation.py +++ b/mcdc/object_/simulation.py @@ -192,6 +192,8 @@ class Simulation(MCDCBase): # GPU metadata gpu_meta: GPUMeta source_seed: int + gen_count: Annotated[NDArray[int64], (1,)] + def __init__(self, name: str = "") -> None: self.compiled = False @@ -276,6 +278,7 @@ def __init__(self, name: str = "") -> None: # GPU metadata self.gpu_meta = GPUMeta() self.source_seed = 0 + self.gen_count = np.zeros(1, dtype=int64) def _reset_model(self) -> None: # Physics diff --git a/mcdc/transport/particle_bank.py b/mcdc/transport/particle_bank.py index 9736c2aa2..29630f250 100644 --- a/mcdc/transport/particle_bank.py +++ b/mcdc/transport/particle_bank.py @@ -26,6 +26,7 @@ @njit def get_bank_size(bank): return bank["size"][0] + #!return add_bank_size(bank,0) @njit @@ -35,7 +36,7 @@ def set_bank_size(bank, value): @njit def add_bank_size(bank, value): - util.atomic_add(bank["size"], 0, value) + return util.atomic_add(bank["size"], 0, value) # ============================================================================= @@ -50,7 +51,9 @@ def _bank_particle(particle_container, bank): report_full_bank(bank) # Set particle data - idx = get_bank_size(bank) + #idx = get_bank_size(bank) + idx = add_bank_size(bank,1) + particle_module.copy(bank["particle_data"][idx : idx + 1], particle_container) @@ -61,7 +64,7 @@ def bank_active_particle(particle_container, program): _bank_particle(particle_container, bank) # Increment bank size - add_bank_size(bank, 1) + #add_bank_size(bank, 1) @njit @@ -71,7 +74,7 @@ def bank_census_particle(particle_container, program): _bank_particle(particle_container, bank) # Increment bank size - add_bank_size(bank, 1) + #add_bank_size(bank, 1) @njit @@ -81,7 +84,7 @@ def bank_future_particle(particle_container, program): _bank_particle(particle_container, bank) # Increment bank size - add_bank_size(bank, 1) + #add_bank_size(bank, 1) @njit @@ -92,7 +95,7 @@ def bank_source_particle(particle_container, simulation): # Increment bank size # Note that we don't use the atomic operation in add_bank_size function # as source particle banking is not thread-parallelized - bank["size"][0] += 1 + #bank["size"][0] += 1 @njit @@ -101,12 +104,10 @@ def pop_particle(particle_container, bank): if get_bank_size(bank) == 0: report_empty_bank(bank) - # Set particle data - idx = get_bank_size(bank) - 1 + # Decrement bank size + idx = add_bank_size(bank, -1) - 1 particle_module.copy(particle_container, bank["particle_data"][idx : idx + 1]) - # Decrement bank size - add_bank_size(bank, -1) # Set default IDs and event for the live particle particle = particle_container[0] @@ -163,10 +164,11 @@ def promote_future_particles(program, data): if particle["t"] < next_census_time: bank_census_particle(particle_container, program) - add_bank_size(future_bank, -1) + #add_bank_size(future_bank, -1) + j = add_bank_size(future_bank,-1) # Consolidate the emptied space in the future bank - j = get_bank_size(future_bank) + #j = get_bank_size(future_bank) particle_module.copy( future_bank["particle_data"][idx : idx + 1], future_bank["particle_data"][j : j + 1], @@ -183,6 +185,11 @@ def manage_particle_banks(simulation): master = simulation["mpi_master"] serial = simulation["mpi_size"] == 1 + with objmode(): + print("Bank census has size: ",get_bank_size(simulation["bank_census"]),flush=True) + print("Bank source has size: ",get_bank_size(simulation["bank_source"]),flush=True) + print("Bank future has size: ",get_bank_size(simulation["bank_future"]),flush=True) + # TIMER: bank management time_start = 0.0 if master: diff --git a/mcdc/transport/simulation.py b/mcdc/transport/simulation.py index 80f3055d2..9323fb7b2 100644 --- a/mcdc/transport/simulation.py +++ b/mcdc/transport/simulation.py @@ -170,6 +170,8 @@ def source_loop(seed, simulation, data): work_start = simulation["mpi_work_start"] work_size = simulation["mpi_work_size"] + print("Gen count before: ",simulation["gen_count"][0]) + simulation["gen_count"][0] = 0 for idx_work in range(work_size): simulation["idx_work"] = work_start + idx_work generate_source_particle(work_start, idx_work, seed, simulation, data) @@ -178,6 +180,7 @@ def source_loop(seed, simulation, data): exhaust_active_bank(simulation, data) source_closeout(simulation, idx_work, N_prog, data) + print("\nGen count after: ",simulation["gen_count"][0]) @njit @@ -202,6 +205,8 @@ def generate_source_particle(work_start, idx_work, seed, program, data): ] particle = particle_container[0] + particle["step_count"] = 0 + # Skip if beyond time boundary if particle["t"] > settings["time_boundary"]: return @@ -281,17 +286,28 @@ def step_particle(particle_container, program, data): simulation = util.access_simulation(program) particle = particle_container[0] + particle["step_count"] += 1 + # Determine and move to event move_to_event(particle_container, simulation, data) + # Execute events if particle["event"] == EVENT_LOST: return + + # In first step of first phase: 10 CPU alive, 0 GPU alive + if (particle["alive"]) and (particle["step_count"] <= 1) : + util.atomic_add(simulation["gen_count"],0,1) # Collision if particle["event"] & EVENT_COLLISION: collision_data_container = util.local_array(1, type_.collision_data) + # In first step of first phase: 3 CPU alive, 0 GPU alive + #if (particle["alive"]) and (particle["step_count"] <= 1) : + # util.atomic_add(simulation["gen_count"],0,1) + # Execute the physics physics.collision(particle_container, collision_data_container, program, data) diff --git a/mcdc/transport/util.py b/mcdc/transport/util.py index 45ecdb1fa..768e73a9b 100644 --- a/mcdc/transport/util.py +++ b/mcdc/transport/util.py @@ -153,8 +153,9 @@ def log_interpolation(x, x1, x2, y1, y2): @njit def atomic_add(array, idx, value): + result = array[idx] array[idx] += value - + return result @njit def local_array(shape, dtype): From c360dd18f44e0709c0b9a55a1d18b5a330386d42 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 11 Aug 2026 14:17:06 -0700 Subject: [PATCH 11/34] Partial fix to data struction initialization ordering Multi-dimensional array returns still need to be added. --- mcdc/code_factory/array_return.py | 7 ++--- mcdc/code_factory/numba_layers_generator.py | 35 ++++++++------------- mcdc/numba_types.py | 12 +++---- 3 files changed, 22 insertions(+), 32 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 4cfe6c308..758155468 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -1,9 +1,8 @@ -from numba import njit, jit, objmode, literal_unroll, types -from numba.extending import intrinsic +import cffi import numba as nb import numpy as np - -import cffi +from numba import jit, literal_unroll, njit, objmode, types +from numba.extending import intrinsic ffi = cffi.FFI() diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index ce24e5b1b..af7b385f9 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -1,23 +1,21 @@ from __future__ import annotations #### - import importlib +from pathlib import Path + import numba as nb import numpy as np from numba import njit from numba.extending import intrinsic -from pathlib import Path #### - import mcdc import mcdc.code_factory.gpu.program_builder as gpu_builder import mcdc.config as config import mcdc.object_ as object_module import mcdc.object_.base as base - from mcdc.object_.base import ( MCDCBase, MCDCObject, @@ -50,7 +48,7 @@ np.bool_: 1, np.float64: 8, np.int64: 8, - np.uint64:8, + np.uint64: 8, np.str_: 32, } @@ -297,11 +295,6 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data) set_object(simulation, annotations, structures, records, data) - print("\n\n\nA\n\n\n",flush=True) - # Allocate the flattened data and re-set the objects - data["array"], data["pointer"] = create_data_array(data["size"], type_map[float],size_map[float]) - print("\n\n\nB\n\n\n",flush=True) - data["size"] = 0 records = {} for mcdc_class in mcdc_classes: @@ -311,10 +304,6 @@ def generate_numba_layers(simulation): records[mcdc_class.label] = {} records["simulation"] = records.pop("simulation") - for object_ in objects: - set_object(object_, annotations, structures, records, data, set_data=True) - set_object(simulation, annotations, structures, records, data, set_data=True) - # ================================================================================== # Finalize the simulation object structure and set record # ================================================================================== @@ -443,7 +432,7 @@ def generate_numba_layers(simulation): # Manually set particle bank attributes for name in bank_names: mcdc_simulation[name]["tag"] = getattr(simulation, name).tag - + mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] @@ -796,19 +785,19 @@ def create_data_array(size): data = np.zeros(size, dtype=np.float64) return data, 0 else: - return create_data_array_on_gpu(nb.types.float64,size,size*8) + return create_data_array_on_gpu(nb.types.float64, size, size * 8) @njit -def create_data_array_on_gpu(dtype,size,byte_size): +def create_data_array_on_gpu(dtype, size, byte_size): if config.gpu_state_storage == "managed": data_tally_ptr = gpu_builder.alloc_managed_bytes(byte_size) else: data_tally_ptr = gpu_builder.alloc_device_bytes(byte_size) data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) - + if config.gpu_state_storage == "separate": - data_tally = np.zeros( (size,),dtype=dtype) + data_tally = np.zeros((size,), dtype=dtype) else: data_tally = nb.carray(data_tally_ptr, (size,), dtype) return data_tally, data_tally_uint @@ -831,7 +820,7 @@ def create_simulation_container_on_gpu(dtype, size): mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) if config.gpu_state_storage == "separate": - mcdc_container = np.zeros((1,),dtype=dtype) + mcdc_container = np.zeros((1,), dtype=dtype) else: mcdc_container = nb.carray(mcdc_ptr, (1,), dtype) return mcdc_container, mcdc_uint @@ -907,8 +896,10 @@ def align(field_list): pad_id = 0 for field in field_list: if len(field) > 3: - print_error("Unexpected struct field specification. Specifications \ - usually only consist of 3 or fewer members") + print_error( + "Unexpected struct field specification. Specifications \ + usually only consist of 3 or fewer members" + ) multiplier = 1 if len(field) == 3: field = (field[0], field[1], fixup_dims(field[2])) diff --git a/mcdc/numba_types.py b/mcdc/numba_types.py index 4b4ed807a..b487f7f36 100644 --- a/mcdc/numba_types.py +++ b/mcdc/numba_types.py @@ -99,9 +99,9 @@ ]) collision_tally = into_dtype([ - ('cell_filtered', bool), + ('cell_filtered', bool_), ('cell_filter_ID', int64), - ('mesh_filtered', bool), + ('mesh_filtered', bool_), ('mesh_filter_type', int64), ('mesh_filter_ID', int64), ('mesh_stride_z', int64), @@ -112,9 +112,9 @@ ]) tracklength_tally = into_dtype([ - ('cell_filtered', bool), + ('cell_filtered', bool_), ('cell_filter_ID', int64), - ('mesh_filtered', bool), + ('mesh_filtered', bool_), ('mesh_filter_type', int64), ('mesh_filter_ID', int64), ('mesh_stride_z', int64), @@ -677,9 +677,9 @@ ]) surface_crossing_tally = into_dtype([ - ('surface_filtered', bool), + ('surface_filtered', bool_), ('surface_filter_ID', int64), - ('cell_filtered', bool), + ('cell_filtered', bool_), ('cell_filter_ID', int64), ('ID', int64), ('base_ID', int64), From 47a777ee7bd9be567d4aed36def63d4a3c71cefa Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Wed, 12 Aug 2026 16:35:34 -0700 Subject: [PATCH 12/34] Updated array return for multi-dim returns. Investigating load_state breakage --- mcdc/code_factory/array_return.py | 36 ++-- mcdc/code_factory/gpu/program_builder.py | 158 +++++++++++++++--- mcdc/code_factory/gpu/transport/simulation.py | 110 +----------- mcdc/code_factory/numba_layers_generator.py | 18 +- 4 files changed, 176 insertions(+), 146 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 758155468..9611726c0 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -132,7 +132,7 @@ def array_result_overload(array): ) def impl(array): - return (into_voidptr(array), len(array)) + return (into_voidptr(array), array.shape) return impl @@ -148,16 +148,16 @@ def context_guard(context): raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") -def array_return_typing(fn, elem_type): +def array_return_typing(fn, elem_type, ndim): from inspect import signature arg_list = ",".join([param for param in signature(fn).parameters]) - template = "def typer({arg_list}):\n return nb.types.Array(dtype=elem_type,ndim=1,layout='C')({arg_list})" + template = "def typer({arg_list}):\n return nb.types.Array(dtype=elem_type,ndim={ndim},layout='C')({arg_list})" gns = globals() | {"elem_type": elem_type} lns = {} - exec(template.format(arg_list=arg_list), gns, lns) + exec(template.format(arg_list=arg_list, ndim=ndim), gns, lns) typer = lns["typer"] def typer_factory(context): @@ -170,20 +170,22 @@ def typer_factory(context): nb.extending.type_callable(fn)(typer_factory) -def array_return_lowering(fn, elem_type): +def array_return_lowering(fn, elem_type, ndim): from inspect import signature param_count = len(signature(fn).parameters) - retty = nb.types.Array(dtype=elem_type, ndim=1, layout="C") + retty = nb.types.Array(dtype=elem_type, ndim=ndim, layout="C") sig = retty(*([nb.types.Any] * param_count)) jit_fn = nb.njit(fn) def builtin(context, builder, sig, args): - thing, data = args - thing_type, data_type = sig.args + # print(f"\n\n\nARGS ARE: {args}\n\n\n", flush=True) + # print(f"\n\n\nARG TYPES ARE: {sig.args}\n\n\n", flush=True) + # thing, data = args + # thing_type, data_type = sig.args import llvmlite.binding as ll from llvmlite import ir @@ -206,14 +208,16 @@ def builtin(context, builder, sig, args): CUDA_AVAILABLE = False lmod = builder.module - retty = nb.types.Tuple([nb.types.voidptr, nb.types.uintp]) + retty = nb.types.Tuple( + [nb.types.voidptr, nb.types.Tuple([nb.types.uintp] * ndim)] + ) ptr_sig = retty(*sig.args) res = context.compile_internal(builder, jit_fn.py_func, ptr_sig, args) ptr_res = builder.extract_value(res, 0) size_res = builder.extract_value(res, 1) - shape = [size_res] - dtype = data_type.dtype + shape = size_res + dtype = elem_type if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): targetdata = ll.create_target_data(nb.hip.amdgcn.DATA_LAYOUT) @@ -233,14 +237,14 @@ def builtin(context, builder, sig, args): kstrides = [context.get_constant(types.intp, itemsize)] - aryty = types.Array(dtype=elem_type, ndim=1, layout="C") + aryty = types.Array(dtype=elem_type, ndim=ndim, layout="C") ary = context.make_array(aryty)(context, builder) dataptr = builder.addrspacecast( ptr_res, ir.PointerType(ir.IntType(8)), "generic" ) - kshape = [size_res] + kshape = size_res context.populate_array( ary, data=builder.bitcast(dataptr, ary.data.type), @@ -254,10 +258,10 @@ def builtin(context, builder, sig, args): nb.extending.lower_builtin(fn, *sig.args)(builtin) -def array_return(sig): +def array_return(sig, ndim=1): def array_return_true_decorator(fn): - array_return_typing(fn, sig) - array_return_lowering(fn, sig) + array_return_typing(fn, sig, ndim) + array_return_lowering(fn, sig, ndim) return fn return array_return_true_decorator diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index 6e3d37868..e83831e13 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -1,11 +1,9 @@ import numba as nb import numba.extending as nbxt import numpy as np - from mpi4py import MPI #### - import mcdc.config as config # ====================================================================================== @@ -98,11 +96,19 @@ def _prepare_gpu_program(simulation_dtype, data_size): def forward_declare_gpu_program(simulation_dtype): import harmonize + import mcdc.numba_types as type_ # Get to set the globals global none_type, simulation_type, data_type - global state_spec, access_simulation, access_data_ptr, access_group, access_thread, particle_gpu, particle_record_gpu + global \ + state_spec, \ + access_simulation, \ + access_data_ptr, \ + access_group, \ + access_thread, \ + particle_gpu, \ + particle_record_gpu global step_async, find_cell_async global alloc_managed_bytes, alloc_device_bytes @@ -178,6 +184,13 @@ def find_cell(program: nb.uintp, particle: particle_gpu): store_state_device_data = None store_pointer_state_device_data = None +load_global = None +store_global = None +store_pointer_global = None +load_data = None +store_data = None +store_pointer_data = None + init_program = None exec_program = None complete = None @@ -188,20 +201,123 @@ def find_cell(program: nb.uintp, particle: particle_gpu): BLOCK_COUNT = 0 + + +def build_gpu_progs(input_deck): + + STRAT = config.args.gpu_strategy + + src_spec = gpu_sources_spec() + + adapt.harm.RuntimeSpec.bind_specs() + + rank = MPI.COMM_WORLD.Get_rank() + device_id = rank % config.args.gpu_share_stride + + if MPI.COMM_WORLD.Get_size() > 1: + MPI.COMM_WORLD.Barrier() + + adapt.harm.RuntimeSpec.load_specs() + + if STRAT == "async": + config.args.gpu_arena_size = config.args.gpu_arena_size // 32 + src_fns = src_spec.async_functions() + pre_fns = pre_spec.async_functions() + else: + src_fns = src_spec.event_functions() + pre_fns = pre_spec.event_functions() + + ARENA_SIZE = config.args.gpu_arena_size + BLOCK_COUNT = config.args.gpu_block_count + + global alloc_state, free_state + alloc_state = src_fns["alloc_state"] + free_state = src_fns["free_state"] + + global src_alloc_program, src_free_program + global src_load_global, src_store_global, src_load_data, src_store_data, src_store_pointer_data + global src_init_program, src_exec_program, src_complete, src_clear_flags + src_alloc_program = src_fns["alloc_program"] + src_free_program = src_fns["free_program"] + src_load_global = src_fns["load_state_device_global"] + src_store_global = src_fns["store_state_device_global"] + src_store_pointer_global = src_fns["store_pointer_state_device_global"] + src_load_data = src_fns["load_state_device_data"] + src_store_data = src_fns["store_state_device_data"] + src_store_pointer_data = src_fns["store_pointer_state_device_data"] + src_init_program = src_fns["init_program"] + src_exec_program = src_fns["exec_program"] + src_complete = src_fns["complete"] + src_clear_flags = src_fns["clear_flags"] + src_set_device = src_fns["set_device"] + + global pre_alloc_program, pre_free_program + global pre_load_global, pre_store_global, pre_load_data, pre_store_data + global pre_init_program, pre_exec_program, pre_complete, pre_clear_flags + pre_alloc_state = pre_fns["alloc_state"] + pre_free_state = pre_fns["free_state"] + pre_alloc_program = pre_fns["alloc_program"] + pre_free_program = pre_fns["free_program"] + pre_load_global = pre_fns["load_state_device_global"] + pre_store_global = pre_fns["store_state_device_global"] + pre_load_data = pre_fns["load_state_device_data"] + pre_store_data = pre_fns["store_state_device_data"] + pre_init_program = pre_fns["init_program"] + pre_exec_program = pre_fns["exec_program"] + pre_complete = pre_fns["complete"] + pre_clear_flags = pre_fns["clear_flags"] + + @njit + def real_setup_gpu(mcdc_array, data_tally): + mcdc = mcdc_array[0] + + print("STATE POINTER {mcdc['gpu_meta']['state_pointer']}") + print("GLOBAL POINTER {mcdc['gpu_meta']['global_pointer']}") + print("TALLY POINTER {mcdc['gpu_meta']['tally_pointer']}") + src_set_device(device_id) + arena_size = ARENA_SIZE + mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) + # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) + if config.gpu_state_storage == "separate": + print("LOADING!") + harmonize.memcpy_device_to_host( + simulation, simulation["gpu_meta"]["simulation_pointer"] + ) + harmonize.memcpy_device_to_host( + data, simulation["gpu_meta"]["data_pointer"] + ) + + gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) + + simulation["mpi_work_size"] = full_work_size + + particle_bank_module.set_bank_size(simulation["bank_active"], 0) + + source_closeout(simulation, 1, 1, data) + print("\nGen count after: ",simulation["gen_count"][0]) + + + def build_gpu_program(data_size): import harmonize + import mcdc.numba_types as type_ import mcdc.transport.util as util - from mcdc.transport.simulation import generate_source_particle, step_particle global alloc_state, free_state global alloc_program, free_program - global load_state_device_simulation, store_state_device_simulation, store_pointer_state_device_simulation + global \ + load_state_device_simulation, \ + store_state_device_simulation, \ + store_pointer_state_device_simulation - global load_state_device_data, store_state_device_data, store_pointer_state_device_data + global \ + load_state_device_data, \ + store_state_device_data, \ + store_pointer_state_device_data global init_program, exec_program, complete, clear_flags, set_device global ARENA_SIZE, BLOCK_COUNT @@ -300,20 +416,14 @@ def step(program: nb.uintp, particle_input: particle_gpu): complete = src_fns["complete"] clear_flags = src_fns["clear_flags"] set_device = src_fns["set_device"] - - src_alloc_program = src_fns["alloc_program"] - src_free_program = src_fns["free_program"] - src_load_global = src_fns["load_state_device_global"] - src_store_global = src_fns["store_state_device_global"] - src_store_pointer_global = src_fns["store_pointer_state_device_global"] - src_load_data = src_fns["load_state_device_data"] - src_store_data = src_fns["store_state_device_data"] - src_store_pointer_data = src_fns["store_pointer_state_device_data"] - src_init_program = src_fns["init_program"] - src_exec_program = src_fns["exec_program"] - src_complete = src_fns["complete"] - src_clear_flags = src_fns["clear_flags"] - src_set_device = src_fns["set_device"] + + alloc_program = src_fns["alloc_program"] + free_program = src_fns["free_program"] + init_program = src_fns["init_program"] + exec_program = src_fns["exec_program"] + complete = src_fns["complete"] + clear_flags = src_fns["clear_flags"] + set_device = src_fns["set_device"] # ================================================================================== # @@ -380,13 +490,13 @@ def teardown_gpu_program(simulation): # ====================================================================================== -#def create_data_array(size, dtype): +# def create_data_array(size, dtype): # if config.gpu_state_storage == "managed": # data_tally_ptr = harmonize.alloc_managed_bytes(size) # else: # data_tally_ptr = harmonize.alloc_device_bytes(size) # data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) -# +# # if config.gpu_state_storage == "separate": # data_tally = nb.zeros( (size,),dtype=dtype) # else: @@ -394,12 +504,12 @@ def teardown_gpu_program(simulation): # return data_tally, data_tally_uint -#def create_mcdc_container(dtype): +# def create_mcdc_container(dtype): # if config.gpu_state_storage == "managed": # mcdc_ptr = harmonize.alloc_managed_bytes(dtype.itemsize) # else: # mcdc_ptr = harmonize.alloc_device_bytes(dtype.itemsize) -# +# # mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) # if config.gpu_state_storage == "separate": # mcdc_tally = nb.zeros((size,),dtype=dtype) diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index fedb03820..d9b853d43 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -71,113 +71,19 @@ def source_loop(seed, simulation, data): gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) # Recover the original program state - src_load_constant(mcdc, mcdc["gpu_state_pointer"]) - src_load_data(data, mcdc["gpu_state_pointer"]) - src_clear_flags(mcdc["source_program_pointer"]) + gpu_module.load_state_device_simulation(simulation, simulation["gpu_meta"]["state_pointer"]) + gpu_module.load_state_device_data(data, simulation["gpu_state_pointer"]) + gpu_module.clear_flags(simulation["source_program_pointer"]) - mcdc["mpi_work_size"] = full_work_size + simulation["mpi_work_size"] = full_work_size - particle_bank_module.set_bank_size(mcdc["bank_active"], 0) + particle_bank_module.set_bank_size(simulation["bank_active"], 0) # ===================================================================== # Closeout (Moved out of the typical particle loop) # ===================================================================== - source_closeout(mcdc, 1, 1, data) - - if mcdc["technique"]["domain_decomposition"]: - source_dd_resolution(data, mcdc) - - -def build_gpu_progs(input_deck, args): - - STRAT = args.gpu_strategy - - src_spec = gpu_sources_spec() - - adapt.harm.RuntimeSpec.bind_specs() - - rank = MPI.COMM_WORLD.Get_rank() - device_id = rank % args.gpu_share_stride - - if MPI.COMM_WORLD.Get_size() > 1: - MPI.COMM_WORLD.Barrier() - - adapt.harm.RuntimeSpec.load_specs() - - if STRAT == "async": - args.gpu_arena_size = args.gpu_arena_size // 32 - src_fns = src_spec.async_functions() - pre_fns = pre_spec.async_functions() - else: - src_fns = src_spec.event_functions() - pre_fns = pre_spec.event_functions() - - ARENA_SIZE = args.gpu_arena_size - BLOCK_COUNT = args.gpu_block_count - - global alloc_state, free_state - alloc_state = src_fns["alloc_state"] - free_state = src_fns["free_state"] - - global src_alloc_program, src_free_program - global src_load_global, src_store_global, src_load_data, src_store_data, src_store_pointer_data - global src_init_program, src_exec_program, src_complete, src_clear_flags - src_alloc_program = src_fns["alloc_program"] - src_free_program = src_fns["free_program"] - src_load_global = src_fns["load_state_device_global"] - src_store_global = src_fns["store_state_device_global"] - src_store_pointer_global = src_fns["store_pointer_state_device_global"] - src_load_data = src_fns["load_state_device_data"] - src_store_data = src_fns["store_state_device_data"] - src_store_pointer_data = src_fns["store_pointer_state_device_data"] - src_init_program = src_fns["init_program"] - src_exec_program = src_fns["exec_program"] - src_complete = src_fns["complete"] - src_clear_flags = src_fns["clear_flags"] - src_set_device = src_fns["set_device"] - - global pre_alloc_program, pre_free_program - global pre_load_global, pre_store_global, pre_load_data, pre_store_data - global pre_init_program, pre_exec_program, pre_complete, pre_clear_flags - pre_alloc_state = pre_fns["alloc_state"] - pre_free_state = pre_fns["free_state"] - pre_alloc_program = pre_fns["alloc_program"] - pre_free_program = pre_fns["free_program"] - pre_load_global = pre_fns["load_state_device_global"] - pre_store_global = pre_fns["store_state_device_global"] - pre_load_data = pre_fns["load_state_device_data"] - pre_store_data = pre_fns["store_state_device_data"] - pre_init_program = pre_fns["init_program"] - pre_exec_program = pre_fns["exec_program"] - pre_complete = pre_fns["complete"] - pre_clear_flags = pre_fns["clear_flags"] - - @njit - def real_setup_gpu(mcdc_array, data_tally): - mcdc = mcdc_array[0] - - print("STATE POINTER {mcdc['gpu_meta']['state_pointer']}") - print("GLOBAL POINTER {mcdc['gpu_meta']['global_pointer']}") - print("TALLY POINTER {mcdc['gpu_meta']['tally_pointer']}") - src_set_device(device_id) - arena_size = ARENA_SIZE - mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) - # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) - if config.gpu_state_storage == "separate": - print("LOADING!") - harmonize.memcpy_device_to_host( - simulation, simulation["gpu_meta"]["simulation_pointer"] - ) - harmonize.memcpy_device_to_host( - data, simulation["gpu_meta"]["data_pointer"] - ) - - gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) - - simulation["mpi_work_size"] = full_work_size - - particle_bank_module.set_bank_size(simulation["bank_active"], 0) - source_closeout(simulation, 1, 1, data) - print("\nGen count after: ",simulation["gen_count"][0]) + + #if simulation["technique"]["domain_decomposition"]: + # source_dd_resolution(data, simulation) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index af7b385f9..e8d48ebf2 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -295,7 +295,7 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data) set_object(simulation, annotations, structures, records, data) - data["size"] = 0 + # data["size"] = 0 records = {} for mcdc_class in mcdc_classes: if issubclass(mcdc_class, ObjectNonSingleton): @@ -378,6 +378,7 @@ def generate_numba_layers(simulation): # Allocate the flattened data and re-set the objects # ================================================================================== + print("\n\nDATA SIZE SHOULD BE: ", data["size"]) data["array"], data["pointer"] = create_data_array(data["size"]) data["size"] = 0 @@ -394,7 +395,7 @@ def generate_numba_layers(simulation): simulation_dtype ) mcdc_simulation = mcdc_simulation_container[0] - mcdc_simulation["gpu_meta"]["global_pointer"] = mcdc_simulation_pointer + mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] record = records["simulation"] @@ -701,6 +702,13 @@ def set_object( record[f"{attribute_name}_offset"] = data["size"] record[f"{attribute_name}_length"] = len(attribute_flatten) if set_data: + print( + "\n\n\nSHAPE IS : ", + f"(size: {data['size']} : {data['size']} + {len(attribute_flatten)} -- {len(data['array'])} )", + data["array"][ + data["size"] : data["size"] + len(attribute_flatten) + ].shape, + ) data["array"][data["size"] : data["size"] + len(attribute_flatten)] = ( attribute_flatten[:] ) @@ -790,6 +798,7 @@ def create_data_array(size): @njit def create_data_array_on_gpu(dtype, size, byte_size): + print(f"STORAGE TYPE IS {config.gpu_state_storage}") if config.gpu_state_storage == "managed": data_tally_ptr = gpu_builder.alloc_managed_bytes(byte_size) else: @@ -1287,10 +1296,11 @@ def _accessor_2d_element( def _accessor_2d_vector(object_name, attribute_name, stride, setter=False): - text = f"@njit\n" if setter: + text = f"@njit\n" text += f"def {attribute_name}_vector(index_1, {object_name}, data, value):\n" else: + text = f"@array_return(nb.types.float64,1)\n" text += f"def {attribute_name}_vector(index_1, {object_name}, data):\n" text += f' offset = {object_name}["{attribute_name}_offset"]\n' text += accessor_dimension("stride", stride, object_name) @@ -1299,7 +1309,7 @@ def _accessor_2d_vector(object_name, attribute_name, stride, setter=False): if setter: text += f" data[start:end] = value\n\n\n" else: - text += f" return data[start:end]\n\n\n" + text += f" return array_result(data[start:end])\n\n\n" return text From 0ae6f9c3bb957898bee1caac72a58760046afd07 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 14 Aug 2026 01:31:24 -0700 Subject: [PATCH 13/34] Reached close CPU/GPU parity on Kobayashi-TD --- mcdc/code_factory/gpu/transport/simulation.py | 31 ++++++++++---- mcdc/code_factory/numba_layers_generator.py | 41 ++++++++++--------- mcdc/main.py | 2 + mcdc/object_/tally.py | 2 + mcdc/output.py | 6 +++ 5 files changed, 54 insertions(+), 28 deletions(-) diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index d9b853d43..ee85bc495 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -42,16 +42,19 @@ def source_loop(seed, simulation, data): # Store the global state to the GPU if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: print("STORING!") - harmonize.memcpy_host_to_device( - simulation["gpu_meta"]["simulation_pointer"], simulation - ) - harmonize.memcpy_host_to_device( - simulation["gpu_meta"]["data_pointer"], data - ) + gpu_module.store_state_device_simulation(simulation["gpu_meta"]["state_pointer"],simulation) + gpu_module.store_state_device_data(simulation["gpu_meta"]["state_pointer"],data) + #harmonize.memcpy_host_to_device( + # simulation["gpu_meta"]["simulation_pointer"], simulation + #) + #harmonize.memcpy_host_to_device( + # simulation["gpu_meta"]["data_pointer"], data + #) # Execute the program, and continue to do so until it is done block_count = gpu_module.BLOCK_COUNT + print("Beginning loop!") if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: gpu_module.exec_program( simulation["gpu_meta"]["program_pointer"], block_count, iter_count @@ -68,12 +71,22 @@ def source_loop(seed, simulation, data): gpu_module.exec_program( simulation["gpu_meta"]["program_pointer"], block_count, batch_size ) + print("Loop done!") gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) + print("Cleared flags!") # Recover the original program state - gpu_module.load_state_device_simulation(simulation, simulation["gpu_meta"]["state_pointer"]) - gpu_module.load_state_device_data(data, simulation["gpu_state_pointer"]) - gpu_module.clear_flags(simulation["source_program_pointer"]) + if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: + print("STORING!") + gpu_module.load_state_device_simulation(simulation, simulation["gpu_meta"]["state_pointer"]) + gpu_module.load_state_device_data(data, simulation["gpu_meta"]["state_pointer"]) + #harmonize.memcpy_device_to_host( + # simulation,simulation["gpu_meta"]["simulation_pointer"] + #) + #harmonize.memcpy_device_to_host( + # data,simulation["gpu_meta"]["data_pointer"] + #) + gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) simulation["mpi_work_size"] = full_work_size diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index e8d48ebf2..9dae71bbb 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -295,14 +295,14 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data) set_object(simulation, annotations, structures, records, data) - # data["size"] = 0 - records = {} - for mcdc_class in mcdc_classes: - if issubclass(mcdc_class, ObjectNonSingleton): - records[mcdc_class.label] = [] - else: - records[mcdc_class.label] = {} - records["simulation"] = records.pop("simulation") + ## data["size"] = 0 + #records = {} + #for mcdc_class in mcdc_classes: + # if issubclass(mcdc_class, ObjectNonSingleton): + # records[mcdc_class.label] = [] + # else: + # records[mcdc_class.label] = {} + #records["simulation"] = records.pop("simulation") # ================================================================================== # Finalize the simulation object structure and set record @@ -386,6 +386,7 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data, set_data=True) set_object(simulation, annotations, structures, records, data, set_data=True) + # ================================================================================== # Set with records # ================================================================================== @@ -436,6 +437,8 @@ def generate_numba_layers(simulation): mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] + + print(f"\n\n{mcdc_simulation}\n\n") # GPU program setup if config.target == "gpu": @@ -702,13 +705,13 @@ def set_object( record[f"{attribute_name}_offset"] = data["size"] record[f"{attribute_name}_length"] = len(attribute_flatten) if set_data: - print( - "\n\n\nSHAPE IS : ", - f"(size: {data['size']} : {data['size']} + {len(attribute_flatten)} -- {len(data['array'])} )", - data["array"][ - data["size"] : data["size"] + len(attribute_flatten) - ].shape, - ) + #print( + # "\n\n\nSHAPE IS : ", + # f"(size: {data['size']} : {data['size']} + {len(attribute_flatten)} -- {len(data['array'])} )", + # data["array"][ + # data["size"] : data["size"] + len(attribute_flatten) + # ].shape, + #) data["array"][data["size"] : data["size"] + len(attribute_flatten)] = ( attribute_flatten[:] ) @@ -793,7 +796,7 @@ def create_data_array(size): data = np.zeros(size, dtype=np.float64) return data, 0 else: - return create_data_array_on_gpu(nb.types.float64, size, size * 8) + return create_data_array_on_gpu(nb.types.float64, size, size * 16) @njit @@ -823,9 +826,9 @@ def create_simulation_container(dtype): @njit def create_simulation_container_on_gpu(dtype, size): if config.gpu_state_storage == "managed": - mcdc_ptr = gpu_builder.alloc_managed_bytes(size) + mcdc_ptr = gpu_builder.alloc_managed_bytes(size*8) else: - mcdc_ptr = gpu_builder.alloc_device_bytes(size) + mcdc_ptr = gpu_builder.alloc_device_bytes(size*8) mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) if config.gpu_state_storage == "separate": @@ -1300,7 +1303,7 @@ def _accessor_2d_vector(object_name, attribute_name, stride, setter=False): text = f"@njit\n" text += f"def {attribute_name}_vector(index_1, {object_name}, data, value):\n" else: - text = f"@array_return(nb.types.float64,1)\n" + text = f"@array_return(nb.types.float64)\n" text += f"def {attribute_name}_vector(index_1, {object_name}, data):\n" text += f' offset = {object_name}["{attribute_name}_offset"]\n' text += accessor_dimension("stride", stride, object_name) diff --git a/mcdc/main.py b/mcdc/main.py index 0d2e915cb..16c508ceb 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -64,6 +64,8 @@ def run_simulation(simulationPy: Simulation): # Run simulation import mcdc.transport.simulation as simulation_module + + print(f"\n\n{simulation_container}\n\n") if settings.neutron_eigenvalue_mode: print("Eigenvalue!!!!") simulation_module.eigenvalue_simulation(simulation_container, data) diff --git a/mcdc/object_/tally.py b/mcdc/object_/tally.py index 91717b7f3..82980d178 100644 --- a/mcdc/object_/tally.py +++ b/mcdc/object_/tally.py @@ -372,6 +372,8 @@ def _set_bin_shape_and_strides(self, shape: tuple): # Set bins self.bin_shape = list(shape) + print("SHAPE IS: ",shape) + # Set strides self.stride_time = reduce(operator.mul, shape[4:]) self.stride_energy = reduce(operator.mul, shape[3:]) diff --git a/mcdc/output.py b/mcdc/output.py index 71dde4a9c..7b8d3d9d5 100644 --- a/mcdc/output.py +++ b/mcdc/output.py @@ -202,6 +202,12 @@ def create_tally_dataset(file, mcdc, data): start_sdev = tally["bin_sum_square_offset"] mean = data[start_mean : start_mean + N_bin] sdev = data[start_sdev : start_sdev + N_bin] + print("TALLY : ",tally) + print("BIN SHAPE OFFSET : ",tally["bin_shape_offset"]) + print("BIN_SHAPE_LENGTH : ",tally["bin_shape_length"]) + print("DATA_SHAPE : ",data.shape) + print("DATA : ",data) + print("RESULT : ",mcdc_get.tally.bin_shape_all(tally, data)) shape = tuple([int(x) for x in mcdc_get.tally.bin_shape_all(tally, data)]) mean = mean.reshape(shape) sdev = sdev.reshape(shape) From c3d8077ba2f064f950a16ac45a01ffbd0f0cd3fa Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 20 Aug 2026 00:59:51 -0700 Subject: [PATCH 14/34] GPU/CPU Regression Test Parity --- mcdc/code_factory/gpu/program_builder.py | 14 ++++++++++---- mcdc/code_factory/gpu/transport/simulation.py | 1 - mcdc/transport/simulation.py | 3 +++ mcdc/transport/util.py | 9 ++++++++- test/regression/conftest.py | 13 ++++++++++++- 5 files changed, 33 insertions(+), 7 deletions(-) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index e83831e13..695ce117b 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -12,13 +12,10 @@ def adapt_transport_functions(): - global access_simulation import mcdc.code_factory.gpu.transport as gpu_transport import mcdc.transport as transport - transport.util.access_simulation = access_simulation - # TODO: Make the following automatic transport.geometry.interface.report_lost_particle = ( gpu_transport.geometry.interface.report_lost_particle @@ -165,6 +162,12 @@ def find_cell(program: nb.uintp, particle: particle_gpu): alloc_managed_bytes = harmonize.alloc_managed_bytes alloc_device_bytes = harmonize.alloc_device_bytes + from mcdc.transport import util + @nb.extending.overload(util.access_simulation,target="hip") + def access_simulation_gpu_overload(program): + def impl(program): + return access_simulation(program) + return impl # ====================================================================================== # Program builder @@ -305,7 +308,9 @@ def build_gpu_program(data_size): import mcdc.transport.util as util from mcdc.transport.simulation import generate_source_particle, step_particle - global alloc_state, free_state + global access_simulation, alloc_state, free_state + + print(f"Access simulation is {access_simulation}") global alloc_program, free_program @@ -381,6 +386,7 @@ def step(program: nb.uintp, particle_input: particle_gpu): base_fns = (initialize, finalize, make_work) async_fns = [step] src_spec = harmonize.RuntimeSpec("mcdc_source", state_spec, base_fns, async_fns) + print(f"ACCESS SIMULATION IS NOW SET TO {access_simulation}") harmonize.RuntimeSpec.bind_specs() # Load the specs diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index ee85bc495..758918e69 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -86,7 +86,6 @@ def source_loop(seed, simulation, data): #harmonize.memcpy_device_to_host( # data,simulation["gpu_meta"]["data_pointer"] #) - gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) simulation["mpi_work_size"] = full_work_size diff --git a/mcdc/transport/simulation.py b/mcdc/transport/simulation.py index 9323fb7b2..b636046a2 100644 --- a/mcdc/transport/simulation.py +++ b/mcdc/transport/simulation.py @@ -261,6 +261,9 @@ def source_closeout(simulation, idx_work, N_prog, data): tally_module.closeout.accumulate(simulation, data) # Progress printout + if simulation["mpi_work_size"] == 0: + return + percent = (idx_work + 1.0) / simulation["mpi_work_size"] if simulation["settings"]["use_progress_bar"] and int(percent * 100.0) > N_prog: N_prog += 1 diff --git a/mcdc/transport/util.py b/mcdc/transport/util.py index 768e73a9b..f9b3f4b5f 100644 --- a/mcdc/transport/util.py +++ b/mcdc/transport/util.py @@ -1,6 +1,7 @@ import math import numpy as np +import numba as nb from numba import njit from typing import Sequence @@ -162,6 +163,12 @@ def local_array(shape, dtype): return np.zeros(shape, dtype=dtype) -@njit def access_simulation(program): return program + +@nb.extending.overload(access_simulation,target="cpu") +def access_simulation_cpu_overload(program): + def impl(program): + return program + return impl + diff --git a/test/regression/conftest.py b/test/regression/conftest.py index 707d0f4ca..ec831c43c 100644 --- a/test/regression/conftest.py +++ b/test/regression/conftest.py @@ -136,11 +136,20 @@ def build_command(config): target = config.getoption("--target") mpiexec = config.getoption("--mpiexec") srun = config.getoption("--srun") + + state = "" + if target == "gpu": + state = "--gpu_state_storage=united" + mode = "numba" + command = [ sys.executable, "input.py", + f"--clear_cache", + f"--caching", f"--mode={mode}", f"--target={target}", + state, "--output=output", "--no-progress-bar", ] @@ -174,6 +183,7 @@ def compare_outputs(output_path, answer_path, target): def compare_tallies(output, answer, target, errors): + gpu_pass_list = ["uq_var","sdev"] name_root = "tallies" for tally in answer[name_root].keys(): name_tally = f"{name_root}/{tally}" @@ -182,7 +192,8 @@ def compare_tallies(output, answer, target, errors): continue name_score = f"{name_tally}/{score}" for result in answer[name_score].keys(): - if "uq_var" in result and target == "gpu": + should_pass = any([x in result for x in gpu_pass_list]) + if should_pass and target == "gpu": continue name = f"{name_score}/{result}" assert_allclose(output[name][()], answer[name][()], name, errors) From c78a438b145f9fca2721fa9d05eb70aca07d80c4 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 21 Aug 2026 12:51:41 -0700 Subject: [PATCH 15/34] Fixed issues with root solver --- mcdc/transport/geometry/root_solve.py | 72 ++++++++++++++-------- mcdc/transport/geometry/surface/torus.py | 8 +-- mcdc/transport/geometry/surface/torus_x.py | 12 ++-- mcdc/transport/geometry/surface/torus_y.py | 12 ++-- mcdc/transport/geometry/surface/torus_z.py | 12 ++-- 5 files changed, 65 insertions(+), 51 deletions(-) diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py index fa7aa0546..16aef6bcf 100644 --- a/mcdc/transport/geometry/root_solve.py +++ b/mcdc/transport/geometry/root_solve.py @@ -1,4 +1,5 @@ import math +import cmath import numpy as np import numba as nb @@ -14,8 +15,14 @@ def modulus(x): @njit() def sqrt(x): - return math.sqrt(x.real) * (x + x.real) / modulus(x + x.real) - + r = modulus(x) + real_part = math.sqrt((r + x.real) / 2.0) + if x.imag < 0: + imag_part = -math.sqrt((r - x.real) / 2.0) + else: + imag_part = math.sqrt((r - x.real) / 2.0) + + return complex(real_part, imag_part) @njit() def power(x, n): @@ -43,7 +50,9 @@ def nth_root(x, n, index): @njit() def principal_nth_root(x, n): - return nth_root(x, n, 0) + result = nth_root(x, n, 0) + return result + @njit() @@ -70,10 +79,10 @@ def solve_biquadratic(coeff, roots): # Yield roots for x by taking square roots of the x^2 # solution. - roots[4] = sqrt(roots[1]) - roots[3] = -roots[2] - roots[0] = sqrt(roots[0]) - roots[1] = -roots[0] + roots[3] = sqrt(roots[1]) + roots[2] = -roots[3] + roots[1] = sqrt(roots[0]) + roots[0] = -roots[1] # Restore the original positions of the coefficients coeff[4] = coeff[2] @@ -97,11 +106,11 @@ def solve_depressed_quartic(coeff, roots): # To solve the depressed quartic, one must first find one # root of a cubic polynomial. - p = (-power(a, 2) / 12) - c - q = (-power(a, 3) / 108) + (a * c / 3) - (power(b, 2) / 8) + p = (-power(a, 2) / 12.0) - c + q = (-power(a, 3) / 108.0) + (a * c / 3.0) - (power(b, 2) / 8.0) - cube_const = -q / 2 - sqrt_body = (power(q, 2) / 4) + (power(p, 3) / 27) + cube_const = -q / 2.0 + sqrt_body = (power(q, 2) / 4.0) + (power(p, 3) / 27.0) w_pos = principal_nth_root(cube_const + sqrt(sqrt_body), 3) w_neg = principal_nth_root(cube_const - sqrt(sqrt_body), 3) @@ -113,21 +122,21 @@ def solve_depressed_quartic(coeff, roots): w = w_neg # A root of the cubic - y = (a / 6) + w - (p / (3 * w)) + y = (a / 6.0) + w - (p / (3.0 * w)) # The different roots are found by flipping the signs # of some terms in a formula. There are three sections # unaffected by these flips, represented below by # alpha, beta, and gamma - alpha = sqrt(2 * y - a) - beta = -2 * y - a - gamma = (2 * b) / sqrt(2 * y - a) + alpha = sqrt(2.0 * y - a) + beta = -2.0 * y - a + gamma = (2.0 * b) / sqrt(2.0 * y - a) - roots[0] = (-alpha) + sqrt(beta + gamma) # - + + - roots[1] = (-alpha) - sqrt(beta + gamma) # - - + - roots[2] = (alpha) + sqrt(beta - gamma) # + + - - roots[3] = (alpha) - sqrt(beta - gamma) # + - - + roots[0] = ((-alpha) + sqrt(beta + gamma)) / 2.0 # - + + + roots[1] = ((-alpha) - sqrt(beta + gamma)) / 2.0 # - - + + roots[2] = ((alpha) + sqrt(beta - gamma)) / 2.0 # + + - + roots[3] = ((alpha) - sqrt(beta - gamma)) / 2.0 # + - - @njit() @@ -152,15 +161,15 @@ def solve_quartic(coeff, roots): # they can be plugged into this equation to yeild the # roots for x. - sub_coeff = util.local_array(4, np.complex128) + sub_coeff = util.local_array(5, np.complex128) sub_coeff[4] = 1.0 + 0.0j sub_coeff[3] = 0.0j - sub_coeff[2] = (-3 * power(b, 2)) / (8 * power(a, 2)) + c / a - sub_coeff[1] = power(b, 3) / (8 * power(a, 3)) - (b * c) / (2 * power(a, 2)) + d / a + sub_coeff[2] = (-3.0 * power(b, 2)) / (8.0 * power(a, 2)) + c / a + sub_coeff[1] = power(b, 3) / (8.0 * power(a, 3)) - (b * c) / (2.0 * power(a, 2)) + d / a sub_coeff[0] = ( - (-3 * power(b, 4)) / (256 * power(a, 4)) - + (c * power(b, 2)) / (16 * power(a, 3)) - - (b * d) / (4 * power(a, 2)) + (-3.0 * power(b, 4)) / (256.0 * power(a, 4)) + + (c * power(b, 2)) / (16.0 * power(a, 3)) + - (b * d) / (4.0 * power(a, 2)) + e / a ) @@ -177,4 +186,15 @@ def solve_quartic(coeff, roots): solve_depressed_quartic(sub_coeff, sub_roots) for idx in range(4): - roots[idx] = sub_roots[idx] - b / (4 * a) + roots[idx] = sub_roots[idx] - b / (4.0 * a) + + ans_roots = np.roots(np.flip(coeff)) + + + for idx in range(4): + y = 0 + ans_y = 0 + for i in range(5): + y += (roots[idx]**i) * coeff[i] + ans_y += (ans_roots[idx]**i) * coeff[i] + diff --git a/mcdc/transport/geometry/surface/torus.py b/mcdc/transport/geometry/surface/torus.py index 6c3e5c096..7a2e7fbf5 100644 --- a/mcdc/transport/geometry/surface/torus.py +++ b/mcdc/transport/geometry/surface/torus.py @@ -168,11 +168,11 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; coefficients = util.local_array(5, np.complex128) - coefficients[0] = a4 + 0.0j - coefficients[1] = a3 + 0.0j + coefficients[0] = a0 + 0.0j + coefficients[1] = a1 + 0.0j coefficients[2] = a2 + 0.0j - coefficients[3] = a1 + 0.0j - coefficients[4] = a0 + 0.0j + coefficients[3] = a3 + 0.0j + coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) root_solve.solve_quartic(coefficients, roots) diff --git a/mcdc/transport/geometry/surface/torus_x.py b/mcdc/transport/geometry/surface/torus_x.py index 178e11e3e..b342a0f7b 100644 --- a/mcdc/transport/geometry/surface/torus_x.py +++ b/mcdc/transport/geometry/surface/torus_x.py @@ -193,15 +193,13 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. coefficients = util.local_array(5, np.complex128) - coefficients[0] = a4 + 0.0j - coefficients[1] = a3 + 0.0j + coefficients[0] = a0 + 0.0j + coefficients[1] = a1 + 0.0j coefficients[2] = a2 + 0.0j - coefficients[3] = a1 + 0.0j - coefficients[4] = a0 + 0.0j + coefficients[3] = a3 + 0.0j + coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - # root_solve.solve_quartic(coefficients,roots) - for idx in range(5): - roots[idx] = 0 + root_solve.solve_quartic(coefficients,roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_y.py b/mcdc/transport/geometry/surface/torus_y.py index 75a9c18a2..a44f93878 100644 --- a/mcdc/transport/geometry/surface/torus_y.py +++ b/mcdc/transport/geometry/surface/torus_y.py @@ -193,15 +193,13 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. coefficients = util.local_array(5, np.complex128) - coefficients[0] = a4 + 0.0j - coefficients[1] = a3 + 0.0j + coefficients[0] = a0 + 0.0j + coefficients[1] = a1 + 0.0j coefficients[2] = a2 + 0.0j - coefficients[3] = a1 + 0.0j - coefficients[4] = a0 + 0.0j + coefficients[3] = a3 + 0.0j + coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - # root_solve.solve_quartic(coefficients,roots) - for idx in range(5): - roots[idx] = 0 + root_solve.solve_quartic(coefficients,roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_z.py b/mcdc/transport/geometry/surface/torus_z.py index 283592eeb..0766d88ec 100644 --- a/mcdc/transport/geometry/surface/torus_z.py +++ b/mcdc/transport/geometry/surface/torus_z.py @@ -193,15 +193,13 @@ def get_distance(particle_container, surface): # TODO: May replace with a fully numba-native quartic solver if torus performance becomes important; # np.roots is sufficient for now. coefficients = util.local_array(5, np.complex128) - coefficients[0] = a4 + 0.0j - coefficients[1] = a3 + 0.0j + coefficients[0] = a0 + 0.0j + coefficients[1] = a1 + 0.0j coefficients[2] = a2 + 0.0j - coefficients[3] = a1 + 0.0j - coefficients[4] = a0 + 0.0j + coefficients[3] = a3 + 0.0j + coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - # root_solve.solve_quartic(coefficients,roots) - for idx in range(5): - roots[idx] = 0 + root_solve.solve_quartic(coefficients,roots) min_t = INF From b746564b11b3c27dd6e94c384c24c3dd1409715f Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 21 Aug 2026 18:07:10 -0700 Subject: [PATCH 16/34] Rebased on mcdc-project/dev --- mcdc/main.py | 1 + mcdc/transport/geometry/root_solve.py | 10 ---------- test/unit/test_annotation_shape.py | 5 +++-- test/unit/test_numba_layers_generator.py | 5 +++-- 4 files changed, 7 insertions(+), 14 deletions(-) diff --git a/mcdc/main.py b/mcdc/main.py index 16c508ceb..daae1517a 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -1,4 +1,5 @@ from mcdc.object_.simulation import Simulation +import mcdc.config as config # ====================================================================================== # Run Simulation diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py index 16aef6bcf..eb1aba113 100644 --- a/mcdc/transport/geometry/root_solve.py +++ b/mcdc/transport/geometry/root_solve.py @@ -187,14 +187,4 @@ def solve_quartic(coeff, roots): for idx in range(4): roots[idx] = sub_roots[idx] - b / (4.0 * a) - - ans_roots = np.roots(np.flip(coeff)) - - - for idx in range(4): - y = 0 - ans_y = 0 - for i in range(5): - y += (roots[idx]**i) * coeff[i] - ans_y += (ans_roots[idx]**i) * coeff[i] diff --git a/test/unit/test_annotation_shape.py b/test/unit/test_annotation_shape.py index f2f058de5..1d408cdbe 100644 --- a/test/unit/test_annotation_shape.py +++ b/test/unit/test_annotation_shape.py @@ -2,6 +2,7 @@ from typing import Annotated +import numba as nb import numpy as np import pytest from numpy.typing import NDArray @@ -66,8 +67,8 @@ def test_stringified_annotation_rejects_incorrect_offset_shape(capsys): def test_generated_accessor_resolves_dimension_offset(): - all_source = _accessor_1d_all("mgxs", "energy", "G+1") - last_source = _accessor_1d_last("mgxs", "energy", "N - 2") + all_source = _accessor_1d_all("mgxs", "energy", "G+1",nb.types.float64) + last_source = _accessor_1d_last("mgxs", "energy", "N - 2",nb.types.float64) assert 'size = mgxs["G"] + 1' in all_source assert 'size = mgxs["N"] - 2' in last_source diff --git a/test/unit/test_numba_layers_generator.py b/test/unit/test_numba_layers_generator.py index 21531dced..687da7d1f 100644 --- a/test/unit/test_numba_layers_generator.py +++ b/test/unit/test_numba_layers_generator.py @@ -1,5 +1,6 @@ from typing import Annotated +import numba as nb import numpy as np import pytest from numpy.typing import NDArray @@ -129,8 +130,8 @@ def test_scalar_integer_getters_cast_values_from_data(): def test_float_and_bulk_getters_remain_zero_copy_views(): assert "return data[offset + index]" in _accessor_1d_element("example", "values") - assert "return data[start:end]" in _accessor_1d_all( - "example", "values", "values_length" + assert "return array_result(data[start:end])" in _accessor_1d_all( + "example", "values", "values_length",nb.types.float64 ) From 560e974025e189df8d909574a8ae878c44c1eb52 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 25 Aug 2026 12:37:44 -0700 Subject: [PATCH 17/34] Cleaned up changes --- mcdc/code_factory/array_return.py | 87 +++++++++++-------- mcdc/code_factory/gpu/program_builder.py | 52 +---------- mcdc/code_factory/gpu/transport/simulation.py | 21 ----- mcdc/code_factory/numba_layers_generator.py | 19 ---- mcdc/config.py | 21 +++++ mcdc/main.py | 3 - mcdc/numba_types.py | 54 ++++++------ mcdc/object_/particle.py | 2 - mcdc/object_/simulation.py | 2 - mcdc/object_/tally.py | 2 - mcdc/output.py | 6 -- mcdc/transport/particle_bank.py | 18 ---- mcdc/transport/rng.py | 2 - mcdc/transport/simulation.py | 16 ---- test/regression/conftest.py | 2 +- 15 files changed, 101 insertions(+), 206 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 9611726c0..6be3028da 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -3,12 +3,13 @@ import numpy as np from numba import jit, literal_unroll, njit, objmode, types from numba.extending import intrinsic +import mcdc.config as config ffi = cffi.FFI() # ============================================================================= -# uintp/voidptr casters +# uintp/voidptr casting helper functions - for internal use only # ============================================================================= @@ -67,12 +68,6 @@ def codegen(context, builder, signature, args): return sig, codegen -@njit() -def uintp_to_voidptr(value): - val = nb.uintp(value) - return cast_uintp_to_voidptr(val) - - @njit() def voidptr_to_uintp(value): return cast_voidptr_to_uintp(value) @@ -83,10 +78,26 @@ def into_voidptr(value): return into_voidptr_python(value) + +# ============================================================================= +# uintp/voidptr casting utility functions +# ============================================================================= + + +# Converts a pointer-sized integer to a void* +@njit() +def uintp_to_voidptr(value): + val = nb.uintp(value) + return cast_uintp_to_voidptr(val) + + +# Placeholder function for casting to void*. There currently is no use case +# for void* values in python mode for mcdc. def into_voidptr_python(value): raise RuntimeError("`into_voidptr` is only supported in nopython mode.") + @nb.extending.overload(into_voidptr_python) def into_voidptr_overload(value): @@ -115,10 +126,10 @@ def impl(value): ############################################################################### -# New Code +# Helper decorators, functions, and builtins for returning arrays ############################################################################### - +# Overload target def array_result(array): return array @@ -136,7 +147,7 @@ def impl(array): return impl - +# Raises an error if the context is not recognized def context_guard(context): if isinstance(context, nb.core.typing.context.Context): pass @@ -147,7 +158,7 @@ def context_guard(context): else: raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") - +# Typing for the `array_return` builtin. def array_return_typing(fn, elem_type, ndim): from inspect import signature @@ -170,42 +181,28 @@ def typer_factory(context): nb.extending.type_callable(fn)(typer_factory) +# The logic forthe `array_return` builtin def array_return_lowering(fn, elem_type, ndim): from inspect import signature + # The builtin returns an array with the given element + # type and the given dimensionality (default 1) param_count = len(signature(fn).parameters) retty = nb.types.Array(dtype=elem_type, ndim=ndim, layout="C") sig = retty(*([nb.types.Any] * param_count)) jit_fn = nb.njit(fn) + # This builtin effectively replaces the original decorated function whenever it + # is referenced in code. The original functions still exists, but it is called through + # this builtin which converts the pointer/shape tuple that the function (should) + # generate with `array_result` and return. def builtin(context, builder, sig, args): - # print(f"\n\n\nARGS ARE: {args}\n\n\n", flush=True) - # print(f"\n\n\nARG TYPES ARE: {sig.args}\n\n\n", flush=True) - # thing, data = args - # thing_type, data_type = sig.args - import llvmlite.binding as ll from llvmlite import ir - try: - import numba.hip as hip - - ROCM_AVAILABLE = True - except: - ROCM_AVAILABLE = False - - if not ROCM_AVAILABLE: - try: - import numba.cuda as cuda - - CUDA_AVAILABLE = True - except: - CUDA_AVAILABLE = False - else: - CUDA_AVAILABLE = False lmod = builder.module retty = nb.types.Tuple( @@ -219,31 +216,42 @@ def builtin(context, builder, sig, args): shape = size_res dtype = elem_type - if ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + # GPU platforms require a `targetdata` for array construction, which is created + # slightly differently depending upon the platform. + if config.ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): targetdata = ll.create_target_data(nb.hip.amdgcn.DATA_LAYOUT) - elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + elif config.CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) lldtype = context.get_data_type(dtype) + + # The size of the item is derived either from the lldtype or `targetdata` depending + # upon platform if isinstance(context, nb.core.cpu.CPUContext): itemsize = context.get_abi_sizeof(lldtype) - elif ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + elif config.ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): itemsize = lldtype.get_abi_size(targetdata) - elif CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + elif config.CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): itemsize = lldtype.get_abi_size(targetdata) else: raise nb.core.errors.UnsupportedError( f"Unsupported target context {context}." ) + # The number of elements-worth of bytes that must be skipped to advance by 1 element + # in a given dimension kstrides = [context.get_constant(types.intp, itemsize)] + # Create array structure based on the supplied type information aryty = types.Array(dtype=elem_type, ndim=ndim, layout="C") ary = context.make_array(aryty)(context, builder) + # Array populating logic expects pointers to the array buffer to be expressed as + # a pointer to a byte in a generic address space. dataptr = builder.addrspacecast( ptr_res, ir.PointerType(ir.IntType(8)), "generic" ) + # Initialize the array structure with the data pointer, shape, and strides kshape = size_res context.populate_array( ary, @@ -255,9 +263,13 @@ def builtin(context, builder, sig, args): ) return ary._getvalue() + # To complete the illusion of the decorated function acting just like a normal + # `njit` function, the decorated function is overloaded as the builtin that + # was defined above. nb.extending.lower_builtin(fn, *sig.args)(builtin) - +# A function decorated with `array_return` may return an array by passing +# it through the `array_result` function and returning the output def array_return(sig, ndim=1): def array_return_true_decorator(fn): array_return_typing(fn, sig, ndim) @@ -265,3 +277,4 @@ def array_return_true_decorator(fn): return fn return array_return_true_decorator + diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index 695ce117b..bc56af070 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -10,7 +10,7 @@ # Transport function adapter # ====================================================================================== - +# Overwrites global symbols in other modules with gpu-compatible counterparts def adapt_transport_functions(): import mcdc.code_factory.gpu.transport as gpu_transport @@ -187,13 +187,6 @@ def impl(program): store_state_device_data = None store_pointer_state_device_data = None -load_global = None -store_global = None -store_pointer_global = None -load_data = None -store_data = None -store_pointer_data = None - init_program = None exec_program = None complete = None @@ -205,7 +198,7 @@ def impl(program): - +# Compiles gpu kernels and loads in the functions that call into said kernels def build_gpu_progs(input_deck): STRAT = config.args.gpu_strategy @@ -274,15 +267,11 @@ def build_gpu_progs(input_deck): def real_setup_gpu(mcdc_array, data_tally): mcdc = mcdc_array[0] - print("STATE POINTER {mcdc['gpu_meta']['state_pointer']}") - print("GLOBAL POINTER {mcdc['gpu_meta']['global_pointer']}") - print("TALLY POINTER {mcdc['gpu_meta']['tally_pointer']}") src_set_device(device_id) arena_size = ARENA_SIZE mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) if config.gpu_state_storage == "separate": - print("LOADING!") harmonize.memcpy_device_to_host( simulation, simulation["gpu_meta"]["simulation_pointer"] ) @@ -297,7 +286,6 @@ def real_setup_gpu(mcdc_array, data_tally): particle_bank_module.set_bank_size(simulation["bank_active"], 0) source_closeout(simulation, 1, 1, data) - print("\nGen count after: ",simulation["gen_count"][0]) @@ -310,8 +298,6 @@ def build_gpu_program(data_size): global access_simulation, alloc_state, free_state - print(f"Access simulation is {access_simulation}") - global alloc_program, free_program global \ @@ -386,7 +372,6 @@ def step(program: nb.uintp, particle_input: particle_gpu): base_fns = (initialize, finalize, make_work) async_fns = [step] src_spec = harmonize.RuntimeSpec("mcdc_source", state_spec, base_fns, async_fns) - print(f"ACCESS SIMULATION IS NOW SET TO {access_simulation}") harmonize.RuntimeSpec.bind_specs() # Load the specs @@ -491,39 +476,6 @@ def teardown_gpu_program(simulation): free_state(cast_uintp_to_voidptr(simulation["gpu_meta"]["state_pointer"])) -# ====================================================================================== -# Simulation structure and data creators -# ====================================================================================== - - -# def create_data_array(size, dtype): -# if config.gpu_state_storage == "managed": -# data_tally_ptr = harmonize.alloc_managed_bytes(size) -# else: -# data_tally_ptr = harmonize.alloc_device_bytes(size) -# data_tally_uint = cast_voidptr_to_uintp(data_tally_ptr) -# -# if config.gpu_state_storage == "separate": -# data_tally = nb.zeros( (size,),dtype=dtype) -# else: -# data_tally = nb.carray(data_tally_ptr, (size,), dtype) -# return data_tally, data_tally_uint - - -# def create_mcdc_container(dtype): -# if config.gpu_state_storage == "managed": -# mcdc_ptr = harmonize.alloc_managed_bytes(dtype.itemsize) -# else: -# mcdc_ptr = harmonize.alloc_device_bytes(dtype.itemsize) -# -# mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) -# if config.gpu_state_storage == "separate": -# mcdc_tally = nb.zeros((size,),dtype=dtype) -# else: -# mcdc_tally = nb.carray(mcdc_ptr, (size,), dtype) -# mcdc_container = nb.carray(mcdc_ptr, (1,), dtype) -# return mcdc_container, mcdc_uint - # ====================================================================================== # Type casters diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 758918e69..e4485a9a0 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -25,7 +25,6 @@ def source_loop(seed, simulation, data): full_work_size = simulation["mpi_work_size"] - print("Gen count before: ",simulation["gen_count"][0]) simulation["gen_count"][0] = 0 if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: phase_size = 1000000000 @@ -41,20 +40,12 @@ def source_loop(seed, simulation, data): # Store the global state to the GPU if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: - print("STORING!") gpu_module.store_state_device_simulation(simulation["gpu_meta"]["state_pointer"],simulation) gpu_module.store_state_device_data(simulation["gpu_meta"]["state_pointer"],data) - #harmonize.memcpy_host_to_device( - # simulation["gpu_meta"]["simulation_pointer"], simulation - #) - #harmonize.memcpy_host_to_device( - # simulation["gpu_meta"]["data_pointer"], data - #) # Execute the program, and continue to do so until it is done block_count = gpu_module.BLOCK_COUNT - print("Beginning loop!") if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: gpu_module.exec_program( simulation["gpu_meta"]["program_pointer"], block_count, iter_count @@ -71,21 +62,12 @@ def source_loop(seed, simulation, data): gpu_module.exec_program( simulation["gpu_meta"]["program_pointer"], block_count, batch_size ) - print("Loop done!") gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) - print("Cleared flags!") # Recover the original program state if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: - print("STORING!") gpu_module.load_state_device_simulation(simulation, simulation["gpu_meta"]["state_pointer"]) gpu_module.load_state_device_data(data, simulation["gpu_meta"]["state_pointer"]) - #harmonize.memcpy_device_to_host( - # simulation,simulation["gpu_meta"]["simulation_pointer"] - #) - #harmonize.memcpy_device_to_host( - # data,simulation["gpu_meta"]["data_pointer"] - #) simulation["mpi_work_size"] = full_work_size @@ -96,6 +78,3 @@ def source_loop(seed, simulation, data): # ===================================================================== source_closeout(simulation, 1, 1, data) - - #if simulation["technique"]["domain_decomposition"]: - # source_dd_resolution(data, simulation) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 9dae71bbb..5d0530547 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -295,15 +295,6 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data) set_object(simulation, annotations, structures, records, data) - ## data["size"] = 0 - #records = {} - #for mcdc_class in mcdc_classes: - # if issubclass(mcdc_class, ObjectNonSingleton): - # records[mcdc_class.label] = [] - # else: - # records[mcdc_class.label] = {} - #records["simulation"] = records.pop("simulation") - # ================================================================================== # Finalize the simulation object structure and set record # ================================================================================== @@ -378,7 +369,6 @@ def generate_numba_layers(simulation): # Allocate the flattened data and re-set the objects # ================================================================================== - print("\n\nDATA SIZE SHOULD BE: ", data["size"]) data["array"], data["pointer"] = create_data_array(data["size"]) data["size"] = 0 @@ -438,7 +428,6 @@ def generate_numba_layers(simulation): mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] - print(f"\n\n{mcdc_simulation}\n\n") # GPU program setup if config.target == "gpu": @@ -705,13 +694,6 @@ def set_object( record[f"{attribute_name}_offset"] = data["size"] record[f"{attribute_name}_length"] = len(attribute_flatten) if set_data: - #print( - # "\n\n\nSHAPE IS : ", - # f"(size: {data['size']} : {data['size']} + {len(attribute_flatten)} -- {len(data['array'])} )", - # data["array"][ - # data["size"] : data["size"] + len(attribute_flatten) - # ].shape, - #) data["array"][data["size"] : data["size"] + len(attribute_flatten)] = ( attribute_flatten[:] ) @@ -801,7 +783,6 @@ def create_data_array(size): @njit def create_data_array_on_gpu(dtype, size, byte_size): - print(f"STORAGE TYPE IS {config.gpu_state_storage}") if config.gpu_state_storage == "managed": data_tally_ptr = gpu_builder.alloc_managed_bytes(byte_size) else: diff --git a/mcdc/config.py b/mcdc/config.py index d841ea567..db607395c 100644 --- a/mcdc/config.py +++ b/mcdc/config.py @@ -114,6 +114,27 @@ def _build_parser() -> argparse.ArgumentParser: clear_cache = args.clear_cache +# ====================================================================================== +# Flags for GPU platform availability +# ====================================================================================== + +try: + import numba.hip as hip + ROCM_AVAILABLE = True +except: + ROCM_AVAILABLE = False + +if not ROCM_AVAILABLE: + try: + import numba.cuda as cuda + CUDA_AVAILABLE = True + except: + CUDA_AVAILABLE = False +else: + CUDA_AVAILABLE = False + + + # ====================================================================================== # Simulation-setting overrides # ====================================================================================== diff --git a/mcdc/main.py b/mcdc/main.py index daae1517a..ca600ae01 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -66,12 +66,9 @@ def run_simulation(simulationPy: Simulation): import mcdc.transport.simulation as simulation_module - print(f"\n\n{simulation_container}\n\n") if settings.neutron_eigenvalue_mode: - print("Eigenvalue!!!!") simulation_module.eigenvalue_simulation(simulation_container, data) else: - print("Fixed source!!!!") simulation_module.fixed_source_simulation(simulation_container, data) # TIMER: simulation diff --git a/mcdc/numba_types.py b/mcdc/numba_types.py index b487f7f36..b14904c85 100644 --- a/mcdc/numba_types.py +++ b/mcdc/numba_types.py @@ -54,8 +54,8 @@ ('surface_IDs_offset', int64), ('fill_type', int64), ('fill_ID', int64), - ('fill_translated', bool), - ('fill_rotated', bool), + ('fill_translated', bool_), + ('fill_rotated', bool_), ('translation', float64, (3,)), ('rotation', float64, (3,)), ('N_collision_tally', int64), @@ -84,8 +84,8 @@ material = into_dtype([ ('name', 'U32'), ('temperature', float64), - ('fissionable', bool), - ('has_neutron_multigroup', bool), + ('fissionable', bool_), + ('has_neutron_multigroup', bool_), ('neutron_multigroup_ID', int64), ('N_nuclide', int64), ('nuclide_IDs_offset', int64), @@ -376,7 +376,7 @@ ('chi_p_length', int64), ('chi_d_offset', int64), ('chi_d_length', int64), - ('fissionable', bool), + ('fissionable', bool_), ('ID', int64), ]) @@ -537,17 +537,17 @@ ('census_bank_buffer_ratio', float64), ('source_bank_buffer_ratio', float64), ('future_bank_buffer_ratio', float64), - ('neutron_transport', bool), - ('electron_transport', bool), - ('proton_transport', bool), - ('neutron_eigenvalue_mode', bool), + ('neutron_transport', bool_), + ('electron_transport', bool_), + ('proton_transport', bool_), + ('neutron_eigenvalue_mode', bool_), ('gpu_strategy', int64), ('gpu_async_type', int64), ('gpu_storage', int64), ]) neutron_multigroup = into_dtype([ - ('hybrid', bool), + ('hybrid', bool_), ]) implicit_capture = into_dtype([ @@ -555,12 +555,12 @@ ]) weighted_emission = into_dtype([ - ('active', bool), + ('active', bool_), ('weight_target', float64), ]) global_weight_roulette = into_dtype([ - ('active', bool), + ('active', bool_), ('weight_threshold', float64), ('weight_target', float64), ]) @@ -583,7 +583,7 @@ ]) population_control = into_dtype([ - ('active', bool), + ('active', bool_), ]) technique = into_dtype([ @@ -597,10 +597,10 @@ source = into_dtype([ ('name', 'U32'), - ('point_source', bool), - ('uniform_x', bool), - ('uniform_y', bool), - ('uniform_z', bool), + ('point_source', bool_), + ('uniform_x', bool_), + ('uniform_y', bool_), + ('uniform_z', bool_), ('point', float64, (3,)), ('x', float64, (2,)), ('y', float64, (2,)), @@ -608,18 +608,18 @@ ('x_pdf_ID', int64), ('y_pdf_ID', int64), ('z_pdf_ID', int64), - ('isotropic_direction', bool), - ('mono_direction', bool), - ('white_direction', bool), + ('isotropic_direction', bool_), + ('mono_direction', bool_), + ('white_direction', bool_), ('direction', float64, (3,)), ('polar_cosine', float64, (2,)), ('azimuthal', float64, (2,)), - ('mono_energetic', bool), - ('discrete_energy', bool), + ('mono_energetic', bool_), + ('discrete_energy', bool_), ('energy', float64), ('energy_pdf_ID', int64), ('energy_pmf_ID', int64), - ('discrete_time', bool), + ('discrete_time', bool_), ('time', float64), ('time_range', float64, (2,)), ('particle_type', int64), @@ -690,9 +690,9 @@ ('scores_offset', int64), ('scores_length', int64), ('particle_type', int64), - ('filter_direction', bool), - ('filter_energy', bool), - ('filter_time', bool), + ('filter_direction', bool_), + ('filter_energy', bool_), + ('filter_time', bool_), ('mu_offset', int64), ('mu_length', int64), ('azi_offset', int64), @@ -851,7 +851,7 @@ def make_simulation_type(N: dict): ('eigenvalue_tally_C', float64, (1,)), ('gyration_radius_offset', int64), ('gyration_radius_length', int64), - ('cycle_active', bool), + ('cycle_active', bool_), ('mpi_size', int64), ('mpi_rank', int64), ('mpi_master', bool_), diff --git a/mcdc/object_/particle.py b/mcdc/object_/particle.py index 50d65fe36..c28b925f0 100644 --- a/mcdc/object_/particle.py +++ b/mcdc/object_/particle.py @@ -29,7 +29,6 @@ class ParticleData(MCDCBase): w: float = 0.0 particle_type: int = PARTICLE_NEUTRON rng_seed: uint64 = uint64(1) - step_count: uint64 = uint64(0) @dataclass @@ -55,7 +54,6 @@ class Particle(ParticleData): alive: bool = False fresh: bool = False event: int = -1 - step_count: uint64 = uint64(0) class ParticleBank(MCDCBase): diff --git a/mcdc/object_/simulation.py b/mcdc/object_/simulation.py index 8a46d278d..09356a2e4 100644 --- a/mcdc/object_/simulation.py +++ b/mcdc/object_/simulation.py @@ -192,7 +192,6 @@ class Simulation(MCDCBase): # GPU metadata gpu_meta: GPUMeta source_seed: int - gen_count: Annotated[NDArray[int64], (1,)] def __init__(self, name: str = "") -> None: @@ -278,7 +277,6 @@ def __init__(self, name: str = "") -> None: # GPU metadata self.gpu_meta = GPUMeta() self.source_seed = 0 - self.gen_count = np.zeros(1, dtype=int64) def _reset_model(self) -> None: # Physics diff --git a/mcdc/object_/tally.py b/mcdc/object_/tally.py index 82980d178..91717b7f3 100644 --- a/mcdc/object_/tally.py +++ b/mcdc/object_/tally.py @@ -372,8 +372,6 @@ def _set_bin_shape_and_strides(self, shape: tuple): # Set bins self.bin_shape = list(shape) - print("SHAPE IS: ",shape) - # Set strides self.stride_time = reduce(operator.mul, shape[4:]) self.stride_energy = reduce(operator.mul, shape[3:]) diff --git a/mcdc/output.py b/mcdc/output.py index 7b8d3d9d5..71dde4a9c 100644 --- a/mcdc/output.py +++ b/mcdc/output.py @@ -202,12 +202,6 @@ def create_tally_dataset(file, mcdc, data): start_sdev = tally["bin_sum_square_offset"] mean = data[start_mean : start_mean + N_bin] sdev = data[start_sdev : start_sdev + N_bin] - print("TALLY : ",tally) - print("BIN SHAPE OFFSET : ",tally["bin_shape_offset"]) - print("BIN_SHAPE_LENGTH : ",tally["bin_shape_length"]) - print("DATA_SHAPE : ",data.shape) - print("DATA : ",data) - print("RESULT : ",mcdc_get.tally.bin_shape_all(tally, data)) shape = tuple([int(x) for x in mcdc_get.tally.bin_shape_all(tally, data)]) mean = mean.reshape(shape) sdev = sdev.reshape(shape) diff --git a/mcdc/transport/particle_bank.py b/mcdc/transport/particle_bank.py index 29630f250..8ab613335 100644 --- a/mcdc/transport/particle_bank.py +++ b/mcdc/transport/particle_bank.py @@ -26,7 +26,6 @@ @njit def get_bank_size(bank): return bank["size"][0] - #!return add_bank_size(bank,0) @njit @@ -51,7 +50,6 @@ def _bank_particle(particle_container, bank): report_full_bank(bank) # Set particle data - #idx = get_bank_size(bank) idx = add_bank_size(bank,1) particle_module.copy(bank["particle_data"][idx : idx + 1], particle_container) @@ -63,9 +61,6 @@ def bank_active_particle(particle_container, program): bank = simulation["bank_active"] _bank_particle(particle_container, bank) - # Increment bank size - #add_bank_size(bank, 1) - @njit def bank_census_particle(particle_container, program): @@ -73,9 +68,6 @@ def bank_census_particle(particle_container, program): bank = simulation["bank_census"] _bank_particle(particle_container, bank) - # Increment bank size - #add_bank_size(bank, 1) - @njit def bank_future_particle(particle_container, program): @@ -83,20 +75,12 @@ def bank_future_particle(particle_container, program): bank = simulation["bank_future"] _bank_particle(particle_container, bank) - # Increment bank size - #add_bank_size(bank, 1) - @njit def bank_source_particle(particle_container, simulation): bank = simulation["bank_source"] _bank_particle(particle_container, bank) - # Increment bank size - # Note that we don't use the atomic operation in add_bank_size function - # as source particle banking is not thread-parallelized - #bank["size"][0] += 1 - @njit def pop_particle(particle_container, bank): @@ -164,11 +148,9 @@ def promote_future_particles(program, data): if particle["t"] < next_census_time: bank_census_particle(particle_container, program) - #add_bank_size(future_bank, -1) j = add_bank_size(future_bank,-1) # Consolidate the emptied space in the future bank - #j = get_bank_size(future_bank) particle_module.copy( future_bank["particle_data"][idx : idx + 1], future_bank["particle_data"][j : j + 1], diff --git a/mcdc/transport/rng.py b/mcdc/transport/rng.py index afb799d16..4a8415e1e 100644 --- a/mcdc/transport/rng.py +++ b/mcdc/transport/rng.py @@ -3,8 +3,6 @@ from numba import uint64, njit -# from mcdc.code_factory.jit import njit - # ====================================================================================== # Random number generator # LCG with hash seed-split diff --git a/mcdc/transport/simulation.py b/mcdc/transport/simulation.py index b636046a2..5c7fd652d 100644 --- a/mcdc/transport/simulation.py +++ b/mcdc/transport/simulation.py @@ -170,8 +170,6 @@ def source_loop(seed, simulation, data): work_start = simulation["mpi_work_start"] work_size = simulation["mpi_work_size"] - print("Gen count before: ",simulation["gen_count"][0]) - simulation["gen_count"][0] = 0 for idx_work in range(work_size): simulation["idx_work"] = work_start + idx_work generate_source_particle(work_start, idx_work, seed, simulation, data) @@ -180,7 +178,6 @@ def source_loop(seed, simulation, data): exhaust_active_bank(simulation, data) source_closeout(simulation, idx_work, N_prog, data) - print("\nGen count after: ",simulation["gen_count"][0]) @njit @@ -205,8 +202,6 @@ def generate_source_particle(work_start, idx_work, seed, program, data): ] particle = particle_container[0] - particle["step_count"] = 0 - # Skip if beyond time boundary if particle["t"] > settings["time_boundary"]: return @@ -289,28 +284,17 @@ def step_particle(particle_container, program, data): simulation = util.access_simulation(program) particle = particle_container[0] - particle["step_count"] += 1 - # Determine and move to event move_to_event(particle_container, simulation, data) - # Execute events if particle["event"] == EVENT_LOST: return - - # In first step of first phase: 10 CPU alive, 0 GPU alive - if (particle["alive"]) and (particle["step_count"] <= 1) : - util.atomic_add(simulation["gen_count"],0,1) # Collision if particle["event"] & EVENT_COLLISION: collision_data_container = util.local_array(1, type_.collision_data) - # In first step of first phase: 3 CPU alive, 0 GPU alive - #if (particle["alive"]) and (particle["step_count"] <= 1) : - # util.atomic_add(simulation["gen_count"],0,1) - # Execute the physics physics.collision(particle_container, collision_data_container, program, data) diff --git a/test/regression/conftest.py b/test/regression/conftest.py index ec831c43c..031555888 100644 --- a/test/regression/conftest.py +++ b/test/regression/conftest.py @@ -139,7 +139,7 @@ def build_command(config): state = "" if target == "gpu": - state = "--gpu_state_storage=united" + state = "--gpu_state_storage=managed" mode = "numba" command = [ From f60e704b1664b82d09df2796b3f49eb372b2576c Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 25 Aug 2026 12:38:38 -0700 Subject: [PATCH 18/34] Back in black --- mcdc/code_factory/array_return.py | 24 ++++++++++------ mcdc/code_factory/gpu/program_builder.py | 28 ++++++------------- mcdc/code_factory/gpu/transport/simulation.py | 16 ++++++++--- mcdc/code_factory/gpu/transport/util.py | 14 ++++++---- mcdc/code_factory/numba_layers_generator.py | 12 +++----- mcdc/config.py | 3 +- mcdc/main.py | 2 -- mcdc/object_/simulation.py | 5 ++-- mcdc/transport/geometry/root_solve.py | 17 +++++------ mcdc/transport/geometry/surface/torus_x.py | 2 +- mcdc/transport/geometry/surface/torus_y.py | 2 +- mcdc/transport/geometry/surface/torus_z.py | 2 +- mcdc/transport/particle_bank.py | 23 +++++++++++---- mcdc/transport/util.py | 6 ++-- test/regression/conftest.py | 2 +- test/unit/test_annotation_shape.py | 4 +-- test/unit/test_numba_layers_generator.py | 2 +- 17 files changed, 90 insertions(+), 74 deletions(-) diff --git a/mcdc/code_factory/array_return.py b/mcdc/code_factory/array_return.py index 6be3028da..b793a4969 100644 --- a/mcdc/code_factory/array_return.py +++ b/mcdc/code_factory/array_return.py @@ -78,7 +78,6 @@ def into_voidptr(value): return into_voidptr_python(value) - # ============================================================================= # uintp/voidptr casting utility functions # ============================================================================= @@ -97,7 +96,6 @@ def into_voidptr_python(value): raise RuntimeError("`into_voidptr` is only supported in nopython mode.") - @nb.extending.overload(into_voidptr_python) def into_voidptr_overload(value): @@ -129,6 +127,7 @@ def impl(value): # Helper decorators, functions, and builtins for returning arrays ############################################################################### + # Overload target def array_result(array): return array @@ -147,6 +146,7 @@ def impl(array): return impl + # Raises an error if the context is not recognized def context_guard(context): if isinstance(context, nb.core.typing.context.Context): @@ -158,6 +158,7 @@ def context_guard(context): else: raise nb.core.errors.UnsupportedError(f"Unsupported target context {context}.") + # Typing for the `array_return` builtin. def array_return_typing(fn, elem_type, ndim): @@ -203,7 +204,6 @@ def builtin(context, builder, sig, args): import llvmlite.binding as ll from llvmlite import ir - lmod = builder.module retty = nb.types.Tuple( [nb.types.voidptr, nb.types.Tuple([nb.types.uintp] * ndim)] @@ -218,9 +218,13 @@ def builtin(context, builder, sig, args): # GPU platforms require a `targetdata` for array construction, which is created # slightly differently depending upon the platform. - if config.ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + if config.ROCM_AVAILABLE and isinstance( + context, nb.hip.target.HIPTargetContext + ): targetdata = ll.create_target_data(nb.hip.amdgcn.DATA_LAYOUT) - elif config.CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + elif config.CUDA_AVAILABLE and isinstance( + context, nb.cuda.target.CUDATargetContext + ): targetdata = ll.create_target_data(nb.cuda.cudadrv.nvvm.NVVM().data_layout) lldtype = context.get_data_type(dtype) @@ -228,9 +232,13 @@ def builtin(context, builder, sig, args): # upon platform if isinstance(context, nb.core.cpu.CPUContext): itemsize = context.get_abi_sizeof(lldtype) - elif config.ROCM_AVAILABLE and isinstance(context, nb.hip.target.HIPTargetContext): + elif config.ROCM_AVAILABLE and isinstance( + context, nb.hip.target.HIPTargetContext + ): itemsize = lldtype.get_abi_size(targetdata) - elif config.CUDA_AVAILABLE and isinstance(context, nb.cuda.target.CUDATargetContext): + elif config.CUDA_AVAILABLE and isinstance( + context, nb.cuda.target.CUDATargetContext + ): itemsize = lldtype.get_abi_size(targetdata) else: raise nb.core.errors.UnsupportedError( @@ -268,6 +276,7 @@ def builtin(context, builder, sig, args): # was defined above. nb.extending.lower_builtin(fn, *sig.args)(builtin) + # A function decorated with `array_return` may return an array by passing # it through the `array_result` function and returning the output def array_return(sig, ndim=1): @@ -277,4 +286,3 @@ def array_return_true_decorator(fn): return fn return array_return_true_decorator - diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index bc56af070..851ea8416 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -10,6 +10,7 @@ # Transport function adapter # ====================================================================================== + # Overwrites global symbols in other modules with gpu-compatible counterparts def adapt_transport_functions(): @@ -98,14 +99,7 @@ def forward_declare_gpu_program(simulation_dtype): # Get to set the globals global none_type, simulation_type, data_type - global \ - state_spec, \ - access_simulation, \ - access_data_ptr, \ - access_group, \ - access_thread, \ - particle_gpu, \ - particle_record_gpu + global state_spec, access_simulation, access_data_ptr, access_group, access_thread, particle_gpu, particle_record_gpu global step_async, find_cell_async global alloc_managed_bytes, alloc_device_bytes @@ -163,12 +157,15 @@ def find_cell(program: nb.uintp, particle: particle_gpu): alloc_device_bytes = harmonize.alloc_device_bytes from mcdc.transport import util - @nb.extending.overload(util.access_simulation,target="hip") + + @nb.extending.overload(util.access_simulation, target="hip") def access_simulation_gpu_overload(program): def impl(program): return access_simulation(program) + return impl + # ====================================================================================== # Program builder # ====================================================================================== @@ -197,7 +194,6 @@ def impl(program): BLOCK_COUNT = 0 - # Compiles gpu kernels and loads in the functions that call into said kernels def build_gpu_progs(input_deck): @@ -288,7 +284,6 @@ def real_setup_gpu(mcdc_array, data_tally): source_closeout(simulation, 1, 1, data) - def build_gpu_program(data_size): import harmonize @@ -300,15 +295,9 @@ def build_gpu_program(data_size): global alloc_program, free_program - global \ - load_state_device_simulation, \ - store_state_device_simulation, \ - store_pointer_state_device_simulation + global load_state_device_simulation, store_state_device_simulation, store_pointer_state_device_simulation - global \ - load_state_device_data, \ - store_state_device_data, \ - store_pointer_state_device_data + global load_state_device_data, store_state_device_data, store_pointer_state_device_data global init_program, exec_program, complete, clear_flags, set_device global ARENA_SIZE, BLOCK_COUNT @@ -476,7 +465,6 @@ def teardown_gpu_program(simulation): free_state(cast_uintp_to_voidptr(simulation["gpu_meta"]["state_pointer"])) - # ====================================================================================== # Type casters # ====================================================================================== diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index e4485a9a0..2e22923b5 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -40,8 +40,12 @@ def source_loop(seed, simulation, data): # Store the global state to the GPU if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: - gpu_module.store_state_device_simulation(simulation["gpu_meta"]["state_pointer"],simulation) - gpu_module.store_state_device_data(simulation["gpu_meta"]["state_pointer"],data) + gpu_module.store_state_device_simulation( + simulation["gpu_meta"]["state_pointer"], simulation + ) + gpu_module.store_state_device_data( + simulation["gpu_meta"]["state_pointer"], data + ) # Execute the program, and continue to do so until it is done block_count = gpu_module.BLOCK_COUNT @@ -66,8 +70,12 @@ def source_loop(seed, simulation, data): # Recover the original program state if settings["gpu_storage"] == GPU_STORAGE_SEPARATE: - gpu_module.load_state_device_simulation(simulation, simulation["gpu_meta"]["state_pointer"]) - gpu_module.load_state_device_data(data, simulation["gpu_meta"]["state_pointer"]) + gpu_module.load_state_device_simulation( + simulation, simulation["gpu_meta"]["state_pointer"] + ) + gpu_module.load_state_device_data( + data, simulation["gpu_meta"]["state_pointer"] + ) simulation["mpi_work_size"] = full_work_size diff --git a/mcdc/code_factory/gpu/transport/util.py b/mcdc/code_factory/gpu/transport/util.py index 321307832..74faddab1 100644 --- a/mcdc/code_factory/gpu/transport/util.py +++ b/mcdc/code_factory/gpu/transport/util.py @@ -5,26 +5,30 @@ from numba import njit, types -def atomic_add(array,idx,value): +def atomic_add(array, idx, value): result = array[idx] array[idx] += value return result -@nb.extending.overload(atomic_add,target="gpu") -def overload_atomic_add_gpu(array,idx,value): +@nb.extending.overload(atomic_add, target="gpu") +def overload_atomic_add_gpu(array, idx, value): def impl(array, idx, value): return harmonize.array_atomic_add(array, idx, value) + return impl -@nb.extending.overload(atomic_add,target="cpu") -def overload_atomic_add_cpu(array,idx,value): + +@nb.extending.overload(atomic_add, target="cpu") +def overload_atomic_add_cpu(array, idx, value): def impl(array, idx, value): result = array[idx] array[idx] += value return result + return impl + # ============================================================================= # Generic GPU/CPU local array variable constructors # ============================================================================= diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index 5d0530547..f609cd881 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -376,7 +376,6 @@ def generate_numba_layers(simulation): set_object(object_, annotations, structures, records, data, set_data=True) set_object(simulation, annotations, structures, records, data, set_data=True) - # ================================================================================== # Set with records # ================================================================================== @@ -427,7 +426,6 @@ def generate_numba_layers(simulation): mcdc_simulation["gpu_meta"]["simulation_pointer"] = mcdc_simulation_pointer mcdc_simulation["gpu_meta"]["data_pointer"] = data["pointer"] - # GPU program setup if config.target == "gpu": @@ -807,9 +805,9 @@ def create_simulation_container(dtype): @njit def create_simulation_container_on_gpu(dtype, size): if config.gpu_state_storage == "managed": - mcdc_ptr = gpu_builder.alloc_managed_bytes(size*8) + mcdc_ptr = gpu_builder.alloc_managed_bytes(size * 8) else: - mcdc_ptr = gpu_builder.alloc_device_bytes(size*8) + mcdc_ptr = gpu_builder.alloc_device_bytes(size * 8) mcdc_uint = cast_voidptr_to_uintp(mcdc_ptr) if config.gpu_state_storage == "separate": @@ -889,10 +887,8 @@ def align(field_list): pad_id = 0 for field in field_list: if len(field) > 3: - print_error( - "Unexpected struct field specification. Specifications \ - usually only consist of 3 or fewer members" - ) + print_error("Unexpected struct field specification. Specifications \ + usually only consist of 3 or fewer members") multiplier = 1 if len(field) == 3: field = (field[0], field[1], fixup_dims(field[2])) diff --git a/mcdc/config.py b/mcdc/config.py index db607395c..e0d9067f4 100644 --- a/mcdc/config.py +++ b/mcdc/config.py @@ -120,6 +120,7 @@ def _build_parser() -> argparse.ArgumentParser: try: import numba.hip as hip + ROCM_AVAILABLE = True except: ROCM_AVAILABLE = False @@ -127,6 +128,7 @@ def _build_parser() -> argparse.ArgumentParser: if not ROCM_AVAILABLE: try: import numba.cuda as cuda + CUDA_AVAILABLE = True except: CUDA_AVAILABLE = False @@ -134,7 +136,6 @@ def _build_parser() -> argparse.ArgumentParser: CUDA_AVAILABLE = False - # ====================================================================================== # Simulation-setting overrides # ====================================================================================== diff --git a/mcdc/main.py b/mcdc/main.py index ca600ae01..d2c191309 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -65,7 +65,6 @@ def run_simulation(simulationPy: Simulation): # Run simulation import mcdc.transport.simulation as simulation_module - if settings.neutron_eigenvalue_mode: simulation_module.eigenvalue_simulation(simulation_container, data) else: @@ -179,7 +178,6 @@ def prepare(simulationPy: Simulation): from mcdc.code_factory.gpu.program_builder import setup_gpu_program setup_gpu_program(simulation_container, data) - # ================================================================================== # Finalize diff --git a/mcdc/object_/simulation.py b/mcdc/object_/simulation.py index 09356a2e4..e473d829c 100644 --- a/mcdc/object_/simulation.py +++ b/mcdc/object_/simulation.py @@ -193,7 +193,6 @@ class Simulation(MCDCBase): gpu_meta: GPUMeta source_seed: int - def __init__(self, name: str = "") -> None: self.compiled = False @@ -436,16 +435,16 @@ def _finalize_compilation(self) -> None: self.bank_census.size[0] = int(settings.census_bank_buffer_ratio * N_work) self.bank_source.size[0] = int(settings.source_bank_buffer_ratio * N_work) self.bank_future.size[0] = int(settings.future_bank_buffer_ratio * N_work) - + # ================================================================================== # Platform targeting, adapters, and toggles for portability # ================================================================================== import mcdc.config as config + # Build GPU program if desired if config.target == "gpu": from mcdc.code_factory.gpu.program_builder import build_gpu_program - # Initialize run state derived from the compiled settings self.k_eff = settings.k_init diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/root_solve.py index eb1aba113..5f72a8483 100644 --- a/mcdc/transport/geometry/root_solve.py +++ b/mcdc/transport/geometry/root_solve.py @@ -21,9 +21,10 @@ def sqrt(x): imag_part = -math.sqrt((r - x.real) / 2.0) else: imag_part = math.sqrt((r - x.real) / 2.0) - + return complex(real_part, imag_part) + @njit() def power(x, n): result = 1 @@ -54,7 +55,6 @@ def principal_nth_root(x, n): return result - @njit() def solve_quadratic(coeff, roots): a = coeff[2] @@ -133,10 +133,10 @@ def solve_depressed_quartic(coeff, roots): beta = -2.0 * y - a gamma = (2.0 * b) / sqrt(2.0 * y - a) - roots[0] = ((-alpha) + sqrt(beta + gamma)) / 2.0 # - + + - roots[1] = ((-alpha) - sqrt(beta + gamma)) / 2.0 # - - + - roots[2] = ((alpha) + sqrt(beta - gamma)) / 2.0 # + + - - roots[3] = ((alpha) - sqrt(beta - gamma)) / 2.0 # + - - + roots[0] = ((-alpha) + sqrt(beta + gamma)) / 2.0 # - + + + roots[1] = ((-alpha) - sqrt(beta + gamma)) / 2.0 # - - + + roots[2] = ((alpha) + sqrt(beta - gamma)) / 2.0 # + + - + roots[3] = ((alpha) - sqrt(beta - gamma)) / 2.0 # + - - @njit() @@ -165,7 +165,9 @@ def solve_quartic(coeff, roots): sub_coeff[4] = 1.0 + 0.0j sub_coeff[3] = 0.0j sub_coeff[2] = (-3.0 * power(b, 2)) / (8.0 * power(a, 2)) + c / a - sub_coeff[1] = power(b, 3) / (8.0 * power(a, 3)) - (b * c) / (2.0 * power(a, 2)) + d / a + sub_coeff[1] = ( + power(b, 3) / (8.0 * power(a, 3)) - (b * c) / (2.0 * power(a, 2)) + d / a + ) sub_coeff[0] = ( (-3.0 * power(b, 4)) / (256.0 * power(a, 4)) + (c * power(b, 2)) / (16.0 * power(a, 3)) @@ -187,4 +189,3 @@ def solve_quartic(coeff, roots): for idx in range(4): roots[idx] = sub_roots[idx] - b / (4.0 * a) - diff --git a/mcdc/transport/geometry/surface/torus_x.py b/mcdc/transport/geometry/surface/torus_x.py index b342a0f7b..3f9152d8f 100644 --- a/mcdc/transport/geometry/surface/torus_x.py +++ b/mcdc/transport/geometry/surface/torus_x.py @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients,roots) + root_solve.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_y.py b/mcdc/transport/geometry/surface/torus_y.py index a44f93878..86c1892d4 100644 --- a/mcdc/transport/geometry/surface/torus_y.py +++ b/mcdc/transport/geometry/surface/torus_y.py @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients,roots) + root_solve.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_z.py b/mcdc/transport/geometry/surface/torus_z.py index 0766d88ec..d65c0fd33 100644 --- a/mcdc/transport/geometry/surface/torus_z.py +++ b/mcdc/transport/geometry/surface/torus_z.py @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients,roots) + root_solve.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/particle_bank.py b/mcdc/transport/particle_bank.py index 8ab613335..519598491 100644 --- a/mcdc/transport/particle_bank.py +++ b/mcdc/transport/particle_bank.py @@ -50,7 +50,7 @@ def _bank_particle(particle_container, bank): report_full_bank(bank) # Set particle data - idx = add_bank_size(bank,1) + idx = add_bank_size(bank, 1) particle_module.copy(bank["particle_data"][idx : idx + 1], particle_container) @@ -92,7 +92,6 @@ def pop_particle(particle_container, bank): idx = add_bank_size(bank, -1) - 1 particle_module.copy(particle_container, bank["particle_data"][idx : idx + 1]) - # Set default IDs and event for the live particle particle = particle_container[0] particle["alive"] = True @@ -148,7 +147,7 @@ def promote_future_particles(program, data): if particle["t"] < next_census_time: bank_census_particle(particle_container, program) - j = add_bank_size(future_bank,-1) + j = add_bank_size(future_bank, -1) # Consolidate the emptied space in the future bank particle_module.copy( @@ -168,9 +167,21 @@ def manage_particle_banks(simulation): serial = simulation["mpi_size"] == 1 with objmode(): - print("Bank census has size: ",get_bank_size(simulation["bank_census"]),flush=True) - print("Bank source has size: ",get_bank_size(simulation["bank_source"]),flush=True) - print("Bank future has size: ",get_bank_size(simulation["bank_future"]),flush=True) + print( + "Bank census has size: ", + get_bank_size(simulation["bank_census"]), + flush=True, + ) + print( + "Bank source has size: ", + get_bank_size(simulation["bank_source"]), + flush=True, + ) + print( + "Bank future has size: ", + get_bank_size(simulation["bank_future"]), + flush=True, + ) # TIMER: bank management time_start = 0.0 diff --git a/mcdc/transport/util.py b/mcdc/transport/util.py index f9b3f4b5f..4d1e7e7d1 100644 --- a/mcdc/transport/util.py +++ b/mcdc/transport/util.py @@ -158,6 +158,7 @@ def atomic_add(array, idx, value): array[idx] += value return result + @njit def local_array(shape, dtype): return np.zeros(shape, dtype=dtype) @@ -166,9 +167,10 @@ def local_array(shape, dtype): def access_simulation(program): return program -@nb.extending.overload(access_simulation,target="cpu") + +@nb.extending.overload(access_simulation, target="cpu") def access_simulation_cpu_overload(program): def impl(program): return program - return impl + return impl diff --git a/test/regression/conftest.py b/test/regression/conftest.py index 031555888..144efcbb5 100644 --- a/test/regression/conftest.py +++ b/test/regression/conftest.py @@ -183,7 +183,7 @@ def compare_outputs(output_path, answer_path, target): def compare_tallies(output, answer, target, errors): - gpu_pass_list = ["uq_var","sdev"] + gpu_pass_list = ["uq_var", "sdev"] name_root = "tallies" for tally in answer[name_root].keys(): name_tally = f"{name_root}/{tally}" diff --git a/test/unit/test_annotation_shape.py b/test/unit/test_annotation_shape.py index 1d408cdbe..dbc0686f0 100644 --- a/test/unit/test_annotation_shape.py +++ b/test/unit/test_annotation_shape.py @@ -67,8 +67,8 @@ def test_stringified_annotation_rejects_incorrect_offset_shape(capsys): def test_generated_accessor_resolves_dimension_offset(): - all_source = _accessor_1d_all("mgxs", "energy", "G+1",nb.types.float64) - last_source = _accessor_1d_last("mgxs", "energy", "N - 2",nb.types.float64) + all_source = _accessor_1d_all("mgxs", "energy", "G+1", nb.types.float64) + last_source = _accessor_1d_last("mgxs", "energy", "N - 2", nb.types.float64) assert 'size = mgxs["G"] + 1' in all_source assert 'size = mgxs["N"] - 2' in last_source diff --git a/test/unit/test_numba_layers_generator.py b/test/unit/test_numba_layers_generator.py index 687da7d1f..64f786135 100644 --- a/test/unit/test_numba_layers_generator.py +++ b/test/unit/test_numba_layers_generator.py @@ -131,7 +131,7 @@ def test_scalar_integer_getters_cast_values_from_data(): def test_float_and_bulk_getters_remain_zero_copy_views(): assert "return data[offset + index]" in _accessor_1d_element("example", "values") assert "return array_result(data[start:end])" in _accessor_1d_all( - "example", "values", "values_length",nb.types.float64 + "example", "values", "values_length", nb.types.float64 ) From 598da8b4ecdc105311a95cacc349c2f4907781e2 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 25 Aug 2026 13:42:05 -0700 Subject: [PATCH 19/34] Removed debug fields --- mcdc/code_factory/gpu/transport/simulation.py | 1 - mcdc/numba_types.py | 3 --- mcdc/transport/particle_bank.py | 17 ----------------- 3 files changed, 21 deletions(-) diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 2e22923b5..a95520b2a 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -25,7 +25,6 @@ def source_loop(seed, simulation, data): full_work_size = simulation["mpi_work_size"] - simulation["gen_count"][0] = 0 if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: phase_size = 1000000000 else: diff --git a/mcdc/numba_types.py b/mcdc/numba_types.py index b14904c85..fd2c161e2 100644 --- a/mcdc/numba_types.py +++ b/mcdc/numba_types.py @@ -22,7 +22,6 @@ ('w', float64), ('particle_type', int64), ('rng_seed', uint64), - ('step_count', uint64), ]) particle = into_dtype([ @@ -32,7 +31,6 @@ ('alive', bool_), ('fresh', bool_), ('event', int64), - ('step_count', uint64), ('x', float64), ('y', float64), ('z', float64), @@ -865,6 +863,5 @@ def make_simulation_type(N: dict): ('runtime_output', float64), ('runtime_bank_management', float64), ('source_seed', int64), - ('gen_count', int64, (1,)), ]) diff --git a/mcdc/transport/particle_bank.py b/mcdc/transport/particle_bank.py index 519598491..2213d87d5 100644 --- a/mcdc/transport/particle_bank.py +++ b/mcdc/transport/particle_bank.py @@ -166,23 +166,6 @@ def manage_particle_banks(simulation): master = simulation["mpi_master"] serial = simulation["mpi_size"] == 1 - with objmode(): - print( - "Bank census has size: ", - get_bank_size(simulation["bank_census"]), - flush=True, - ) - print( - "Bank source has size: ", - get_bank_size(simulation["bank_source"]), - flush=True, - ) - print( - "Bank future has size: ", - get_bank_size(simulation["bank_future"]), - flush=True, - ) - # TIMER: bank management time_start = 0.0 if master: From cf7479b41a35c423059f16e9db112eb9e54ad7b7 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 25 Aug 2026 13:50:23 -0700 Subject: [PATCH 20/34] Removed regression test slab_isobeam_td_census --- .../slab_isobeam_td_census/answer.h5 | Bin 71256 -> 0 bytes .../slab_isobeam_td_census/input.py | 60 ------------------ 2 files changed, 60 deletions(-) delete mode 100644 test/regression/slab_isobeam_td_census/answer.h5 delete mode 100644 test/regression/slab_isobeam_td_census/input.py diff --git a/test/regression/slab_isobeam_td_census/answer.h5 b/test/regression/slab_isobeam_td_census/answer.h5 deleted file mode 100644 index b6a2d5cb30508b7a4a16880e252fbca4cdf69d3e..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 71256 zcmeHQ1z1%}*FH2N-OZu9yUwD!yW!9vofe8o2x5Rz5&}x8CQ8Ics9=*)#82Gkf;z*)J(ZdInS!>=cNZgPa^ejG+A-#Q%K4 z9Zt*ypOf%$e7+YhfXzi9l5z2JLTrElG4qE&u;S9`aOKa}v$N6HM_@GZ<>MvpG=vN> zHbjn3Sd_mN0!I4URx=HD!nIlkCtGlShKnNN{Gy_R!ov{JelamYVFA(j^cZwVNRVGN zA}Tg4CMeYJvro?CiQpB!APEp8SY4ch6Cki2i1&9QjAze;*$4{ zaY90DGC8g%5aA1&*WVC`>@U>h$4Iudu+$^Kj1qXyOuC#SI8GxO610dG$R1tygY{jkg$MB^Ff#0!f5B;U^W}_AA=Y4nvD+Ft*KFpIbWT zPlTVLobxAI9HMxH@W2$|13$F<(tTpSpWw%VdVIgfNi9um^A_TYg@pjqKW7e{yUj{` z|D{91EY1o`5elDh{`F!o+at4nXADcO55Hi%YH2aJcuH#}+ zo5M{;N>)w^(_}pScyGgxxA<;J^0`C=A_VR27lH}nZr_L{xb(2_Fwd~qP;WFYj@QTA zCpI-6)vu&hm8;87aVZs zEs7r#0sJ~Je*I@TegWqT)x|ANA`&^ipnyw8;^a2Yuj}Fcu88aFdOv#Z!u3ilqWEq= zgn2(O6aMq-z^q7YoAH*xBT$47{Lu36#{-GPHh2+_861cRL>BSD;4t^zMCK>Qv;(zh$`}p;Ed^%qG;sSVyi)!KXv2qyxJDnVtPm7a`I61r) zr#Fc|0T&;JnQn`JK`~P)f|)l6_=K6a3Ntqze0=8Z!+*Zfn90MBq?q8$fxRUDgt_L; zq~leG&wo}$3gku?a1|uBL5uV7861cRw0?0su>0b^e_xH!7jVg* zf$fJq{{H@Ag~T@Y=j<=|^&uo~PF=3_b<@e<=iaJc*VDr3(W-r255KO6RR6l(&zg(J zR|VIA8t$*R#0ZzoD)MzbCtS9O=-2h|>%1yrU)TH5bHCkhMLs_kvSOY+XIyXQ!4d0% z^sv$5fs-{iUA4-on`9)ERQPeEdvhXL@vj^Axu z-;TsKe?AUH|Aoqr8o`e{+i(>$690Nl%W>J)=2lw%@Ps;isx&_E$Fc>NOpAX3@W*OD zeCj{f&vKC1#?SY!pS_>`)$c5dad^;TVVV8?+41@TdJ z>xJLfKeIGEbKoU@U;p=imvaR9xgZ2iU?Z5{9}5rsI6Q!>A5wO4H-MaM?ZrTeV#3Gu_E2ZjCT85W)q=%|>WnXJzTHqX-=9k$HVJJ#RdFAAHE zE$u(kJ)_VuLE-rFarrZ!yRPu`kMfI*^$YV^fk5EXXTE)%IsAeG{KDeUA+dg*q2a!M z2>;la*eJg*mm9F+^JT&?SMv>ujrNRL5%Ia8WtesZVZPJCf9Dk%8xj)~5EUL9fh~9b zcV5`~^G`({4!_4E#p z4f91uVVaMKjrPM?0@KE5KVQ$!2w~IbwQUBT&-xKj;Q>*8(b1SQaOh*x|KpmoPxJMk z&-(d({%n8#{(00W5{t&KcK`dU;~Nn#7LT86eEYV4{9X6Ir8sP4dC9@<;hyfmKL9^w4Vc+v#yxcE)+L!i? z8(CnR9Ej`R-x>#0Y_a;^JHGs~!2V_pn?DX{;o6^!^B+7Mp|RKydS%e(fM1wjQ~+iH z=QC*=CgtZ02Rmiz0FHB8Pby^P(cSc#Y${jCoNRT)f_CUdFtr3NC)ld0xg}jQ)Fm#q(46 z)xU>_^e=%2{&hB9;@{VP|F{D`7x$(g72f@p~ zjz_ov77zT0|8x2W9tXVq>v*jFN$~jN?_=lB`@Z%3QseQ**2jCnqWrZHnE(8OpU)ZL zo*VG35A{Qb&4i@9hIEGx{sZyf8;l>I^=}F9>!Rx% z;ov8Y{kyG)Q9nSHe(NfnEFSmcO7Vno+1wq{KK6_K@_Q0*}(F z3Om>|R;gmR(Gc#cpS;q#)DnDSWYc$UHwMi$m%Nef5{q{EhwWmr`TX#^JrUp|${j#< z%M-4J^l9tS>4N-99k-QIR)F|;fOND@12c|ICq<4M!%NlD_;MBTMZ5e%b{Wn}Haov8 z8VGn%m3~5LI+nh>1qN8#}do7MZMa zfDp=iT~DI`yPj=HkgPWW@5AoJb#x+&cKMs^A}pB0?o_=3gsQn>tR4o!)0Aa3NOyb4 zJiVb{+RhG+W=alhd#DYGVcjBrfmX0mEPWm0E>Td4A3?;g6^E;vR~2dq8$u=5_AHhd z{zbd|KkdR!O)membtN36H`5e&5CJ4y=Zpl7IDtY8yI$TVN8r|0^?Xle0+LlptyAkQ zU=p#svY=WEE;cCWr0kM{@%tuPglmQ1%&VYHO^JqZdtJ|mQ7XPgyZlAFl-kM2WW`2+ zrT&yLU3@%DK1tfof6xo$Mj4*nBXotXLvq?DH`zisx@^PCOS+JnJtU%W(GWaq6g@un zD?^fG7{S`hvJh!dvNnAUFI>yru+{y%0c=i}p$ojixoDT~v`dBo{{fYaL2$HrS(AcN zG#J6dQ$lKykSo3f#Z9CS(>27?N(CMe@!Vv}!o&dXr+^!j(=gR2TxuMF_H4HFWTi#+Qp?ui?l)C3lJMqldFuP zpuKY7>O=its6z8R?28NlhVdiL;RdGQW1ke}d&C*2oTFVHXGj9QA9>dBB~v(hbdPm% zv@j^Xso{CoC=0Kem>NbeD?zb?)l|kf6*Njz9zRi`2Vio5y)=Sl(JtSxOLo$}NVB96 zkdo6Oe;sEJW*0xmnbbwWD`HjBo3{SIH?$*NyW0s4y&zd97-I`#?_HKkA+6w|YRL<* z(}fP3{e_Ggg77w)Z`+nENs#8hhZyyfgw>-r@*OWK!lv}w+xMmoqKMG+;~PHdL4Nr% z9-RipMZ0{#E`vLa5qFTwpd>u#Zb_>XNZGrW=GFQE@&UEa25WzCA5yc~*XRv0#%_TR z*4n{Me`1!X@SXQ#@m_JFHzw=ZJkY<1mLZYQ^RXr7TDIe zR=oGH>@5CaD00~{!)`icUeduaKqYpa|eWRyA82pL$HGyIC? z!rXd_{!|ue>q`M3=f09uu!<1`)~~A!xg#@+zksnfwJDD>tZ%(${l4wv;(HgI=S{2= zA7F8e2U3a?57UPoAi~(+=%#=ts!1N+!(xOS~ZL;1RWYHMWd z;C5R_hCn9(y5UL63QiB;?C-f)HQ)kUBHXrlY&QdT+BFtRx@Mqcazd`8TpJ3HUC@4u zkbuJvOu|>A=kDK|>e%($sl_2rQSyYqRwkfPz4~#-(z$go={=T9&K=6vL$570|zyqFY$hK*bN502z1fLpc!T70h-Y?YEP z|0s?IrHQp-#ZOJaZICLvnA8SNTx$$!aMyw(S5_-#lN*DY?=x-M{WP#|&zjZIljO7g z#Z0%k#IAuKc6V8Zg^+Q=hTWyRUq;W3bGsU@v>l6P1BMaN60g3eznxv`2LrWJo1$TS zIHmt$Z6J))4N#U`M1x01*eNleBJwKfFVJu)E=7U_WI+={yEo{J<^OBd6_xG`jO!HmQl*njU zc}c2&*Mks9xXEsnL~jqPNZ5Lb;=MqCOik^=O<%~NHr}LWtN@o&3b)*DQ3j&E0X=~n zYiLPn@4w$|4QaZTuCXY62-giwiyYR6M)&2O_dV6YJ2bKgB_jmmeJ{`4(d7iYW=PX( zPeMI6%G#*YJ&KadmA7p-x5)%TpHtPNQ{V?zf`gb)B3>ZK?~rwaRnQBsGVpYbGv5FnWW8y;o%8YB#8`xawu% zzz8D6<`(=}hH$W8Ros4UJ1|t6JQoP6F!7F$*zLX7YsD=a3)k~{m3~bP*X@~ zzHi9~PZ9DDE?Ca3qY!O=A27m82)lK+GSozK19896o)jlus8#OX5HZ9wJI)DS)LRvM z8Z(YPjpXs!`{b9kOJr>$#q;ICAoAwS&Z6)LxE~tXcp}gTg>mGODEI(_IkFum&1gavHxOla5p-!?y{sZgmM#7N!03t-R}JbQp_UIS1*xt+e8UeDD(xyTagg@ z_)}Z?O`2J}n+0l|+l=Vo9{TR9`|8Y)pVYg`^`H>Q_4}pgY$Bc2%V*EbGkL)VuVp7V zPx1dccJZ!jdsXNV1!oAFMVJMBKt%7+#t%#pVBvP|c=1D17-D0oJmSj)Y>`B+ z<2>|meCg_=8+P%6v{p*di`r+i{i3<)w74P>J51g>ndj&GXg2+4eXt*KRC?cI8wLK( zcPN(CM}x|NZ2{iOeo)EG+(UcE2MFFhX7i#kgX0z_pJqooLNCS2?44Ftz-(dI`Q8hl zPCEX660HN6-XYlPbVnZ)TYbv7+m+xpXD#V+J8d}YeUJNy89nf@q2F1HYXO5Le|Jo+ z6d#e3_L&W_EcD~JWVND_0x<3oyay!}UJ z@qZKX_%sg(I}jwUW~VEAIGc_y=Vy^`ueh^H9t{VbzEhnQbm2goyQA|Rdl-~&Y_ZDG z4uT3-i_4|!t-<*Kn}EKeHIVadB+rv@0>52XKJ+is2D9f(uj0*gAl2rGf`g9@JPPMa zqTg-=O8dF{m%Nn$SM5rE+dO4(d9`k(e7*o|ynLvrg9qSlv*yPukNCh$r}Ry9ff!J> zA=7So&D|%yh2Q*n0bZY-bEz4HJUa{tLe81~zMsx6%)4*UrYc1MqEBNo7qc%ell(D5 zjzJ*gSM5=a&0Gc)>Qv4xCcd!#WqvTM)`fD>sDO$yHlR;sL|nPo1~fNrB|aCc0#7ZwN=BmK-FpRNa`F4L^;mbsj{hphS`x#snvck|g_9q)vB!+H?bZa@ zXusO%AZF0IGqvePi3AW7?^aA?Bbgny@Ho$xx4YiimP|iFX>6{I2_0rzpdWNe?(SO? zcDT?+z#1s@^VlVI7mqeiU@*iSlOsxR^#;Ydro!gSb};b4vGt^H5GcLX+R$Yg1lgiA zX)DjT1FO_4^tv1B5JJR|e8$EK601LIUC6Y6;|CvXrQNRyD3NKNGzwj4mn1*J_`n`q zI8L`%Rw%*Qr1ypaVroEiFOB{2L4Ig^V%V9oTnZWv88h0~$-zA%(*%Ps5>S`4XrX4{ z1NHuOtfNv>-}GbN+Djj_%8e2Cz2n3o?KZbwg3DD zfAIGxzFXYo3!gTRp5Y1bglftu<0(N4n4;|{uA~eCHLrmFHKyKhVuK7BA=(~DH4;w* z?-hjR^t3mt15BWYEM@G4r4g8IeqclQL0N;! zV(eOfjKT}93Oc^*p)!F?VK{Wg-W&=tqCTD98UW^_bluryZjdH)a*x&qV|dR3JFGvd zLgU`2@eMV)&~k*%C+mPdyb9BCVt9;#xb@eI+;-}KLdh|AsdimZv=BLPsY4E^HAwDdWaae^$gIF8i3{ia=f zxY;BwOi=)fASb1w_1ruFKi<6;Q5R{@=LKSho$0BJKi+(d|6*@^kT4pyJzRhJ-5wvv zW07LsbQKNx>MfIgudU%Dg{|y?N_n`MQ4=Jv!V%USOGqLQ3WO6Pdp@mgwFbl%Hv6eH zYM`zl+rH9A2iWc$c-p)|7t*UmFKjT>gS;~_?Od z8d=H!$tHeY&)Z@kDd~Hu!bB9-CEpw{6&8eBsWR;~tNG#jD&}!@8Y=J(+rg8Y&I&Fe z%O>7Lj(pQD)00++inp{7e!8bPq?cjAIEHdMecjxi7u=XmIP)=8dg0B=(YBCV zm8JJq-XE^fUX*0E@`d*VQ(Mx@Y=PAcO^FUOhsrDWhS_8_;4~MhvEEGwXlri1eW<}7 z1d@+N)HWJHU%p~fl#L8*P2;8cn4u39^5vP5Q5wL0clGmpUL$yQ#kOxMKn}K*W%>p? zp&-;S-i`1bKyB6;y?fnq5QW&SwN_0C9&jY?Bq?VF^2`fwK48W}`IAdFH(sZQPs^^U zJa|hFPh*};vDPqxP!9F(H*(+perUUQR&ZnpBji}`R4zp?oWCU|sVDUM^1=;~bWJAy z|NVV`vTApt=71yA7c)z*xnvJ(Ph8?Mu=%l)86cr-i~apU_9oUVO8^h_v^1t3F|Y z>}z{+PjoF=&|h93KP~JR!wahd7{!m--T!8J`0?-WmMPiP$c^?6phsV)zc$wbl35w= z8_RftTeA!z$lVi$Ldq$Fx0r%&y16~6kQFp_7KGTQXh0RUJ2H{o6dn>spyTM>pr-Qb z;mYM!!0%IVD*Lb&h&gnhmoAWo$~8@`f+E_WcQFaIa=8rfNXV__Qqclw7sgJC0D3^D ztkpW8V*nKOsoI`8q98JMWHNk(3cS*M`G`4$8K~bXZ?MT>1g~U&lUog}!0R`ZmwAR3 zH@>mKa=a$i#34lt?KUvmWo1N?WVYv z!*Y<#+^N%gPzi)wFRY1m2OxGqnubk?fb2`DUYh|%2;tqHfHXA#n)RIm1 zuvc&vF9O85ItcLgoH;IAiIo{+H~5Pj@Kv zu4wDxh9O(l$KuZ?7U1s2bY}JHoxE@+k9fwDjW1}5Vx4aQFDl-PRM0Dk|OI9!(waQs4S_iZk#n$Vl z8iGslrTt+&R?z>+aj67mzpsQf^RfgkLulL8az2n$3=|qW&Qd;BhOLL(@9$d23xf8P zBu4Z)Fui`KPj5IEFpb-v;@!=QnWs4n@sep{=55y>`&NiSfGzPc2XEe4oZXnueY_pd z4Am-p;>kuN@Su0=U5{)gp!O}#uW_QE)%&wJWo+z9X%*y!-7ZuyjEoDPPxk0(2$Pra zLi%Asb^F9y|F~T)hhONDZ;u4u+yDj*Qp|T!h%(z_)uzC<;rWRT(Z;Z4+#;Oc-T+$i z#wIjLTp|3LGsj9-RUk&9UG+xHVfyv{iqaA-_{bfD&((#qmSf1r7scQJWB=RuK_w`t+)DFMKn@Bv_PabjtqeBVsiONZ&(qh_d$uuq zi-Gp;R=-Ft6vT`f7uuE~fk(cJ2vNrdPuxRgnwC96*&QSm-@!%;)oGecCp^gDuDd~5 z@-!=~UDkZ|{>?w%Pw;m8;5;*a&b9CZ$SJNw zLd+0(`wez)NNugSm9*3g%8<@CyW{nNJ#_1%B0~e%B(#Oo|B@caWh?v}xlGbV82a?eI^(}$9WWq$8os)DE_b=UO|x-d9l zuEWnS3c|GgN%^TtuyxCbr{HT%kQi1u_B2l(&ay{eVy9sQA0}3d+YG{>e^>Igm4FHy z*U;wN@SYy-cEo*rM9BbG7^8G#H;$ut6a2E{oe&UjXnA0G2{pLr2UmBTMgp4Cx>7Y@ zZhan)>z5>DpLTBA$7Yn@^7ih$M+=`P@#l5Dby1Tpp#Wy9r~#gr=l*uPyxq6Q*OMj? zlvFMXk{|Vh^YW)sZr*T)5%K^fl3flUz!LCr2(u5j>aN!_Cv`*EnlyABjoCknDAKNI zTq5@BCn)*%M4ME;k&sFWU5xCkf%jhYQhrsH97lJlwu-f1; zDW9(npx$oZ#YZj?}0N^{jg0I}ETIin0L+2?soHA)tmvRCmFOD_S|l0>rgF3`>$%|i$7V1K*CvKhM zgWHrZx>x_f{Q-~fd@1!>tkiz%APP}(dhpf;>MxFe9sZW;9yb;*5J8XDkjU%n%{lny#M*1vghj{@2P(}|YsI)S>J_cYjI z6&*-qs^~mAiCI@BE3A+cX!+CroO0|D?Z$)8P$j*7imwlTdwugu_VG*Kj= zNDPnpOnbsX(=NBOqt?JhI^=yg+Y7RCb3dIFG=c-3efN7=tl?ARouHLElCbIQ*~XJ# z0DkF%$TJ5FKr_#f`RF-w(B)vuxw}sm7<>5llK2=x*~vVoT>{3ysX;X;QifSq7rBAr z%D}AmYP23R7FLF`iqx8b5jCiK&4LtMt_;*QjCOq0N)UWI+GSlE1MqZyBqCYL2gQgL z!MlP#p#E@7J@EK>t#$H(TrpAvmc(bzAP-{o^SMY3z#Q z^bl#&k#ujhz!&}fOSqkz>e)X;@dj0ydZeF~Y3_a5m-PR?={BBwTBDMpp^~02;k2nQ zh$V1z-d|z`oi$`j^`h+{N^*L0*_a2YU45}OSwjtwLF_h-@#c_wzgsGhMiPop+KrN_ zMvzaHoxhhu3z}WJ9Z@-kz{wX}8F^h63|_M5s(2WIQ(StJlCmMN*oRrZU5|v(@wf-_ z$CW{08T0#62{oWNUwejFPzI*l-^SW`17sH$vNL1WpH$?Jq$(t_10#p$iERhw=2_^8 zJIjq)Md3{!`@ygjK{))()uUUK7|!hIGutb=1d4)o^}3}H&*Fc@uQ{bYh!xDl!#}8E z=1bp;Tla31VIFE6C8Yi-OJ^10yU+i#*X!zgpK9C~3=XMAceAr1A)m!wt9p|&>_$F) zxP06UWDkGHmv3?d&Ii3acHOmvi&qSfa=4i^I3JS%GCo;W^%@7rx{#d_U%>*anwIbK ze>Z}9La|Np&=V184-p~T>MQ_-MTzx->O`=o!D6|BJ7)fM;j#Bf+}yhTo`>B;@y8ee zxz|YkP4(RVl|PHuRi(10hEeZOojtqLZZ*97)8pTtE&t>v@52w{LZCo@=^mBXP_R3U zh`;$V2=<)XJi&tLC!8jy$cWSxpeRfA{g|pZFy2qT*{`7h`L!`wLH)Y$#$abM3%x!- z)+#i6vodIZIHBWuP7k6e4J@dXFzbnKe#;v?WS~3kMBbV@LpXlhaLMb!Grx>TL&p@d8Zv!kjDptlZsvkMEQZpxr{H~j1WrHCEIojG6Jio zqM*a^xpj8)Ym=E$H>iNhsbjQ&xw8^@AN}y|W$RZ3hFBSiK%}4OO|{c%lgcmWKKfmt3LOsB$Go z3qRcKso@iCG=w3eri+2Q)!{)n`I=n_eHhuHSb?P0gldC?EAel2fIfEmh4X|wG$+3Z ztDH84a-Y(z&qtKOdS6)8X*3aJEMt7(`B4nSy?k`j0i0(Tk@M4L>r@I%mVRgFrIN41L zILOpAwc>~o6pJXwiqNw{8>&+G%~s5JgC`v~vUbYDV|mqku8rofuW@aLwFo<~4u+h% zvtJJkp-5YAw+sxpQ+w`NWd;84PP->&B4_>6QF&;!AzBvPtfXlLIWhZ&N1h!^64ilC z&YP4gV#GkW)N9TN8yKCuu$$gU2Wq;C1ijW!!Ks0*MCs+~aFJ1s+uutGsEVR1!nsYM zacx<%R31HCQy{!0n_vX5d5T#-o>Kt&td|=e2%+F{sN5MkeXd!%P^#Jxd+I8JdCgv{ zu$L0R&*+o5>%InT2ud3DN)iG_!#nmnPXKykq z@BoVj4ecH&3Xq5lW>Suvn}^~1ON%>pevOZRue83!BbR_Ahq(AL_3UeRzS}PUJKjOo zXd`wc6B=ZB38pi&X{zlSoS+Gv@3WA1zM-w=}jX_GtlGSZr} zXu#tKwDL@t{SW7P8*>k5(ShRWZETI13-@RKPb>QyT)cawOz(C>Q|R-~7Nv^h29Aia z@evgvSpQscDrXxbMDoU0=$n53I`(e>-9OUA5F!sAHd{aNtv5Wv+h7g4Tx{yYKB{nZ zUGaLZ^rbM8%w-u(!?Z72{^xp z-ZwLt8)7C^2G>aP!%o@EM*naW?39vSxls?ZevadZT@k~3z0ialUg6{-lz$HbAzn7m zrbhih!th`*IY?JU@jtS{9ojUb)6dOPv7(c^^k32ec{vcK&FLnCL z7;2)L_{Fwd8L%yvvFhQnoaKkr%RkPLyjSUq{IIkQF7nz9l%OX`WU0PZ08%Oitt!{KeN%t3 zdcU2TQ09gvXA*y4efBzaQcw?EB0uFSRNDbFfpYwrL_3I1?OFP8`W?!MKZncZnF5^d zmvLy#(}(++)>EsF$O93%X751%d(^l1VNIu$0?sB0g4$DqU2%!u{!Z|>GX$=|mj`Z{ zs&B_U4~{!scaN~c@W(xN?Rabks`A7qOLZ{c3l!=R=Nz#Gxgz(oT8YB&ASouN?2sie z6T6TzpJn@Qe%Rdi+!O00u>XM*nX0%vAl#c%S% z+A=Wua(Q&1=JUfAd=lT;Xvez94~uiGv8Bzr8^f$Y5uq3}UWk^uxi#=4#@_KSkcDTJ2cz~2yhGvnVY`?PQ0Jdk0^hXC zx0QRO;5v=lbG|Fq;5Wg-s@TmBCy2VdZ*&TPmt;-wsoM;oUY%luoWF-<_sogYzaQ~L(BM~x?^I$ACGm8YCXja*=ubf7?c|`k-cCjw1ReN5Uw9q z#D?I8W*NN#t$G14S#vm@qe&GW$2~7S$;Ajo?R&U+k14@B)%>IYVJ?uYZ(ON3Ed-}u z+&JOQP6)!mA8q#D=Ue25U0?}pLR#rI*4BPBR6cic%R3{`(vD95*e(l=NCzd2Fj)`?w|Vr8o)$(+{8&Z%)nH%ubuv$D zm05n+_B6kvg1JnQkf6pM7?ArzQ~uBK)sgUguuqhn;5^+e=PorZ4-#yHll+ zCof~Z^Nb*%ILD;{5eBbfc{8-(;IXVnde_HMJGmXzw!e@Ec7fMc#nIYuN%|~@p{@iR zBhU(W!SHRoYe}qPkd+5vq1N+OUdfLJYhH!#TujMScX`Ul&9>>+oy z6-P1u_ipRsZuihW)a7RhYfYFGfY*IAD{m_;)H?>qi`WW7Bz0qLH0C?n75Qs|uX$=h zLbGX;-dV|6eppG1i#t@WQiDgzAw%C&>@Xp9{#2MLGaS3_S3`G}2J$K?AKkwE=huIK zReso$jSc$k=Vd`9Wt`J&jVHLPywN~Z$iZzF&dgI+wc#Y?7zVD+w;R(ub^pCrX_<=N%CG7O>;05^H`=9 zWe!6J(=2yL7y^=HQ3IMyV zM%S>ifCj6bs<)#O1XX)e1Pn;R*uIqhGX<;*_+dG;OT}ygg+TX7bi-P%rzku>tch&U z7~uo5-7`#s*AW&v`gBmEpS${_C{ zl@!6jw}2m3L$`28M=KM^rd>)T+RHS{51ZfPwG+|y4)we7!xo?0`#Q|Z21FuHU+7LY z0}h7g#ePcaz;2Nyu@SRB@vhwyMY*l4a1w3m;J{`DsNGHiHziFWn?d&IeNILw$Pq9| zB~*nwCfBr1UQ>dV&rG*eG0H=vNXKnuWhvNmfL!;^Wfma0)ZW9=B?~ta8(bV*7#8ru z7CjJO&NICPk~XnW)wJ=>^26?wIs5vh;{(+1!Vl|vw7oC+y)s0Hit0|gVfg&mB*f*l zbYW7GvHuJ=8SL)dpoC`P1XXrc0|iqGFbbi_ov2cSW}cXH&BUrex8}tb{%K~+|L+Vh zB(_3<`51l4ot*+uS>H8w(OeZ2MSLtT2Z+PohuTGFpJ8~FvTlx)q2HmTdqzqg~08#n{po>&;n@x$7!C01tC7(o5%{IKcVq73TH7~Z7YsnZWW=zuz<+3_YX zH%L9@K&iD~35w`0ZsA9YgXVJqR1caCWZP|}k1EhXeZpl&0$*BSrt~-9Bcp^vg_}e^ z35}x2_zzv3j>62vdd4=bIzE&p&ZhQEh!591)37hYa0i7_e? z0@~t)5mkL-(0OQVltw5ID|w=MYP2+<%iw;8;1&YNFdr&UP?iK$iYJ+NdcqJM>*0T; zi3>bSnxFQU6G3a*ksVvB0Vu0A4h!b;LzQ|QhjFbYth*t^*J?}x38%+O!tyzRm&ZSY zTi|hrOz=9G+A>f#N?%T~Nrzuz(*{yGqWFf{5amx zBU7H`hxOriEgwji1#c6RLz$yO;OqAG&N2g`1^lqj`(9i*k^2$#tMkJi?cL2e_Rtx& zUYforhWWo86;BEXZz7Qp({`cf5$1ak=RFy>rxewp!tp?*zgr(lpH%(9omxpK$>zBj zs7DAUod6P*?s7D}`Hu(^#p@49W31V_#{qcy%5o-6_v*BdWI z;ZZx2Bk@ilc$PD0E+JBsJ;_rKunL}|X%ipp@_Wo!75N< z9s2B;zTzxDtk0lM#!)JIs93i)@SGwUQ1n~26YNxgRAr4f_en`+`C<1+_!M@Qvj9K6 z^N3R`=WoLgYwfh`o`sS#tT}PcvA@t7b{^bKxKEaLmLE1jHsifxt^&w6NatjFBOz@| z-+C17yeo@ARx2UAo=ED7~;z(Q55&4H$ecZQ>Fp$&Qv&JlI-%^?w>Kkjm9GC&df zrfznWVtAR%mBS{Bma)Ll6~TZO%>P;MV5Qp~j@d7?yLx@5?_F^)W$XDAaG!jZAC|=- zMg2B2|8K(&>k*8y&*xTvsJep&dUxo7YUQ#@9w|rAv(M6uW>kTux3}`%?GcB9q5>(J zFg7Tcf3rd4ybd5a-IB?jk?^82-eBb?Rw!Pb+>+HI2NyUFqVzWNE#QZhD~=l)ua^Pg zr$S5fp2`5f`wsL`^Xn+`Hks;^x#F|@ur&-DE~W`nE#QZ(C0UuQc~#(-=Z9T|z7s;A zU=5)U&o-i4bm6hT0@*$zQ`nxC{-#9A8m3g#GYeM;gXYv^noO1ewCCJMJ-j3WyE7(V z6G-SoTME+&?Uf2JA@{t8_JbsNPiBZoK2!ykm#YUWkH7+c*mI@Y+k~zPLR0e6+Ex!) z(AYgGUY#_GVyB47T)#;LlGm7TC3K~SjXY@q`WSvzJU^@w1v#lgiZGlTR)||QBJ#`f z!=i;k`&KsDK^N+nNs%1#S}D;sM1 zZVQ48L$zop338SnmV%Men&eJDD*m8kT6G@BEI({n>Fdb0ZK5FW++UW{Bl^qn!!9Q{ zKi>6D15dGT zuzRzpDU;0d!zxtmtKTlc2V*@w7vok)&ho=@Cg&zUA$@pCaF)O3tz`)d^9b<%b=CrlC`cov8WzuvfAQ2k%ew{c`-UCld=)kMP^W zaMPJxfu7#5N|X4uo|y^e|Jex27+cA}*zH62o0mDla2|#B5it#5wcH%5bdvxERjke! z(`qc>hYgMj8}hRmM7_SXqr5vs5@?bnR3jOb!DD=As9#TImLHaFr=@sW^;;C4A69Ma zaUGHXM(C|!?v>#tp5=$VNq4+!>|sAj>A_*PH#Z0u@Wa;a&85}U=Kf{)VGrCp&SF*N z44+;pG!{i#!M#l=2hvGL=nmSrC)!K{m|x!?f6UAXR|{1TR_tyd5x-3S!e$EaZ$FoB z8czvRgFANcWh*Y=hh1km)zs-F2$}b&hr7r)AYD0JMvzGc^Bqs_;6cp)kHquCerl|+ zoLs{J+^i&*n2t%nfL%Q;_t8;QIbpBjidSNPnjiLK#pB(%Nn}4gKkUVD!v4E0UO={& zJ4~I&80tBd9lBH$!S(#k{u%}oKv8#(_YX?U^23&O56T-p#{7TJTECFjAGl!Mi=yjv z?(86Fc?zx6A-jMd7PvyJP6unk)62o8RG!Sx7xjv5r|wmhIl;j4W1>n6_+cMB><%tJ ziv<2=kxiRU6Mi>8EV=tS{WUN8P(Ld_?EbUne_WDA3H3!?%g^P*8Z3l*VG`{4m%4pSgFS$pBa1L0YISjv9m`U_#A5H_mvY!btJ zJf9!7rrp!$#LZo5vWjN%bM@u$=DNrkY8I1OT8IPRB-h%V!|Wv z#}$;pC95!gT@R19)snC4A#vq-QG8twkK*J!|M=rt{`;x_o?mJG6n^#Z;bHVs;PLN! z Date: Tue, 25 Aug 2026 18:01:22 -0700 Subject: [PATCH 21/34] Removed dead code --- mcdc/code_factory/gpu/program_builder.py | 103 ------------------ mcdc/code_factory/gpu/transport/simulation.py | 4 - mcdc/code_factory/numba_layers_generator.py | 12 -- mcdc/object_/simulation.py | 10 -- 4 files changed, 129 deletions(-) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index 851ea8416..669c1308f 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -94,7 +94,6 @@ def _prepare_gpu_program(simulation_dtype, data_size): def forward_declare_gpu_program(simulation_dtype): import harmonize - import mcdc.numba_types as type_ # Get to set the globals @@ -194,99 +193,10 @@ def impl(program): BLOCK_COUNT = 0 -# Compiles gpu kernels and loads in the functions that call into said kernels -def build_gpu_progs(input_deck): - - STRAT = config.args.gpu_strategy - - src_spec = gpu_sources_spec() - - adapt.harm.RuntimeSpec.bind_specs() - - rank = MPI.COMM_WORLD.Get_rank() - device_id = rank % config.args.gpu_share_stride - - if MPI.COMM_WORLD.Get_size() > 1: - MPI.COMM_WORLD.Barrier() - - adapt.harm.RuntimeSpec.load_specs() - - if STRAT == "async": - config.args.gpu_arena_size = config.args.gpu_arena_size // 32 - src_fns = src_spec.async_functions() - pre_fns = pre_spec.async_functions() - else: - src_fns = src_spec.event_functions() - pre_fns = pre_spec.event_functions() - - ARENA_SIZE = config.args.gpu_arena_size - BLOCK_COUNT = config.args.gpu_block_count - - global alloc_state, free_state - alloc_state = src_fns["alloc_state"] - free_state = src_fns["free_state"] - - global src_alloc_program, src_free_program - global src_load_global, src_store_global, src_load_data, src_store_data, src_store_pointer_data - global src_init_program, src_exec_program, src_complete, src_clear_flags - src_alloc_program = src_fns["alloc_program"] - src_free_program = src_fns["free_program"] - src_load_global = src_fns["load_state_device_global"] - src_store_global = src_fns["store_state_device_global"] - src_store_pointer_global = src_fns["store_pointer_state_device_global"] - src_load_data = src_fns["load_state_device_data"] - src_store_data = src_fns["store_state_device_data"] - src_store_pointer_data = src_fns["store_pointer_state_device_data"] - src_init_program = src_fns["init_program"] - src_exec_program = src_fns["exec_program"] - src_complete = src_fns["complete"] - src_clear_flags = src_fns["clear_flags"] - src_set_device = src_fns["set_device"] - - global pre_alloc_program, pre_free_program - global pre_load_global, pre_store_global, pre_load_data, pre_store_data - global pre_init_program, pre_exec_program, pre_complete, pre_clear_flags - pre_alloc_state = pre_fns["alloc_state"] - pre_free_state = pre_fns["free_state"] - pre_alloc_program = pre_fns["alloc_program"] - pre_free_program = pre_fns["free_program"] - pre_load_global = pre_fns["load_state_device_global"] - pre_store_global = pre_fns["store_state_device_global"] - pre_load_data = pre_fns["load_state_device_data"] - pre_store_data = pre_fns["store_state_device_data"] - pre_init_program = pre_fns["init_program"] - pre_exec_program = pre_fns["exec_program"] - pre_complete = pre_fns["complete"] - pre_clear_flags = pre_fns["clear_flags"] - - @njit - def real_setup_gpu(mcdc_array, data_tally): - mcdc = mcdc_array[0] - - src_set_device(device_id) - arena_size = ARENA_SIZE - mcdc["gpu_meta"]["state_pointer"] = adapt.cast_voidptr_to_uintp(alloc_state()) - # src_store_global(mcdc["gpu_meta"]["state_pointer"], mcdc_array[0]) - if config.gpu_state_storage == "separate": - harmonize.memcpy_device_to_host( - simulation, simulation["gpu_meta"]["simulation_pointer"] - ) - harmonize.memcpy_device_to_host( - data, simulation["gpu_meta"]["data_pointer"] - ) - - gpu_module.clear_flags(simulation["gpu_meta"]["program_pointer"]) - - simulation["mpi_work_size"] = full_work_size - - particle_bank_module.set_bank_size(simulation["bank_active"], 0) - - source_closeout(simulation, 1, 1, data) def build_gpu_program(data_size): import harmonize - import mcdc.numba_types as type_ import mcdc.transport.util as util from mcdc.transport.simulation import generate_source_particle, step_particle @@ -405,19 +315,6 @@ def step(program: nb.uintp, particle_input: particle_gpu): clear_flags = src_fns["clear_flags"] set_device = src_fns["set_device"] - # ================================================================================== - # - # ================================================================================== - - """ - global loop_source - loop_source = gpu_loop_source - # - # Overwrite function - for impl in target_rosters["cpu"].values(): - overwrite_func(impl, impl) - """ - alloc_managed_bytes = harmonize.alloc_managed_bytes alloc_device_bytes = harmonize.alloc_device_bytes diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index a95520b2a..027b9cca9 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -80,8 +80,4 @@ def source_loop(seed, simulation, data): particle_bank_module.set_bank_size(simulation["bank_active"], 0) - # ===================================================================== - # Closeout (Moved out of the typical particle loop) - # ===================================================================== - source_closeout(simulation, 1, 1, data) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index f609cd881..c380632c2 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -40,18 +40,6 @@ np.uintp: np.uintp, } -size_map = { - bool: 1, - float: 8, - int: 8, - str: 32, - np.bool_: 1, - np.float64: 8, - np.int64: 8, - np.uint64: 8, - np.str_: 32, -} - bank_names = ["bank_active", "bank_census", "bank_source", "bank_future"] diff --git a/mcdc/object_/simulation.py b/mcdc/object_/simulation.py index e473d829c..bd9cbd06d 100644 --- a/mcdc/object_/simulation.py +++ b/mcdc/object_/simulation.py @@ -436,16 +436,6 @@ def _finalize_compilation(self) -> None: self.bank_source.size[0] = int(settings.source_bank_buffer_ratio * N_work) self.bank_future.size[0] = int(settings.future_bank_buffer_ratio * N_work) - # ================================================================================== - # Platform targeting, adapters, and toggles for portability - # ================================================================================== - - import mcdc.config as config - - # Build GPU program if desired - if config.target == "gpu": - from mcdc.code_factory.gpu.program_builder import build_gpu_program - # Initialize run state derived from the compiled settings self.k_eff = settings.k_init self.cycle_active = ( From fd2bf0a23107bf9de3b798a04b1ce2fdd307a873 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Tue, 25 Aug 2026 19:00:29 -0700 Subject: [PATCH 22/34] Back in Black --- mcdc/code_factory/gpu/program_builder.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/mcdc/code_factory/gpu/program_builder.py b/mcdc/code_factory/gpu/program_builder.py index 669c1308f..1129bfb73 100644 --- a/mcdc/code_factory/gpu/program_builder.py +++ b/mcdc/code_factory/gpu/program_builder.py @@ -193,8 +193,6 @@ def impl(program): BLOCK_COUNT = 0 - - def build_gpu_program(data_size): import harmonize import mcdc.numba_types as type_ From f6dd0d751225e3b9abd00de6a869bae659918adc Mon Sep 17 00:00:00 2001 From: Ilham Variansyah Date: Thu, 27 Aug 2026 13:22:37 +0700 Subject: [PATCH 23/34] add a comment --- mcdc/main.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/mcdc/main.py b/mcdc/main.py index d2c191309..b192c981e 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -137,9 +137,9 @@ def prepare(simulationPy: Simulation): simulation = simulation_container[0] # Pick Python-version RNG if needed - import mcdc.transport.rng as rng - if config.mode == "python": + import mcdc.transport.rng as rng + rng.wrapping_add = rng.wrapping_add_python rng.wrapping_mul = rng.wrapping_mul_python @@ -171,7 +171,7 @@ def prepare(simulationPy: Simulation): # MPI.COMM_WORLD.Barrier() # ================================================================================== - # Setup GPU-Related Data Structures, if Necessary + # Setup GPU-Related Data Structures # ================================================================================== if config.target == "gpu": From faa3e4174a65b57af47d66a5d1519f7bf860a7cd Mon Sep 17 00:00:00 2001 From: Ilham Variansyah Date: Thu, 27 Aug 2026 13:40:31 +0700 Subject: [PATCH 24/34] fix future-to-census bank promotion --- mcdc/transport/particle_bank.py | 3 +- .../slab_isobeam_td_census/answer.h5 | Bin 0 -> 71256 bytes .../slab_isobeam_td_census/input.py | 60 ++++++++++++++++++ 3 files changed, 62 insertions(+), 1 deletion(-) create mode 100644 test/regression/slab_isobeam_td_census/answer.h5 create mode 100644 test/regression/slab_isobeam_td_census/input.py diff --git a/mcdc/transport/particle_bank.py b/mcdc/transport/particle_bank.py index 2213d87d5..c697305e3 100644 --- a/mcdc/transport/particle_bank.py +++ b/mcdc/transport/particle_bank.py @@ -35,6 +35,7 @@ def set_bank_size(bank, value): @njit def add_bank_size(bank, value): + # Perform atomic increment to the bank size; return the initial size return util.atomic_add(bank["size"], 0, value) @@ -147,7 +148,7 @@ def promote_future_particles(program, data): if particle["t"] < next_census_time: bank_census_particle(particle_container, program) - j = add_bank_size(future_bank, -1) + j = add_bank_size(future_bank, -1) - 1 # Consolidate the emptied space in the future bank particle_module.copy( diff --git a/test/regression/slab_isobeam_td_census/answer.h5 b/test/regression/slab_isobeam_td_census/answer.h5 new file mode 100644 index 0000000000000000000000000000000000000000..b6a2d5cb30508b7a4a16880e252fbca4cdf69d3e GIT binary patch literal 71256 zcmeHQ1z1%}*FH2N-OZu9yUwD!yW!9vofe8o2x5Rz5&}x8CQ8Ics9=*)#82Gkf;z*)J(ZdInS!>=cNZgPa^ejG+A-#Q%K4 z9Zt*ypOf%$e7+YhfXzi9l5z2JLTrElG4qE&u;S9`aOKa}v$N6HM_@GZ<>MvpG=vN> zHbjn3Sd_mN0!I4URx=HD!nIlkCtGlShKnNN{Gy_R!ov{JelamYVFA(j^cZwVNRVGN zA}Tg4CMeYJvro?CiQpB!APEp8SY4ch6Cki2i1&9QjAze;*$4{ zaY90DGC8g%5aA1&*WVC`>@U>h$4Iudu+$^Kj1qXyOuC#SI8GxO610dG$R1tygY{jkg$MB^Ff#0!f5B;U^W}_AA=Y4nvD+Ft*KFpIbWT zPlTVLobxAI9HMxH@W2$|13$F<(tTpSpWw%VdVIgfNi9um^A_TYg@pjqKW7e{yUj{` z|D{91EY1o`5elDh{`F!o+at4nXADcO55Hi%YH2aJcuH#}+ zo5M{;N>)w^(_}pScyGgxxA<;J^0`C=A_VR27lH}nZr_L{xb(2_Fwd~qP;WFYj@QTA zCpI-6)vu&hm8;87aVZs zEs7r#0sJ~Je*I@TegWqT)x|ANA`&^ipnyw8;^a2Yuj}Fcu88aFdOv#Z!u3ilqWEq= zgn2(O6aMq-z^q7YoAH*xBT$47{Lu36#{-GPHh2+_861cRL>BSD;4t^zMCK>Qv;(zh$`}p;Ed^%qG;sSVyi)!KXv2qyxJDnVtPm7a`I61r) zr#Fc|0T&;JnQn`JK`~P)f|)l6_=K6a3Ntqze0=8Z!+*Zfn90MBq?q8$fxRUDgt_L; zq~leG&wo}$3gku?a1|uBL5uV7861cRw0?0su>0b^e_xH!7jVg* zf$fJq{{H@Ag~T@Y=j<=|^&uo~PF=3_b<@e<=iaJc*VDr3(W-r255KO6RR6l(&zg(J zR|VIA8t$*R#0ZzoD)MzbCtS9O=-2h|>%1yrU)TH5bHCkhMLs_kvSOY+XIyXQ!4d0% z^sv$5fs-{iUA4-on`9)ERQPeEdvhXL@vj^Axu z-;TsKe?AUH|Aoqr8o`e{+i(>$690Nl%W>J)=2lw%@Ps;isx&_E$Fc>NOpAX3@W*OD zeCj{f&vKC1#?SY!pS_>`)$c5dad^;TVVV8?+41@TdJ z>xJLfKeIGEbKoU@U;p=imvaR9xgZ2iU?Z5{9}5rsI6Q!>A5wO4H-MaM?ZrTeV#3Gu_E2ZjCT85W)q=%|>WnXJzTHqX-=9k$HVJJ#RdFAAHE zE$u(kJ)_VuLE-rFarrZ!yRPu`kMfI*^$YV^fk5EXXTE)%IsAeG{KDeUA+dg*q2a!M z2>;la*eJg*mm9F+^JT&?SMv>ujrNRL5%Ia8WtesZVZPJCf9Dk%8xj)~5EUL9fh~9b zcV5`~^G`({4!_4E#p z4f91uVVaMKjrPM?0@KE5KVQ$!2w~IbwQUBT&-xKj;Q>*8(b1SQaOh*x|KpmoPxJMk z&-(d({%n8#{(00W5{t&KcK`dU;~Nn#7LT86eEYV4{9X6Ir8sP4dC9@<;hyfmKL9^w4Vc+v#yxcE)+L!i? z8(CnR9Ej`R-x>#0Y_a;^JHGs~!2V_pn?DX{;o6^!^B+7Mp|RKydS%e(fM1wjQ~+iH z=QC*=CgtZ02Rmiz0FHB8Pby^P(cSc#Y${jCoNRT)f_CUdFtr3NC)ld0xg}jQ)Fm#q(46 z)xU>_^e=%2{&hB9;@{VP|F{D`7x$(g72f@p~ zjz_ov77zT0|8x2W9tXVq>v*jFN$~jN?_=lB`@Z%3QseQ**2jCnqWrZHnE(8OpU)ZL zo*VG35A{Qb&4i@9hIEGx{sZyf8;l>I^=}F9>!Rx% z;ov8Y{kyG)Q9nSHe(NfnEFSmcO7Vno+1wq{KK6_K@_Q0*}(F z3Om>|R;gmR(Gc#cpS;q#)DnDSWYc$UHwMi$m%Nef5{q{EhwWmr`TX#^JrUp|${j#< z%M-4J^l9tS>4N-99k-QIR)F|;fOND@12c|ICq<4M!%NlD_;MBTMZ5e%b{Wn}Haov8 z8VGn%m3~5LI+nh>1qN8#}do7MZMa zfDp=iT~DI`yPj=HkgPWW@5AoJb#x+&cKMs^A}pB0?o_=3gsQn>tR4o!)0Aa3NOyb4 zJiVb{+RhG+W=alhd#DYGVcjBrfmX0mEPWm0E>Td4A3?;g6^E;vR~2dq8$u=5_AHhd z{zbd|KkdR!O)membtN36H`5e&5CJ4y=Zpl7IDtY8yI$TVN8r|0^?Xle0+LlptyAkQ zU=p#svY=WEE;cCWr0kM{@%tuPglmQ1%&VYHO^JqZdtJ|mQ7XPgyZlAFl-kM2WW`2+ zrT&yLU3@%DK1tfof6xo$Mj4*nBXotXLvq?DH`zisx@^PCOS+JnJtU%W(GWaq6g@un zD?^fG7{S`hvJh!dvNnAUFI>yru+{y%0c=i}p$ojixoDT~v`dBo{{fYaL2$HrS(AcN zG#J6dQ$lKykSo3f#Z9CS(>27?N(CMe@!Vv}!o&dXr+^!j(=gR2TxuMF_H4HFWTi#+Qp?ui?l)C3lJMqldFuP zpuKY7>O=its6z8R?28NlhVdiL;RdGQW1ke}d&C*2oTFVHXGj9QA9>dBB~v(hbdPm% zv@j^Xso{CoC=0Kem>NbeD?zb?)l|kf6*Njz9zRi`2Vio5y)=Sl(JtSxOLo$}NVB96 zkdo6Oe;sEJW*0xmnbbwWD`HjBo3{SIH?$*NyW0s4y&zd97-I`#?_HKkA+6w|YRL<* z(}fP3{e_Ggg77w)Z`+nENs#8hhZyyfgw>-r@*OWK!lv}w+xMmoqKMG+;~PHdL4Nr% z9-RipMZ0{#E`vLa5qFTwpd>u#Zb_>XNZGrW=GFQE@&UEa25WzCA5yc~*XRv0#%_TR z*4n{Me`1!X@SXQ#@m_JFHzw=ZJkY<1mLZYQ^RXr7TDIe zR=oGH>@5CaD00~{!)`icUeduaKqYpa|eWRyA82pL$HGyIC? z!rXd_{!|ue>q`M3=f09uu!<1`)~~A!xg#@+zksnfwJDD>tZ%(${l4wv;(HgI=S{2= zA7F8e2U3a?57UPoAi~(+=%#=ts!1N+!(xOS~ZL;1RWYHMWd z;C5R_hCn9(y5UL63QiB;?C-f)HQ)kUBHXrlY&QdT+BFtRx@Mqcazd`8TpJ3HUC@4u zkbuJvOu|>A=kDK|>e%($sl_2rQSyYqRwkfPz4~#-(z$go={=T9&K=6vL$570|zyqFY$hK*bN502z1fLpc!T70h-Y?YEP z|0s?IrHQp-#ZOJaZICLvnA8SNTx$$!aMyw(S5_-#lN*DY?=x-M{WP#|&zjZIljO7g z#Z0%k#IAuKc6V8Zg^+Q=hTWyRUq;W3bGsU@v>l6P1BMaN60g3eznxv`2LrWJo1$TS zIHmt$Z6J))4N#U`M1x01*eNleBJwKfFVJu)E=7U_WI+={yEo{J<^OBd6_xG`jO!HmQl*njU zc}c2&*Mks9xXEsnL~jqPNZ5Lb;=MqCOik^=O<%~NHr}LWtN@o&3b)*DQ3j&E0X=~n zYiLPn@4w$|4QaZTuCXY62-giwiyYR6M)&2O_dV6YJ2bKgB_jmmeJ{`4(d7iYW=PX( zPeMI6%G#*YJ&KadmA7p-x5)%TpHtPNQ{V?zf`gb)B3>ZK?~rwaRnQBsGVpYbGv5FnWW8y;o%8YB#8`xawu% zzz8D6<`(=}hH$W8Ros4UJ1|t6JQoP6F!7F$*zLX7YsD=a3)k~{m3~bP*X@~ zzHi9~PZ9DDE?Ca3qY!O=A27m82)lK+GSozK19896o)jlus8#OX5HZ9wJI)DS)LRvM z8Z(YPjpXs!`{b9kOJr>$#q;ICAoAwS&Z6)LxE~tXcp}gTg>mGODEI(_IkFum&1gavHxOla5p-!?y{sZgmM#7N!03t-R}JbQp_UIS1*xt+e8UeDD(xyTagg@ z_)}Z?O`2J}n+0l|+l=Vo9{TR9`|8Y)pVYg`^`H>Q_4}pgY$Bc2%V*EbGkL)VuVp7V zPx1dccJZ!jdsXNV1!oAFMVJMBKt%7+#t%#pVBvP|c=1D17-D0oJmSj)Y>`B+ z<2>|meCg_=8+P%6v{p*di`r+i{i3<)w74P>J51g>ndj&GXg2+4eXt*KRC?cI8wLK( zcPN(CM}x|NZ2{iOeo)EG+(UcE2MFFhX7i#kgX0z_pJqooLNCS2?44Ftz-(dI`Q8hl zPCEX660HN6-XYlPbVnZ)TYbv7+m+xpXD#V+J8d}YeUJNy89nf@q2F1HYXO5Le|Jo+ z6d#e3_L&W_EcD~JWVND_0x<3oyay!}UJ z@qZKX_%sg(I}jwUW~VEAIGc_y=Vy^`ueh^H9t{VbzEhnQbm2goyQA|Rdl-~&Y_ZDG z4uT3-i_4|!t-<*Kn}EKeHIVadB+rv@0>52XKJ+is2D9f(uj0*gAl2rGf`g9@JPPMa zqTg-=O8dF{m%Nn$SM5rE+dO4(d9`k(e7*o|ynLvrg9qSlv*yPukNCh$r}Ry9ff!J> zA=7So&D|%yh2Q*n0bZY-bEz4HJUa{tLe81~zMsx6%)4*UrYc1MqEBNo7qc%ell(D5 zjzJ*gSM5=a&0Gc)>Qv4xCcd!#WqvTM)`fD>sDO$yHlR;sL|nPo1~fNrB|aCc0#7ZwN=BmK-FpRNa`F4L^;mbsj{hphS`x#snvck|g_9q)vB!+H?bZa@ zXusO%AZF0IGqvePi3AW7?^aA?Bbgny@Ho$xx4YiimP|iFX>6{I2_0rzpdWNe?(SO? zcDT?+z#1s@^VlVI7mqeiU@*iSlOsxR^#;Ydro!gSb};b4vGt^H5GcLX+R$Yg1lgiA zX)DjT1FO_4^tv1B5JJR|e8$EK601LIUC6Y6;|CvXrQNRyD3NKNGzwj4mn1*J_`n`q zI8L`%Rw%*Qr1ypaVroEiFOB{2L4Ig^V%V9oTnZWv88h0~$-zA%(*%Ps5>S`4XrX4{ z1NHuOtfNv>-}GbN+Djj_%8e2Cz2n3o?KZbwg3DD zfAIGxzFXYo3!gTRp5Y1bglftu<0(N4n4;|{uA~eCHLrmFHKyKhVuK7BA=(~DH4;w* z?-hjR^t3mt15BWYEM@G4r4g8IeqclQL0N;! zV(eOfjKT}93Oc^*p)!F?VK{Wg-W&=tqCTD98UW^_bluryZjdH)a*x&qV|dR3JFGvd zLgU`2@eMV)&~k*%C+mPdyb9BCVt9;#xb@eI+;-}KLdh|AsdimZv=BLPsY4E^HAwDdWaae^$gIF8i3{ia=f zxY;BwOi=)fASb1w_1ruFKi<6;Q5R{@=LKSho$0BJKi+(d|6*@^kT4pyJzRhJ-5wvv zW07LsbQKNx>MfIgudU%Dg{|y?N_n`MQ4=Jv!V%USOGqLQ3WO6Pdp@mgwFbl%Hv6eH zYM`zl+rH9A2iWc$c-p)|7t*UmFKjT>gS;~_?Od z8d=H!$tHeY&)Z@kDd~Hu!bB9-CEpw{6&8eBsWR;~tNG#jD&}!@8Y=J(+rg8Y&I&Fe z%O>7Lj(pQD)00++inp{7e!8bPq?cjAIEHdMecjxi7u=XmIP)=8dg0B=(YBCV zm8JJq-XE^fUX*0E@`d*VQ(Mx@Y=PAcO^FUOhsrDWhS_8_;4~MhvEEGwXlri1eW<}7 z1d@+N)HWJHU%p~fl#L8*P2;8cn4u39^5vP5Q5wL0clGmpUL$yQ#kOxMKn}K*W%>p? zp&-;S-i`1bKyB6;y?fnq5QW&SwN_0C9&jY?Bq?VF^2`fwK48W}`IAdFH(sZQPs^^U zJa|hFPh*};vDPqxP!9F(H*(+perUUQR&ZnpBji}`R4zp?oWCU|sVDUM^1=;~bWJAy z|NVV`vTApt=71yA7c)z*xnvJ(Ph8?Mu=%l)86cr-i~apU_9oUVO8^h_v^1t3F|Y z>}z{+PjoF=&|h93KP~JR!wahd7{!m--T!8J`0?-WmMPiP$c^?6phsV)zc$wbl35w= z8_RftTeA!z$lVi$Ldq$Fx0r%&y16~6kQFp_7KGTQXh0RUJ2H{o6dn>spyTM>pr-Qb z;mYM!!0%IVD*Lb&h&gnhmoAWo$~8@`f+E_WcQFaIa=8rfNXV__Qqclw7sgJC0D3^D ztkpW8V*nKOsoI`8q98JMWHNk(3cS*M`G`4$8K~bXZ?MT>1g~U&lUog}!0R`ZmwAR3 zH@>mKa=a$i#34lt?KUvmWo1N?WVYv z!*Y<#+^N%gPzi)wFRY1m2OxGqnubk?fb2`DUYh|%2;tqHfHXA#n)RIm1 zuvc&vF9O85ItcLgoH;IAiIo{+H~5Pj@Kv zu4wDxh9O(l$KuZ?7U1s2bY}JHoxE@+k9fwDjW1}5Vx4aQFDl-PRM0Dk|OI9!(waQs4S_iZk#n$Vl z8iGslrTt+&R?z>+aj67mzpsQf^RfgkLulL8az2n$3=|qW&Qd;BhOLL(@9$d23xf8P zBu4Z)Fui`KPj5IEFpb-v;@!=QnWs4n@sep{=55y>`&NiSfGzPc2XEe4oZXnueY_pd z4Am-p;>kuN@Su0=U5{)gp!O}#uW_QE)%&wJWo+z9X%*y!-7ZuyjEoDPPxk0(2$Pra zLi%Asb^F9y|F~T)hhONDZ;u4u+yDj*Qp|T!h%(z_)uzC<;rWRT(Z;Z4+#;Oc-T+$i z#wIjLTp|3LGsj9-RUk&9UG+xHVfyv{iqaA-_{bfD&((#qmSf1r7scQJWB=RuK_w`t+)DFMKn@Bv_PabjtqeBVsiONZ&(qh_d$uuq zi-Gp;R=-Ft6vT`f7uuE~fk(cJ2vNrdPuxRgnwC96*&QSm-@!%;)oGecCp^gDuDd~5 z@-!=~UDkZ|{>?w%Pw;m8;5;*a&b9CZ$SJNw zLd+0(`wez)NNugSm9*3g%8<@CyW{nNJ#_1%B0~e%B(#Oo|B@caWh?v}xlGbV82a?eI^(}$9WWq$8os)DE_b=UO|x-d9l zuEWnS3c|GgN%^TtuyxCbr{HT%kQi1u_B2l(&ay{eVy9sQA0}3d+YG{>e^>Igm4FHy z*U;wN@SYy-cEo*rM9BbG7^8G#H;$ut6a2E{oe&UjXnA0G2{pLr2UmBTMgp4Cx>7Y@ zZhan)>z5>DpLTBA$7Yn@^7ih$M+=`P@#l5Dby1Tpp#Wy9r~#gr=l*uPyxq6Q*OMj? zlvFMXk{|Vh^YW)sZr*T)5%K^fl3flUz!LCr2(u5j>aN!_Cv`*EnlyABjoCknDAKNI zTq5@BCn)*%M4ME;k&sFWU5xCkf%jhYQhrsH97lJlwu-f1; zDW9(npx$oZ#YZj?}0N^{jg0I}ETIin0L+2?soHA)tmvRCmFOD_S|l0>rgF3`>$%|i$7V1K*CvKhM zgWHrZx>x_f{Q-~fd@1!>tkiz%APP}(dhpf;>MxFe9sZW;9yb;*5J8XDkjU%n%{lny#M*1vghj{@2P(}|YsI)S>J_cYjI z6&*-qs^~mAiCI@BE3A+cX!+CroO0|D?Z$)8P$j*7imwlTdwugu_VG*Kj= zNDPnpOnbsX(=NBOqt?JhI^=yg+Y7RCb3dIFG=c-3efN7=tl?ARouHLElCbIQ*~XJ# z0DkF%$TJ5FKr_#f`RF-w(B)vuxw}sm7<>5llK2=x*~vVoT>{3ysX;X;QifSq7rBAr z%D}AmYP23R7FLF`iqx8b5jCiK&4LtMt_;*QjCOq0N)UWI+GSlE1MqZyBqCYL2gQgL z!MlP#p#E@7J@EK>t#$H(TrpAvmc(bzAP-{o^SMY3z#Q z^bl#&k#ujhz!&}fOSqkz>e)X;@dj0ydZeF~Y3_a5m-PR?={BBwTBDMpp^~02;k2nQ zh$V1z-d|z`oi$`j^`h+{N^*L0*_a2YU45}OSwjtwLF_h-@#c_wzgsGhMiPop+KrN_ zMvzaHoxhhu3z}WJ9Z@-kz{wX}8F^h63|_M5s(2WIQ(StJlCmMN*oRrZU5|v(@wf-_ z$CW{08T0#62{oWNUwejFPzI*l-^SW`17sH$vNL1WpH$?Jq$(t_10#p$iERhw=2_^8 zJIjq)Md3{!`@ygjK{))()uUUK7|!hIGutb=1d4)o^}3}H&*Fc@uQ{bYh!xDl!#}8E z=1bp;Tla31VIFE6C8Yi-OJ^10yU+i#*X!zgpK9C~3=XMAceAr1A)m!wt9p|&>_$F) zxP06UWDkGHmv3?d&Ii3acHOmvi&qSfa=4i^I3JS%GCo;W^%@7rx{#d_U%>*anwIbK ze>Z}9La|Np&=V184-p~T>MQ_-MTzx->O`=o!D6|BJ7)fM;j#Bf+}yhTo`>B;@y8ee zxz|YkP4(RVl|PHuRi(10hEeZOojtqLZZ*97)8pTtE&t>v@52w{LZCo@=^mBXP_R3U zh`;$V2=<)XJi&tLC!8jy$cWSxpeRfA{g|pZFy2qT*{`7h`L!`wLH)Y$#$abM3%x!- z)+#i6vodIZIHBWuP7k6e4J@dXFzbnKe#;v?WS~3kMBbV@LpXlhaLMb!Grx>TL&p@d8Zv!kjDptlZsvkMEQZpxr{H~j1WrHCEIojG6Jio zqM*a^xpj8)Ym=E$H>iNhsbjQ&xw8^@AN}y|W$RZ3hFBSiK%}4OO|{c%lgcmWKKfmt3LOsB$Go z3qRcKso@iCG=w3eri+2Q)!{)n`I=n_eHhuHSb?P0gldC?EAel2fIfEmh4X|wG$+3Z ztDH84a-Y(z&qtKOdS6)8X*3aJEMt7(`B4nSy?k`j0i0(Tk@M4L>r@I%mVRgFrIN41L zILOpAwc>~o6pJXwiqNw{8>&+G%~s5JgC`v~vUbYDV|mqku8rofuW@aLwFo<~4u+h% zvtJJkp-5YAw+sxpQ+w`NWd;84PP->&B4_>6QF&;!AzBvPtfXlLIWhZ&N1h!^64ilC z&YP4gV#GkW)N9TN8yKCuu$$gU2Wq;C1ijW!!Ks0*MCs+~aFJ1s+uutGsEVR1!nsYM zacx<%R31HCQy{!0n_vX5d5T#-o>Kt&td|=e2%+F{sN5MkeXd!%P^#Jxd+I8JdCgv{ zu$L0R&*+o5>%InT2ud3DN)iG_!#nmnPXKykq z@BoVj4ecH&3Xq5lW>Suvn}^~1ON%>pevOZRue83!BbR_Ahq(AL_3UeRzS}PUJKjOo zXd`wc6B=ZB38pi&X{zlSoS+Gv@3WA1zM-w=}jX_GtlGSZr} zXu#tKwDL@t{SW7P8*>k5(ShRWZETI13-@RKPb>QyT)cawOz(C>Q|R-~7Nv^h29Aia z@evgvSpQscDrXxbMDoU0=$n53I`(e>-9OUA5F!sAHd{aNtv5Wv+h7g4Tx{yYKB{nZ zUGaLZ^rbM8%w-u(!?Z72{^xp z-ZwLt8)7C^2G>aP!%o@EM*naW?39vSxls?ZevadZT@k~3z0ialUg6{-lz$HbAzn7m zrbhih!th`*IY?JU@jtS{9ojUb)6dOPv7(c^^k32ec{vcK&FLnCL z7;2)L_{Fwd8L%yvvFhQnoaKkr%RkPLyjSUq{IIkQF7nz9l%OX`WU0PZ08%Oitt!{KeN%t3 zdcU2TQ09gvXA*y4efBzaQcw?EB0uFSRNDbFfpYwrL_3I1?OFP8`W?!MKZncZnF5^d zmvLy#(}(++)>EsF$O93%X751%d(^l1VNIu$0?sB0g4$DqU2%!u{!Z|>GX$=|mj`Z{ zs&B_U4~{!scaN~c@W(xN?Rabks`A7qOLZ{c3l!=R=Nz#Gxgz(oT8YB&ASouN?2sie z6T6TzpJn@Qe%Rdi+!O00u>XM*nX0%vAl#c%S% z+A=Wua(Q&1=JUfAd=lT;Xvez94~uiGv8Bzr8^f$Y5uq3}UWk^uxi#=4#@_KSkcDTJ2cz~2yhGvnVY`?PQ0Jdk0^hXC zx0QRO;5v=lbG|Fq;5Wg-s@TmBCy2VdZ*&TPmt;-wsoM;oUY%luoWF-<_sogYzaQ~L(BM~x?^I$ACGm8YCXja*=ubf7?c|`k-cCjw1ReN5Uw9q z#D?I8W*NN#t$G14S#vm@qe&GW$2~7S$;Ajo?R&U+k14@B)%>IYVJ?uYZ(ON3Ed-}u z+&JOQP6)!mA8q#D=Ue25U0?}pLR#rI*4BPBR6cic%R3{`(vD95*e(l=NCzd2Fj)`?w|Vr8o)$(+{8&Z%)nH%ubuv$D zm05n+_B6kvg1JnQkf6pM7?ArzQ~uBK)sgUguuqhn;5^+e=PorZ4-#yHll+ zCof~Z^Nb*%ILD;{5eBbfc{8-(;IXVnde_HMJGmXzw!e@Ec7fMc#nIYuN%|~@p{@iR zBhU(W!SHRoYe}qPkd+5vq1N+OUdfLJYhH!#TujMScX`Ul&9>>+oy z6-P1u_ipRsZuihW)a7RhYfYFGfY*IAD{m_;)H?>qi`WW7Bz0qLH0C?n75Qs|uX$=h zLbGX;-dV|6eppG1i#t@WQiDgzAw%C&>@Xp9{#2MLGaS3_S3`G}2J$K?AKkwE=huIK zReso$jSc$k=Vd`9Wt`J&jVHLPywN~Z$iZzF&dgI+wc#Y?7zVD+w;R(ub^pCrX_<=N%CG7O>;05^H`=9 zWe!6J(=2yL7y^=HQ3IMyV zM%S>ifCj6bs<)#O1XX)e1Pn;R*uIqhGX<;*_+dG;OT}ygg+TX7bi-P%rzku>tch&U z7~uo5-7`#s*AW&v`gBmEpS${_C{ zl@!6jw}2m3L$`28M=KM^rd>)T+RHS{51ZfPwG+|y4)we7!xo?0`#Q|Z21FuHU+7LY z0}h7g#ePcaz;2Nyu@SRB@vhwyMY*l4a1w3m;J{`DsNGHiHziFWn?d&IeNILw$Pq9| zB~*nwCfBr1UQ>dV&rG*eG0H=vNXKnuWhvNmfL!;^Wfma0)ZW9=B?~ta8(bV*7#8ru z7CjJO&NICPk~XnW)wJ=>^26?wIs5vh;{(+1!Vl|vw7oC+y)s0Hit0|gVfg&mB*f*l zbYW7GvHuJ=8SL)dpoC`P1XXrc0|iqGFbbi_ov2cSW}cXH&BUrex8}tb{%K~+|L+Vh zB(_3<`51l4ot*+uS>H8w(OeZ2MSLtT2Z+PohuTGFpJ8~FvTlx)q2HmTdqzqg~08#n{po>&;n@x$7!C01tC7(o5%{IKcVq73TH7~Z7YsnZWW=zuz<+3_YX zH%L9@K&iD~35w`0ZsA9YgXVJqR1caCWZP|}k1EhXeZpl&0$*BSrt~-9Bcp^vg_}e^ z35}x2_zzv3j>62vdd4=bIzE&p&ZhQEh!591)37hYa0i7_e? z0@~t)5mkL-(0OQVltw5ID|w=MYP2+<%iw;8;1&YNFdr&UP?iK$iYJ+NdcqJM>*0T; zi3>bSnxFQU6G3a*ksVvB0Vu0A4h!b;LzQ|QhjFbYth*t^*J?}x38%+O!tyzRm&ZSY zTi|hrOz=9G+A>f#N?%T~Nrzuz(*{yGqWFf{5amx zBU7H`hxOriEgwji1#c6RLz$yO;OqAG&N2g`1^lqj`(9i*k^2$#tMkJi?cL2e_Rtx& zUYforhWWo86;BEXZz7Qp({`cf5$1ak=RFy>rxewp!tp?*zgr(lpH%(9omxpK$>zBj zs7DAUod6P*?s7D}`Hu(^#p@49W31V_#{qcy%5o-6_v*BdWI z;ZZx2Bk@ilc$PD0E+JBsJ;_rKunL}|X%ipp@_Wo!75N< z9s2B;zTzxDtk0lM#!)JIs93i)@SGwUQ1n~26YNxgRAr4f_en`+`C<1+_!M@Qvj9K6 z^N3R`=WoLgYwfh`o`sS#tT}PcvA@t7b{^bKxKEaLmLE1jHsifxt^&w6NatjFBOz@| z-+C17yeo@ARx2UAo=ED7~;z(Q55&4H$ecZQ>Fp$&Qv&JlI-%^?w>Kkjm9GC&df zrfznWVtAR%mBS{Bma)Ll6~TZO%>P;MV5Qp~j@d7?yLx@5?_F^)W$XDAaG!jZAC|=- zMg2B2|8K(&>k*8y&*xTvsJep&dUxo7YUQ#@9w|rAv(M6uW>kTux3}`%?GcB9q5>(J zFg7Tcf3rd4ybd5a-IB?jk?^82-eBb?Rw!Pb+>+HI2NyUFqVzWNE#QZhD~=l)ua^Pg zr$S5fp2`5f`wsL`^Xn+`Hks;^x#F|@ur&-DE~W`nE#QZ(C0UuQc~#(-=Z9T|z7s;A zU=5)U&o-i4bm6hT0@*$zQ`nxC{-#9A8m3g#GYeM;gXYv^noO1ewCCJMJ-j3WyE7(V z6G-SoTME+&?Uf2JA@{t8_JbsNPiBZoK2!ykm#YUWkH7+c*mI@Y+k~zPLR0e6+Ex!) z(AYgGUY#_GVyB47T)#;LlGm7TC3K~SjXY@q`WSvzJU^@w1v#lgiZGlTR)||QBJ#`f z!=i;k`&KsDK^N+nNs%1#S}D;sM1 zZVQ48L$zop338SnmV%Men&eJDD*m8kT6G@BEI({n>Fdb0ZK5FW++UW{Bl^qn!!9Q{ zKi>6D15dGT zuzRzpDU;0d!zxtmtKTlc2V*@w7vok)&ho=@Cg&zUA$@pCaF)O3tz`)d^9b<%b=CrlC`cov8WzuvfAQ2k%ew{c`-UCld=)kMP^W zaMPJxfu7#5N|X4uo|y^e|Jex27+cA}*zH62o0mDla2|#B5it#5wcH%5bdvxERjke! z(`qc>hYgMj8}hRmM7_SXqr5vs5@?bnR3jOb!DD=As9#TImLHaFr=@sW^;;C4A69Ma zaUGHXM(C|!?v>#tp5=$VNq4+!>|sAj>A_*PH#Z0u@Wa;a&85}U=Kf{)VGrCp&SF*N z44+;pG!{i#!M#l=2hvGL=nmSrC)!K{m|x!?f6UAXR|{1TR_tyd5x-3S!e$EaZ$FoB z8czvRgFANcWh*Y=hh1km)zs-F2$}b&hr7r)AYD0JMvzGc^Bqs_;6cp)kHquCerl|+ zoLs{J+^i&*n2t%nfL%Q;_t8;QIbpBjidSNPnjiLK#pB(%Nn}4gKkUVD!v4E0UO={& zJ4~I&80tBd9lBH$!S(#k{u%}oKv8#(_YX?U^23&O56T-p#{7TJTECFjAGl!Mi=yjv z?(86Fc?zx6A-jMd7PvyJP6unk)62o8RG!Sx7xjv5r|wmhIl;j4W1>n6_+cMB><%tJ ziv<2=kxiRU6Mi>8EV=tS{WUN8P(Ld_?EbUne_WDA3H3!?%g^P*8Z3l*VG`{4m%4pSgFS$pBa1L0YISjv9m`U_#A5H_mvY!btJ zJf9!7rrp!$#LZo5vWjN%bM@u$=DNrkY8I1OT8IPRB-h%V!|Wv z#}$;pC95!gT@R19)snC4A#vq-QG8twkK*J!|M=rt{`;x_o?mJG6n^#Z;bHVs;PLN! z Date: Thu, 27 Aug 2026 13:49:36 +0700 Subject: [PATCH 25/34] only import config once at the beginning --- mcdc/main.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/mcdc/main.py b/mcdc/main.py index b192c981e..816a66eac 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -192,8 +192,6 @@ def prepare(simulationPy: Simulation): def finalize(simulation): - import mcdc.config as config - # GPU teardowns if needed if config.target == "gpu": from mcdc.code_factory.gpu.program_builder import teardown_gpu_program From 31870141afe03846b4e7116c19b20bd1ab47c552 Mon Sep 17 00:00:00 2001 From: Ilham Variansyah Date: Thu, 27 Aug 2026 14:08:12 +0700 Subject: [PATCH 26/34] remove unused import --- mcdc/transport/geometry/surface/interface.py | 1 - 1 file changed, 1 deletion(-) diff --git a/mcdc/transport/geometry/surface/interface.py b/mcdc/transport/geometry/surface/interface.py index 04e00e47d..62e31604b 100644 --- a/mcdc/transport/geometry/surface/interface.py +++ b/mcdc/transport/geometry/surface/interface.py @@ -21,7 +21,6 @@ import mcdc.transport.geometry.surface.torus_y as torus_y import mcdc.transport.geometry.surface.torus_z as torus_z import mcdc.transport.geometry.surface.torus as torus -import mcdc.transport.util as util from mcdc.constant import ( COINCIDENCE_TOLERANCE, From ac8fdad71542c597d9808ed8383e800229b2e20a Mon Sep 17 00:00:00 2001 From: Ilham Variansyah Date: Thu, 27 Aug 2026 15:16:28 +0700 Subject: [PATCH 27/34] organize torus-specific solver --- mcdc/transport/geometry/surface/torus.py | 4 ++-- .../geometry/{root_solve.py => surface/torus_root_solver.py} | 0 mcdc/transport/geometry/surface/torus_x.py | 4 ++-- mcdc/transport/geometry/surface/torus_y.py | 4 ++-- mcdc/transport/geometry/surface/torus_z.py | 4 ++-- 5 files changed, 8 insertions(+), 8 deletions(-) rename mcdc/transport/geometry/{root_solve.py => surface/torus_root_solver.py} (100%) diff --git a/mcdc/transport/geometry/surface/torus.py b/mcdc/transport/geometry/surface/torus.py index 7a2e7fbf5..239c16d84 100644 --- a/mcdc/transport/geometry/surface/torus.py +++ b/mcdc/transport/geometry/surface/torus.py @@ -30,7 +30,7 @@ from numba import njit import mcdc.transport.util as util -import mcdc.transport.geometry.root_solve as root_solve +import mcdc.transport.geometry.surface.torus_root_solver as torus_root_solver from mcdc.constant import ( COINCIDENCE_TOLERANCE, @@ -174,7 +174,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients, roots) + torus_root_solver.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/root_solve.py b/mcdc/transport/geometry/surface/torus_root_solver.py similarity index 100% rename from mcdc/transport/geometry/root_solve.py rename to mcdc/transport/geometry/surface/torus_root_solver.py diff --git a/mcdc/transport/geometry/surface/torus_x.py b/mcdc/transport/geometry/surface/torus_x.py index 3f9152d8f..f238b6be5 100644 --- a/mcdc/transport/geometry/surface/torus_x.py +++ b/mcdc/transport/geometry/surface/torus_x.py @@ -16,7 +16,7 @@ from numba import njit import mcdc.transport.util as util -import mcdc.transport.geometry.root_solve as root_solve +import mcdc.transport.geometry.surface.torus_root_solver as torus_root_solver from mcdc.constant import ( COINCIDENCE_TOLERANCE, @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients, roots) + torus_root_solver.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_y.py b/mcdc/transport/geometry/surface/torus_y.py index 86c1892d4..ed7594e4d 100644 --- a/mcdc/transport/geometry/surface/torus_y.py +++ b/mcdc/transport/geometry/surface/torus_y.py @@ -16,7 +16,7 @@ from numba import njit import mcdc.transport.util as util -import mcdc.transport.geometry.root_solve as root_solve +import mcdc.transport.geometry.surface.torus_root_solver as torus_root_solver from mcdc.constant import ( COINCIDENCE_TOLERANCE, @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients, roots) + torus_root_solver.solve_quartic(coefficients, roots) min_t = INF diff --git a/mcdc/transport/geometry/surface/torus_z.py b/mcdc/transport/geometry/surface/torus_z.py index d65c0fd33..a64ab1e5a 100644 --- a/mcdc/transport/geometry/surface/torus_z.py +++ b/mcdc/transport/geometry/surface/torus_z.py @@ -16,7 +16,7 @@ from numba import njit import mcdc.transport.util as util -import mcdc.transport.geometry.root_solve as root_solve +import mcdc.transport.geometry.surface.torus_root_solver as torus_root_solver from mcdc.constant import ( COINCIDENCE_TOLERANCE, @@ -199,7 +199,7 @@ def get_distance(particle_container, surface): coefficients[3] = a3 + 0.0j coefficients[4] = a4 + 0.0j roots = util.local_array(4, np.complex128) - root_solve.solve_quartic(coefficients, roots) + torus_root_solver.solve_quartic(coefficients, roots) min_t = INF From 4f855929d410c908eafafe0965329d553df03d2e Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 27 Aug 2026 17:38:37 -0700 Subject: [PATCH 28/34] Early exit for GPU source loop --- mcdc/code_factory/gpu/transport/simulation.py | 3 +++ mcdc/transport/simulation.py | 3 --- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 027b9cca9..11750ae29 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -24,6 +24,9 @@ def source_loop(seed, simulation, data): settings = simulation["settings"] full_work_size = simulation["mpi_work_size"] + + if full_work_size == 0: + return if settings["gpu_strategy"] == GPU_STRATEGY_ASYNC: phase_size = 1000000000 diff --git a/mcdc/transport/simulation.py b/mcdc/transport/simulation.py index 5c7fd652d..80f3055d2 100644 --- a/mcdc/transport/simulation.py +++ b/mcdc/transport/simulation.py @@ -256,9 +256,6 @@ def source_closeout(simulation, idx_work, N_prog, data): tally_module.closeout.accumulate(simulation, data) # Progress printout - if simulation["mpi_work_size"] == 0: - return - percent = (idx_work + 1.0) / simulation["mpi_work_size"] if simulation["settings"]["use_progress_bar"] and int(percent * 100.0) > N_prog: N_prog += 1 From d08b63f2120439c55c1add36c5a286120f05fb74 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Thu, 27 Aug 2026 17:39:53 -0700 Subject: [PATCH 29/34] Removed `mcdc/code_factory/jit.py` --- mcdc/code_factory/jit.py | 54 ---------------------------------------- 1 file changed, 54 deletions(-) delete mode 100644 mcdc/code_factory/jit.py diff --git a/mcdc/code_factory/jit.py b/mcdc/code_factory/jit.py deleted file mode 100644 index bf36e963d..000000000 --- a/mcdc/code_factory/jit.py +++ /dev/null @@ -1,54 +0,0 @@ -import numba as nb -from inspect import getfullargspec - -ARG_CHECK = True - -njit_trace_template = """ -def arg_check_{name}({args}): - return fn({args}) - -@nb.extending.overload(arg_check_{name}) -def arg_check_{name}_overload({args}): - arg_list = [{args}] - arg_name_list = [{arg_names}] - for idx in range(len(arg_list)): - arg = arg_list[idx] - arg_name = arg_name_list[idx] - if isinstance(arg,nb.types.Record): - raise RuntimeError(f"Argument {{arg_name}} has a Record type. Records should be passed in an array.") - elif isinstance(arg,nb.types.Optional): - raise RuntimeError(f"Argument {{arg_name}} has a Record type. Records should be passed in an array.") - return fn -""" - - -def wrap_with_check(fn, njit_fn): - nb.extending.register_jitable(fn) - arg_names = getfullargspec(fn).args - args_str = ",".join(arg_names) - arg_names_str = ",".join(f'"{n}"' for n in arg_names) - name = fn.__name__ - print(f"wrapping {name}") - gns = {"nb": nb, "fn": fn, "njit_fn": njit_fn} - lns = {} - code = njit_trace_template.format(args=args_str, arg_names=arg_names_str, name=name) - exec(code, gns, lns) - return lns[f"arg_check_{name}"] - - -def njit(*args, **kwargs): - - if (len(args) == 1) and (len(kwargs) == 0): - if not ARG_CHECK: - return njit(args[0]) - fn = args[0] - njit_fn = nb.njit(args[0]) - return wrap_with_check(fn, njit_fn) - - else: - if not ARG_CHECK: - return nb.njit(*args, **kwargs) - - def wrapper(fn): - njit_fn = nb.njit(*args, **kwargs)(fn) - return wrap_with_check(fn, njit_fn) From ce618bc00dbcbba0c77755719c6fb84e3663e3a3 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 28 Aug 2026 01:51:31 -0700 Subject: [PATCH 30/34] Restored usage of `mcdc_get` for accessing vector data --- mcdc/code_factory/numba_layers_generator.py | 5 ++- mcdc/transport/data.py | 20 ++------- mcdc/transport/distribution.py | 43 +++++--------------- mcdc/transport/geometry/surface/interface.py | 19 ++------- mcdc/transport/mesh/structured.py | 22 ++-------- mcdc/transport/physics/neutron/multigroup.py | 25 +++--------- mcdc/transport/physics/neutron/native.py | 10 ++--- mcdc/transport/physics/util.py | 10 +---- mcdc/transport/source.py | 15 ++----- mcdc/transport/tally/filter.py | 16 ++------ 10 files changed, 44 insertions(+), 141 deletions(-) diff --git a/mcdc/code_factory/numba_layers_generator.py b/mcdc/code_factory/numba_layers_generator.py index c380632c2..673da8db8 100644 --- a/mcdc/code_factory/numba_layers_generator.py +++ b/mcdc/code_factory/numba_layers_generator.py @@ -1228,19 +1228,20 @@ def _accessor_1d_last( def _accessor_chunk(object_name, attribute_name, setter=False): - text = f"@njit\n" if setter: + text = f"@njit\n" text += ( f"def {attribute_name}_chunk(start, length, {object_name}, data, value):\n" ) else: + text = f"@array_return(nb.types.float64)\n" text += f"def {attribute_name}_chunk(start, length, {object_name}, data):\n" text += f' start += {object_name}["{attribute_name}_offset"]\n' text += f" end = start + length\n" if setter: text += f" data[start:end] = value\n\n\n" else: - text += f" return data[start:end]\n\n\n" + text += f" return array_result(data[start:end])\n\n\n" return text diff --git a/mcdc/transport/data.py b/mcdc/transport/data.py index f0894d1c3..4e4ba387e 100644 --- a/mcdc/transport/data.py +++ b/mcdc/transport/data.py @@ -39,10 +39,7 @@ def evaluate_data(x, data_, simulation, data): @njit def evaluate_table(x, table, data): - offset = table["x_offset"] - length = table["x_length"] - grid = data[offset : offset + length] - # Above is equivalent to: grid = mcdc_get.table_data.x_all(table, data) + grid = mcdc_get.table_data.x_all(table, data) idx = find_bin(x, grid) x1 = grid[idx] @@ -69,15 +66,9 @@ def evaluate_table(x, table, data): @njit def get_table_interpolation_law(idx, table, data) -> int: """Return the interpolation law for interval [idx, idx + 1].""" - offset = table["interpolation_boundaries_offset"] - length = table["interpolation_boundaries_length"] - boundaries = data[offset : offset + length] - # Above is equivalent to: boundaries = mcdc_get.table_data.interpolation_boundaries_all(table, data) + boundaries = mcdc_get.table_data.interpolation_boundaries_all(table, data) - offset = table["interpolations_offset"] - length = table["interpolations_length"] - interpolations = data[offset : offset + length] - # Above is equivalent to: interpolations = mcdc_get.table_data.interpolations_all(table, data) + interpolations = mcdc_get.table_data.interpolations_all(table, data) # Boundaries are exclusive upper point indices. upper_point = idx + 1 @@ -91,10 +82,7 @@ def get_table_interpolation_law(idx, table, data) -> int: @njit def evaluate_polynomial(x, polynomial, data): - offset = polynomial["coefficients_offset"] - length = polynomial["coefficients_length"] - coeffs = data[offset : offset + length] - # Above is equivalent to: coeffs = mcdc_get.polynomial_data.coefficients_all(polynomial, data) + coeffs = mcdc_get.polynomial_data.coefficients_all(polynomial, data) total = 0.0 for i in range(len(coeffs)): diff --git a/mcdc/transport/distribution.py b/mcdc/transport/distribution.py index 190f3ae04..d32d0f2d5 100644 --- a/mcdc/transport/distribution.py +++ b/mcdc/transport/distribution.py @@ -247,11 +247,7 @@ def invert_tabulated_segment(xi, c0, v0, v1, p0, p1, interpolation): @njit def sample_pmf(pmf, rng_state, data): xi = rng.lcg(rng_state) - - offset = pmf["cmf_offset"] - length = pmf["cmf_length"] - cmf = data[offset : offset + length] - # Above is equivalent to: cmf = mcdc_get.pmf_distribution.cmf_all(pmf, data) + cmf = mcdc_get.pmf_distribution.cmf_all(pmf, data) idx = find_bin(xi, cmf) return mcdc_get.pmf_distribution.value(idx, pmf, data) @@ -297,10 +293,7 @@ def _sample_multi_table(E, rng_state, multi_table, simulation, data, scale): """Sample from a multi-table distribution.""" # Get the grid - offset = multi_table["grid_offset"] - length = multi_table["grid_length"] - grid = data[offset : offset + length] - # Above is equivalent to: grid = mcdc_get.multi_table_distribution.grid_all(multi_table, data) + grid = mcdc_get.multi_table_distribution.grid_all(multi_table, data) # Helper flag for scaling later use_next_table = False @@ -440,10 +433,7 @@ def sample_evaporation(E, rng_state, evaporation, simulation, data): @njit def sample_kalbach_mann(E, rng_state, kalbach_mann, data): - offset = kalbach_mann["energy_offset"] - length = kalbach_mann["energy_length"] - grid = data[offset : offset + length] - # Above is equivalent to: grid = mcdc_get.kalbach_mann_distribution.energy_all(kalbach_mann, data) + grid = mcdc_get.kalbach_mann_distribution.energy_all(kalbach_mann, data) # Random numbers xi1 = rng.lcg(rng_state) @@ -493,9 +483,7 @@ def sample_kalbach_mann(E, rng_state, kalbach_mann, data): size = end - start # The CDF - offset = kalbach_mann["cdf_offset"] - cdf = data[start + offset : start + offset + size] - # Above is equivalent to: cdf = mcdc_get.kalbach_mann_distribution.cdf_chunk(start, size, kalbach_mann, data) + cdf = mcdc_get.kalbach_mann_distribution.cdf_chunk(start, size, kalbach_mann, data) # Sample bin index idx = find_bin(xi2, cdf) @@ -544,10 +532,7 @@ def sample_kalbach_mann(E, rng_state, kalbach_mann, data): @njit def sample_tabulated_energy_angle(E, rng_state, table, data): - offset = table["energy_offset"] - length = table["energy_length"] - grid = data[offset : offset + length] - # Above is equivalent to: grid = mcdc_get.tabulated_energy_angle_distribution.energy_all(table, data) + grid = mcdc_get.tabulated_energy_angle_distribution.energy_all(table, data) # Random numbers xi1 = rng.lcg(rng_state) @@ -600,12 +585,9 @@ def sample_tabulated_energy_angle(E, rng_state, table, data): size = end - start # The CDF - offset = table["cdf_offset"] - cdf = data[start + offset : start + offset + size] - # Above is equivalent to: - # cdf = mcdc_get.tabulated_energy_angle_distribution.cdf_chunk( - # start, size, table, data - # ) + cdf = mcdc_get.tabulated_energy_angle_distribution.cdf_chunk( + start, size, table, data + ) # Sample bin index idx = find_bin(xi2, cdf) @@ -653,12 +635,9 @@ def sample_tabulated_energy_angle(E, rng_state, table, data): size = end - start # The CDF - offset = table["cosine_cdf_offset"] - cdf = data[start + offset : start + offset + size] - # Above is equivalent to: - # cdf = mcdc_get.tabulated_energy_angle_distribution.cosine_cdf_chunk( - # start, size, table, data - # ) + cdf = mcdc_get.tabulated_energy_angle_distribution.cosine_cdf_chunk( + start, size, table, data + ) # Sample bin index idx = find_bin(xi3, cdf) diff --git a/mcdc/transport/geometry/surface/interface.py b/mcdc/transport/geometry/surface/interface.py index 62e31604b..02f7836ce 100644 --- a/mcdc/transport/geometry/surface/interface.py +++ b/mcdc/transport/geometry/surface/interface.py @@ -394,12 +394,7 @@ def _get_move_idx(t, surface, data): """ Get moving interval index wrt the given time """ - time_grid = data[ - surface["move_time_grid_offset"] : ( - surface["move_time_grid_offset"] + surface["N_move_grid"] - ) - ] - # Above is equivalent to: time_grid = mcdc_get.surface.move_time_grid_all(surface, data) + time_grid = mcdc_get.surface.move_time_grid_all(surface, data) tolerance = COINCIDENCE_TOLERANCE_TIME go_lower = False idx = find_bin_with_rules(t, time_grid, tolerance, go_lower) @@ -419,14 +414,10 @@ def _translate_particle_position(particle_container, surface, idx, data): particle = particle_container[0] # Surface move translations - start = surface["move_translations_offset"] + idx * 3 - trans_0 = data[start : start + 3] - # Above is equivalent to: trans_0 = mcdc_get.surface.move_translations_vector(idx, surface, data) + trans_0 = mcdc_get.surface.move_translations_vector(idx, surface, data) # Surface move velocities - start = surface["move_velocities_offset"] + idx * 3 - V = data[start : start + 3] - # Above is equivalent to: V = mcdc_get.surface.move_velocities_vector(idx, surface, data) + V = mcdc_get.surface.move_velocities_vector(idx, surface, data) # Surface move time grid time_0 = mcdc_get.surface.move_time_grid(idx, surface, data) @@ -446,9 +437,7 @@ def _translate_particle_direction(particle_container, speed, surface, idx, data) particle = particle_container[0] # Surface move velocities - start = surface["move_velocities_offset"] + idx * 3 - V = data[start : start + 3] - # Above is equivalent to: V = mcdc_get.surface.move_velocities_vector(idx, surface, data) + V = mcdc_get.surface.move_velocities_vector(idx, surface, data) # Translate the particle particle["ux"] -= V[0] / speed diff --git a/mcdc/transport/mesh/structured.py b/mcdc/transport/mesh/structured.py index 7e3ad989f..3a41810b0 100644 --- a/mcdc/transport/mesh/structured.py +++ b/mcdc/transport/mesh/structured.py @@ -2,6 +2,7 @@ #### +import mcdc.mcdc_get as mcdc_get from mcdc.constant import COINCIDENCE_TOLERANCE, INF from mcdc.transport.util import find_bin_with_rules @@ -21,24 +22,9 @@ def get_indices(particle_container, structured_mesh, data): uy = particle["uy"] uz = particle["uz"] - grid_x = data[ - structured_mesh["x_offset"] : ( - structured_mesh["x_offset"] + structured_mesh["x_length"] - ) - ] - # Above is equivalent to: grid_x = mcdc_get.structured_mesh.x_all(structured_mesh, data) - grid_y = data[ - structured_mesh["y_offset"] : ( - structured_mesh["y_offset"] + structured_mesh["y_length"] - ) - ] - # Above is equivalent to: grid_y = mcdc_get.structured_structured_mesh.y_all(structured_mesh, data) - grid_z = data[ - structured_mesh["z_offset"] : ( - structured_mesh["z_offset"] + structured_mesh["z_length"] - ) - ] - # Above is equivalent to: grid_z = mcdc_get.structured_structured_mesh.z_all(structured_mesh, data) + grid_x = mcdc_get.structured_mesh.x_all(structured_mesh, data) + grid_y = mcdc_get.structured_mesh.y_all(structured_mesh, data) + grid_z = mcdc_get.structured_mesh.z_all(structured_mesh, data) tolerance = COINCIDENCE_TOLERANCE ux_go_lower = ux < 0.0 diff --git a/mcdc/transport/physics/neutron/multigroup.py b/mcdc/transport/physics/neutron/multigroup.py index e1d48d495..aa0433918 100644 --- a/mcdc/transport/physics/neutron/multigroup.py +++ b/mcdc/transport/physics/neutron/multigroup.py @@ -251,10 +251,7 @@ def scattering(particle_container, program, data): particle_new["uz"] = uz_new # Get outgoing spectrum - stride = mgxs["G"] - start = mgxs["chi_s_offset"] + group * stride - chi_s = data[start : start + stride] - # Above is equivalent to: chi_s = mcdc_get.neutron_multigroup_data.chi_s_vector(group, mgxs, data) + chi_s = mcdc_get.neutron_multigroup_data.chi_s_vector(group, mgxs, data) # Sample outgoing energy xi = rng.lcg(particle_container_new) @@ -311,10 +308,7 @@ def fission(particle_container, program, data): nu = mcdc_get.neutron_multigroup_data.nu_f(group, mgxs, data) nu_p = mcdc_get.neutron_multigroup_data.nu_p(group, mgxs, data) if J > 0: - stride = mgxs["J"] - start = mgxs["nu_d_offset"] + group * stride - nu_d = data[start : start + stride] - # Above is equivalent to: nu_d = mcdc_get.neutron_multigroup_data.nu_d_vector(group, mgxs, data) + nu_d = mcdc_get.neutron_multigroup_data.nu_d_vector(group, mgxs, data) # Get number of secondaries N = int( @@ -346,10 +340,7 @@ def fission(particle_container, program, data): total = nu_p if xi < total: prompt = True - stride = mgxs["G"] - start = mgxs["chi_p_offset"] + group * stride - spectrum = data[start : start + stride] - # Above is equivalent to: spectrum = mcdc_get.neutron_multigroup_data.chi_p_vector(group, mgxs, data) + spectrum = mcdc_get.neutron_multigroup_data.chi_p_vector(group, mgxs, data) else: prompt = False @@ -357,13 +348,9 @@ def fission(particle_container, program, data): for j in range(J): total += nu_d[j] if xi < total: - stride = mgxs["G"] - start = mgxs["chi_d_offset"] + j * stride - spectrum = data[start : start + stride] - # Above is equivalent to: - # spectrum = mcdc_get.neutron_multigroup_data.chi_d_vector( - # j, mgxs, data - # ) + spectrum = mcdc_get.neutron_multigroup_data.chi_d_vector( + j, mgxs, data + ) decay = mcdc_get.neutron_multigroup_data.decay_rate(j, mgxs, data) break diff --git a/mcdc/transport/physics/neutron/native.py b/mcdc/transport/physics/neutron/native.py index bb69d7e5e..6c938a2e9 100644 --- a/mcdc/transport/physics/neutron/native.py +++ b/mcdc/transport/physics/neutron/native.py @@ -654,13 +654,9 @@ def sample_inelastic_scattering( ) spectrum = simulation["distributions"][ID] else: - offset = inelastic_scattering["spectrum_probability_grid_offset"] - length = inelastic_scattering["spectrum_probability_grid_length"] - probability_grid = data[offset : offset + length] - # Above is equivalent to: - # probability_grid = mcdc_get.neutron_inelastic_scattering_reaction.spectrum_probability_grid_all( - # inelastic_scattering, data - # ) + probability_grid = mcdc_get.neutron_inelastic_scattering_reaction.spectrum_probability_grid_all( + inelastic_scattering, data + ) probability_idx = find_bin(E, probability_grid) xi = rng.lcg(particle_container_new) total = 0.0 diff --git a/mcdc/transport/physics/util.py b/mcdc/transport/physics/util.py index 12db31f87..487b9091e 100644 --- a/mcdc/transport/physics/util.py +++ b/mcdc/transport/physics/util.py @@ -11,10 +11,7 @@ @njit def evaluate_neutron_xs_energy_grid(e, nuclide, data): - offset = nuclide["neutron_xs_energy_grid_offset"] - length = nuclide["neutron_xs_energy_grid_length"] - energy_grid = data[offset : offset + length] - # Above is equivalent to: energy_grid = mcdc_get.nuclide.neutron_xs_energy_grid_all(nuclide, data) + energy_grid = mcdc_get.nuclide.neutron_xs_energy_grid_all(nuclide, data) idx = find_bin(e, energy_grid) e0 = energy_grid[idx] @@ -24,10 +21,7 @@ def evaluate_neutron_xs_energy_grid(e, nuclide, data): @njit def evaluate_electron_xs_energy_grid(e, element, data): - offset = element["electron_xs_energy_grid_offset"] - length = element["electron_xs_energy_grid_length"] - energy_grid = data[offset : offset + length] - # Above is equivalent to: energy_grid = mcdc_get.element.electron_xs_energy_grid_all(element, data) + energy_grid = mcdc_get.element.electron_xs_energy_grid_all(element, data) idx = find_bin(e, energy_grid) e0 = energy_grid[idx] e1 = energy_grid[idx + 1] diff --git a/mcdc/transport/source.py b/mcdc/transport/source.py index c89cf4ef2..12c6a7a76 100644 --- a/mcdc/transport/source.py +++ b/mcdc/transport/source.py @@ -108,12 +108,7 @@ def source_particle(particle_container, seed, simulation, data): # Motion translation if source["moving"]: # Get moving interval index wrt the given time - time_grid = data[ - source["move_time_grid_offset"] : ( - source["move_time_grid_offset"] + source["N_move_grid"] - ) - ] - # Above is equivalent to: time_grid = mcdc_get.source.move_time_grid_all(source, data) + time_grid = mcdc_get.source.move_time_grid_all(source, data) tolerance = COINCIDENCE_TOLERANCE_TIME go_lower = False @@ -124,14 +119,10 @@ def source_particle(particle_container, seed, simulation, data): idx += 1 # Source move translations - start = source["move_translations_offset"] + idx * 3 - trans_0 = data[start : start + 3] - # Above is equivalent to: trans_0 = mcdc_get.source.move_translations_vector(idx, source, data) + trans_0 = mcdc_get.source.move_translations_vector(idx, source, data) # Source move velocities - start = source["move_velocities_offset"] + idx * 3 - V = data[start : start + 3] - # Above is equivalent to: V = mcdc_get.source.move_velocities_vector(idx, source, data) + V = mcdc_get.source.move_velocities_vector(idx, source, data) # Source move time grid time_0 = mcdc_get.source.move_time_grid(idx, source, data) diff --git a/mcdc/transport/tally/filter.py b/mcdc/transport/tally/filter.py index d976b1bee..824ab5a12 100644 --- a/mcdc/transport/tally/filter.py +++ b/mcdc/transport/tally/filter.py @@ -56,10 +56,8 @@ def get_direction_index(particle_container, tally, data): tolerance = COINCIDENCE_TOLERANCE_DIRECTION - grid_mu = data[tally["mu_offset"] : (tally["mu_offset"] + tally["mu_length"])] - # Above is equivalent to: grid_mu = mcdc_get.tally.mu_all(tally, data) - grid_azi = data[tally["azi_offset"] : (tally["azi_offset"] + tally["azi_length"])] - # Above is equivalent to: grid_azi = mcdc_get.tally.azi_all(tally, data) + grid_mu = mcdc_get.tally.mu_all(tally, data) + grid_azi = mcdc_get.tally.azi_all(tally, data) i_mu = find_bin_with_tolerance(mu, grid_mu, tolerance) i_azi = find_bin_with_tolerance(azi, grid_azi, tolerance) @@ -73,10 +71,7 @@ def get_energy_index(particle_container, tally, data): E = particle["E"] tolerance = COINCIDENCE_TOLERANCE_ENERGY - grid_energy = data[ - tally["energy_offset"] : (tally["energy_offset"] + tally["energy_length"]) - ] - # Above is equivalent to: grid_energy = mcdc_get.tally.energy_all(tally, data) + grid_energy = mcdc_get.tally.energy_all(tally, data) return find_bin_with_tolerance(E, grid_energy, tolerance) @@ -88,10 +83,7 @@ def get_time_index(particle_container, tally, data): # Particle properties time = particle["t"] - grid_time = data[ - tally["time_offset"] : (tally["time_offset"] + tally["time_length"]) - ] - # Above is equivalent to: grid_time = mcdc_get.tally.time_all(tally, data) + grid_time = mcdc_get.tally.time_all(tally, data) tolerance = COINCIDENCE_TOLERANCE_TIME go_lower = False From 5a2f2189548f00a3a04bc8a1f3fc6ea77249cd03 Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 28 Aug 2026 01:52:06 -0700 Subject: [PATCH 31/34] Back in Black --- mcdc/code_factory/gpu/transport/simulation.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mcdc/code_factory/gpu/transport/simulation.py b/mcdc/code_factory/gpu/transport/simulation.py index 11750ae29..b9ba444c7 100644 --- a/mcdc/code_factory/gpu/transport/simulation.py +++ b/mcdc/code_factory/gpu/transport/simulation.py @@ -24,7 +24,7 @@ def source_loop(seed, simulation, data): settings = simulation["settings"] full_work_size = simulation["mpi_work_size"] - + if full_work_size == 0: return From 479bd40875965d92075e1c816b112b73b2faf02b Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 28 Aug 2026 02:27:46 -0700 Subject: [PATCH 32/34] Removed extra call to `setup_gpu_program` --- mcdc/main.py | 9 --------- 1 file changed, 9 deletions(-) diff --git a/mcdc/main.py b/mcdc/main.py index 816a66eac..0b2e85fa2 100644 --- a/mcdc/main.py +++ b/mcdc/main.py @@ -170,15 +170,6 @@ def prepare(simulationPy: Simulation): # simulation["bank_source"]["size"] = N_local # MPI.COMM_WORLD.Barrier() - # ================================================================================== - # Setup GPU-Related Data Structures - # ================================================================================== - - if config.target == "gpu": - from mcdc.code_factory.gpu.program_builder import setup_gpu_program - - setup_gpu_program(simulation_container, data) - # ================================================================================== # Finalize # ================================================================================== From 28b1a7abdd56af9c07d46657acd4b419e4c0c30d Mon Sep 17 00:00:00 2001 From: Braxton Cuneo Date: Fri, 28 Aug 2026 04:04:15 -0700 Subject: [PATCH 33/34] Reverted local array sizing in `_check_cell` to generated literal --- mcdc/transport/geometry/interface.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mcdc/transport/geometry/interface.py b/mcdc/transport/geometry/interface.py index 0050383af..b904dbe3f 100644 --- a/mcdc/transport/geometry/interface.py +++ b/mcdc/transport/geometry/interface.py @@ -369,7 +369,7 @@ def _check_cell(particle_container, speed, cell, simulation, data): return True # Create local value array - value = util.local_array(100, np.bool_) + value = util.local_array(literals.rpn_evaluation_buffer_size(), np.bool_) N_value = 0 # March forward through RPN tokens From 923151e3aaf9cf641eaa89bf69242976dc9c01dc Mon Sep 17 00:00:00 2001 From: Ilham Variansyah Date: Sat, 29 Aug 2026 11:08:35 +0700 Subject: [PATCH 34/34] update CHANGELOG --- CHANGELOG.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 083037ee9..76c232824 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -13,6 +13,7 @@ The format is based on [Keep a Changelog](https://keepachangelog.com/en/2.0.0/), ### Changed +- Update GPU transport support for the current MC/DC data model and Harmonize runtime, including GPU-compatible state access, particle-bank operations, array accessors, and torus intersections, from [@braxtoncuneo]. - Show previously published documentation versions in the documentation version switcher, from [@ilhamv] ### Deprecated @@ -203,3 +204,4 @@ The pre-refactor implementation remains available in the `cement` branch as a re [@gunnarrl]: https://github.com/gunnarrl [@Talen-Ayers]: https://github.com/Talen-Ayers [@steps-re]: https://github.com/steps-re +[@braxtoncuneo]: https://github.com/braxtoncuneo