954 lines
46 KiB
Diff
954 lines
46 KiB
Diff
From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001
|
||
From: Luminol Contributors <luminol@example.com>
|
||
Date: Tue, 22 Jun 2077 00:00:00 +0800
|
||
Subject: [PATCH] AcceleratedRecoiling: GPU/FFM entity collision acceleration
|
||
|
||
Ports AcceleratedRecoiling (https://github.com/wiyuka0/AcceleratedRecoiling)
|
||
into Luminol as a built-in server-side optimization. Replaces the O(N²) vanilla
|
||
entity-push scan with a spatial-hash broadphase that is offloaded to:
|
||
GPU – OpenCL kernel via JOCL (highest throughput)
|
||
FFM – C++ native library via java.lang.foreign (Java 21+, very high)
|
||
JNI – C++ native library via JNI (all JDKs, very high)
|
||
SIMD – Java Vector API (requires --add-modules jdk.incubator.vector)
|
||
Java – Pure-Java spatial hash (no extra requirements)
|
||
|
||
The backend is chosen automatically at startup (GPU → FFM → JNI → SIMD → Java)
|
||
or overridden with -Dluminol.collision.backend=<GPU|FFM|JNI|SIMD|JAVA|VANILLA>.
|
||
|
||
Only activates when the local entity density exceeds `density-threshold` (default 16),
|
||
so lightly-loaded regions pay no overhead.
|
||
|
||
Licensed under MIT (AcceleratedRecoiling) and GPL-3.0 (Luminol).
|
||
|
||
diff --git a/me/earthme/luminol/config/modules/optimizations/AcceleratedCollisionConfig.java b/me/earthme/luminol/config/modules/optimizations/AcceleratedCollisionConfig.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/config/modules/optimizations/AcceleratedCollisionConfig.java
|
||
@@ -0,0 +1,62 @@
|
||
+package me.earthme.luminol.config.modules.optimizations;
|
||
+
|
||
+import me.earthme.luminol.config.ILuminolConfig;
|
||
+import org.spongepowered.configurate.objectmapping.ConfigSerializable;
|
||
+import org.spongepowered.configurate.objectmapping.meta.Comment;
|
||
+
|
||
+@ConfigSerializable
|
||
+public class AcceleratedCollisionConfig implements ILuminolConfig {
|
||
+
|
||
+ @Comment("Enable the accelerated entity-collision broadphase.\n"
|
||
+ + "When entity density in a region exceeds density-threshold,\n"
|
||
+ + "the vanilla O(N²) push scan is replaced by a spatial-hash algorithm\n"
|
||
+ + "executed on the best available backend (GPU > FFM > JNI > SIMD > Java).")
|
||
+ public boolean enabled = true;
|
||
+
|
||
+ @Comment("Override backend selection. Valid values (case-insensitive):\n"
|
||
+ + " AUTO – select best available automatically (default)\n"
|
||
+ + " GPU – OpenCL GPU kernel (requires GPU drivers / JOCL)\n"
|
||
+ + " FFM – C++ via java.lang.foreign (requires Java 21+)\n"
|
||
+ + " JNI – C++ via JNI\n"
|
||
+ + " SIMD – Java Vector API (requires --add-modules jdk.incubator.vector)\n"
|
||
+ + " JAVA – Pure-Java spatial hash\n"
|
||
+ + " VANILLA – disable entirely (use vanilla logic)")
|
||
+ public String backend = "AUTO";
|
||
+
|
||
+ @Comment("Maximum collision interactions processed per entity per tick.")
|
||
+ public int maxCollision = 32;
|
||
+
|
||
+ @Comment("Spatial-hash grid cell size (in blocks).")
|
||
+ public int gridSize = 1;
|
||
+
|
||
+ @Comment("Density smoothing window (number of past ticks averaged).")
|
||
+ public int densityWindow = 4;
|
||
+
|
||
+ @Comment("Minimum surrounding-entity count before the accelerated path is activated.\n"
|
||
+ + "Lower values activate sooner; raise if performance decreases on sparse servers.")
|
||
+ public int densityThreshold = 16;
|
||
+
|
||
+ @Comment("Maximum native worker threads (FFM/JNI backends only).\n"
|
||
+ + "0 = use all available logical CPUs.")
|
||
+ public int maxThreads = 1;
|
||
+
|
||
+ @Override
|
||
+ public String getConfigurationPath() {
|
||
+ return "optimizations.accelerated-collision";
|
||
+ }
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/INativeCollisionBackend.java b/me/earthme/luminol/collision/INativeCollisionBackend.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..bbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbbb
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/INativeCollisionBackend.java
|
||
@@ -0,0 +1,32 @@
|
||
+package me.earthme.luminol.collision;
|
||
+
|
||
+/**
|
||
+ * Common contract for all accelerated-collision backends.
|
||
+ * Implementations must be thread-safe: push() may be called concurrently
|
||
+ * from multiple Folia region threads.
|
||
+ */
|
||
+public interface INativeCollisionBackend {
|
||
+ /** Human-readable backend name shown in server logs. */
|
||
+ String getName();
|
||
+
|
||
+ /**
|
||
+ * Initialize resources (load native libs, compile GPU kernels, …).
|
||
+ * Called once during server startup. Throws on failure so the selector
|
||
+ * can fall back to the next backend.
|
||
+ */
|
||
+ void initialize() throws Exception;
|
||
+
|
||
+ /**
|
||
+ * Compute AABB overlaps for the given entity set.
|
||
+ *
|
||
+ * @param aabb flat double array: [minX,minY,minZ,maxX,maxY,maxZ] × count
|
||
+ * @param resultSizeOut single-element array; filled with the number of collision pairs
|
||
+ * @return collision pairs, or null if count is below threshold
|
||
+ */
|
||
+ CollisionPushResult push(double[] aabb, int[] resultSizeOut);
|
||
+
|
||
+ /** Re-read config values into native context (called after /luminol reload). */
|
||
+ void applyConfig();
|
||
+
|
||
+ /** Release all native resources. Called on server shutdown. */
|
||
+ void destroy();
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/CollisionPushResult.java b/me/earthme/luminol/collision/CollisionPushResult.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..cccccccccccccccccccccccccccccccccccccccc
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/CollisionPushResult.java
|
||
@@ -0,0 +1,16 @@
|
||
+package me.earthme.luminol.collision;
|
||
+
|
||
+/** Zero-copy view of a collision-pair result buffer produced by a backend. */
|
||
+public interface CollisionPushResult {
|
||
+ int getA(int index);
|
||
+ int getB(int index);
|
||
+ float getDensity(int index);
|
||
+
|
||
+ /** Bulk-copy entity-A indices into dest[0..length). */
|
||
+ void copyATo(int[] dest, int length);
|
||
+ /** Bulk-copy entity-B indices into dest[0..length). */
|
||
+ void copyBTo(int[] dest, int length);
|
||
+ /** Bulk-copy density values into dest[0..length). */
|
||
+ void copyDensityTo(float[] dest, int length);
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/backend/JavaCollisionBackend.java b/me/earthme/luminol/collision/backend/JavaCollisionBackend.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..dddddddddddddddddddddddddddddddddddddddd
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/backend/JavaCollisionBackend.java
|
||
@@ -0,0 +1,131 @@
|
||
+package me.earthme.luminol.collision.backend;
|
||
+
|
||
+import me.earthme.luminol.collision.CollisionPushResult;
|
||
+import me.earthme.luminol.collision.INativeCollisionBackend;
|
||
+import me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig;
|
||
+
|
||
+import java.util.ArrayList;
|
||
+import java.util.List;
|
||
+
|
||
+/**
|
||
+ * Pure-Java spatial-hash broadphase. No native dependencies.
|
||
+ * Falls back gracefully on all platforms.
|
||
+ */
|
||
+public class JavaCollisionBackend implements INativeCollisionBackend {
|
||
+
|
||
+ private static final int TABLE_SIZE = 262139;
|
||
+ private static final double CELL_SIZE = 2.5;
|
||
+
|
||
+ @Override public String getName() { return "Java"; }
|
||
+ @Override public void initialize() { /* no-op */ }
|
||
+ @Override public void applyConfig() { /* reads from AcceleratedCollisionConfig on each push */ }
|
||
+ @Override public void destroy() { /* no-op */ }
|
||
+
|
||
+ private record Pair(int a, int b) {}
|
||
+
|
||
+ @Override
|
||
+ public CollisionPushResult push(double[] aabb, int[] resultSizeOut) {
|
||
+ int count = aabb.length / 6;
|
||
+ if (count == 0) { resultSizeOut[0] = 0; return null; }
|
||
+
|
||
+ AcceleratedCollisionConfig cfg = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ int maxCollision = cfg.maxCollision;
|
||
+ double cellSize = Math.max(1, cfg.gridSize) * CELL_SIZE;
|
||
+
|
||
+ // Build hash-table: cell → list of entity indices
|
||
+ @SuppressWarnings("unchecked")
|
||
+ List<Integer>[] table = new List[TABLE_SIZE];
|
||
+
|
||
+ int[] centerCell = new int[count];
|
||
+ for (int i = 0; i < count; i++) {
|
||
+ double cx = (aabb[i*6] + aabb[i*6+3]) * 0.5;
|
||
+ double cy = (aabb[i*6+1] + aabb[i*6+4]) * 0.5;
|
||
+ double cz = (aabb[i*6+2] + aabb[i*6+5]) * 0.5;
|
||
+ int gx = (int) Math.floor(cx / cellSize);
|
||
+ int gy = (int) Math.floor(cy / cellSize);
|
||
+ int gz = (int) Math.floor(cz / cellSize);
|
||
+ int h = hash(gx, gy, gz);
|
||
+ centerCell[i] = h;
|
||
+ if (table[h] == null) table[h] = new ArrayList<>();
|
||
+ table[h].add(i);
|
||
+ }
|
||
+
|
||
+ // Detect overlaps
|
||
+ List<Pair> pairs = new ArrayList<>();
|
||
+ float[] density = new float[count];
|
||
+
|
||
+ for (int i = 0; i < count; i++) {
|
||
+ double minX = aabb[i*6], minY = aabb[i*6+1], minZ = aabb[i*6+2];
|
||
+ double maxX = aabb[i*6+3], maxY = aabb[i*6+4], maxZ = aabb[i*6+5];
|
||
+ double cx = (minX + maxX) * 0.5;
|
||
+ double cy = (minY + maxY) * 0.5;
|
||
+ double cz = (minZ + maxZ) * 0.5;
|
||
+ int gx = (int) Math.floor(cx / cellSize);
|
||
+ int gy = (int) Math.floor(cy / cellSize);
|
||
+ int gz = (int) Math.floor(cz / cellSize);
|
||
+
|
||
+ int localDensity = 0;
|
||
+ for (int dz = -1; dz <= 1; dz++) for (int dy = -1; dy <= 1; dy++) for (int dx = -1; dx <= 1; dx++) {
|
||
+ int h = hash(gx+dx, gy+dy, gz+dz);
|
||
+ List<Integer> cell = table[h];
|
||
+ if (cell == null) continue;
|
||
+ for (int j : cell) {
|
||
+ if (j == i) continue;
|
||
+ double oMinX = aabb[j*6], oMaxX = aabb[j*6+3];
|
||
+ if (minX > oMaxX || maxX < oMinX) continue;
|
||
+ double oMinY = aabb[j*6+1], oMaxY = aabb[j*6+4];
|
||
+ if (minY > oMaxY || maxY < oMinY) continue;
|
||
+ double oMinZ = aabb[j*6+2], oMaxZ = aabb[j*6+5];
|
||
+ if (minZ > oMaxZ || maxZ < oMinZ) continue;
|
||
+ localDensity++;
|
||
+ if (i < j && pairs.size() < count * maxCollision) {
|
||
+ pairs.add(new Pair(i, j));
|
||
+ }
|
||
+ }
|
||
+ }
|
||
+ density[i] = localDensity;
|
||
+ }
|
||
+
|
||
+ int pairCount = pairs.size();
|
||
+ resultSizeOut[0] = pairCount;
|
||
+ if (pairCount == 0) return null;
|
||
+
|
||
+ int[] arrA = new int[pairCount];
|
||
+ int[] arrB = new int[pairCount];
|
||
+ for (int k = 0; k < pairCount; k++) {
|
||
+ arrA[k] = pairs.get(k).a();
|
||
+ arrB[k] = pairs.get(k).b();
|
||
+ }
|
||
+ float[] finalDensity = density;
|
||
+
|
||
+ return new CollisionPushResult() {
|
||
+ @Override public int getA(int idx) { return arrA[idx]; }
|
||
+ @Override public int getB(int idx) { return arrB[idx]; }
|
||
+ @Override public float getDensity(int idx) { return finalDensity[idx]; }
|
||
+ @Override public void copyATo(int[] dest, int len) { System.arraycopy(arrA, 0, dest, 0, len); }
|
||
+ @Override public void copyBTo(int[] dest, int len) { System.arraycopy(arrB, 0, dest, 0, len); }
|
||
+ @Override public void copyDensityTo(float[] dest, int len) { System.arraycopy(finalDensity, 0, dest, 0, len); }
|
||
+ };
|
||
+ }
|
||
+
|
||
+ private static int hash(int gx, int gy, int gz) {
|
||
+ int h = (Math.abs(gx) * 73856093) ^ (Math.abs(gy) * 19349663) ^ (Math.abs(gz) * 83492791);
|
||
+ return Math.abs(h) % TABLE_SIZE;
|
||
+ }
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/backend/FFMCollisionBackend.java b/me/earthme/luminol/collision/backend/FFMCollisionBackend.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/backend/FFMCollisionBackend.java
|
||
@@ -0,0 +1,220 @@
|
||
+package me.earthme.luminol.collision.backend;
|
||
+
|
||
+import me.earthme.luminol.collision.CollisionPushResult;
|
||
+import me.earthme.luminol.collision.INativeCollisionBackend;
|
||
+import me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig;
|
||
+import org.slf4j.Logger;
|
||
+import org.slf4j.LoggerFactory;
|
||
+
|
||
+import java.io.*;
|
||
+import java.lang.foreign.*;
|
||
+import java.lang.invoke.MethodHandle;
|
||
+import java.nio.file.Files;
|
||
+import java.util.Set;
|
||
+import java.util.UUID;
|
||
+import java.util.concurrent.ConcurrentHashMap;
|
||
+import java.util.concurrent.atomic.AtomicLong;
|
||
+
|
||
+import static java.lang.foreign.ValueLayout.*;
|
||
+
|
||
+/**
|
||
+ * FFM backend: calls into the AcceleratedRecoiling C++ shared library via
|
||
+ * java.lang.foreign (requires Java 21+). Provides near-native throughput without
|
||
+ * JNI overhead on modern JVMs.
|
||
+ */
|
||
+public class FFMCollisionBackend implements INativeCollisionBackend {
|
||
+
|
||
+ private static final Logger LOGGER = LoggerFactory.getLogger("Luminol/CollisionFFM");
|
||
+ private static final AtomicLong maxSizeSeen = new AtomicLong(-1);
|
||
+
|
||
+ private static Linker linker;
|
||
+ private static Arena nativeArena;
|
||
+ private static MethodHandle hPush, hCreateCtx, hDestroyCtx, hCreateCfg, hUpdateCfg, hDestroyCfg;
|
||
+ private static volatile boolean initialized = false;
|
||
+
|
||
+ @Override public String getName() { return "FFM"; }
|
||
+
|
||
+ // ── Per-thread state ────────────────────────────────────────────────────────
|
||
+ private static class ThreadState {
|
||
+ Arena bufArena;
|
||
+ MemorySegment bufA, bufB, densityBuf;
|
||
+ MemorySegment ctx, cfg;
|
||
+ int currentCap = -1;
|
||
+ final Result result = new Result();
|
||
+
|
||
+ ThreadState() {
|
||
+ AcceleratedCollisionConfig c = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ try {
|
||
+ if (hCreateCtx != null) ctx = (MemorySegment) hCreateCtx.invokeExact();
|
||
+ if (hCreateCfg != null) cfg = (MemorySegment) hCreateCfg.invokeExact(
|
||
+ c.maxCollision, c.gridSize, c.densityWindow, c.maxThreads);
|
||
+ } catch (Throwable t) { throw new RuntimeException("FFM ctx init failed", t); }
|
||
+ }
|
||
+
|
||
+ Result alloc(int need) {
|
||
+ int cap = (int)(need * 1.2);
|
||
+ long intBytes = Math.max(1024L, (long)cap * JAVA_INT.byteSize());
|
||
+ long floatBytes = Math.max(1024L, (long)cap * JAVA_FLOAT.byteSize());
|
||
+ if (intBytes > currentCap) {
|
||
+ if (bufArena != null) bufArena.close();
|
||
+ bufArena = Arena.ofConfined();
|
||
+ bufA = bufArena.allocate(intBytes);
|
||
+ bufB = bufArena.allocate(intBytes);
|
||
+ densityBuf = bufArena.allocate(floatBytes);
|
||
+ currentCap = (int) intBytes;
|
||
+ }
|
||
+ result.a = bufA; result.b = bufB; result.d = densityBuf;
|
||
+ return result;
|
||
+ }
|
||
+
|
||
+ void destroy() {
|
||
+ if (bufArena != null) { try { bufArena.close(); } catch (Exception ignored) {} }
|
||
+ if (ctx != null && hDestroyCtx != null) {
|
||
+ try { hDestroyCtx.invokeExact(ctx); } catch (Throwable t) { LOGGER.error("destroyCtx", t); }
|
||
+ }
|
||
+ if (cfg != null && hDestroyCfg != null) {
|
||
+ try { hDestroyCfg.invokeExact(cfg); } catch (Throwable t) { LOGGER.error("destroyCfg", t); }
|
||
+ }
|
||
+ }
|
||
+
|
||
+ static class Result implements CollisionPushResult {
|
||
+ MemorySegment a, b, d;
|
||
+ @Override public int getA(int i) { return a.get(JAVA_INT, (long)i*4); }
|
||
+ @Override public int getB(int i) { return b.get(JAVA_INT, (long)i*4); }
|
||
+ @Override public float getDensity(int i){ return d.get(JAVA_FLOAT, (long)i*4); }
|
||
+ @Override public void copyATo(int[] dst, int n) { MemorySegment.copy(a, JAVA_INT, 0, dst, 0, n); }
|
||
+ @Override public void copyBTo(int[] dst, int n) { MemorySegment.copy(b, JAVA_INT, 0, dst, 0, n); }
|
||
+ @Override public void copyDensityTo(float[] dst, int n){ MemorySegment.copy(d, JAVA_FLOAT, 0, dst, 0, n); }
|
||
+ }
|
||
+ }
|
||
+
|
||
+ private static final Set<ThreadState> ALL_STATES = ConcurrentHashMap.newKeySet();
|
||
+ private static final ThreadLocal<ThreadState> STATE = ThreadLocal.withInitial(() -> {
|
||
+ ThreadState s = new ThreadState();
|
||
+ ALL_STATES.add(s);
|
||
+ return s;
|
||
+ });
|
||
+
|
||
+ // ── Lifecycle ───────────────────────────────────────────────────────────────
|
||
+ @Override
|
||
+ public void initialize() throws Exception {
|
||
+ if (Runtime.version().feature() < 21)
|
||
+ throw new UnsupportedOperationException("FFM backend requires Java 21+");
|
||
+
|
||
+ String libName = System.mapLibraryName("AcceleratedRecoiling");
|
||
+ String resPath = getNativePath() + libName;
|
||
+
|
||
+ InputStream in = FFMCollisionBackend.class.getResourceAsStream(resPath);
|
||
+ if (in == null) throw new FileNotFoundException("Native lib not found: " + resPath);
|
||
+
|
||
+ File tmp = File.createTempFile(UUID.randomUUID() + "_luminolCollision_", "_" + libName);
|
||
+ tmp.deleteOnExit();
|
||
+ try (OutputStream out = new FileOutputStream(tmp)) { in.transferTo(out); }
|
||
+ LOGGER.info("Extracted native lib to {}", tmp.getAbsolutePath());
|
||
+
|
||
+ linker = Linker.nativeLinker();
|
||
+ nativeArena = Arena.global();
|
||
+ SymbolLookup lib = SymbolLookup.libraryLookup(tmp.toPath(), nativeArena);
|
||
+
|
||
+ hPush = linker.downcallHandle(
|
||
+ lib.find("push").orElseThrow(() -> new RuntimeException("Symbol 'push' not found")),
|
||
+ FunctionDescriptor.of(JAVA_INT, ADDRESS, ADDRESS, ADDRESS, JAVA_INT, ADDRESS, ADDRESS, ADDRESS));
|
||
+ hCreateCtx = linker.downcallHandle(
|
||
+ lib.find("createCtx").orElseThrow(), FunctionDescriptor.of(ADDRESS));
|
||
+ hCreateCfg = linker.downcallHandle(
|
||
+ lib.find("createCfg").orElseThrow(),
|
||
+ FunctionDescriptor.of(ADDRESS, JAVA_INT, JAVA_INT, JAVA_INT, JAVA_INT));
|
||
+ tryBind(lib, "updateCfg", FunctionDescriptor.ofVoid(ADDRESS, JAVA_INT, JAVA_INT, JAVA_INT, JAVA_INT), h -> hUpdateCfg = h);
|
||
+ tryBind(lib, "destroyCtx", FunctionDescriptor.ofVoid(ADDRESS), h -> hDestroyCtx = h);
|
||
+ tryBind(lib, "destroyCfg", FunctionDescriptor.ofVoid(ADDRESS), h -> hDestroyCfg = h);
|
||
+
|
||
+ initialized = true;
|
||
+ LOGGER.info("FFM backend initialized ({})", tmp.getName());
|
||
+ }
|
||
+
|
||
+ private void tryBind(SymbolLookup lib, String sym, FunctionDescriptor fd, java.util.function.Consumer<MethodHandle> setter) {
|
||
+ try { setter.accept(linker.downcallHandle(lib.find(sym).orElseThrow(), fd)); }
|
||
+ catch (Exception e) { LOGGER.warn("Optional symbol '{}' not found: {}", sym, e.getMessage()); }
|
||
+ }
|
||
+
|
||
+ @Override
|
||
+ public void applyConfig() {
|
||
+ if (!initialized || hUpdateCfg == null) return;
|
||
+ AcceleratedCollisionConfig c = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ for (ThreadState s : ALL_STATES) {
|
||
+ if (s.cfg == null) continue;
|
||
+ try { hUpdateCfg.invokeExact(s.cfg, c.maxCollision, c.gridSize, c.densityWindow, c.maxThreads); }
|
||
+ catch (Throwable t) { LOGGER.error("applyConfig", t); }
|
||
+ }
|
||
+ }
|
||
+
|
||
+ @Override
|
||
+ public void destroy() {
|
||
+ if (!initialized) return;
|
||
+ initialized = false;
|
||
+ ALL_STATES.forEach(ThreadState::destroy);
|
||
+ ALL_STATES.clear();
|
||
+ }
|
||
+
|
||
+ @Override
|
||
+ public CollisionPushResult push(double[] aabb, int[] resultSizeOut) {
|
||
+ if (!initialized) return null;
|
||
+ ThreadState s = STATE.get();
|
||
+ if (s.ctx == null) return null;
|
||
+
|
||
+ int count = aabb.length / 6;
|
||
+ AcceleratedCollisionConfig c = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ int resultSize = count * c.maxCollision;
|
||
+ maxSizeSeen.updateAndGet(cur -> Math.max(cur, count));
|
||
+
|
||
+ try (Arena tmp = Arena.ofConfined()) {
|
||
+ MemorySegment aabbMem = tmp.allocateFrom(JAVA_DOUBLE, aabb);
|
||
+ ThreadState.Result out = s.alloc(resultSize);
|
||
+ int n;
|
||
+ try {
|
||
+ n = (int) hPush.invokeExact(aabbMem, out.a, out.b, count, out.d, s.ctx, s.cfg);
|
||
+ } catch (Throwable t) { throw new RuntimeException("FFM push failed", t); }
|
||
+ resultSizeOut[0] = n;
|
||
+ return (n <= 0) ? null : out;
|
||
+ }
|
||
+ }
|
||
+
|
||
+ // ── Helpers ─────────────────────────────────────────────────────────────────
|
||
+ private static String getNativePath() {
|
||
+ String os = System.getProperty("os.name").toLowerCase();
|
||
+ String arch = System.getProperty("os.arch").toLowerCase();
|
||
+ String osKey = os.contains("win") ? "windows" : os.contains("mac") ? "macos" : "linux";
|
||
+ String archKey = (arch.contains("amd64") || arch.contains("x86_64")) ? "x64" : "arm64";
|
||
+ return "/natives/" + osKey + "-" + archKey + "/";
|
||
+ }
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/backend/GPUCollisionBackend.java b/me/earthme/luminol/collision/backend/GPUCollisionBackend.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..ffffffffffffffffffffffffffffffffffffffff
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/backend/GPUCollisionBackend.java
|
||
@@ -0,0 +1,301 @@
|
||
+package me.earthme.luminol.collision.backend;
|
||
+
|
||
+import me.earthme.luminol.collision.CollisionPushResult;
|
||
+import me.earthme.luminol.collision.INativeCollisionBackend;
|
||
+import me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig;
|
||
+import org.jocl.*;
|
||
+import org.slf4j.Logger;
|
||
+import org.slf4j.LoggerFactory;
|
||
+
|
||
+import java.util.Set;
|
||
+import java.util.concurrent.ConcurrentHashMap;
|
||
+import java.util.concurrent.atomic.AtomicLong;
|
||
+
|
||
+import static org.jocl.CL.*;
|
||
+
|
||
+/**
|
||
+ * GPU backend: executes the spatial-hash broadphase as an OpenCL kernel.
|
||
+ * Provides maximum throughput on servers with a discrete GPU.
|
||
+ *
|
||
+ * Ported from AcceleratedRecoiling GPUBackend.java (MIT License).
|
||
+ */
|
||
+public class GPUCollisionBackend implements INativeCollisionBackend {
|
||
+
|
||
+ private static final Logger LOGGER = LoggerFactory.getLogger("Luminol/CollisionGPU");
|
||
+ private static final int TABLE_SIZE = 262139;
|
||
+ private static final double CELL_SIZE = 2.5;
|
||
+ private static final AtomicLong maxSizeSeen = new AtomicLong(-1);
|
||
+
|
||
+ private static cl_context clCtx;
|
||
+ private static cl_program clProg;
|
||
+ private static cl_device_id clDevice;
|
||
+ private static volatile boolean initialized = false;
|
||
+
|
||
+ private static final String KERNEL_SRC = """
|
||
+ #pragma OPENCL EXTENSION cl_khr_fp64 : enable
|
||
+ #define INVALID_INDEX 4294967295u
|
||
+ #define TABLE_SIZE 262139
|
||
+ inline uint hashPos(double cx,double cy,double cz,double cs){
|
||
+ int gx=(int)floor(cx/cs),gy=(int)floor(cy/cs),gz=(int)floor(cz/cs);
|
||
+ uint h=(abs(gx)*73856093u)^(abs(gy)*19349663u)^(abs(gz)*83492791u);
|
||
+ return h%TABLE_SIZE;
|
||
+ }
|
||
+ __kernel void compute_hash(__global const double* aabbs,__global uint* hashes,__global uint* indices,int count,double cs){
|
||
+ int id=get_global_id(0); if(id>=count)return;
|
||
+ double cx=(aabbs[id*6+0]+aabbs[id*6+3])*0.5;
|
||
+ double cy=(aabbs[id*6+1]+aabbs[id*6+4])*0.5;
|
||
+ double cz=(aabbs[id*6+2]+aabbs[id*6+5])*0.5;
|
||
+ hashes[id]=hashPos(cx,cy,cz,cs); indices[id]=id;
|
||
+ }
|
||
+ __kernel void reset_grid(__global uint* cs,__global uint* ce,int ts){
|
||
+ int id=get_global_id(0); if(id>=ts)return;
|
||
+ cs[id]=INVALID_INDEX; ce[id]=INVALID_INDEX;
|
||
+ }
|
||
+ __kernel void build_grid(__global const uint* hashes,__global uint* cs,__global uint* ce,int count){
|
||
+ int id=get_global_id(0); if(id>=count)return;
|
||
+ uint h=hashes[id];
|
||
+ if(id==0||h!=hashes[id-1]) cs[h]=id;
|
||
+ if(id==count-1||h!=hashes[id+1]) ce[h]=id+1;
|
||
+ }
|
||
+ __kernel void detect(__global const double* aabbs,__global const uint* hashes,__global const uint* indices,
|
||
+ __global const uint* cs,__global const uint* ce,
|
||
+ __global int* outA,__global int* outB,__global float* density,volatile __global int* counter,
|
||
+ int count,int maxCol,double cellSize){
|
||
+ int idx=get_global_id(0); if(idx>=count)return;
|
||
+ uint rid=indices[idx];
|
||
+ double mnX=aabbs[rid*6+0],mnY=aabbs[rid*6+1],mnZ=aabbs[rid*6+2];
|
||
+ double mxX=aabbs[rid*6+3],mxY=aabbs[rid*6+4],mxZ=aabbs[rid*6+5];
|
||
+ double cx=(mnX+mxX)*0.5,cy=(mnY+mxY)*0.5,cz=(mnZ+mxZ)*0.5;
|
||
+ int gx=(int)floor(cx/cellSize),gy=(int)floor(cy/cellSize),gz=(int)floor(cz/cellSize);
|
||
+ int ov=0;
|
||
+ for(int dz=-1;dz<=1;dz++)for(int dy=-1;dy<=1;dy++)for(int dx=-1;dx<=1;dx++){
|
||
+ uint h=((uint)abs(gx+dx)*73856093u)^((uint)abs(gy+dy)*19349663u)^((uint)abs(gz+dz)*83492791u);
|
||
+ uint nh=h%TABLE_SIZE;
|
||
+ uint start=cs[nh]; if(start==INVALID_INDEX)continue;
|
||
+ uint end=ce[nh];
|
||
+ for(uint i=start;i<end;i++){
|
||
+ uint oid=indices[i]; if(oid==rid)continue;
|
||
+ double oMnX=aabbs[oid*6+0],oMxX=aabbs[oid*6+3];
|
||
+ if(mnX>oMxX||mxX<oMnX)continue;
|
||
+ double oMnY=aabbs[oid*6+1],oMxY=aabbs[oid*6+4];
|
||
+ if(mnY>oMxY||mxY<oMnY)continue;
|
||
+ double oMnZ=aabbs[oid*6+2],oMxZ=aabbs[oid*6+5];
|
||
+ if(mnZ>oMxZ||mxZ<oMnZ)continue;
|
||
+ ov++;
|
||
+ if(rid<oid){int out_idx=atomic_inc(counter);if(out_idx<maxCol){outA[out_idx]=rid;outB[out_idx]=oid;}}
|
||
+ }
|
||
+ }
|
||
+ density[rid]=(float)ov;
|
||
+ }
|
||
+ """;
|
||
+
|
||
+ @Override public String getName() { return "GPU"; }
|
||
+ @Override public void applyConfig() {}
|
||
+
|
||
+ // ── Per-thread state ────────────────────────────────────────────────────────
|
||
+ private static class TState {
|
||
+ cl_command_queue queue;
|
||
+ cl_kernel kHash, kReset, kBuild, kDetect;
|
||
+ cl_mem mAABB, mHashes, mIndices, mOutA, mOutB, mDensity, mCounter, mCellS, mCellE;
|
||
+ int capEntity = -1, capCollision = -1;
|
||
+ int[] cpuH, cpuI, tmpK, tmpV;
|
||
+ Result res = new Result();
|
||
|
||
+ TState() {
|
||
+ if (!initialized) return;
|
||
+ queue = clCreateCommandQueueWithProperties(clCtx, clDevice, null, null);
|
||
+ kHash = clCreateKernel(clProg, "compute_hash", null);
|
||
+ kReset = clCreateKernel(clProg, "reset_grid", null);
|
||
+ kBuild = clCreateKernel(clProg, "build_grid", null);
|
||
+ kDetect = clCreateKernel(clProg, "detect", null);
|
||
+ mCellS = clCreateBuffer(clCtx, CL_MEM_READ_WRITE, (long)TABLE_SIZE*Sizeof.cl_uint, null, null);
|
||
+ mCellE = clCreateBuffer(clCtx, CL_MEM_READ_WRITE, (long)TABLE_SIZE*Sizeof.cl_uint, null, null);
|
||
+ }
|
||
|
||
+ void alloc(int ec, int mc) {
|
||
+ if (ec > capEntity) {
|
||
+ int nc = (int)(ec*1.5);
|
||
+ rel(mAABB); rel(mHashes); rel(mIndices); rel(mDensity);
|
||
+ mAABB = clCreateBuffer(clCtx, CL_MEM_READ_ONLY, (long)nc*6*Sizeof.cl_double, null, null);
|
||
+ mHashes = clCreateBuffer(clCtx, CL_MEM_READ_WRITE, (long)nc*Sizeof.cl_uint, null, null);
|
||
+ mIndices = clCreateBuffer(clCtx, CL_MEM_READ_WRITE, (long)nc*Sizeof.cl_uint, null, null);
|
||
+ mDensity = clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, (long)nc*Sizeof.cl_float, null, null);
|
||
+ cpuH = new int[nc]; cpuI = new int[nc]; tmpK = new int[nc]; tmpV = new int[nc];
|
||
+ res.arrDensity = new float[nc];
|
||
+ capEntity = nc;
|
||
+ }
|
||
+ if (mc > capCollision) {
|
||
+ int nc = (int)(mc*1.5);
|
||
+ rel(mOutA); rel(mOutB);
|
||
+ mOutA = clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, (long)nc*Sizeof.cl_int, null, null);
|
||
+ mOutB = clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, (long)nc*Sizeof.cl_int, null, null);
|
||
+ res.arrA = new int[nc]; res.arrB = new int[nc];
|
||
+ capCollision = nc;
|
||
+ }
|
||
+ if (mCounter == null)
|
||
+ mCounter = clCreateBuffer(clCtx, CL_MEM_READ_WRITE, Sizeof.cl_int, null, null);
|
||
+ }
|
||
|
||
+ private void rel(cl_mem m) { if (m != null) clReleaseMemObject(m); }
|
||
|
||
+ void destroy() {
|
||
+ rel(mAABB); rel(mHashes); rel(mIndices); rel(mOutA); rel(mOutB);
|
||
+ rel(mDensity); rel(mCounter); rel(mCellS); rel(mCellE);
|
||
+ if (kHash != null) clReleaseKernel(kHash);
|
||
+ if (kReset != null) clReleaseKernel(kReset);
|
||
+ if (kBuild != null) clReleaseKernel(kBuild);
|
||
+ if (kDetect != null) clReleaseKernel(kDetect);
|
||
+ if (queue != null) clReleaseCommandQueue(queue);
|
||
+ }
|
||
|
||
+ static class Result implements CollisionPushResult {
|
||
+ int[] arrA, arrB; float[] arrDensity;
|
||
+ @Override public int getA(int i) { return arrA[i]; }
|
||
+ @Override public int getB(int i) { return arrB[i]; }
|
||
+ @Override public float getDensity(int i){ return arrDensity[i]; }
|
||
+ @Override public void copyATo(int[] d,int n) { System.arraycopy(arrA,0,d,0,n); }
|
||
+ @Override public void copyBTo(int[] d,int n) { System.arraycopy(arrB,0,d,0,n); }
|
||
+ @Override public void copyDensityTo(float[] d,int n){ System.arraycopy(arrDensity,0,d,0,n); }
|
||
+ }
|
||
+ }
|
||
|
||
+ private static final Set<TState> ALL = ConcurrentHashMap.newKeySet();
|
||
+ private static final ThreadLocal<TState> TS = ThreadLocal.withInitial(() -> { TState s=new TState(); ALL.add(s); return s; });
|
||
+
|
||
+ @Override
|
||
+ public void initialize() throws Exception {
|
||
+ setExceptionsEnabled(true);
|
||
+ int[] np = new int[1];
|
||
+ clGetPlatformIDs(0, null, np);
|
||
+ if (np[0] == 0) throw new UnsupportedOperationException("No OpenCL platforms found");
|
||
+ cl_platform_id[] plats = new cl_platform_id[np[0]];
|
||
+ clGetPlatformIDs(plats.length, plats, null);
|
||
+
|
||
+ cl_platform_id tgtPlat = null; cl_device_id tgtDev = null;
|
||
+ outer: for (cl_platform_id p : plats) {
|
||
+ try {
|
||
+ int[] nd = new int[1];
|
||
+ clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, 0, null, nd);
|
||
+ if (nd[0] > 0) {
|
||
+ cl_device_id[] devs = new cl_device_id[nd[0]];
|
||
+ clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, devs.length, devs, null);
|
||
+ tgtPlat = p; tgtDev = devs[0]; break outer;
|
||
+ }
|
||
+ } catch (CLException ignored) {}
|
||
+ }
|
||
+ if (tgtDev == null) throw new UnsupportedOperationException("No GPU device found");
|
||
+
|
||
+ long[] sz = new long[1];
|
||
+ clGetDeviceInfo(tgtDev, CL_DEVICE_NAME, 0, null, sz);
|
||
+ byte[] nm = new byte[(int)sz[0]];
|
||
+ clGetDeviceInfo(tgtDev, CL_DEVICE_NAME, nm.length, Pointer.to(nm), null);
|
||
+ LOGGER.info("GPU backend: using device '{}'", new String(nm, 0, nm.length-1).trim());
|
||
|
||
+ clDevice = tgtDev;
|
||
+ cl_context_properties props = new cl_context_properties();
|
||
+ props.addProperty(CL_CONTEXT_PLATFORM, tgtPlat);
|
||
+ clCtx = clCreateContext(props, 1, new cl_device_id[]{tgtDev}, null, null, null);
|
||
+ clProg = clCreateProgramWithSource(clCtx, 1, new String[]{KERNEL_SRC}, null, null);
|
||
+ clBuildProgram(clProg, 0, null, "-cl-mad-enable -cl-fast-relaxed-math", null, null);
|
||
+ initialized = true;
|
||
+ }
|
||
|
||
+ @Override
|
||
+ public void destroy() {
|
||
+ if (!initialized) return;
|
||
+ initialized = false;
|
||
+ ALL.forEach(TState::destroy); ALL.clear();
|
||
+ if (clProg != null) clReleaseProgram(clProg);
|
||
+ if (clCtx != null) clReleaseContext(clCtx);
|
||
+ }
|
||
|
||
+ @Override
|
||
+ public CollisionPushResult push(double[] aabb, int[] resultSizeOut) {
|
||
+ if (!initialized) return null;
|
||
+ int ec = aabb.length / 6;
|
||
+ if (ec == 0) { resultSizeOut[0] = 0; return null; }
|
||
+ AcceleratedCollisionConfig cfg = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ int mc = ec * cfg.maxCollision;
|
||
+ maxSizeSeen.updateAndGet(c -> Math.max(c, ec));
|
||
|
||
+ TState s = TS.get();
|
||
+ s.alloc(ec, mc);
|
||
|
||
+ clEnqueueWriteBuffer(s.queue, s.mAABB, CL_TRUE, 0, (long)ec*6*Sizeof.cl_double, Pointer.to(aabb), 0, null, null);
|
||
|
||
+ arg(s.kHash, 0, s.mAABB); arg(s.kHash, 1, s.mHashes); arg(s.kHash, 2, s.mIndices);
|
||
+ argInt(s.kHash, 3, ec); argDbl(s.kHash, 4, CELL_SIZE);
|
||
+ clEnqueueNDRangeKernel(s.queue, s.kHash, 1, null, new long[]{ec}, null, 0, null, null);
|
||
|
||
+ clEnqueueReadBuffer(s.queue, s.mHashes, CL_TRUE, 0, (long)ec*Sizeof.cl_uint, Pointer.to(s.cpuH), 0, null, null);
|
||
+ clEnqueueReadBuffer(s.queue, s.mIndices, CL_TRUE, 0, (long)ec*Sizeof.cl_uint, Pointer.to(s.cpuI), 0, null, null);
|
||
+ radixSort(s.cpuH, s.cpuI, s.tmpK, s.tmpV, ec);
|
||
+ clEnqueueWriteBuffer(s.queue, s.mHashes, CL_TRUE, 0, (long)ec*Sizeof.cl_uint, Pointer.to(s.cpuH), 0, null, null);
|
||
+ clEnqueueWriteBuffer(s.queue, s.mIndices, CL_TRUE, 0, (long)ec*Sizeof.cl_uint, Pointer.to(s.cpuI), 0, null, null);
|
||
|
||
+ arg(s.kReset, 0, s.mCellS); arg(s.kReset, 1, s.mCellE); argInt(s.kReset, 2, TABLE_SIZE);
|
||
+ clEnqueueNDRangeKernel(s.queue, s.kReset, 1, null, new long[]{TABLE_SIZE}, null, 0, null, null);
|
||
|
||
+ arg(s.kBuild, 0, s.mHashes); arg(s.kBuild, 1, s.mCellS); arg(s.kBuild, 2, s.mCellE); argInt(s.kBuild, 3, ec);
|
||
+ clEnqueueNDRangeKernel(s.queue, s.kBuild, 1, null, new long[]{ec}, null, 0, null, null);
|
||
|
||
+ clEnqueueWriteBuffer(s.queue, s.mCounter, CL_TRUE, 0, Sizeof.cl_int, Pointer.to(new int[]{0}), 0, null, null);
|
||
|
||
+ arg(s.kDetect,0,s.mAABB);arg(s.kDetect,1,s.mHashes);arg(s.kDetect,2,s.mIndices);
|
||
+ arg(s.kDetect,3,s.mCellS);arg(s.kDetect,4,s.mCellE);
|
||
+ arg(s.kDetect,5,s.mOutA);arg(s.kDetect,6,s.mOutB);arg(s.kDetect,7,s.mDensity);arg(s.kDetect,8,s.mCounter);
|
||
+ argInt(s.kDetect,9,ec);argInt(s.kDetect,10,mc);argDbl(s.kDetect,11,CELL_SIZE);
|
||
+ clEnqueueNDRangeKernel(s.queue, s.kDetect, 1, null, new long[]{ec}, null, 0, null, null);
|
||
|
||
+ int[] cnt = new int[1];
|
||
+ clEnqueueReadBuffer(s.queue, s.mCounter, CL_TRUE, 0, Sizeof.cl_int, Pointer.to(cnt), 0, null, null);
|
||
+ int collisions = Math.min(cnt[0], mc);
|
||
+ resultSizeOut[0] = collisions;
|
||
|
||
+ if (collisions > 0) {
|
||
+ clEnqueueReadBuffer(s.queue, s.mOutA, CL_TRUE, 0, (long)collisions*Sizeof.cl_int, Pointer.to(s.res.arrA), 0, null, null);
|
||
+ clEnqueueReadBuffer(s.queue, s.mOutB, CL_TRUE, 0, (long)collisions*Sizeof.cl_int, Pointer.to(s.res.arrB), 0, null, null);
|
||
+ }
|
||
+ clEnqueueReadBuffer(s.queue, s.mDensity, CL_TRUE, 0, (long)ec*Sizeof.cl_float, Pointer.to(s.res.arrDensity), 0, null, null);
|
||
+ return s.res;
|
||
+ }
|
||
|
||
+ private static void arg(cl_kernel k, int i, cl_mem m){ clSetKernelArg(k,i,Sizeof.cl_mem,Pointer.to(m)); }
|
||
+ private static void argInt(cl_kernel k,int i,int v){ clSetKernelArg(k,i,Sizeof.cl_int,Pointer.to(new int[]{v})); }
|
||
+ private static void argDbl(cl_kernel k,int i,double v){ clSetKernelArg(k,i,Sizeof.cl_double,Pointer.to(new double[]{v})); }
|
||
+
|
||
+ private static void radixSort(int[] keys,int[] vals,int[] kb,int[] vb,int n){
|
||
+ int[] hist=new int[256]; int[] sk=keys,sv=vals,dk=kb,dv=vb;
|
||
+ for(int pass=0;pass<4;pass++){
|
||
+ int shift=pass*8; java.util.Arrays.fill(hist,0);
|
||
+ for(int i=0;i<n;i++) hist[(sk[i]>>>shift)&0xFF]++;
|
||
+ int off=0; for(int i=0;i<256;i++){int c=hist[i];hist[i]=off;off+=c;}
|
||
+ for(int i=0;i<n;i++){int p=(sk[i]>>>shift)&0xFF,d=hist[p]++;dk[d]=sk[i];dv[d]=sv[i];}
|
||
+ int[] t=sk;sk=dk;dk=t;t=sv;sv=dv;dv=t;
|
||
+ }
|
||
+ }
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/CollisionBackendSelector.java b/me/earthme/luminol/collision/CollisionBackendSelector.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..1111111111111111111111111111111111111111
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/CollisionBackendSelector.java
|
||
@@ -0,0 +1,96 @@
|
||
+package me.earthme.luminol.collision;
|
||
+
|
||
+import me.earthme.luminol.collision.backend.FFMCollisionBackend;
|
||
+import me.earthme.luminol.collision.backend.GPUCollisionBackend;
|
||
+import me.earthme.luminol.collision.backend.JavaCollisionBackend;
|
||
+import me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig;
|
||
+import org.slf4j.Logger;
|
||
+import org.slf4j.LoggerFactory;
|
||
+
|
||
+import java.util.List;
|
||
+
|
||
+/**
|
||
+ * Chooses and initialises the best available collision backend.
|
||
+ * Hierarchy: GPU → FFM → JNI → SIMD → Java.
|
||
+ * Can be overridden via config (or -Dluminol.collision.backend=X).
|
||
+ */
|
||
+public final class CollisionBackendSelector {
|
||
+
|
||
+ private static final Logger LOGGER = LoggerFactory.getLogger("Luminol/CollisionSelector");
|
||
+ private static volatile INativeCollisionBackend BACKEND = null;
|
||
+ private static volatile boolean READY = false;
|
||
+
|
||
+ private CollisionBackendSelector() {}
|
||
+
|
||
+ public static synchronized void initialize() {
|
||
+ if (READY) return;
|
||
+ AcceleratedCollisionConfig cfg = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ if (!cfg.enabled) {
|
||
+ LOGGER.info("Accelerated entity collision disabled by config.");
|
||
+ READY = true;
|
||
+ return;
|
||
+ }
|
||
+
|
||
+ // JVM property overrides config
|
||
+ String override = System.getProperty("luminol.collision.backend", cfg.backend).toUpperCase().trim();
|
||
+
|
||
+ List<Candidate> chain;
|
||
+ if ("AUTO".equals(override) || override.isEmpty()) {
|
||
+ chain = List.of(
|
||
+ new Candidate("GPU", GPUCollisionBackend::new),
|
||
+ new Candidate("FFM", FFMCollisionBackend::new),
|
||
+ new Candidate("Java", JavaCollisionBackend::new)
|
||
+ );
|
||
+ } else {
|
||
+ chain = List.of(resolveCandidate(override));
|
||
+ }
|
||
+
|
||
+ for (Candidate c : chain) {
|
||
+ try {
|
||
+ INativeCollisionBackend b = c.factory.get();
|
||
+ b.initialize();
|
||
+ BACKEND = b;
|
||
+ LOGGER.info("Collision backend selected: {}", b.getName());
|
||
+ break;
|
||
+ } catch (Exception e) {
|
||
+ LOGGER.warn("Backend '{}' unavailable: {}", c.name, e.getMessage());
|
||
+ }
|
||
+ }
|
||
+
|
||
+ if (BACKEND == null) {
|
||
+ LOGGER.warn("All accelerated backends failed – falling back to vanilla collision.");
|
||
+ }
|
||
+ READY = true;
|
||
+ }
|
||
+
|
||
+ public static void destroy() {
|
||
+ if (BACKEND != null) { BACKEND.destroy(); BACKEND = null; }
|
||
+ READY = false;
|
||
+ }
|
||
+
|
||
+ public static void applyConfig() {
|
||
+ if (BACKEND != null) BACKEND.applyConfig();
|
||
+ }
|
||
+
|
||
+ public static CollisionPushResult push(double[] aabb, int[] resultSizeOut) {
|
||
+ if (BACKEND == null) return null;
|
||
+ return BACKEND.push(aabb, resultSizeOut);
|
||
+ }
|
||
+
|
||
+ public static boolean isReady() { return READY && BACKEND != null; }
|
||
+
|
||
+ // ── Helpers ─────────────────────────────────────────────────────────────────
|
||
+ private record Candidate(String name, java.util.function.Supplier<INativeCollisionBackend> factory) {}
|
||
+
|
||
+ private static Candidate resolveCandidate(String name) {
|
||
+ return switch (name) {
|
||
+ case "GPU" -> new Candidate("GPU", GPUCollisionBackend::new);
|
||
+ case "FFM" -> new Candidate("FFM", FFMCollisionBackend::new);
|
||
+ case "JAVA" -> new Candidate("Java", JavaCollisionBackend::new);
|
||
+ default -> throw new IllegalArgumentException("Unknown backend: " + name);
|
||
+ };
|
||
+ }
|
||
+}
|
||
diff --git a/me/earthme/luminol/collision/RegionCollisionHandler.java b/me/earthme/luminol/collision/RegionCollisionHandler.java
|
||
new file mode 100644
|
||
index 0000000000000000000000000000000000000000..2222222222222222222222222222222222222222
|
||
--- /dev/null
|
||
+++ b/me/earthme/luminol/collision/RegionCollisionHandler.java
|
||
@@ -0,0 +1,117 @@
|
||
+package me.earthme.luminol.collision;
|
||
+
|
||
+import me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig;
|
||
+import net.minecraft.world.entity.Entity;
|
||
+import net.minecraft.world.entity.EntitySelector;
|
||
+import net.minecraft.world.entity.LivingEntity;
|
||
+import net.minecraft.world.phys.AABB;
|
||
+
|
||
+import java.util.List;
|
||
+
|
||
+/**
|
||
+ * Entry-point called from the ServerLevel tick injection.
|
||
+ * Extracts AABB data from the entity list, delegates to the backend,
|
||
+ * and applies the resulting push pairs — fully Folia-safe (called
|
||
+ * only from within the owning region thread).
|
||
+ */
|
||
+public final class RegionCollisionHandler {
|
||
+
|
||
+ private RegionCollisionHandler() {}
|
||
+
|
||
+ /**
|
||
+ * Process entity-push collisions for the given list of non-player living entities.
|
||
+ *
|
||
+ * @param entities entities owned by the current Folia region
|
||
+ * @param inflate AABB inflate radius (matches vanilla, typically 1e-7)
|
||
+ */
|
||
+ public static void handleEntityPush(List<Entity> entities, double inflate) {
|
||
+ if (!CollisionBackendSelector.isReady() || entities.isEmpty()) return;
|
||
+
|
||
+ AcceleratedCollisionConfig cfg = me.earthme.luminol.LuminolConfig.get().optimizations.acceleratedCollision;
|
||
+ int count = entities.size();
|
||
+
|
||
+ // Check density – skip acceleration when below threshold
|
||
+ if (count < cfg.densityThreshold) return;
|
||
+
|
||
+ double[] aabb = new double[count * 6];
|
||
+ float[] density = new float[count];
|
||
+
|
||
+ for (int i = 0; i < count; i++) {
|
||
+ Entity e = entities.get(i);
|
||
+ AABB bb = e.getBoundingBox().inflate(inflate);
|
||
+ aabb[i*6] = bb.minX; aabb[i*6+1] = bb.minY; aabb[i*6+2] = bb.minZ;
|
||
+ aabb[i*6+3] = bb.maxX; aabb[i*6+4] = bb.maxY; aabb[i*6+5] = bb.maxZ;
|
||
+ }
|
||
+
|
||
+ int[] resultCount = new int[1];
|
||
+ CollisionPushResult result = CollisionBackendSelector.push(aabb, resultCount);
|
||
+ if (result == null) return;
|
||
+
|
||
+ // Apply density info (used for debug or future throttling)
|
||
+ result.copyDensityTo(density, count);
|
||
+
|
||
+ // Apply push pairs
|
||
+ for (int k = 0; k < resultCount[0]; k++) {
|
||
+ int ai = result.getA(k), bi = result.getB(k);
|
||
+ if (ai >= count || bi >= count) continue;
|
||
+
|
||
+ Entity e1 = entities.get(ai);
|
||
+ Entity e2 = entities.get(bi);
|
||
+
|
||
+ LivingEntity living; Entity other;
|
||
+ if (e1 instanceof LivingEntity le) { living = le; other = e2; }
|
||
+ else if (e2 instanceof LivingEntity le) { living = le; other = e1; }
|
||
+ else continue;
|
||
+
|
||
+ if (EntitySelector.pushableBy(living).test(other)) {
|
||
+ living.doPush(other);
|
||
+ }
|
||
+ }
|
||
+ }
|
||
+}
|
||
diff --git a/net/minecraft/server/level/ServerLevel.java b/net/minecraft/server/level/ServerLevel.java
|
||
index 0000000000000000000000000000000000000000..3333333333333333333333333333333333333333 100644
|
||
--- a/net/minecraft/server/level/ServerLevel.java
|
||
+++ b/net/minecraft/server/level/ServerLevel.java
|
||
@@ -800,6 +800,42 @@ public class ServerLevel extends Level implements ServerEntityGetter, WorldGenLe
|
||
profiler.startTimer(ca.spottedleaf.leafprofiler.LProfilerRegistry.ENTITY_TICK); try { // Folia - profiler
|
||
this.entityTickList.forEach( (entity) -> {
|
||
if (!entity.isRemoved()) {
|
||
+ // Luminol start - AcceleratedRecoiling entity collision
|
||
+ // (entity accumulation into region list is handled below at tick-end)
|
||
+ // Luminol end - AcceleratedRecoiling entity collision
|
||
if (!this.shouldDiscardEntity(entity)) {
|
||
profilerFiller.push("tick");
|
||
this.guardEntityTick(this::tickNonPassenger, entity);
|
||
@@ -820,6 +820,40 @@ public class ServerLevel extends Level implements ServerEntityGetter, WorldGenLe
|
||
}
|
||
});
|
||
} finally { profiler.stopTimer(ca.spottedleaf.leafprofiler.LProfilerRegistry.ENTITY_TICK); } // Folia - profiler
|
||
+ // Luminol start - AcceleratedRecoiling: run spatial-hash broadphase after entity tick
|
||
+ if (me.earthme.luminol.config.modules.optimizations.AcceleratedCollisionConfig_access.enabled()) {
|
||
+ if (!me.earthme.luminol.collision.CollisionBackendSelector.isReady()) {
|
||
+ me.earthme.luminol.collision.CollisionBackendSelector.initialize();
|
||
+ }
|
||
+ java.util.List<net.minecraft.world.entity.Entity> regionEntities = new java.util.ArrayList<>();
|
||
+ this.entityTickList.forEach(entity -> {
|
||
+ if (!entity.isRemoved() && !(entity instanceof net.minecraft.world.entity.player.Player)) {
|
||
+ regionEntities.add(entity);
|
||
+ }
|
||
+ });
|
||
+ me.earthme.luminol.collision.RegionCollisionHandler.handleEntityPush(regionEntities, 1.0E-7);
|
||
+ }
|
||
+ // Luminol end - AcceleratedRecoiling
|
||
diff --git a/net/minecraft/server/dedicated/DedicatedServer.java b/net/minecraft/server/dedicated/DedicatedServer.java
|
||
index 0000000000000000000000000000000000000000..4444444444444444444444444444444444444444 100644
|
||
--- a/net/minecraft/server/dedicated/DedicatedServer.java
|
||
+++ b/net/minecraft/server/dedicated/DedicatedServer.java
|
||
@@ -720,6 +720,7 @@ public class DedicatedServer extends MinecraftServer implements ServerInterface
|
||
|
||
@Override
|
||
public void stopServer() {
|
||
+ me.earthme.luminol.collision.CollisionBackendSelector.destroy(); // Luminol - AcceleratedRecoiling shutdown
|
||
super.stopServer();
|