diff --git a/Kha/Backends/Kore-HL/kha/korehl/graphics4/Graphics.hx b/Kha/Backends/Kore-HL/kha/korehl/graphics4/Graphics.hx index bc10fe54..e664054d 100644 --- a/Kha/Backends/Kore-HL/kha/korehl/graphics4/Graphics.hx +++ b/Kha/Backends/Kore-HL/kha/korehl/graphics4/Graphics.hx @@ -266,7 +266,6 @@ class Graphics implements kha.graphics4.Graphics { } 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); } diff --git a/Kha/Backends/Krom/kha/graphics4/ComputeShader.hx b/Kha/Backends/Krom/kha/graphics4/ComputeShader.hx index b86268fc..4b624ddb 100644 --- a/Kha/Backends/Krom/kha/graphics4/ComputeShader.hx +++ b/Kha/Backends/Krom/kha/graphics4/ComputeShader.hx @@ -1,22 +1,25 @@ -package kha.graphics4; - -import haxe.io.Bytes; -import kha.Blob; - -class ComputeShader { - public function new(sources: Array, files: Array) { - - } - - public function delete(): Void { - - } - - public function getConstantLocation(name: String): ConstantLocation { - return null; - } - - public function getTextureUnit(name: String): TextureUnit { - return null; - } -} +package kha.graphics4; + +import haxe.io.Bytes; +import kha.Blob; + +class ComputeShader { + public var shader_: Dynamic; + + public function new(sources: Array, files: Array) { + shader_ = Krom.createShaderCompute(sources[0].toBytes().getData()); + } + + public function delete(): Void { + Krom.deleteShaderCompute(shader_); + shader_ = null; + } + + public function getConstantLocation(name: String): ConstantLocation { + return Krom.getConstantLocationCompute(shader_, name); + } + + public function getTextureUnit(name: String): TextureUnit { + return Krom.getTextureUnitCompute(shader_, name); + } +} diff --git a/Kha/Backends/Krom/kha/krom/Graphics.hx b/Kha/Backends/Krom/kha/krom/Graphics.hx index fe0c796f..3ea614fe 100644 --- a/Kha/Backends/Krom/kha/krom/Graphics.hx +++ b/Kha/Backends/Krom/kha/krom/Graphics.hx @@ -261,10 +261,10 @@ class Graphics implements kha.graphics4.Graphics { } public function setComputeShader(shader: ComputeShader) { - + Krom.setShaderCompute(shader.shader_); } public function compute(x: Int, y: Int, z: Int) { - + Krom.compute(x, y, z); } } diff --git a/Kha/Kore/Tools/freebsd_x64/icon.png b/Kha/Kore/Tools/freebsd_x64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/freebsd_x64/icon.png and b/Kha/Kore/Tools/freebsd_x64/icon.png differ diff --git a/Kha/Kore/Tools/linux_arm/icon.png b/Kha/Kore/Tools/linux_arm/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/linux_arm/icon.png and b/Kha/Kore/Tools/linux_arm/icon.png differ diff --git a/Kha/Kore/Tools/linux_arm64/icon.png b/Kha/Kore/Tools/linux_arm64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/linux_arm64/icon.png and b/Kha/Kore/Tools/linux_arm64/icon.png differ diff --git a/Kha/Kore/Tools/linux_x64/icon.png b/Kha/Kore/Tools/linux_x64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/linux_x64/icon.png and b/Kha/Kore/Tools/linux_x64/icon.png differ diff --git a/Kha/Kore/Tools/macos_arm64/icon.png b/Kha/Kore/Tools/macos_arm64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/macos_arm64/icon.png and b/Kha/Kore/Tools/macos_arm64/icon.png differ diff --git a/Kha/Kore/Tools/macos_x64/icon.png b/Kha/Kore/Tools/macos_x64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/macos_x64/icon.png and b/Kha/Kore/Tools/macos_x64/icon.png differ diff --git a/Kha/Kore/Tools/windows_x64/icon.png b/Kha/Kore/Tools/windows_x64/icon.png index 33cc8469..5dd0e8d7 100644 Binary files a/Kha/Kore/Tools/windows_x64/icon.png and b/Kha/Kore/Tools/windows_x64/icon.png differ diff --git a/leenkx/Sources/iron/RenderPath.hx b/leenkx/Sources/iron/RenderPath.hx index 18a85cee..e79e6e41 100644 --- a/leenkx/Sources/iron/RenderPath.hx +++ b/leenkx/Sources/iron/RenderPath.hx @@ -672,7 +672,7 @@ class RenderPath { } #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"); } #end diff --git a/leenkx/Sources/leenkx/renderpath/Inc.hx b/leenkx/Sources/leenkx/renderpath/Inc.hx index 2e6e76eb..f83d0fb7 100644 --- a/leenkx/Sources/leenkx/renderpath/Inc.hx +++ b/leenkx/Sources/leenkx/renderpath/Inc.hx @@ -22,99 +22,99 @@ class Inc { #end #if (rp_voxels != "Off") - static var voxel_sh0:kha.compute.Shader = null; - static var voxel_sh1:kha.compute.Shader = null; - static var voxel_ta0:kha.compute.TextureUnit; - static var voxel_tb0:kha.compute.TextureUnit; - static var voxel_ca0:kha.compute.ConstantLocation; - static var voxel_cb0:kha.compute.ConstantLocation; - static var voxel_ta1:kha.compute.TextureUnit; - static var voxel_tb1:kha.compute.TextureUnit; - static var voxel_tc1:kha.compute.TextureUnit; + static var voxel_sh0:kha.graphics4.ComputeShader = null; + static var voxel_sh1:kha.graphics4.ComputeShader = null; + static var voxel_ta0:kha.graphics4.TextureUnit; + static var voxel_tb0:kha.graphics4.TextureUnit; + static var voxel_ca0:kha.graphics4.ConstantLocation; + static var voxel_cb0:kha.graphics4.ConstantLocation; + static var voxel_ta1:kha.graphics4.TextureUnit; + static var voxel_tb1:kha.graphics4.TextureUnit; + static var voxel_tc1:kha.graphics4.TextureUnit; static var m = iron.math.Mat4.identity(); - static var voxel_ca1:kha.compute.ConstantLocation; - static var voxel_cb1:kha.compute.ConstantLocation; + static var voxel_ca1:kha.graphics4.ConstantLocation; + static var voxel_cb1:kha.graphics4.ConstantLocation; #if (rp_voxels == "Voxel GI") - static var voxel_td1:kha.compute.TextureUnit; - static var voxel_te1:kha.compute.TextureUnit; - static var voxel_tf1:kha.compute.TextureUnit; - static var voxel_cc1:kha.compute.ConstantLocation; + static var voxel_td1:kha.graphics4.TextureUnit; + static var voxel_te1:kha.graphics4.TextureUnit; + static var voxel_tf1:kha.graphics4.TextureUnit; + static var voxel_cc1:kha.graphics4.ConstantLocation; #else #if lnx_voxelgi_shadows - static var voxel_te1:kha.compute.TextureUnit; + static var voxel_te1:kha.graphics4.TextureUnit; #end #end #if (lnx_voxelgi_shadows || rp_voxels == "Voxel GI") - static var voxel_sh2:kha.compute.Shader = null; - static var voxel_ta2:kha.compute.TextureUnit; - static var voxel_tb2:kha.compute.TextureUnit; - static var voxel_ca2:kha.compute.ConstantLocation; - static var voxel_cb2:kha.compute.ConstantLocation; - static var voxel_cc2:kha.compute.ConstantLocation; + static var voxel_sh2:kha.graphics4.ComputeShader = null; + static var voxel_ta2:kha.graphics4.TextureUnit; + static var voxel_tb2:kha.graphics4.TextureUnit; + static var voxel_ca2:kha.graphics4.ConstantLocation; + static var voxel_cb2:kha.graphics4.ConstantLocation; + static var voxel_cc2:kha.graphics4.ConstantLocation; #end - static var voxel_sh3:kha.compute.Shader = null; - static var voxel_ta3:kha.compute.TextureUnit; - static var voxel_tb3:kha.compute.TextureUnit; - static var voxel_tc3:kha.compute.TextureUnit; - static var voxel_td3:kha.compute.TextureUnit; - static var voxel_te3:kha.compute.TextureUnit; - static var voxel_tf3:kha.compute.TextureUnit; + static var voxel_sh3:kha.graphics4.ComputeShader = null; + static var voxel_ta3:kha.graphics4.TextureUnit; + static var voxel_tb3:kha.graphics4.TextureUnit; + static var voxel_tc3:kha.graphics4.TextureUnit; + static var voxel_td3:kha.graphics4.TextureUnit; + static var voxel_te3:kha.graphics4.TextureUnit; + static var voxel_tf3:kha.graphics4.TextureUnit; #if lnx_brdf - static var voxel_tg3:kha.compute.TextureUnit; + static var voxel_tg3:kha.graphics4.TextureUnit; #end #if lnx_radiance - static var voxel_th3:kha.compute.TextureUnit; + static var voxel_th3:kha.graphics4.TextureUnit; #end - static var voxel_ca3:kha.compute.ConstantLocation; - static var voxel_cb3:kha.compute.ConstantLocation; - static var voxel_cc3:kha.compute.ConstantLocation; - static var voxel_cd3:kha.compute.ConstantLocation; - static var voxel_ce3:kha.compute.ConstantLocation; + static var voxel_ca3:kha.graphics4.ConstantLocation; + static var voxel_cb3:kha.graphics4.ConstantLocation; + static var voxel_cc3:kha.graphics4.ConstantLocation; + static var voxel_cd3:kha.graphics4.ConstantLocation; + static var voxel_ce3:kha.graphics4.ConstantLocation; #if lnx_irradiance - static var voxel_cf3:kha.compute.ConstantLocation; + static var voxel_cf3:kha.graphics4.ConstantLocation; #end #if lnx_radiance - static var voxel_cg3:kha.compute.ConstantLocation; + static var voxel_cg3:kha.graphics4.ConstantLocation; #end #if lnx_envcol - static var voxel_ch3:kha.compute.ConstantLocation; + static var voxel_ch3:kha.graphics4.ConstantLocation; #end #if (rp_voxels == "Voxel GI") - static var voxel_sh4:kha.compute.Shader = null; - static var voxel_ta4:kha.compute.TextureUnit; - static var voxel_tb4:kha.compute.TextureUnit; - static var voxel_tc4:kha.compute.TextureUnit; - static var voxel_td4:kha.compute.TextureUnit; - static var voxel_te4:kha.compute.TextureUnit; - static var voxel_tf4:kha.compute.TextureUnit; - static var voxel_ca4:kha.compute.ConstantLocation; - static var voxel_cb4:kha.compute.ConstantLocation; - static var voxel_cc4:kha.compute.ConstantLocation; - static var voxel_cd4:kha.compute.ConstantLocation; + static var voxel_sh4:kha.graphics4.ComputeShader = null; + static var voxel_ta4:kha.graphics4.TextureUnit; + static var voxel_tb4:kha.graphics4.TextureUnit; + static var voxel_tc4:kha.graphics4.TextureUnit; + static var voxel_td4:kha.graphics4.TextureUnit; + static var voxel_te4:kha.graphics4.TextureUnit; + static var voxel_tf4:kha.graphics4.TextureUnit; + static var voxel_ca4:kha.graphics4.ConstantLocation; + static var voxel_cb4:kha.graphics4.ConstantLocation; + static var voxel_cc4:kha.graphics4.ConstantLocation; + static var voxel_cd4:kha.graphics4.ConstantLocation; - static var voxel_sh5:kha.compute.Shader = null; - static var voxel_ta5:kha.compute.TextureUnit; - static var voxel_te5:kha.compute.TextureUnit; - static var voxel_tf5:kha.compute.TextureUnit; - static var voxel_tg5:kha.compute.TextureUnit; - static var voxel_ca5:kha.compute.ConstantLocation; - static var voxel_cb5:kha.compute.ConstantLocation; - static var voxel_cc5:kha.compute.ConstantLocation; - static var voxel_cd5:kha.compute.ConstantLocation; - static var voxel_ce5:kha.compute.ConstantLocation; - static var voxel_cf5:kha.compute.ConstantLocation; - static var voxel_cg5:kha.compute.ConstantLocation; + static var voxel_sh5:kha.graphics4.ComputeShader = null; + static var voxel_ta5:kha.graphics4.TextureUnit; + static var voxel_te5:kha.graphics4.TextureUnit; + static var voxel_tf5:kha.graphics4.TextureUnit; + static var voxel_tg5:kha.graphics4.TextureUnit; + static var voxel_ca5:kha.graphics4.ConstantLocation; + static var voxel_cb5:kha.graphics4.ConstantLocation; + static var voxel_cc5:kha.graphics4.ConstantLocation; + static var voxel_cd5:kha.graphics4.ConstantLocation; + static var voxel_ce5:kha.graphics4.ConstantLocation; + static var voxel_cf5:kha.graphics4.ConstantLocation; + static var voxel_cg5:kha.graphics4.ConstantLocation; #if rp_shadowmap - static var voxel_tb5:kha.compute.TextureUnit; - static var voxel_tc5:kha.compute.TextureUnit; - static var voxel_td5:kha.compute.TextureUnit; - static var voxel_ch5:kha.compute.ConstantLocation; - static var voxel_ci5:kha.compute.ConstantLocation; - static var voxel_cj5:kha.compute.ConstantLocation; - static var voxel_ck5:kha.compute.ConstantLocation; + static var voxel_tb5:kha.graphics4.TextureUnit; + static var voxel_tc5:kha.graphics4.TextureUnit; + static var voxel_td5:kha.graphics4.TextureUnit; + static var voxel_ch5:kha.graphics4.ConstantLocation; + static var voxel_ci5:kha.graphics4.ConstantLocation; + static var voxel_cj5:kha.graphics4.ConstantLocation; + static var voxel_ck5:kha.graphics4.ConstantLocation; #if lnx_shadowmap_atlas - static var voxel_cl5:kha.compute.ConstantLocation; - static var voxel_cm5:kha.compute.ConstantLocation; + static var voxel_cl5:kha.graphics4.ConstantLocation; + static var voxel_cm5:kha.graphics4.ConstantLocation; #end #end #end //rp_voxels == "Voxel GI" @@ -852,7 +852,8 @@ class Inc { #end #if (rp_voxels != "Off") - public static function computeVoxelsBegin() { + // TODO: set global graphic + public static function computeVoxelsBegin(g: kha.graphics4.Graphics) { if (voxel_sh0 == null) { voxel_sh0 = path.getComputeShader("voxel_offsetprev"); @@ -984,16 +985,16 @@ class Inc { #end } - public static function computeVoxelsOffsetPrev() { + public static function computeVoxelsOffsetPrev(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var clipmaps = iron.RenderPath.clipmaps; 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); - kha.compute.Compute.setTexture(voxel_tb0, rts.get("voxelsOutB").image, kha.compute.Access.Write); + g.setImageTexture(voxel_ta0, rts.get("voxelsOut").image); + g.setImageTexture(voxel_tb0, rts.get("voxelsOutB").image); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); for (i in 0...Main.voxelgiClipmapCount) { @@ -1009,14 +1010,14 @@ class Inc { 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 res = iron.RenderPath.getVoxelRes(); var camera = iron.Scene.active.camera; @@ -1037,19 +1038,19 @@ class Inc { // Check again after init 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); - kha.compute.Compute.setTexture(voxel_tb1, rts.get("voxelsOutB").image, kha.compute.Access.Read); - kha.compute.Compute.setTexture(voxel_tc1, rts.get("voxelsOut").image, kha.compute.Access.Write); + g.setImageTexture(voxel_ta1, rts.get("voxels").image); + g.setImageTexture(voxel_tb1, rts.get("voxelsOutB").image); + g.setImageTexture(voxel_tc1, rts.get("voxelsOut").image); #if (rp_voxels == "Voxel GI") - kha.compute.Compute.setSampledTexture(voxel_td1, rts.get("voxelsOutB").image); - kha.compute.Compute.setTexture(voxel_te1, rts.get("voxelsLight").image, kha.compute.Access.Read); - kha.compute.Compute.setTexture(voxel_tf1, rts.get("voxelsSDF").image, kha.compute.Access.Write); - kha.compute.Compute.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); + g.setTexture(voxel_td1, rts.get("voxelsOutB").image); + g.setImageTexture(voxel_te1, rts.get("voxelsLight").image); + g.setImageTexture(voxel_tf1, rts.get("voxelsSDF").image); + g.setFloat(voxel_cc1, iron.Scene.active.world == null ? 0.0 : iron.Scene.active.world.probe.raw.strength); #else #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 @@ -1067,17 +1068,17 @@ class Inc { 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")) - public static function computeVoxelsSDF() { + public static function computeVoxelsSDF(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var clipmaps = iron.RenderPath.clipmaps; @@ -1089,10 +1090,10 @@ class Inc { var passcount = Std.int(Math.ceil(Math.log(res) / Math.log(2.0))); 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); - kha.compute.Compute.setTexture(voxel_tb2, rts.get(write_sdf).image, kha.compute.Access.Write); + g.setImageTexture(voxel_ta2, rts.get(read_sdf).image); + g.setImageTexture(voxel_tb2, rts.get(write_sdf).image); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); for (i in 0...Main.voxelgiClipmapCount) { @@ -1108,14 +1109,14 @@ class Inc { 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); - 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) { @@ -1127,29 +1128,29 @@ class Inc { #end #if (rp_voxels == "Voxel AO") - public static function resolveAO() { + public static function resolveAO(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var camera = iron.Scene.active.camera; var clipmaps = iron.RenderPath.clipmaps; 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); - kha.compute.Compute.setSampledTexture(voxel_tb3, rts.get("half").image); - kha.compute.Compute.setSampledTexture(voxel_tc3, rts.get("gbuffer0").image); - kha.compute.Compute.setTexture(voxel_td3, rts.get("voxels_ao").image, kha.compute.Access.Write); + g.setTexture(voxel_ta3, rts.get("voxelsOut").image); + g.setTexture(voxel_tb3, rts.get("half").image); + g.setTexture(voxel_tc3, rts.get("gbuffer0").image); + 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 - kha.compute.Compute.setSampledTexture(voxel_tf3, rts.get("gbuffer2").image); + g.setTexture(voxel_tf3, rts.get("gbuffer2").image); #end #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 #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 var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); @@ -1166,7 +1167,7 @@ class Inc { fa[i * 10 + 9] = clipmaps[i].offset_prev.z; } - kha.compute.Compute.setFloats(voxel_ca3, fa); + g.setFloats(voxel_ca3, fa); #if lnx_centerworld m.setFrom(vmat(camera.V)); @@ -1176,9 +1177,9 @@ class Inc { m.multmat(camera.P); 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 height = iron.App.h(); @@ -1193,17 +1194,17 @@ class Inc { 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 var irradiance = iron.Scene.active.world == null ? iron.data.WorldData.getEmptyIrradiance() : iron.Scene.active.world.probe.irradiance; - kha.compute.Compute.setFloats(voxel_cf3, irradiance); + g.setFloats(voxel_cf3, irradiance); #end #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 #if lnx_envcol @@ -1217,34 +1218,34 @@ class Inc { z = camera.data.raw.clear_color[2]; } - kha.compute.Compute.setFloat3(voxel_ch3, x, y, z); + g.setFloat3(voxel_ch3, x, y, z); #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 - public static function resolveDiffuse() { + public static function resolveDiffuse(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var camera = iron.Scene.active.camera; var clipmaps = iron.RenderPath.clipmaps; 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); - kha.compute.Compute.setSampledTexture(voxel_tb3, rts.get("half").image); - kha.compute.Compute.setSampledTexture(voxel_tc3, rts.get("gbuffer0").image); - kha.compute.Compute.setTexture(voxel_td3, rts.get("voxels_diffuse").image, kha.compute.Access.Write); - kha.compute.Compute.setSampledTexture(voxel_te3, rts.get("gbuffer1").image); + g.setTexture(voxel_ta3, rts.get("voxelsOut").image); + g.setTexture(voxel_tb3, rts.get("half").image); + g.setTexture(voxel_tc3, rts.get("gbuffer0").image); + g.setImageTexture(voxel_td3, rts.get("voxels_diffuse").image); + g.setTexture(voxel_te3, rts.get("gbuffer1").image); #if rp_gbuffer2 - kha.compute.Compute.setSampledTexture(voxel_tf3, rts.get("gbuffer2").image); + g.setTexture(voxel_tf3, rts.get("gbuffer2").image); #end #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 #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 var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); @@ -1261,7 +1262,7 @@ class Inc { fa[i * 10 + 9] = clipmaps[i].offset_prev.z; } - kha.compute.Compute.setFloats(voxel_ca3, fa); + g.setFloats(voxel_ca3, fa); #if lnx_centerworld m.setFrom(vmat(camera.V)); @@ -1271,9 +1272,9 @@ class Inc { m.multmat(camera.P); 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 height = iron.App.h(); @@ -1288,17 +1289,17 @@ class Inc { 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 var irradiance = iron.Scene.active.world == null ? iron.data.WorldData.getEmptyIrradiance() : iron.Scene.active.world.probe.irradiance; - kha.compute.Compute.setFloats(voxel_cf3, irradiance); + g.setFloats(voxel_cf3, irradiance); #end #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 #if lnx_envcol @@ -1312,29 +1313,29 @@ class Inc { z = camera.data.raw.clear_color[2]; } - kha.compute.Compute.setFloat3(voxel_ch3, x, y, z); + g.setFloat3(voxel_ch3, x, y, z); #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 - public static function resolveSpecular() { + public static function resolveSpecular(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var camera = iron.Scene.active.camera; var clipmaps = iron.RenderPath.clipmaps; 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); - kha.compute.Compute.setSampledTexture(voxel_tb4, rts.get("half").image); - kha.compute.Compute.setSampledTexture(voxel_tc4, rts.get("gbuffer0").image); - kha.compute.Compute.setSampledTexture(voxel_td4, rts.get("voxelsSDF").image); - kha.compute.Compute.setTexture(voxel_te4, rts.get("voxels_specular").image, kha.compute.Access.Write); + g.setTexture(voxel_ta4, rts.get("voxelsOut").image); + g.setTexture(voxel_tb4, rts.get("half").image); + g.setTexture(voxel_tc4, rts.get("gbuffer0").image); + g.setTexture(voxel_td4, rts.get("voxelsSDF").image); + g.setImageTexture(voxel_te4, rts.get("voxels_specular").image); #if rp_gbuffer2 - kha.compute.Compute.setSampledTexture(voxel_tf4, rts.get("gbuffer2").image); + g.setTexture(voxel_tf4, rts.get("gbuffer2").image); #end var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); @@ -1351,7 +1352,7 @@ class Inc { fa[i * 10 + 9] = clipmaps[i].offset_prev.z; } - kha.compute.Compute.setFloats(voxel_ca4, fa); + g.setFloats(voxel_ca4, fa); #if lnx_centerworld m.setFrom(vmat(camera.V)); @@ -1361,9 +1362,9 @@ class Inc { m.multmat(camera.P); 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 height = iron.App.h(); @@ -1378,13 +1379,13 @@ class Inc { 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") - public static function computeVoxelsLight() { + public static function computeVoxelsLight(g: kha.graphics4.Graphics) { var rts = path.renderTargets; var res = iron.RenderPath.getVoxelRes(); var camera = iron.Scene.active.camera; @@ -1398,12 +1399,12 @@ class Inc { if (!l.visible) continue; 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); - kha.compute.Compute.setTexture(voxel_te5, rts.get("voxels").image, kha.compute.Access.Read); - kha.compute.Compute.setSampledTexture(voxel_tf5, rts.get("voxelsOut").image); - kha.compute.Compute.setSampledTexture(voxel_tg5, rts.get("voxelsSDF").image); + g.setImageTexture(voxel_ta5, rts.get("voxelsLight").image); + g.setImageTexture(voxel_te5, rts.get("voxels").image); + g.setTexture(voxel_tf5, rts.get("voxelsOut").image); + g.setTexture(voxel_tg5, rts.get("voxelsSDF").image); var fa:Float32Array = new Float32Array(Main.voxelgiClipmapCount * 10); for (i in 0...Main.voxelgiClipmapCount) { @@ -1419,48 +1420,48 @@ class Inc { 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 (l.data.raw.type == "sun") { #if lnx_shadowmap_atlas #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 - kha.compute.Compute.setSampledTexture(voxel_tb5, rts.get("shadowMapAtlasSun").image); + g.setTexture(voxel_tb5, rts.get("shadowMapAtlasSun").image); #end #else - kha.compute.Compute.setSampledTexture(voxel_tb5, rts.get("shadowMap").image); + g.setTexture(voxel_tb5, rts.get("shadowMap").image); #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") { #if lnx_shadowmap_atlas #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 - kha.compute.Compute.setSampledTexture(voxel_tc5, rts.get("shadowMapAtlasSpot").image); + g.setTexture(voxel_tc5, rts.get("shadowMapAtlasSpot").image); #end #else - kha.compute.Compute.setSampledTexture(voxel_tc5, rts.get("shadowMapSpot[" + lightIndex + "]").image); + g.setTexture(voxel_tc5, rts.get("shadowMapSpot[" + lightIndex + "]").image); #end - kha.compute.Compute.setInt(voxel_ch5, 2); + g.setInt(voxel_ch5, 2); } else { #if lnx_shadowmap_atlas #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 - kha.compute.Compute.setSampledTexture(voxel_td5, rts.get("shadowMapAtlasPoint").image); - kha.compute.Compute.setInt(voxel_cl5, i); - kha.compute.Compute.setFloats(voxel_cm5, iron.object.LightObject.pointLightsData); + g.setTexture(voxel_td5, rts.get("shadowMapAtlasPoint").image); + g.setInt(voxel_cl5, i); + g.setFloats(voxel_cm5, iron.object.LightObject.pointLightsData); #end #else - kha.compute.Compute.setSampledCubeMap(voxel_td5, rts.get("shadowMapPoint[" + lightIndex + "]").cubeMap); + g.setCubeMap(voxel_td5, rts.get("shadowMapPoint[" + lightIndex + "]").cubeMap); #end - kha.compute.Compute.setInt(voxel_ch5, 3); + g.setInt(voxel_ch5, 3); } // lightProj @@ -1472,7 +1473,7 @@ class Inc { var c:kha.FastFloat = f2 * far * near; var vx:kha.FastFloat = a / b; var vy:kha.FastFloat = c / b; - kha.compute.Compute.setFloat2(voxel_ci5, vx, vy); + g.setFloat2(voxel_ci5, vx, vy); // LVP m.setFrom(l.VP); m.multmat(iron.object.Uniforms.biasMat); @@ -1496,29 +1497,29 @@ class Inc { #end } #end - kha.compute.Compute.setMatrix(voxel_cj5, m.self); + g.setMatrix(voxel_cj5, m.self); // shadowsBias - kha.compute.Compute.setFloat(voxel_ck5, l.data.raw.shadows_bias); + g.setFloat(voxel_ck5, l.data.raw.shadows_bias); #end // rp_shadowmap // 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 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 - 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 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 if (l.data.raw.type == "spot") { var vx = l.data.raw.spot_size; 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)) { lightIndex++; diff --git a/leenkx/Sources/leenkx/renderpath/RenderPathDeferred.hx b/leenkx/Sources/leenkx/renderpath/RenderPathDeferred.hx index 34274eb1..9f44f378 100644 --- a/leenkx/Sources/leenkx/renderpath/RenderPathDeferred.hx +++ b/leenkx/Sources/leenkx/renderpath/RenderPathDeferred.hx @@ -676,8 +676,10 @@ class RenderPathDeferred { if (leenkx.data.Config.raw.rp_gi != false) { var path = RenderPath.active; + // TODO: investigate currentTarget + var g = path.frameG; - Inc.computeVoxelsBegin(); + Inc.computeVoxelsBegin(g); if (iron.RenderPath.pre_clear == true) { @@ -693,7 +695,7 @@ class RenderPathDeferred { else { path.clearImage("voxels", 0x00000000); - Inc.computeVoxelsOffsetPrev(); + Inc.computeVoxelsOffsetPrev(g); } path.setTarget(""); @@ -713,11 +715,11 @@ class RenderPathDeferred { path.drawMeshes("voxel"); - Inc.computeVoxelsTemporal(); + Inc.computeVoxelsTemporal(g); #if (rp_voxels == "Voxel GI") - Inc.computeVoxelsLight(); - Inc.computeVoxelsSDF(); + Inc.computeVoxelsLight(g); + Inc.computeVoxelsSDF(g); #end if (iron.RenderPath.res_pre_clear == true) { @@ -777,15 +779,16 @@ class RenderPathDeferred { #if (rp_voxels != "Off") if (leenkx.data.Config.raw.rp_gi != false) { + var g = path.frameG; #if (lnx_config && (rp_voxels == "Voxel AO")) voxelao_pass = true; #end #if (rp_voxels == "Voxel AO") - Inc.resolveAO(); + Inc.resolveAO(g); path.bindTarget("voxels_ao", "voxels_ao"); #else - Inc.resolveDiffuse(); - Inc.resolveSpecular(); + Inc.resolveDiffuse(g); + Inc.resolveSpecular(g); path.bindTarget("voxels_diffuse", "voxels_diffuse"); path.bindTarget("voxels_specular", "voxels_specular"); #end