Onek8/LNXSDK
1
0
Fork 1
forked from LeenkxTeam/LNXSDK
Code Pull Requests Activity

Compute 0

Browse Source
This commit is contained in:
Onek8
2026-06-15 20:07:42 -07:00
parent 3aac63d255
commit 96134404a2
13 changed files with 223 additions and 217 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
Kha/Backends/Kore-HL/kha/korehl/graphics4/Graphics.hx
View File

@ -266,7 +266,6 @@ class Graphics implements kha.graphics4.Graphics {
} }
public function beginFace(face: Int): Void { public function beginFace(face: Int): Void {
// TODO: Hashlink conflict for shadowmap atlasing needs image and probes need cubemap
kinc_graphics_render_to_face((target is CubeMap) ? cast(target, CubeMap)._renderTarget : cast(target, Image)._renderTarget, face); kinc_graphics_render_to_face((target is CubeMap) ? cast(target, CubeMap)._renderTarget : cast(target, Image)._renderTarget, face);
} }

47
Kha/Backends/Krom/kha/graphics4/ComputeShader.hx
View File

@ -1,22 +1,25 @@
package kha.graphics4; package kha.graphics4;
import haxe.io.Bytes; import haxe.io.Bytes;
import kha.Blob; import kha.Blob;
class ComputeShader { class ComputeShader {
public function new(sources: Array<Blob>, files: Array<String>) { public var shader_: Dynamic;
} public function new(sources: Array<Blob>, files: Array<String>) {
shader_ = Krom.createShaderCompute(sources[0].toBytes().getData());
public function delete(): Void { }
} public function delete(): Void {
Krom.deleteShaderCompute(shader_);
public function getConstantLocation(name: String): ConstantLocation { shader_ = null;
return null; }
}
public function getConstantLocation(name: String): ConstantLocation {
public function getTextureUnit(name: String): TextureUnit { return Krom.getConstantLocationCompute(shader_, name);
return null; }
}
} public function getTextureUnit(name: String): TextureUnit {
return Krom.getTextureUnitCompute(shader_, name);
}
}

4
Kha/Backends/Krom/kha/krom/Graphics.hx
View File

@ -261,10 +261,10 @@ class Graphics implements kha.graphics4.Graphics {
} }
public function setComputeShader(shader: ComputeShader) { public function setComputeShader(shader: ComputeShader) {
Krom.setShaderCompute(shader.shader_);
} }
public function compute(x: Int, y: Int, z: Int) { public function compute(x: Int, y: Int, z: Int) {
Krom.compute(x, y, z);
} }
} }

BIN
Kha/Kore/Tools/freebsd_x64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/linux_arm/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/linux_arm64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/linux_x64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/macos_arm64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/macos_x64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

BIN
Kha/Kore/Tools/windows_x64/icon.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 417 KiB

After

Width:  |  Height:  |  Size: 281 KiB

2
leenkx/Sources/iron/RenderPath.hx
View File

@ -672,7 +672,7 @@ class RenderPath {
} }
#if (rp_voxels != "Off") #if (rp_voxels != "Off")
public function getComputeShader(handle: String): kha.compute.Shader { public function getComputeShader(handle: String): kha.graphics4.ComputeShader {
return Reflect.field(kha.Shaders, handle + "_comp"); return Reflect.field(kha.Shaders, handle + "_comp");
} }
#end #end

367
leenkx/Sources/leenkx/renderpath/Inc.hx
View File

@ -22,99 +22,99 @@ class Inc {
#end #end
#if (rp_voxels != "Off") #if (rp_voxels != "Off")
static var voxel_sh0:kha.compute.Shader = null; static var voxel_sh0:kha.graphics4.ComputeShader = null;
static var voxel_sh1:kha.compute.Shader = null; static var voxel_sh1:kha.graphics4.ComputeShader = null;
static var voxel_ta0:kha.compute.TextureUnit; static var voxel_ta0:kha.graphics4.TextureUnit;
static var voxel_tb0:kha.compute.TextureUnit; static var voxel_tb0:kha.graphics4.TextureUnit;
static var voxel_ca0:kha.compute.ConstantLocation; static var voxel_ca0:kha.graphics4.ConstantLocation;
static var voxel_cb0:kha.compute.ConstantLocation; static var voxel_cb0:kha.graphics4.ConstantLocation;
static var voxel_ta1:kha.compute.TextureUnit; static var voxel_ta1:kha.graphics4.TextureUnit;
static var voxel_tb1:kha.compute.TextureUnit; static var voxel_tb1:kha.graphics4.TextureUnit;
static var voxel_tc1:kha.compute.TextureUnit; static var voxel_tc1:kha.graphics4.TextureUnit;
static var m = iron.math.Mat4.identity(); static var m = iron.math.Mat4.identity();
static var voxel_ca1:kha.compute.ConstantLocation; static var voxel_ca1:kha.graphics4.ConstantLocation;
static var voxel_cb1:kha.compute.ConstantLocation; static var voxel_cb1:kha.graphics4.ConstantLocation;
#if (rp_voxels == "Voxel GI") #if (rp_voxels == "Voxel GI")
static var voxel_td1:kha.compute.TextureUnit; static var voxel_td1:kha.graphics4.TextureUnit;
static var voxel_te1:kha.compute.TextureUnit; static var voxel_te1:kha.graphics4.TextureUnit;
static var voxel_tf1:kha.compute.TextureUnit; static var voxel_tf1:kha.graphics4.TextureUnit;
static var voxel_cc1:kha.compute.ConstantLocation; static var voxel_cc1:kha.graphics4.ConstantLocation;
#else #else
#if lnx_voxelgi_shadows #if lnx_voxelgi_shadows
static var voxel_te1:kha.compute.TextureUnit; static var voxel_te1:kha.graphics4.TextureUnit;
#end #end
#end #end
#if (lnx_voxelgi_shadows || rp_voxels == "Voxel GI") #if (lnx_voxelgi_shadows || rp_voxels == "Voxel GI")
static var voxel_sh2:kha.compute.Shader = null; static var voxel_sh2:kha.graphics4.ComputeShader = null;
static var voxel_ta2:kha.compute.TextureUnit; static var voxel_ta2:kha.graphics4.TextureUnit;
static var voxel_tb2:kha.compute.TextureUnit; static var voxel_tb2:kha.graphics4.TextureUnit;
static var voxel_ca2:kha.compute.ConstantLocation; static var voxel_ca2:kha.graphics4.ConstantLocation;
static var voxel_cb2:kha.compute.ConstantLocation; static var voxel_cb2:kha.graphics4.ConstantLocation;
static var voxel_cc2:kha.compute.ConstantLocation; static var voxel_cc2:kha.graphics4.ConstantLocation;
#end #end
static var voxel_sh3:kha.compute.Shader = null; static var voxel_sh3:kha.graphics4.ComputeShader = null;
static var voxel_ta3:kha.compute.TextureUnit; static var voxel_ta3:kha.graphics4.TextureUnit;
static var voxel_tb3:kha.compute.TextureUnit; static var voxel_tb3:kha.graphics4.TextureUnit;
static var voxel_tc3:kha.compute.TextureUnit; static var voxel_tc3:kha.graphics4.TextureUnit;
static var voxel_td3:kha.compute.TextureUnit; static var voxel_td3:kha.graphics4.TextureUnit;
static var voxel_te3:kha.compute.TextureUnit; static var voxel_te3:kha.graphics4.TextureUnit;
static var voxel_tf3:kha.compute.TextureUnit; static var voxel_tf3:kha.graphics4.TextureUnit;
#if lnx_brdf #if lnx_brdf
static var voxel_tg3:kha.compute.TextureUnit; static var voxel_tg3:kha.graphics4.TextureUnit;
#end #end
#if lnx_radiance #if lnx_radiance
static var voxel_th3:kha.compute.TextureUnit; static var voxel_th3:kha.graphics4.TextureUnit;
#end #end
static var voxel_ca3:kha.compute.ConstantLocation; static var voxel_ca3:kha.graphics4.ConstantLocation;
static var voxel_cb3:kha.compute.ConstantLocation; static var voxel_cb3:kha.graphics4.ConstantLocation;
static var voxel_cc3:kha.compute.ConstantLocation; static var voxel_cc3:kha.graphics4.ConstantLocation;
static var voxel_cd3:kha.compute.ConstantLocation; static var voxel_cd3:kha.graphics4.ConstantLocation;
static var voxel_ce3:kha.compute.ConstantLocation; static var voxel_ce3:kha.graphics4.ConstantLocation;
#if lnx_irradiance #if lnx_irradiance
static var voxel_cf3:kha.compute.ConstantLocation; static var voxel_cf3:kha.graphics4.ConstantLocation;
#end #end
#if lnx_radiance #if lnx_radiance
static var voxel_cg3:kha.compute.ConstantLocation; static var voxel_cg3:kha.graphics4.ConstantLocation;
#end #end
#if lnx_envcol #if lnx_envcol
static var voxel_ch3:kha.compute.ConstantLocation; static var voxel_ch3:kha.graphics4.ConstantLocation;
#end #end
#if (rp_voxels == "Voxel GI") #if (rp_voxels == "Voxel GI")
static var voxel_sh4:kha.compute.Shader = null; static var voxel_sh4:kha.graphics4.ComputeShader = null;
static var voxel_ta4:kha.compute.TextureUnit; static var voxel_ta4:kha.graphics4.TextureUnit;
static var voxel_tb4:kha.compute.TextureUnit; static var voxel_tb4:kha.graphics4.TextureUnit;
static var voxel_tc4:kha.compute.TextureUnit; static var voxel_tc4:kha.graphics4.TextureUnit;
static var voxel_td4:kha.compute.TextureUnit; static var voxel_td4:kha.graphics4.TextureUnit;
static var voxel_te4:kha.compute.TextureUnit; static var voxel_te4:kha.graphics4.TextureUnit;
static var voxel_tf4:kha.compute.TextureUnit; static var voxel_tf4:kha.graphics4.TextureUnit;
static var voxel_ca4:kha.compute.ConstantLocation; static var voxel_ca4:kha.graphics4.ConstantLocation;
static var voxel_cb4:kha.compute.ConstantLocation; static var voxel_cb4:kha.graphics4.ConstantLocation;
static var voxel_cc4:kha.compute.ConstantLocation; static var voxel_cc4:kha.graphics4.ConstantLocation;
static var voxel_cd4:kha.compute.ConstantLocation; static var voxel_cd4:kha.graphics4.ConstantLocation;
static var voxel_sh5:kha.compute.Shader = null; static var voxel_sh5:kha.graphics4.ComputeShader = null;
static var voxel_ta5:kha.compute.TextureUnit; static var voxel_ta5:kha.graphics4.TextureUnit;
static var voxel_te5:kha.compute.TextureUnit; static var voxel_te5:kha.graphics4.TextureUnit;
static var voxel_tf5:kha.compute.TextureUnit; static var voxel_tf5:kha.graphics4.TextureUnit;
static var voxel_tg5:kha.compute.TextureUnit; static var voxel_tg5:kha.graphics4.TextureUnit;
static var voxel_ca5:kha.compute.ConstantLocation; static var voxel_ca5:kha.graphics4.ConstantLocation;
static var voxel_cb5:kha.compute.ConstantLocation; static var voxel_cb5:kha.graphics4.ConstantLocation;
static var voxel_cc5:kha.compute.ConstantLocation; static var voxel_cc5:kha.graphics4.ConstantLocation;
static var voxel_cd5:kha.compute.ConstantLocation; static var voxel_cd5:kha.graphics4.ConstantLocation;
static var voxel_ce5:kha.compute.ConstantLocation; static var voxel_ce5:kha.graphics4.ConstantLocation;
static var voxel_cf5:kha.compute.ConstantLocation; static var voxel_cf5:kha.graphics4.ConstantLocation;
static var voxel_cg5:kha.compute.ConstantLocation; static var voxel_cg5:kha.graphics4.ConstantLocation;
#if rp_shadowmap #if rp_shadowmap
static var voxel_tb5:kha.compute.TextureUnit; static var voxel_tb5:kha.graphics4.TextureUnit;
static var voxel_tc5:kha.compute.TextureUnit; static var voxel_tc5:kha.graphics4.TextureUnit;
static var voxel_td5:kha.compute.TextureUnit; static var voxel_td5:kha.graphics4.TextureUnit;
static var voxel_ch5:kha.compute.ConstantLocation; static var voxel_ch5:kha.graphics4.ConstantLocation;
static var voxel_ci5:kha.compute.ConstantLocation; static var voxel_ci5:kha.graphics4.ConstantLocation;
static var voxel_cj5:kha.compute.ConstantLocation; static var voxel_cj5:kha.graphics4.ConstantLocation;
static var voxel_ck5:kha.compute.ConstantLocation; static var voxel_ck5:kha.graphics4.ConstantLocation;
#if lnx_shadowmap_atlas #if lnx_shadowmap_atlas
static var voxel_cl5:kha.compute.ConstantLocation; static var voxel_cl5:kha.graphics4.ConstantLocation;
static var voxel_cm5:kha.compute.ConstantLocation; static var voxel_cm5:kha.graphics4.ConstantLocation;
#end #end
#end #end
#end //rp_voxels == "Voxel GI" #end //rp_voxels == "Voxel GI"
@ -852,7 +852,8 @@ class Inc {
#end #end
#if (rp_voxels != "Off") #if (rp_voxels != "Off")
public static function computeVoxelsBegin() { // TODO: set global graphic
public static function computeVoxelsBegin(g: kha.graphics4.Graphics) {
if (voxel_sh0 == null) if (voxel_sh0 == null)
{ {
voxel_sh0 = path.getComputeShader("voxel_offsetprev"); voxel_sh0 = path.getComputeShader("voxel_offsetprev");
@ -984,16 +985,16 @@ class Inc {
#end #end
} }
public static function computeVoxelsOffsetPrev() { public static function computeVoxelsOffsetPrev(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var clipmaps = iron.RenderPath.clipmaps; var clipmaps = iron.RenderPath.clipmaps;
var clipmap = clipmaps[iron.RenderPath.clipmapLevel]; var clipmap = clipmaps[iron.RenderPath.clipmapLevel];
kha.compute.Compute.setShader(voxel_sh0); g.setComputeShader(voxel_sh0);
kha.compute.Compute.setTexture(voxel_ta0, rts.get("voxelsOut").image, kha.compute.Access.Read); g.setImageTexture(voxel_ta0, rts.get("voxelsOut").image);
kha.compute.Compute.setTexture(voxel_tb0, rts.get("voxelsOutB").image, kha.compute.Access.Write); g.setImageTexture(voxel_tb0, rts.get("voxelsOutB").image);
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
for (i in 0...Main.voxelgiClipmapCount) { for (i in 0...Main.voxelgiClipmapCount) {
@ -1009,14 +1010,14 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca0, fa); g.setFloats(voxel_ca0, fa);
kha.compute.Compute.setInt(voxel_cb0, iron.RenderPath.clipmapLevel); g.setInt(voxel_cb0, iron.RenderPath.clipmapLevel);
kha.compute.Compute.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8)); g.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8));
} }
public static function computeVoxelsTemporal() { public static function computeVoxelsTemporal(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var camera = iron.Scene.active.camera; var camera = iron.Scene.active.camera;
@ -1037,19 +1038,19 @@ class Inc {
// Check again after init // Check again after init
if (rts.get("voxels") == null || rts.get("voxelsOutB") == null || rts.get("voxelsOut") == null) return; if (rts.get("voxels") == null || rts.get("voxelsOutB") == null || rts.get("voxelsOut") == null) return;
kha.compute.Compute.setShader(voxel_sh1); g.setComputeShader(voxel_sh1);
kha.compute.Compute.setTexture(voxel_ta1, rts.get("voxels").image, kha.compute.Access.Read); g.setImageTexture(voxel_ta1, rts.get("voxels").image);
kha.compute.Compute.setTexture(voxel_tb1, rts.get("voxelsOutB").image, kha.compute.Access.Read); g.setImageTexture(voxel_tb1, rts.get("voxelsOutB").image);
kha.compute.Compute.setTexture(voxel_tc1, rts.get("voxelsOut").image, kha.compute.Access.Write); g.setImageTexture(voxel_tc1, rts.get("voxelsOut").image);
#if (rp_voxels == "Voxel GI") #if (rp_voxels == "Voxel GI")
kha.compute.Compute.setSampledTexture(voxel_td1, rts.get("voxelsOutB").image); g.setTexture(voxel_td1, rts.get("voxelsOutB").image);
kha.compute.Compute.setTexture(voxel_te1, rts.get("voxelsLight").image, kha.compute.Access.Read); g.setImageTexture(voxel_te1, rts.get("voxelsLight").image);
kha.compute.Compute.setTexture(voxel_tf1, rts.get("voxelsSDF").image, kha.compute.Access.Write); g.setImageTexture(voxel_tf1, rts.get("voxelsSDF").image);
kha.compute.Compute.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); g.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength);
#else #else
#if lnx_voxelgi_shadows #if lnx_voxelgi_shadows
kha.compute.Compute.setTexture(voxel_te1, rts.get("voxelsSDF").image, kha.compute.Access.Write); g.setImageTexture(voxel_te1, rts.get("voxelsSDF").image);
#end #end
#end #end
@ -1067,17 +1068,17 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca1, fa); g.setFloats(voxel_ca1, fa);
kha.compute.Compute.setInt(voxel_cb1, iron.RenderPath.clipmapLevel); g.setInt(voxel_cb1, iron.RenderPath.clipmapLevel);
kha.compute.Compute.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); g.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength);
kha.compute.Compute.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8)); g.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8));
} }
#if (lnx_voxelgi_shadows || (rp_voxels == "Voxel GI")) #if (lnx_voxelgi_shadows || (rp_voxels == "Voxel GI"))
public static function computeVoxelsSDF() { public static function computeVoxelsSDF(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var clipmaps = iron.RenderPath.clipmaps; var clipmaps = iron.RenderPath.clipmaps;
@ -1089,10 +1090,10 @@ class Inc {
var passcount = Std.int(Math.ceil(Math.log(res) / Math.log(2.0))); var passcount = Std.int(Math.ceil(Math.log(res) / Math.log(2.0)));
for (i in 0...passcount) { for (i in 0...passcount) {
kha.compute.Compute.setShader(voxel_sh2); g.setComputeShader(voxel_sh2);
kha.compute.Compute.setTexture(voxel_ta2, rts.get(read_sdf).image, kha.compute.Access.Read); g.setImageTexture(voxel_ta2, rts.get(read_sdf).image);
kha.compute.Compute.setTexture(voxel_tb2, rts.get(write_sdf).image, kha.compute.Access.Write); g.setImageTexture(voxel_tb2, rts.get(write_sdf).image);
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
for (i in 0...Main.voxelgiClipmapCount) { for (i in 0...Main.voxelgiClipmapCount) {
@ -1108,14 +1109,14 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca2, fa); g.setFloats(voxel_ca2, fa);
kha.compute.Compute.setInt(voxel_cb2, iron.RenderPath.clipmapLevel); g.setInt(voxel_cb2, iron.RenderPath.clipmapLevel);
var jump_size = Math.pow(2.0, passcount - i - 1); var jump_size = Math.pow(2.0, passcount - i - 1);
kha.compute.Compute.setFloat(voxel_cc2, jump_size); g.setFloat(voxel_cc2, jump_size);
kha.compute.Compute.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8)); g.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8));
if (i < passcount - 1) if (i < passcount - 1)
{ {
@ -1127,29 +1128,29 @@ class Inc {
#end #end
#if (rp_voxels == "Voxel AO") #if (rp_voxels == "Voxel AO")
public static function resolveAO() { public static function resolveAO(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var camera = iron.Scene.active.camera; var camera = iron.Scene.active.camera;
var clipmaps = iron.RenderPath.clipmaps; var clipmaps = iron.RenderPath.clipmaps;
var clipmap = clipmaps[iron.RenderPath.clipmapLevel]; var clipmap = clipmaps[iron.RenderPath.clipmapLevel];
kha.compute.Compute.setShader(voxel_sh3); g.setComputeShader(voxel_sh3);
kha.compute.Compute.setSampledTexture(voxel_ta3, rts.get("voxelsOut").image); g.setTexture(voxel_ta3, rts.get("voxelsOut").image);
kha.compute.Compute.setSampledTexture(voxel_tb3, rts.get("half").image); g.setTexture(voxel_tb3, rts.get("half").image);
kha.compute.Compute.setSampledTexture(voxel_tc3, rts.get("gbuffer0").image); g.setTexture(voxel_tc3, rts.get("gbuffer0").image);
kha.compute.Compute.setTexture(voxel_td3, rts.get("voxels_ao").image, kha.compute.Access.Write); g.setImageTexture(voxel_td3, rts.get("voxels_ao").image);
kha.compute.Compute.setSampledTexture(voxel_te3, rts.get("gbuffer1").image); g.setTexture(voxel_te3, rts.get("gbuffer1").image);
#if rp_gbuffer2 #if rp_gbuffer2
kha.compute.Compute.setSampledTexture(voxel_tf3, rts.get("gbuffer2").image); g.setTexture(voxel_tf3, rts.get("gbuffer2").image);
#end #end
#if lnx_brdf #if lnx_brdf
kha.compute.Compute.setSampledTexture(voxel_tg3, iron.Scene.active.embedded.get("brdf.png")); g.setTexture(voxel_tg3, iron.Scene.active.embedded.get("brdf.png"));
#end #end
#if lnx_radiance #if lnx_radiance
kha.compute.Compute.setSampledTexture(voxel_th3, iron.Scene.active.world.probe.radiance); g.setTexture(voxel_th3, iron.Scene.active.world.probe.radiance);
#end #end
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
@ -1166,7 +1167,7 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca3, fa); g.setFloats(voxel_ca3, fa);
#if lnx_centerworld #if lnx_centerworld
m.setFrom(vmat(camera.V)); m.setFrom(vmat(camera.V));
@ -1176,9 +1177,9 @@ class Inc {
m.multmat(camera.P); m.multmat(camera.P);
m.getInverse(m); m.getInverse(m);
kha.compute.Compute.setMatrix(voxel_cb3, m.self); g.setMatrix(voxel_cb3, m.self);
kha.compute.Compute.setFloat3(voxel_cc3, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz()); g.setFloat3(voxel_cc3, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz());
var width = iron.App.w(); var width = iron.App.w();
var height = iron.App.h(); var height = iron.App.h();
@ -1193,17 +1194,17 @@ class Inc {
width = Std.int(dp * Inc.getSuperSampling()); width = Std.int(dp * Inc.getSuperSampling());
} }
} }
kha.compute.Compute.setFloat2(voxel_cd3, width, height); g.setFloat2(voxel_cd3, width, height);
kha.compute.Compute.setFloat(voxel_ce3, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); g.setFloat(voxel_ce3, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength);
#if lnx_irradiance #if lnx_irradiance
var irradiance = iron.Scene.active.world == null ? var irradiance = iron.Scene.active.world == null ?
iron.data.WorldData.getEmptyIrradiance() : iron.data.WorldData.getEmptyIrradiance() :
iron.Scene.active.world.probe.irradiance; iron.Scene.active.world.probe.irradiance;
kha.compute.Compute.setFloats(voxel_cf3, irradiance); g.setFloats(voxel_cf3, irradiance);
#end #end
#if lnx_radiance #if lnx_radiance
kha.compute.Compute.setFloat(voxel_cg3, iron.Scene.active.world != null ? iron.Scene.active.world.probe.raw.radiance_mipmaps + 1 - 2 : 1); g.setFloat(voxel_cg3, iron.Scene.active.world != null ? iron.Scene.active.world.probe.raw.radiance_mipmaps + 1 - 2 : 1);
#end #end
#if lnx_envcol #if lnx_envcol
@ -1217,34 +1218,34 @@ class Inc {
z = camera.data.raw.clear_color[2]; z = camera.data.raw.clear_color[2];
} }
kha.compute.Compute.setFloat3(voxel_ch3, x, y, z); g.setFloat3(voxel_ch3, x, y, z);
#end #end
kha.compute.Compute.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1); g.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1);
} }
#else #else
public static function resolveDiffuse() { public static function resolveDiffuse(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var camera = iron.Scene.active.camera; var camera = iron.Scene.active.camera;
var clipmaps = iron.RenderPath.clipmaps; var clipmaps = iron.RenderPath.clipmaps;
var clipmap = clipmaps[iron.RenderPath.clipmapLevel]; var clipmap = clipmaps[iron.RenderPath.clipmapLevel];
kha.compute.Compute.setShader(voxel_sh3); g.setComputeShader(voxel_sh3);
kha.compute.Compute.setSampledTexture(voxel_ta3, rts.get("voxelsOut").image); g.setTexture(voxel_ta3, rts.get("voxelsOut").image);
kha.compute.Compute.setSampledTexture(voxel_tb3, rts.get("half").image); g.setTexture(voxel_tb3, rts.get("half").image);
kha.compute.Compute.setSampledTexture(voxel_tc3, rts.get("gbuffer0").image); g.setTexture(voxel_tc3, rts.get("gbuffer0").image);
kha.compute.Compute.setTexture(voxel_td3, rts.get("voxels_diffuse").image, kha.compute.Access.Write); g.setImageTexture(voxel_td3, rts.get("voxels_diffuse").image);
kha.compute.Compute.setSampledTexture(voxel_te3, rts.get("gbuffer1").image); g.setTexture(voxel_te3, rts.get("gbuffer1").image);
#if rp_gbuffer2 #if rp_gbuffer2
kha.compute.Compute.setSampledTexture(voxel_tf3, rts.get("gbuffer2").image); g.setTexture(voxel_tf3, rts.get("gbuffer2").image);
#end #end
#if lnx_brdf #if lnx_brdf
kha.compute.Compute.setSampledTexture(voxel_tg3, iron.Scene.active.embedded.get("brdf.png")); g.setTexture(voxel_tg3, iron.Scene.active.embedded.get("brdf.png"));
#end #end
#if lnx_radiance #if lnx_radiance
kha.compute.Compute.setSampledTexture(voxel_th3, iron.Scene.active.world.probe.radiance); g.setTexture(voxel_th3, iron.Scene.active.world.probe.radiance);
#end #end
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
@ -1261,7 +1262,7 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca3, fa); g.setFloats(voxel_ca3, fa);
#if lnx_centerworld #if lnx_centerworld
m.setFrom(vmat(camera.V)); m.setFrom(vmat(camera.V));
@ -1271,9 +1272,9 @@ class Inc {
m.multmat(camera.P); m.multmat(camera.P);
m.getInverse(m); m.getInverse(m);
kha.compute.Compute.setMatrix(voxel_cb3, m.self); g.setMatrix(voxel_cb3, m.self);
kha.compute.Compute.setFloat3(voxel_cc3, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz()); g.setFloat3(voxel_cc3, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz());
var width = iron.App.w(); var width = iron.App.w();
var height = iron.App.h(); var height = iron.App.h();
@ -1288,17 +1289,17 @@ class Inc {
width = Std.int(dp * Inc.getSuperSampling()); width = Std.int(dp * Inc.getSuperSampling());
} }
} }
kha.compute.Compute.setFloat2(voxel_cd3, width, height); g.setFloat2(voxel_cd3, width, height);
kha.compute.Compute.setFloat(voxel_ce3, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); g.setFloat(voxel_ce3, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength);
#if lnx_irradiance #if lnx_irradiance
var irradiance = iron.Scene.active.world == null ? var irradiance = iron.Scene.active.world == null ?
iron.data.WorldData.getEmptyIrradiance() : iron.data.WorldData.getEmptyIrradiance() :
iron.Scene.active.world.probe.irradiance; iron.Scene.active.world.probe.irradiance;
kha.compute.Compute.setFloats(voxel_cf3, irradiance); g.setFloats(voxel_cf3, irradiance);
#end #end
#if lnx_radiance #if lnx_radiance
kha.compute.Compute.setFloat(voxel_cg3, iron.Scene.active.world != null ? iron.Scene.active.world.probe.raw.radiance_mipmaps + 1 - 2 : 1); g.setFloat(voxel_cg3, iron.Scene.active.world != null ? iron.Scene.active.world.probe.raw.radiance_mipmaps + 1 - 2 : 1);
#end #end
#if lnx_envcol #if lnx_envcol
@ -1312,29 +1313,29 @@ class Inc {
z = camera.data.raw.clear_color[2]; z = camera.data.raw.clear_color[2];
} }
kha.compute.Compute.setFloat3(voxel_ch3, x, y, z); g.setFloat3(voxel_ch3, x, y, z);
#end #end
kha.compute.Compute.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1); g.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1);
} }
#end #end
public static function resolveSpecular() { public static function resolveSpecular(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var camera = iron.Scene.active.camera; var camera = iron.Scene.active.camera;
var clipmaps = iron.RenderPath.clipmaps; var clipmaps = iron.RenderPath.clipmaps;
var clipmap = clipmaps[iron.RenderPath.clipmapLevel]; var clipmap = clipmaps[iron.RenderPath.clipmapLevel];
kha.compute.Compute.setShader(voxel_sh4); g.setComputeShader(voxel_sh4);
kha.compute.Compute.setSampledTexture(voxel_ta4, rts.get("voxelsOut").image); g.setTexture(voxel_ta4, rts.get("voxelsOut").image);
kha.compute.Compute.setSampledTexture(voxel_tb4, rts.get("half").image); g.setTexture(voxel_tb4, rts.get("half").image);
kha.compute.Compute.setSampledTexture(voxel_tc4, rts.get("gbuffer0").image); g.setTexture(voxel_tc4, rts.get("gbuffer0").image);
kha.compute.Compute.setSampledTexture(voxel_td4, rts.get("voxelsSDF").image); g.setTexture(voxel_td4, rts.get("voxelsSDF").image);
kha.compute.Compute.setTexture(voxel_te4, rts.get("voxels_specular").image, kha.compute.Access.Write); g.setImageTexture(voxel_te4, rts.get("voxels_specular").image);
#if rp_gbuffer2 #if rp_gbuffer2
kha.compute.Compute.setSampledTexture(voxel_tf4, rts.get("gbuffer2").image); g.setTexture(voxel_tf4, rts.get("gbuffer2").image);
#end #end
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
@ -1351,7 +1352,7 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca4, fa); g.setFloats(voxel_ca4, fa);
#if lnx_centerworld #if lnx_centerworld
m.setFrom(vmat(camera.V)); m.setFrom(vmat(camera.V));
@ -1361,9 +1362,9 @@ class Inc {
m.multmat(camera.P); m.multmat(camera.P);
m.getInverse(m); m.getInverse(m);
kha.compute.Compute.setMatrix(voxel_cb4, m.self); g.setMatrix(voxel_cb4, m.self);
kha.compute.Compute.setFloat3(voxel_cc4, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz()); g.setFloat3(voxel_cc4, camera.transform.worldx(), camera.transform.worldy(), camera.transform.worldz());
var width = iron.App.w(); var width = iron.App.w();
var height = iron.App.h(); var height = iron.App.h();
@ -1378,13 +1379,13 @@ class Inc {
width = Std.int(dp * Inc.getSuperSampling()); width = Std.int(dp * Inc.getSuperSampling());
} }
} }
kha.compute.Compute.setFloat2(voxel_cd4, width, height); g.setFloat2(voxel_cd4, width, height);
kha.compute.Compute.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1); g.compute(Std.int((width + 7) / 8), Std.int((height + 7) / 8), 1);
} }
#if (rp_voxels == "Voxel GI") #if (rp_voxels == "Voxel GI")
public static function computeVoxelsLight() { public static function computeVoxelsLight(g: kha.graphics4.Graphics) {
var rts = path.renderTargets; var rts = path.renderTargets;
var res = iron.RenderPath.getVoxelRes(); var res = iron.RenderPath.getVoxelRes();
var camera = iron.Scene.active.camera; var camera = iron.Scene.active.camera;
@ -1398,12 +1399,12 @@ class Inc {
if (!l.visible) continue; if (!l.visible) continue;
path.light = l; path.light = l;
kha.compute.Compute.setShader(voxel_sh5); g.setComputeShader(voxel_sh5);
kha.compute.Compute.setTexture(voxel_ta5, rts.get("voxelsLight").image, kha.compute.Access.Write); g.setImageTexture(voxel_ta5, rts.get("voxelsLight").image);
kha.compute.Compute.setTexture(voxel_te5, rts.get("voxels").image, kha.compute.Access.Read); g.setImageTexture(voxel_te5, rts.get("voxels").image);
kha.compute.Compute.setSampledTexture(voxel_tf5, rts.get("voxelsOut").image); g.setTexture(voxel_tf5, rts.get("voxelsOut").image);
kha.compute.Compute.setSampledTexture(voxel_tg5, rts.get("voxelsSDF").image); g.setTexture(voxel_tg5, rts.get("voxelsSDF").image);
var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10);
for (i in 0...Main.voxelgiClipmapCount) { for (i in 0...Main.voxelgiClipmapCount) {
@ -1419,48 +1420,48 @@ class Inc {
fa[i * 10 + 9] = clipmaps[i].offset_prev.z; fa[i * 10 + 9] = clipmaps[i].offset_prev.z;
} }
kha.compute.Compute.setFloats(voxel_ca5, fa); g.setFloats(voxel_ca5, fa);
kha.compute.Compute.setInt(voxel_cb5, iron.RenderPath.clipmapLevel); g.setInt(voxel_cb5, iron.RenderPath.clipmapLevel);
#if rp_shadowmap #if rp_shadowmap
if (l.data.raw.type == "sun") { if (l.data.raw.type == "sun") {
#if lnx_shadowmap_atlas #if lnx_shadowmap_atlas
#if lnx_shadowmap_atlas_single_map #if lnx_shadowmap_atlas_single_map
kha.compute.Compute.setSampledTexture(voxel_tb5, rts.get("shadowMapAtlas").image); g.setTexture(voxel_tb5, rts.get("shadowMapAtlas").image);
#else #else
kha.compute.Compute.setSampledTexture(voxel_tb5, rts.get("shadowMapAtlasSun").image); g.setTexture(voxel_tb5, rts.get("shadowMapAtlasSun").image);
#end #end
#else #else
kha.compute.Compute.setSampledTexture(voxel_tb5, rts.get("shadowMap").image); g.setTexture(voxel_tb5, rts.get("shadowMap").image);
#end #end
kha.compute.Compute.setInt(voxel_ch5, 1); // lightShadow g.setInt(voxel_ch5, 1); // lightShadow
} }
else if (l.data.raw.type == "spot" || l.data.raw.type == "area") { else if (l.data.raw.type == "spot" || l.data.raw.type == "area") {
#if lnx_shadowmap_atlas #if lnx_shadowmap_atlas
#if lnx_shadowmap_atlas_single_map #if lnx_shadowmap_atlas_single_map
kha.compute.Compute.setSampledTexture(voxel_tc5, rts.get("shadowMapAtlas").image); g.setTexture(voxel_tc5, rts.get("shadowMapAtlas").image);
#else #else
kha.compute.Compute.setSampledTexture(voxel_tc5, rts.get("shadowMapAtlasSpot").image); g.setTexture(voxel_tc5, rts.get("shadowMapAtlasSpot").image);
#end #end
#else #else
kha.compute.Compute.setSampledTexture(voxel_tc5, rts.get("shadowMapSpot[" + lightIndex + "]").image); g.setTexture(voxel_tc5, rts.get("shadowMapSpot[" + lightIndex + "]").image);
#end #end
kha.compute.Compute.setInt(voxel_ch5, 2); g.setInt(voxel_ch5, 2);
} }
else { else {
#if lnx_shadowmap_atlas #if lnx_shadowmap_atlas
#if lnx_shadowmap_atlas_single_map #if lnx_shadowmap_atlas_single_map
kha.compute.Compute.setSampledTexture(voxel_td5, rts.get("shadowMapAtlas").image); g.setTexture(voxel_td5, rts.get("shadowMapAtlas").image);
#else #else
kha.compute.Compute.setSampledTexture(voxel_td5, rts.get("shadowMapAtlasPoint").image); g.setTexture(voxel_td5, rts.get("shadowMapAtlasPoint").image);
kha.compute.Compute.setInt(voxel_cl5, i); g.setInt(voxel_cl5, i);
kha.compute.Compute.setFloats(voxel_cm5, iron.object.LightObject.pointLightsData); g.setFloats(voxel_cm5, iron.object.LightObject.pointLightsData);
#end #end
#else #else
kha.compute.Compute.setSampledCubeMap(voxel_td5, rts.get("shadowMapPoint[" + lightIndex + "]").cubeMap); g.setCubeMap(voxel_td5, rts.get("shadowMapPoint[" + lightIndex + "]").cubeMap);
#end #end
kha.compute.Compute.setInt(voxel_ch5, 3); g.setInt(voxel_ch5, 3);
} }
// lightProj // lightProj
@ -1472,7 +1473,7 @@ class Inc {
var c:kha.FastFloat = f2 * far * near; var c:kha.FastFloat = f2 * far * near;
var vx:kha.FastFloat = a / b; var vx:kha.FastFloat = a / b;
var vy:kha.FastFloat = c / b; var vy:kha.FastFloat = c / b;
kha.compute.Compute.setFloat2(voxel_ci5, vx, vy); g.setFloat2(voxel_ci5, vx, vy);
// LVP // LVP
m.setFrom(l.VP); m.setFrom(l.VP);
m.multmat(iron.object.Uniforms.biasMat); m.multmat(iron.object.Uniforms.biasMat);
@ -1496,29 +1497,29 @@ class Inc {
#end #end
} }
#end #end
kha.compute.Compute.setMatrix(voxel_cj5, m.self); g.setMatrix(voxel_cj5, m.self);
// shadowsBias // shadowsBias
kha.compute.Compute.setFloat(voxel_ck5, l.data.raw.shadows_bias); g.setFloat(voxel_ck5, l.data.raw.shadows_bias);
#end // rp_shadowmap #end // rp_shadowmap
// lightPos // lightPos
kha.compute.Compute.setFloat3(voxel_cc5, l.transform.worldx(), l.transform.worldy(), l.transform.worldz()); g.setFloat3(voxel_cc5, l.transform.worldx(), l.transform.worldy(), l.transform.worldz());
// lightCol // lightCol
var f = l.data.raw.strength; var f = l.data.raw.strength;
kha.compute.Compute.setFloat3(voxel_cd5, l.data.raw.color[0] * f, l.data.raw.color[1] * f, l.data.raw.color[2] * f); g.setFloat3(voxel_cd5, l.data.raw.color[0] * f, l.data.raw.color[1] * f, l.data.raw.color[2] * f);
// lightType // lightType
kha.compute.Compute.setInt(voxel_ce5, iron.data.LightData.typeToInt(l.data.raw.type)); g.setInt(voxel_ce5, iron.data.LightData.typeToInt(l.data.raw.type));
// lightDir // lightDir
var v = l.look(); var v = l.look();
kha.compute.Compute.setFloat3(voxel_cf5, v.x, v.y, v.z); g.setFloat3(voxel_cf5, v.x, v.y, v.z);
// spotData // spotData
if (l.data.raw.type == "spot") { if (l.data.raw.type == "spot") {
var vx = l.data.raw.spot_size; var vx = l.data.raw.spot_size;
var vy = vx - l.data.raw.spot_blend; var vy = vx - l.data.raw.spot_blend;
kha.compute.Compute.setFloat2(voxel_cg5, vx, vy); g.setFloat2(voxel_cg5, vx, vy);
} }
kha.compute.Compute.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8)); g.compute(Std.int(res / 8), Std.int(res / 8), Std.int(res / 8));
if (!iron.object.LightObject.discardLightCulled(l)) { if (!iron.object.LightObject.discardLightCulled(l)) {
lightIndex++; lightIndex++;

19
leenkx/Sources/leenkx/renderpath/RenderPathDeferred.hx
View File

@ -676,8 +676,10 @@ class RenderPathDeferred {
if (leenkx.data.Config.raw.rp_gi != false) if (leenkx.data.Config.raw.rp_gi != false)
{ {
var path = RenderPath.active; var path = RenderPath.active;
// TODO: investigate currentTarget
var g = path.frameG;
Inc.computeVoxelsBegin(); Inc.computeVoxelsBegin(g);
if (iron.RenderPath.pre_clear == true) if (iron.RenderPath.pre_clear == true)
{ {
@ -693,7 +695,7 @@ class RenderPathDeferred {
else else
{ {
path.clearImage("voxels", 0x00000000); path.clearImage("voxels", 0x00000000);
Inc.computeVoxelsOffsetPrev(); Inc.computeVoxelsOffsetPrev(g);
} }
path.setTarget(""); path.setTarget("");
@ -713,11 +715,11 @@ class RenderPathDeferred {
path.drawMeshes("voxel"); path.drawMeshes("voxel");
Inc.computeVoxelsTemporal(); Inc.computeVoxelsTemporal(g);
#if (rp_voxels == "Voxel GI") #if (rp_voxels == "Voxel GI")
Inc.computeVoxelsLight(); Inc.computeVoxelsLight(g);
Inc.computeVoxelsSDF(); Inc.computeVoxelsSDF(g);
#end #end
if (iron.RenderPath.res_pre_clear == true) { if (iron.RenderPath.res_pre_clear == true) {
@ -777,15 +779,16 @@ class RenderPathDeferred {
#if (rp_voxels != "Off") #if (rp_voxels != "Off")
if (leenkx.data.Config.raw.rp_gi != false) if (leenkx.data.Config.raw.rp_gi != false)
{ {
var g = path.frameG;
#if (lnx_config && (rp_voxels == "Voxel AO")) #if (lnx_config && (rp_voxels == "Voxel AO"))
voxelao_pass = true; voxelao_pass = true;
#end #end
#if (rp_voxels == "Voxel AO") #if (rp_voxels == "Voxel AO")
Inc.resolveAO(); Inc.resolveAO(g);
path.bindTarget("voxels_ao", "voxels_ao"); path.bindTarget("voxels_ao", "voxels_ao");
#else #else
Inc.resolveDiffuse(); Inc.resolveDiffuse(g);
Inc.resolveSpecular(); Inc.resolveSpecular(g);
path.bindTarget("voxels_diffuse", "voxels_diffuse"); path.bindTarget("voxels_diffuse", "voxels_diffuse");
path.bindTarget("voxels_specular", "voxels_specular"); path.bindTarget("voxels_specular", "voxels_specular");
#end #end