From 8d1699bf9273e033eb354d00dc11f4991596c15a Mon Sep 17 00:00:00 2001 From: luboslenco Date: Sat, 22 Aug 2026 13:36:26 +0200 Subject: [PATCH] metal: raytrace fixes --- .../raytrace/raytrace_brute_core.metal | 257 ++++++----- .../raytrace/raytrace_brute_full.metal | 258 ++++++----- ....metal => raytrace_brute_multi_core.metal} | 258 ++++++----- .../raytrace/raytrace_brute_multi_full.metal | 420 ++++++++++++++++++ base/shaders/raytrace/src/build_metal.sh | 6 +- ..._brute_full.metal => raytrace_brute.metal} | 257 ++++++----- base/sources/backends/metal_gpu.h | 2 + base/sources/backends/metal_gpu.m | 99 +++-- paint/sources/render/render_path_raytrace.c | 19 +- 9 files changed, 1073 insertions(+), 503 deletions(-) rename base/shaders/raytrace/{src/raytrace_brute_core.metal => raytrace_brute_multi_core.metal} (58%) create mode 100644 base/shaders/raytrace/raytrace_brute_multi_full.metal rename base/shaders/raytrace/src/{raytrace_brute_full.metal => raytrace_brute.metal} (59%) diff --git a/base/shaders/raytrace/raytrace_brute_core.metal b/base/shaders/raytrace/raytrace_brute_core.metal index 28b94903..1e0de1fa 100644 --- a/base/shaders/raytrace/raytrace_brute_core.metal +++ b/base/shaders/raytrace/raytrace_brute_core.metal @@ -3,10 +3,10 @@ #define _SUBSURFACE #define _TRANSLUCENCY #define _ROULETTE -#define _TRANSPARENCY +// #define _TRANSPARENCY // #define _FRESNEL #endif -#define _RENDER +// #define _RENDER using namespace metal; using namespace raytracing; @@ -18,6 +18,12 @@ struct Vertex { uint tex; }; +#ifdef _MULTI +typedef intersector intersector_t; +#else +typedef intersector intersector_t; +#endif + struct RayGenConstantBuffer { float4 eye; // xyz, frame float4x4 inv_vp; @@ -30,9 +36,9 @@ struct RayPayload { float3 ray_dir; }; -constant int SAMPLES = 2; // 64 +constant int SAMPLES = 8; #ifdef _TRANSLUCENCY -constant int DEPTH = 6; +constant int DEPTH = 16; #else constant int DEPTH = 3; // Opaque hits #endif @@ -100,7 +106,7 @@ float2 s16_to_f32(uint val) { return float2(a, b) / 32767.0f; } -float3 hit_world_position(ray ray, typename intersector::result_type intersection) { +float3 hit_world_position(ray ray, intersector_t::result_type intersection) { return ray.origin + ray.direction * intersection.distance; } @@ -136,19 +142,15 @@ float3 surface_specular(const float3 base_color, const float metalness) { return mix(float3(0.04, 0.04, 0.04), base_color, metalness); } -float3 env_brdf_approx(float3 specular, float roughness, float dotnv) { - const float4 c0 = float4(-1, -0.0275, -0.572, 0.022); - const float4 c1 = float4(1, 0.0425, 1.04, -0.04); - float4 r = roughness * c0 + c1; - float a004 = min(r.x * r.x, exp2(-9.28 * dotnv)) * r.x + r.y; - float2 ab = float2(-1.04, 1.04) * a004 + r.zw; - return specular * ab.x + ab.y; -} - float fresnel(float3 normal, float3 incident) { return mix(0.5, 1.0, pow(1.0 + dot(normal, incident), 5.0)); } +float4 read_texel(texture2d tex, float2 tex_coord) { + uint2 size = uint2(tex.get_width(), tex.get_height()); + return tex.read(uint2(fract(tex_coord) * float2(size)), 0); +} + kernel void raytracingKernel( uint2 tid [[thread_position_in_grid]], constant RayGenConstantBuffer &constant_buffer [[buffer(0)]], @@ -165,24 +167,32 @@ kernel void raytracingKernel( device void *indices [[buffer(2)]], device void *vertices [[buffer(3)]] ) { + uint2 dim = uint2(render_target.get_width(), render_target.get_height()); + if (tid.x >= dim.x || tid.y >= dim.y) { + return; + } + + int frame = int(constant_buffer.eye.w); uint thread_seed = 0; float3 accum = float3(0, 0, 0); for (int j = 0; j < SAMPLES; ++j) { + int sample_index = frame * SAMPLES + j; + // AA float2 xy = float2(tid) + float2(0.5f, 0.5f); - xy.x += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.x += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); thread_seed += 1; - xy.y += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.y += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); - float2 screen_pos = xy / float2(render_target.get_width(), render_target.get_height()) * 2.0 - 1.0; + float2 screen_pos = xy / float2(dim) * 2.0 - 1.0; ray ray; ray.min_distance = 0.0001; - ray.max_distance = 10.0; + ray.max_distance = 100.0; generate_camera_ray(screen_pos, ray.origin, ray.direction, constant_buffer.eye.xyz, constant_buffer.inv_vp); RayPayload payload; - payload.color = float4(1, 1, 1, j); + payload.color = float4(1, 1, 1, sample_index); #ifdef _TRANSPARENCY int transparent_hits = 0; @@ -193,7 +203,7 @@ kernel void raytracingKernel( #ifdef _ROULETTE float rr_factor = 1.0; if (i >= rr_start) { - float f = rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + float f = rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); if (f <= rr_probability) { break; } @@ -201,91 +211,132 @@ kernel void raytracingKernel( } #endif - // #ifdef _SUBSURFACE - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_CULL_BACK_FACING_TRIANGLES, ~0, 0, 1, 0, ray, payload); - // #else - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE, ~0, 0, 1, 0, ray, payload); - // #endif - - intersector in; + intersector_t in; in.assume_geometry_type(geometry_type::triangle); in.force_opacity(forced_opacity::opaque); in.accept_any_intersection(false); - // in.set_triangle_cull_mode(triangle_cull_mode::none); - typename intersector::result_type intersection; + intersector_t::result_type intersection; intersection = in.intersect(ray, scene); + + // Miss if (intersection.type == intersection_type::none) { #ifdef _EMISSION - if (payload.color.a == -2.0) { - return; + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; } #endif - float2 tex_coord = fract(equirect(ray.direction, constant_buffer.params.y)); - float3 texenv = mytexture_env.sample(linear_sampler, tex_coord).rgb * abs(constant_buffer.params.x); - payload.color = float4(payload.color.rgb * texenv.rgb, -1); + if (i == 0 && constant_buffer.params.x < 0.0) { // No envmap + payload.color.rgb = float3(0.0275, 0.0275, 0.0275); + } + else { + float2 tex_coord = equirect(ray.direction, constant_buffer.params.y); + float3 texenv = mytexture_env.sample(linear_sampler, tex_coord, level(0)).rgb * abs(constant_buffer.params.x); + payload.color.rgb *= texenv; + } + + accum += clamp(payload.color.rgb, 0.0, 8.0); + break; } - else { - device uint32_t *inda = (device uint32_t *)(indices); - uint3 indices_sample = uint3( - inda[intersection.primitive_id * 3], - inda[intersection.primitive_id * 3 + 1], - inda[intersection.primitive_id * 3 + 2] - ); - device Vertex *verta = (device Vertex *)(vertices); - float2 vertex_uvs[3] = { - s16_to_f32(verta[indices_sample[0]].tex), - s16_to_f32(verta[indices_sample[1]].tex), - s16_to_f32(verta[indices_sample[2]].tex) - }; - float2 barycentrics = intersection.triangle_barycentric_coord; - float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + device uint32_t *inda = (device uint32_t *)(indices); + uint base_index = intersection.primitive_id * 3; - uint2 size = uint2(mytexture0.get_width(), mytexture0.get_height()); - uint3 utex_coord = uint3(uint2((tex_coord - float2(uint2(tex_coord))) * float2(size)), 0); - float4 texpaint0 = mytexture0.read(utex_coord.xy, utex_coord.z); + #ifdef _MULTI + base_index += intersection.user_instance_id; + #endif - #ifdef _TRANSPARENCY - if (texpaint0.a <= 0.1) { - payload.ray_dir = ray.direction; - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; - payload.color.a = -2; - return; + uint3 indices_sample = uint3( + inda[base_index], + inda[base_index + 1], + inda[base_index + 2] + ); + + device Vertex *verta = (device Vertex *)(vertices); + float2 vertex_uvs[3] = { + s16_to_f32(verta[indices_sample[0]].tex), + s16_to_f32(verta[indices_sample[1]].tex), + s16_to_f32(verta[indices_sample[2]].tex) + }; + float2 barycentrics = intersection.triangle_barycentric_coord; + float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + + float3 hit = hit_world_position(ray, intersection); + float4 texpaint0 = read_texel(mytexture0, tex_coord); + + #ifdef _TRANSPARENCY + if (texpaint0.a <= 0.01) { + ray.origin = hit + ray.direction * 0.0001f; + if (transparent_hits < DEPTH_TRANSPARENT) { + payload.color.a = sample_index; + transparent_hits++; + i--; } - #endif + else { + payload.color.a = -2; + } + continue; + } + #endif - float3 vertex_normals[3] = { - float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), - float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), - float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) - }; - float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); + float3 vertex_normals[3] = { + float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), + float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), + float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) + }; + float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); - float4 texpaint1 = mytexture1.read(utex_coord.xy, utex_coord.z); - float4 texpaint2 = mytexture2.read(utex_coord.xy, utex_coord.z); - float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); + #ifdef _MULTI + float4x3 obj_to_world = intersection.object_to_world_transform; + n = normalize(float3x3(obj_to_world[0], obj_to_world[1], obj_to_world[2]) * n); + #endif - float3 tangent = float3(0, 0, 0); - float3 binormal = float3(0, 0, 0); - create_basis(n, tangent, binormal); + float4 texpaint1 = read_texel(mytexture1, tex_coord); + float4 texpaint2 = read_texel(mytexture2, tex_coord); + float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); - texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); - texpaint1.g = -texpaint1.g; - n = float3x3(tangent, binormal, n) * texpaint1.rgb; + #ifdef _TRANSLUCENCY + if (!intersection.triangle_front_facing) { + float3 absorption = pow(max(texcolor, float3(0.001)), float3(intersection.distance * texpaint0.a)); + payload.color.rgb *= absorption; + } + #endif - float f = rand(tid.x, tid.y, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - thread_seed += 1; + float3 tangent = float3(0, 0, 0); + float3 binormal = float3(0, 0, 0); + create_basis(n, tangent, binormal); + texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); + texpaint1.g = -texpaint1.g; + n = float3x3(tangent, binormal, n) * texpaint1.rgb; + + uint bounce_seed = 0; + + float f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + + bool scatter = false; + + #ifdef _TRANSLUCENCY + if (f > texpaint0.a) { + float roughness = texpaint2.g; + float3 scatter_dir = cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + payload.ray_dir = normalize(mix(ray.direction, scatter_dir, roughness * roughness * 0.5)); + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + scatter = true; + } + #endif + + if (!scatter) { #ifdef _TRANSLUCENCY - float3 diffuse_dir = texpaint0.a < f ? - cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank) : - cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - #else - float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; #endif + float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + #ifdef _FRESNEL float specular_chance = fresnel(n, ray.direction); #else @@ -293,12 +344,7 @@ kernel void raytracingKernel( #endif if (f < specular_chance) { - #ifdef _TRANSLUCENCY - float3 specular_dir = texpaint0.a < f * 2 ? ray.direction : reflect(ray.direction, n); - #else float3 specular_dir = reflect(ray.direction, n); - #endif - payload.ray_dir = mix(specular_dir, diffuse_dir, texpaint2.g * texpaint2.g); float3 specular = surface_specular(texcolor, texpaint2.b); payload.color.xyz *= specular; @@ -314,59 +360,44 @@ kernel void raytracingKernel( payload.color.xyz /= 1.0 - specular_chance; #endif } + #ifdef _FRESNEL payload.color.xyz *= 0.5; #endif - // float dotnv = abs(dot(n, -WorldRayDirection())); - // payload.ray_origin = hit_world_position() + n * mix(0.1f, 0.0001f, dotnv); - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; + payload.ray_origin = hit + payload.ray_dir * 0.0001f; #ifdef _EMISSION if (int(texpaint1.a * 255.0f) % 3 == 1) { // matid payload.color.xyz *= 100.0f; - payload.color.a = -2.0; + payload.color.a = -3.0; } #endif #ifdef _SUBSURFACE if (int(texpaint1.a * 255.0f) % 3 == 2) { - payload.ray_origin += ray.direction * f; + float d = min(1.0 / min(intersection.distance * 2.0, 1.0) / 10.0, 0.5); + payload.color.xyz += payload.color.xyz * d; + if (f < 0.5) { + payload.ray_origin += ray.direction * f * 0.001; + } } #endif } #ifdef _EMISSION - if (payload.color.a == -2) { + if (payload.color.a == -3.0) { accum += payload.color.rgb; break; } #endif - // Miss - if (payload.color.a < 0) { - #ifdef _TRANSPARENCY - if (payload.color.a == -2 && transparent_hits < DEPTH_TRANSPARENT) { - payload.color.a = j; - transparent_hits++; - i--; - } - #endif - - if (i == 0 && constant_buffer.params.x < 0) { // No envmap - payload.color.rgb = float3(0.032, 0.032, 0.032); - } - - accum += clamp(payload.color.rgb, 0.0, 8.0); - break; - } + ray.origin = payload.ray_origin; + ray.direction = payload.ray_dir; #ifdef _ROULETTE payload.color.rgb *= rr_factor; #endif - - ray.origin = payload.ray_origin; - ray.direction = payload.ray_dir; } } @@ -382,6 +413,6 @@ kernel void raytracingKernel( if (constant_buffer.eye.w == 0) { color = accum; } - render_target.write(float4(mix(color, accum, 1.0 / 16.0), 1.0f), tid); + render_target.write(float4(mix(color, accum, 1.0 / 4.0), 1.0f), tid); #endif } diff --git a/base/shaders/raytrace/raytrace_brute_full.metal b/base/shaders/raytrace/raytrace_brute_full.metal index 52cf871c..705736d0 100644 --- a/base/shaders/raytrace/raytrace_brute_full.metal +++ b/base/shaders/raytrace/raytrace_brute_full.metal @@ -1,14 +1,13 @@ #define _FULL - #ifdef _FULL #define _EMISSION #define _SUBSURFACE #define _TRANSLUCENCY #define _ROULETTE -#define _TRANSPARENCY +// #define _TRANSPARENCY // #define _FRESNEL #endif -#define _RENDER +// #define _RENDER using namespace metal; using namespace raytracing; @@ -20,6 +19,12 @@ struct Vertex { uint tex; }; +#ifdef _MULTI +typedef intersector intersector_t; +#else +typedef intersector intersector_t; +#endif + struct RayGenConstantBuffer { float4 eye; // xyz, frame float4x4 inv_vp; @@ -32,9 +37,9 @@ struct RayPayload { float3 ray_dir; }; -constant int SAMPLES = 2; // 64 +constant int SAMPLES = 8; #ifdef _TRANSLUCENCY -constant int DEPTH = 6; +constant int DEPTH = 16; #else constant int DEPTH = 3; // Opaque hits #endif @@ -102,7 +107,7 @@ float2 s16_to_f32(uint val) { return float2(a, b) / 32767.0f; } -float3 hit_world_position(ray ray, typename intersector::result_type intersection) { +float3 hit_world_position(ray ray, intersector_t::result_type intersection) { return ray.origin + ray.direction * intersection.distance; } @@ -138,19 +143,15 @@ float3 surface_specular(const float3 base_color, const float metalness) { return mix(float3(0.04, 0.04, 0.04), base_color, metalness); } -float3 env_brdf_approx(float3 specular, float roughness, float dotnv) { - const float4 c0 = float4(-1, -0.0275, -0.572, 0.022); - const float4 c1 = float4(1, 0.0425, 1.04, -0.04); - float4 r = roughness * c0 + c1; - float a004 = min(r.x * r.x, exp2(-9.28 * dotnv)) * r.x + r.y; - float2 ab = float2(-1.04, 1.04) * a004 + r.zw; - return specular * ab.x + ab.y; -} - float fresnel(float3 normal, float3 incident) { return mix(0.5, 1.0, pow(1.0 + dot(normal, incident), 5.0)); } +float4 read_texel(texture2d tex, float2 tex_coord) { + uint2 size = uint2(tex.get_width(), tex.get_height()); + return tex.read(uint2(fract(tex_coord) * float2(size)), 0); +} + kernel void raytracingKernel( uint2 tid [[thread_position_in_grid]], constant RayGenConstantBuffer &constant_buffer [[buffer(0)]], @@ -167,24 +168,32 @@ kernel void raytracingKernel( device void *indices [[buffer(2)]], device void *vertices [[buffer(3)]] ) { + uint2 dim = uint2(render_target.get_width(), render_target.get_height()); + if (tid.x >= dim.x || tid.y >= dim.y) { + return; + } + + int frame = int(constant_buffer.eye.w); uint thread_seed = 0; float3 accum = float3(0, 0, 0); for (int j = 0; j < SAMPLES; ++j) { + int sample_index = frame * SAMPLES + j; + // AA float2 xy = float2(tid) + float2(0.5f, 0.5f); - xy.x += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.x += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); thread_seed += 1; - xy.y += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.y += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); - float2 screen_pos = xy / float2(render_target.get_width(), render_target.get_height()) * 2.0 - 1.0; + float2 screen_pos = xy / float2(dim) * 2.0 - 1.0; ray ray; ray.min_distance = 0.0001; - ray.max_distance = 10.0; + ray.max_distance = 100.0; generate_camera_ray(screen_pos, ray.origin, ray.direction, constant_buffer.eye.xyz, constant_buffer.inv_vp); RayPayload payload; - payload.color = float4(1, 1, 1, j); + payload.color = float4(1, 1, 1, sample_index); #ifdef _TRANSPARENCY int transparent_hits = 0; @@ -195,7 +204,7 @@ kernel void raytracingKernel( #ifdef _ROULETTE float rr_factor = 1.0; if (i >= rr_start) { - float f = rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + float f = rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); if (f <= rr_probability) { break; } @@ -203,91 +212,132 @@ kernel void raytracingKernel( } #endif - // #ifdef _SUBSURFACE - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_CULL_BACK_FACING_TRIANGLES, ~0, 0, 1, 0, ray, payload); - // #else - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE, ~0, 0, 1, 0, ray, payload); - // #endif - - intersector in; + intersector_t in; in.assume_geometry_type(geometry_type::triangle); in.force_opacity(forced_opacity::opaque); in.accept_any_intersection(false); - // in.set_triangle_cull_mode(triangle_cull_mode::none); - typename intersector::result_type intersection; + intersector_t::result_type intersection; intersection = in.intersect(ray, scene); + + // Miss if (intersection.type == intersection_type::none) { #ifdef _EMISSION - if (payload.color.a == -2.0) { - return; + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; } #endif - float2 tex_coord = fract(equirect(ray.direction, constant_buffer.params.y)); - float3 texenv = mytexture_env.sample(linear_sampler, tex_coord).rgb * abs(constant_buffer.params.x); - payload.color = float4(payload.color.rgb * texenv.rgb, -1); + if (i == 0 && constant_buffer.params.x < 0.0) { // No envmap + payload.color.rgb = float3(0.0275, 0.0275, 0.0275); + } + else { + float2 tex_coord = equirect(ray.direction, constant_buffer.params.y); + float3 texenv = mytexture_env.sample(linear_sampler, tex_coord, level(0)).rgb * abs(constant_buffer.params.x); + payload.color.rgb *= texenv; + } + + accum += clamp(payload.color.rgb, 0.0, 8.0); + break; } - else { - device uint32_t *inda = (device uint32_t *)(indices); - uint3 indices_sample = uint3( - inda[intersection.primitive_id * 3], - inda[intersection.primitive_id * 3 + 1], - inda[intersection.primitive_id * 3 + 2] - ); - device Vertex *verta = (device Vertex *)(vertices); - float2 vertex_uvs[3] = { - s16_to_f32(verta[indices_sample[0]].tex), - s16_to_f32(verta[indices_sample[1]].tex), - s16_to_f32(verta[indices_sample[2]].tex) - }; - float2 barycentrics = intersection.triangle_barycentric_coord; - float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + device uint32_t *inda = (device uint32_t *)(indices); + uint base_index = intersection.primitive_id * 3; - uint2 size = uint2(mytexture0.get_width(), mytexture0.get_height()); - uint3 utex_coord = uint3(uint2((tex_coord - float2(uint2(tex_coord))) * float2(size)), 0); - float4 texpaint0 = mytexture0.read(utex_coord.xy, utex_coord.z); + #ifdef _MULTI + base_index += intersection.user_instance_id; + #endif - #ifdef _TRANSPARENCY - if (texpaint0.a <= 0.1) { - payload.ray_dir = ray.direction; - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; - payload.color.a = -2; - return; + uint3 indices_sample = uint3( + inda[base_index], + inda[base_index + 1], + inda[base_index + 2] + ); + + device Vertex *verta = (device Vertex *)(vertices); + float2 vertex_uvs[3] = { + s16_to_f32(verta[indices_sample[0]].tex), + s16_to_f32(verta[indices_sample[1]].tex), + s16_to_f32(verta[indices_sample[2]].tex) + }; + float2 barycentrics = intersection.triangle_barycentric_coord; + float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + + float3 hit = hit_world_position(ray, intersection); + float4 texpaint0 = read_texel(mytexture0, tex_coord); + + #ifdef _TRANSPARENCY + if (texpaint0.a <= 0.01) { + ray.origin = hit + ray.direction * 0.0001f; + if (transparent_hits < DEPTH_TRANSPARENT) { + payload.color.a = sample_index; + transparent_hits++; + i--; } - #endif + else { + payload.color.a = -2; + } + continue; + } + #endif - float3 vertex_normals[3] = { - float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), - float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), - float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) - }; - float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); + float3 vertex_normals[3] = { + float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), + float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), + float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) + }; + float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); - float4 texpaint1 = mytexture1.read(utex_coord.xy, utex_coord.z); - float4 texpaint2 = mytexture2.read(utex_coord.xy, utex_coord.z); - float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); + #ifdef _MULTI + float4x3 obj_to_world = intersection.object_to_world_transform; + n = normalize(float3x3(obj_to_world[0], obj_to_world[1], obj_to_world[2]) * n); + #endif - float3 tangent = float3(0, 0, 0); - float3 binormal = float3(0, 0, 0); - create_basis(n, tangent, binormal); + float4 texpaint1 = read_texel(mytexture1, tex_coord); + float4 texpaint2 = read_texel(mytexture2, tex_coord); + float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); - texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); - texpaint1.g = -texpaint1.g; - n = float3x3(tangent, binormal, n) * texpaint1.rgb; + #ifdef _TRANSLUCENCY + if (!intersection.triangle_front_facing) { + float3 absorption = pow(max(texcolor, float3(0.001)), float3(intersection.distance * texpaint0.a)); + payload.color.rgb *= absorption; + } + #endif - float f = rand(tid.x, tid.y, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - thread_seed += 1; + float3 tangent = float3(0, 0, 0); + float3 binormal = float3(0, 0, 0); + create_basis(n, tangent, binormal); + texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); + texpaint1.g = -texpaint1.g; + n = float3x3(tangent, binormal, n) * texpaint1.rgb; + + uint bounce_seed = 0; + + float f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + + bool scatter = false; + + #ifdef _TRANSLUCENCY + if (f > texpaint0.a) { + float roughness = texpaint2.g; + float3 scatter_dir = cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + payload.ray_dir = normalize(mix(ray.direction, scatter_dir, roughness * roughness * 0.5)); + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + scatter = true; + } + #endif + + if (!scatter) { #ifdef _TRANSLUCENCY - float3 diffuse_dir = texpaint0.a < f ? - cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank) : - cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - #else - float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; #endif + float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + #ifdef _FRESNEL float specular_chance = fresnel(n, ray.direction); #else @@ -295,12 +345,7 @@ kernel void raytracingKernel( #endif if (f < specular_chance) { - #ifdef _TRANSLUCENCY - float3 specular_dir = texpaint0.a < f * 2 ? ray.direction : reflect(ray.direction, n); - #else float3 specular_dir = reflect(ray.direction, n); - #endif - payload.ray_dir = mix(specular_dir, diffuse_dir, texpaint2.g * texpaint2.g); float3 specular = surface_specular(texcolor, texpaint2.b); payload.color.xyz *= specular; @@ -316,59 +361,44 @@ kernel void raytracingKernel( payload.color.xyz /= 1.0 - specular_chance; #endif } + #ifdef _FRESNEL payload.color.xyz *= 0.5; #endif - // float dotnv = abs(dot(n, -WorldRayDirection())); - // payload.ray_origin = hit_world_position() + n * mix(0.1f, 0.0001f, dotnv); - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; + payload.ray_origin = hit + payload.ray_dir * 0.0001f; #ifdef _EMISSION if (int(texpaint1.a * 255.0f) % 3 == 1) { // matid payload.color.xyz *= 100.0f; - payload.color.a = -2.0; + payload.color.a = -3.0; } #endif #ifdef _SUBSURFACE if (int(texpaint1.a * 255.0f) % 3 == 2) { - payload.ray_origin += ray.direction * f; + float d = min(1.0 / min(intersection.distance * 2.0, 1.0) / 10.0, 0.5); + payload.color.xyz += payload.color.xyz * d; + if (f < 0.5) { + payload.ray_origin += ray.direction * f * 0.001; + } } #endif } #ifdef _EMISSION - if (payload.color.a == -2) { + if (payload.color.a == -3.0) { accum += payload.color.rgb; break; } #endif - // Miss - if (payload.color.a < 0) { - #ifdef _TRANSPARENCY - if (payload.color.a == -2 && transparent_hits < DEPTH_TRANSPARENT) { - payload.color.a = j; - transparent_hits++; - i--; - } - #endif - - if (i == 0 && constant_buffer.params.x < 0) { // No envmap - payload.color.rgb = float3(0.0275, 0.0275, 0.0275); - } - - accum += clamp(payload.color.rgb, 0.0, 8.0); - break; - } + ray.origin = payload.ray_origin; + ray.direction = payload.ray_dir; #ifdef _ROULETTE payload.color.rgb *= rr_factor; #endif - - ray.origin = payload.ray_origin; - ray.direction = payload.ray_dir; } } @@ -384,6 +414,6 @@ kernel void raytracingKernel( if (constant_buffer.eye.w == 0) { color = accum; } - render_target.write(float4(mix(color, accum, 1.0 / 16.0), 1.0f), tid); + render_target.write(float4(mix(color, accum, 1.0 / 4.0), 1.0f), tid); #endif } diff --git a/base/shaders/raytrace/src/raytrace_brute_core.metal b/base/shaders/raytrace/raytrace_brute_multi_core.metal similarity index 58% rename from base/shaders/raytrace/src/raytrace_brute_core.metal rename to base/shaders/raytrace/raytrace_brute_multi_core.metal index 28b94903..2f18a1d3 100644 --- a/base/shaders/raytrace/src/raytrace_brute_core.metal +++ b/base/shaders/raytrace/raytrace_brute_multi_core.metal @@ -1,12 +1,13 @@ +#define _MULTI #ifdef _FULL #define _EMISSION #define _SUBSURFACE #define _TRANSLUCENCY #define _ROULETTE -#define _TRANSPARENCY +// #define _TRANSPARENCY // #define _FRESNEL #endif -#define _RENDER +// #define _RENDER using namespace metal; using namespace raytracing; @@ -18,6 +19,12 @@ struct Vertex { uint tex; }; +#ifdef _MULTI +typedef intersector intersector_t; +#else +typedef intersector intersector_t; +#endif + struct RayGenConstantBuffer { float4 eye; // xyz, frame float4x4 inv_vp; @@ -30,9 +37,9 @@ struct RayPayload { float3 ray_dir; }; -constant int SAMPLES = 2; // 64 +constant int SAMPLES = 8; #ifdef _TRANSLUCENCY -constant int DEPTH = 6; +constant int DEPTH = 16; #else constant int DEPTH = 3; // Opaque hits #endif @@ -100,7 +107,7 @@ float2 s16_to_f32(uint val) { return float2(a, b) / 32767.0f; } -float3 hit_world_position(ray ray, typename intersector::result_type intersection) { +float3 hit_world_position(ray ray, intersector_t::result_type intersection) { return ray.origin + ray.direction * intersection.distance; } @@ -136,19 +143,15 @@ float3 surface_specular(const float3 base_color, const float metalness) { return mix(float3(0.04, 0.04, 0.04), base_color, metalness); } -float3 env_brdf_approx(float3 specular, float roughness, float dotnv) { - const float4 c0 = float4(-1, -0.0275, -0.572, 0.022); - const float4 c1 = float4(1, 0.0425, 1.04, -0.04); - float4 r = roughness * c0 + c1; - float a004 = min(r.x * r.x, exp2(-9.28 * dotnv)) * r.x + r.y; - float2 ab = float2(-1.04, 1.04) * a004 + r.zw; - return specular * ab.x + ab.y; -} - float fresnel(float3 normal, float3 incident) { return mix(0.5, 1.0, pow(1.0 + dot(normal, incident), 5.0)); } +float4 read_texel(texture2d tex, float2 tex_coord) { + uint2 size = uint2(tex.get_width(), tex.get_height()); + return tex.read(uint2(fract(tex_coord) * float2(size)), 0); +} + kernel void raytracingKernel( uint2 tid [[thread_position_in_grid]], constant RayGenConstantBuffer &constant_buffer [[buffer(0)]], @@ -165,24 +168,32 @@ kernel void raytracingKernel( device void *indices [[buffer(2)]], device void *vertices [[buffer(3)]] ) { + uint2 dim = uint2(render_target.get_width(), render_target.get_height()); + if (tid.x >= dim.x || tid.y >= dim.y) { + return; + } + + int frame = int(constant_buffer.eye.w); uint thread_seed = 0; float3 accum = float3(0, 0, 0); for (int j = 0; j < SAMPLES; ++j) { + int sample_index = frame * SAMPLES + j; + // AA float2 xy = float2(tid) + float2(0.5f, 0.5f); - xy.x += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.x += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); thread_seed += 1; - xy.y += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.y += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); - float2 screen_pos = xy / float2(render_target.get_width(), render_target.get_height()) * 2.0 - 1.0; + float2 screen_pos = xy / float2(dim) * 2.0 - 1.0; ray ray; ray.min_distance = 0.0001; - ray.max_distance = 10.0; + ray.max_distance = 100.0; generate_camera_ray(screen_pos, ray.origin, ray.direction, constant_buffer.eye.xyz, constant_buffer.inv_vp); RayPayload payload; - payload.color = float4(1, 1, 1, j); + payload.color = float4(1, 1, 1, sample_index); #ifdef _TRANSPARENCY int transparent_hits = 0; @@ -193,7 +204,7 @@ kernel void raytracingKernel( #ifdef _ROULETTE float rr_factor = 1.0; if (i >= rr_start) { - float f = rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + float f = rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); if (f <= rr_probability) { break; } @@ -201,91 +212,132 @@ kernel void raytracingKernel( } #endif - // #ifdef _SUBSURFACE - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_CULL_BACK_FACING_TRIANGLES, ~0, 0, 1, 0, ray, payload); - // #else - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE, ~0, 0, 1, 0, ray, payload); - // #endif - - intersector in; + intersector_t in; in.assume_geometry_type(geometry_type::triangle); in.force_opacity(forced_opacity::opaque); in.accept_any_intersection(false); - // in.set_triangle_cull_mode(triangle_cull_mode::none); - typename intersector::result_type intersection; + intersector_t::result_type intersection; intersection = in.intersect(ray, scene); + + // Miss if (intersection.type == intersection_type::none) { #ifdef _EMISSION - if (payload.color.a == -2.0) { - return; + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; } #endif - float2 tex_coord = fract(equirect(ray.direction, constant_buffer.params.y)); - float3 texenv = mytexture_env.sample(linear_sampler, tex_coord).rgb * abs(constant_buffer.params.x); - payload.color = float4(payload.color.rgb * texenv.rgb, -1); + if (i == 0 && constant_buffer.params.x < 0.0) { // No envmap + payload.color.rgb = float3(0.0275, 0.0275, 0.0275); + } + else { + float2 tex_coord = equirect(ray.direction, constant_buffer.params.y); + float3 texenv = mytexture_env.sample(linear_sampler, tex_coord, level(0)).rgb * abs(constant_buffer.params.x); + payload.color.rgb *= texenv; + } + + accum += clamp(payload.color.rgb, 0.0, 8.0); + break; } - else { - device uint32_t *inda = (device uint32_t *)(indices); - uint3 indices_sample = uint3( - inda[intersection.primitive_id * 3], - inda[intersection.primitive_id * 3 + 1], - inda[intersection.primitive_id * 3 + 2] - ); - device Vertex *verta = (device Vertex *)(vertices); - float2 vertex_uvs[3] = { - s16_to_f32(verta[indices_sample[0]].tex), - s16_to_f32(verta[indices_sample[1]].tex), - s16_to_f32(verta[indices_sample[2]].tex) - }; - float2 barycentrics = intersection.triangle_barycentric_coord; - float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + device uint32_t *inda = (device uint32_t *)(indices); + uint base_index = intersection.primitive_id * 3; - uint2 size = uint2(mytexture0.get_width(), mytexture0.get_height()); - uint3 utex_coord = uint3(uint2((tex_coord - float2(uint2(tex_coord))) * float2(size)), 0); - float4 texpaint0 = mytexture0.read(utex_coord.xy, utex_coord.z); + #ifdef _MULTI + base_index += intersection.user_instance_id; + #endif - #ifdef _TRANSPARENCY - if (texpaint0.a <= 0.1) { - payload.ray_dir = ray.direction; - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; - payload.color.a = -2; - return; + uint3 indices_sample = uint3( + inda[base_index], + inda[base_index + 1], + inda[base_index + 2] + ); + + device Vertex *verta = (device Vertex *)(vertices); + float2 vertex_uvs[3] = { + s16_to_f32(verta[indices_sample[0]].tex), + s16_to_f32(verta[indices_sample[1]].tex), + s16_to_f32(verta[indices_sample[2]].tex) + }; + float2 barycentrics = intersection.triangle_barycentric_coord; + float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + + float3 hit = hit_world_position(ray, intersection); + float4 texpaint0 = read_texel(mytexture0, tex_coord); + + #ifdef _TRANSPARENCY + if (texpaint0.a <= 0.01) { + ray.origin = hit + ray.direction * 0.0001f; + if (transparent_hits < DEPTH_TRANSPARENT) { + payload.color.a = sample_index; + transparent_hits++; + i--; } - #endif + else { + payload.color.a = -2; + } + continue; + } + #endif - float3 vertex_normals[3] = { - float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), - float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), - float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) - }; - float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); + float3 vertex_normals[3] = { + float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), + float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), + float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) + }; + float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); - float4 texpaint1 = mytexture1.read(utex_coord.xy, utex_coord.z); - float4 texpaint2 = mytexture2.read(utex_coord.xy, utex_coord.z); - float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); + #ifdef _MULTI + float4x3 obj_to_world = intersection.object_to_world_transform; + n = normalize(float3x3(obj_to_world[0], obj_to_world[1], obj_to_world[2]) * n); + #endif - float3 tangent = float3(0, 0, 0); - float3 binormal = float3(0, 0, 0); - create_basis(n, tangent, binormal); + float4 texpaint1 = read_texel(mytexture1, tex_coord); + float4 texpaint2 = read_texel(mytexture2, tex_coord); + float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); - texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); - texpaint1.g = -texpaint1.g; - n = float3x3(tangent, binormal, n) * texpaint1.rgb; + #ifdef _TRANSLUCENCY + if (!intersection.triangle_front_facing) { + float3 absorption = pow(max(texcolor, float3(0.001)), float3(intersection.distance * texpaint0.a)); + payload.color.rgb *= absorption; + } + #endif - float f = rand(tid.x, tid.y, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - thread_seed += 1; + float3 tangent = float3(0, 0, 0); + float3 binormal = float3(0, 0, 0); + create_basis(n, tangent, binormal); + texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); + texpaint1.g = -texpaint1.g; + n = float3x3(tangent, binormal, n) * texpaint1.rgb; + + uint bounce_seed = 0; + + float f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + + bool scatter = false; + + #ifdef _TRANSLUCENCY + if (f > texpaint0.a) { + float roughness = texpaint2.g; + float3 scatter_dir = cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + payload.ray_dir = normalize(mix(ray.direction, scatter_dir, roughness * roughness * 0.5)); + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + scatter = true; + } + #endif + + if (!scatter) { #ifdef _TRANSLUCENCY - float3 diffuse_dir = texpaint0.a < f ? - cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank) : - cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - #else - float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; #endif + float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + #ifdef _FRESNEL float specular_chance = fresnel(n, ray.direction); #else @@ -293,12 +345,7 @@ kernel void raytracingKernel( #endif if (f < specular_chance) { - #ifdef _TRANSLUCENCY - float3 specular_dir = texpaint0.a < f * 2 ? ray.direction : reflect(ray.direction, n); - #else float3 specular_dir = reflect(ray.direction, n); - #endif - payload.ray_dir = mix(specular_dir, diffuse_dir, texpaint2.g * texpaint2.g); float3 specular = surface_specular(texcolor, texpaint2.b); payload.color.xyz *= specular; @@ -314,59 +361,44 @@ kernel void raytracingKernel( payload.color.xyz /= 1.0 - specular_chance; #endif } + #ifdef _FRESNEL payload.color.xyz *= 0.5; #endif - // float dotnv = abs(dot(n, -WorldRayDirection())); - // payload.ray_origin = hit_world_position() + n * mix(0.1f, 0.0001f, dotnv); - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; + payload.ray_origin = hit + payload.ray_dir * 0.0001f; #ifdef _EMISSION if (int(texpaint1.a * 255.0f) % 3 == 1) { // matid payload.color.xyz *= 100.0f; - payload.color.a = -2.0; + payload.color.a = -3.0; } #endif #ifdef _SUBSURFACE if (int(texpaint1.a * 255.0f) % 3 == 2) { - payload.ray_origin += ray.direction * f; + float d = min(1.0 / min(intersection.distance * 2.0, 1.0) / 10.0, 0.5); + payload.color.xyz += payload.color.xyz * d; + if (f < 0.5) { + payload.ray_origin += ray.direction * f * 0.001; + } } #endif } #ifdef _EMISSION - if (payload.color.a == -2) { + if (payload.color.a == -3.0) { accum += payload.color.rgb; break; } #endif - // Miss - if (payload.color.a < 0) { - #ifdef _TRANSPARENCY - if (payload.color.a == -2 && transparent_hits < DEPTH_TRANSPARENT) { - payload.color.a = j; - transparent_hits++; - i--; - } - #endif - - if (i == 0 && constant_buffer.params.x < 0) { // No envmap - payload.color.rgb = float3(0.032, 0.032, 0.032); - } - - accum += clamp(payload.color.rgb, 0.0, 8.0); - break; - } + ray.origin = payload.ray_origin; + ray.direction = payload.ray_dir; #ifdef _ROULETTE payload.color.rgb *= rr_factor; #endif - - ray.origin = payload.ray_origin; - ray.direction = payload.ray_dir; } } @@ -382,6 +414,6 @@ kernel void raytracingKernel( if (constant_buffer.eye.w == 0) { color = accum; } - render_target.write(float4(mix(color, accum, 1.0 / 16.0), 1.0f), tid); + render_target.write(float4(mix(color, accum, 1.0 / 4.0), 1.0f), tid); #endif } diff --git a/base/shaders/raytrace/raytrace_brute_multi_full.metal b/base/shaders/raytrace/raytrace_brute_multi_full.metal new file mode 100644 index 00000000..4baa90ca --- /dev/null +++ b/base/shaders/raytrace/raytrace_brute_multi_full.metal @@ -0,0 +1,420 @@ +#define _MULTI +#define _FULL +#ifdef _FULL +#define _EMISSION +#define _SUBSURFACE +#define _TRANSLUCENCY +#define _ROULETTE +// #define _TRANSPARENCY +// #define _FRESNEL +#endif +// #define _RENDER + +using namespace metal; +using namespace raytracing; + +struct Vertex { + uint posxy; + uint poszw; + uint nor; + uint tex; +}; + +#ifdef _MULTI +typedef intersector intersector_t; +#else +typedef intersector intersector_t; +#endif + +struct RayGenConstantBuffer { + float4 eye; // xyz, frame + float4x4 inv_vp; + float4 params; // envstr, envangle, uvscale +}; + +struct RayPayload { + float4 color; // rgb, frame + float3 ray_origin; + float3 ray_dir; +}; + +constant int SAMPLES = 8; +#ifdef _TRANSLUCENCY +constant int DEPTH = 16; +#else +constant int DEPTH = 3; // Opaque hits +#endif +#ifdef _TRANSPARENCY +constant int DEPTH_TRANSPARENT = 16; // Transparent hits +#endif +#ifdef _ROULETTE +constant int rr_start = 2; +constant float rr_probability = 0.5; // Map to albedo +#endif + +void generate_camera_ray(float2 screen_pos, thread float3 & ray_origin, thread float3 & ray_dir, float3 eye, float4x4 inv_vp) { + screen_pos.y = -screen_pos.y; + float4 world = inv_vp * float4(screen_pos, 0, 1); + world.xyz /= world.w; + ray_origin = eye; + ray_dir = normalize(world.xyz - ray_origin); +} + +float2 equirect(float3 normal, float angle) { + const float PI = 3.1415926535; + const float PI2 = PI * 2.0; + float phi = acos(normal.z); + float theta = atan2(-normal.y, normal.x) + PI + angle; + return float2(theta / PI2, phi / PI); +} + +float rand(int pixel_i, int pixel_j, int sample_index, int sample_dimension, int frame, texture2d sobol, texture2d scramble, texture2d rank) { + pixel_i += frame * 9; + pixel_j += frame * 11; + pixel_i = pixel_i & 127; + pixel_j = pixel_j & 127; + sample_index = sample_index & 255; + sample_dimension = sample_dimension & 255; + + int i = sample_dimension + (pixel_i + pixel_j * 128) * 8; + int ranked_sample_index = sample_index ^ int(rank.read(uint2(i % 128, uint(i / 128)), 0).r * 255); + + i = sample_dimension + ranked_sample_index * 256; + int value = int(sobol.read(uint2(i % 256, uint(i / 256)), 0).r * 255); + + i = (sample_dimension % 8) + (pixel_i + pixel_j * 128) * 8; + value = value ^ int(scramble.read(uint2(i % 128, uint(i / 128)), 0).r * 255); + + float v = (0.5f + value) / 256.0f; + return v; +} + +float3 cos_weighted_hemisphere_direction(uint2 tid, float3 n, uint sample, uint seed, int frame, texture2d sobol, texture2d scramble, texture2d rank) { + const float PI = 3.1415926535; + const float PI2 = PI * 2.0; + float f0 = rand(tid.x, tid.y, sample, seed, frame, sobol, scramble, rank); + float f1 = rand(tid.x, tid.y, sample, seed + 1, frame, sobol, scramble, rank); + float z = f0 * 2.0f - 1.0f; + float a = f1 * PI2; + float r = sqrt(1.0f - z * z); + float x = r * cos(a); + float y = r * sin(a); + return normalize(n + float3(x, y, z)); +} + +float2 s16_to_f32(uint val) { + int a = (int)(val << 16) >> 16; + int b = (int)(val & 0xffff0000) >> 16; + return float2(a, b) / 32767.0f; +} + +float3 hit_world_position(ray ray, intersector_t::result_type intersection) { + return ray.origin + ray.direction * intersection.distance; +} + +float3 hit_attribute(float3 vertex_attribute[3], float2 barycentrics) { + return vertex_attribute[0] + + barycentrics.x * (vertex_attribute[1] - vertex_attribute[0]) + + barycentrics.y * (vertex_attribute[2] - vertex_attribute[0]); +} + +float2 hit_attribute2d(float2 vertex_attribute[3], float2 barycentrics) { + return vertex_attribute[0] + + barycentrics.x * (vertex_attribute[1] - vertex_attribute[0]) + + barycentrics.y * (vertex_attribute[2] - vertex_attribute[0]); +} + +void create_basis(float3 normal, thread float3 & tangent, thread float3 & binormal) { + float3 v = cross(normal, float3(0.0, 0.0, 1.0)); + if (dot(v, v) > 0.0001) { + tangent = normalize(v); + } + else { + v = cross(normal, float3(0.0, 1.0, 0.0)); + tangent = normalize(v); + } + binormal = cross(tangent, normal); +} + +float3 surface_albedo(const float3 base_color, const float metalness) { + return mix(base_color, float3(0.0, 0.0, 0.0), metalness); +} + +float3 surface_specular(const float3 base_color, const float metalness) { + return mix(float3(0.04, 0.04, 0.04), base_color, metalness); +} + +float fresnel(float3 normal, float3 incident) { + return mix(0.5, 1.0, pow(1.0 + dot(normal, incident), 5.0)); +} + +float4 read_texel(texture2d tex, float2 tex_coord) { + uint2 size = uint2(tex.get_width(), tex.get_height()); + return tex.read(uint2(fract(tex_coord) * float2(size)), 0); +} + +kernel void raytracingKernel( + uint2 tid [[thread_position_in_grid]], + constant RayGenConstantBuffer &constant_buffer [[buffer(0)]], + texture2d render_target [[texture(0)]], + texture2d mytexture0 [[texture(1)]], + texture2d mytexture1 [[texture(2)]], + texture2d mytexture2 [[texture(3)]], + texture2d mytexture_env [[texture(4)]], + texture2d mytexture_sobol [[texture(5)]], + texture2d mytexture_scramble [[texture(6)]], + texture2d mytexture_rank [[texture(7)]], + sampler linear_sampler [[sampler(0)]], + instance_acceleration_structure scene [[buffer(1)]], + device void *indices [[buffer(2)]], + device void *vertices [[buffer(3)]] +) { + uint2 dim = uint2(render_target.get_width(), render_target.get_height()); + if (tid.x >= dim.x || tid.y >= dim.y) { + return; + } + + int frame = int(constant_buffer.eye.w); + uint thread_seed = 0; + float3 accum = float3(0, 0, 0); + + for (int j = 0; j < SAMPLES; ++j) { + int sample_index = frame * SAMPLES + j; + + // AA + float2 xy = float2(tid) + float2(0.5f, 0.5f); + xy.x += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + thread_seed += 1; + xy.y += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + + float2 screen_pos = xy / float2(dim) * 2.0 - 1.0; + ray ray; + ray.min_distance = 0.0001; + ray.max_distance = 100.0; + generate_camera_ray(screen_pos, ray.origin, ray.direction, constant_buffer.eye.xyz, constant_buffer.inv_vp); + + RayPayload payload; + payload.color = float4(1, 1, 1, sample_index); + + #ifdef _TRANSPARENCY + int transparent_hits = 0; + #endif + + for (int i = 0; i < DEPTH; ++i) { + + #ifdef _ROULETTE + float rr_factor = 1.0; + if (i >= rr_start) { + float f = rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + if (f <= rr_probability) { + break; + } + rr_factor = 1.0 / (1.0 - rr_probability); + } + #endif + + intersector_t in; + in.assume_geometry_type(geometry_type::triangle); + in.force_opacity(forced_opacity::opaque); + in.accept_any_intersection(false); + + intersector_t::result_type intersection; + intersection = in.intersect(ray, scene); + + // Miss + if (intersection.type == intersection_type::none) { + #ifdef _EMISSION + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; + } + #endif + + if (i == 0 && constant_buffer.params.x < 0.0) { // No envmap + payload.color.rgb = float3(0.0275, 0.0275, 0.0275); + } + else { + float2 tex_coord = equirect(ray.direction, constant_buffer.params.y); + float3 texenv = mytexture_env.sample(linear_sampler, tex_coord, level(0)).rgb * abs(constant_buffer.params.x); + payload.color.rgb *= texenv; + } + + accum += clamp(payload.color.rgb, 0.0, 8.0); + break; + } + + device uint32_t *inda = (device uint32_t *)(indices); + uint base_index = intersection.primitive_id * 3; + + #ifdef _MULTI + base_index += intersection.user_instance_id; + #endif + + uint3 indices_sample = uint3( + inda[base_index], + inda[base_index + 1], + inda[base_index + 2] + ); + + device Vertex *verta = (device Vertex *)(vertices); + float2 vertex_uvs[3] = { + s16_to_f32(verta[indices_sample[0]].tex), + s16_to_f32(verta[indices_sample[1]].tex), + s16_to_f32(verta[indices_sample[2]].tex) + }; + float2 barycentrics = intersection.triangle_barycentric_coord; + float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + + float3 hit = hit_world_position(ray, intersection); + float4 texpaint0 = read_texel(mytexture0, tex_coord); + + #ifdef _TRANSPARENCY + if (texpaint0.a <= 0.01) { + ray.origin = hit + ray.direction * 0.0001f; + if (transparent_hits < DEPTH_TRANSPARENT) { + payload.color.a = sample_index; + transparent_hits++; + i--; + } + else { + payload.color.a = -2; + } + continue; + } + #endif + + float3 vertex_normals[3] = { + float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), + float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), + float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) + }; + float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); + + #ifdef _MULTI + float4x3 obj_to_world = intersection.object_to_world_transform; + n = normalize(float3x3(obj_to_world[0], obj_to_world[1], obj_to_world[2]) * n); + #endif + + float4 texpaint1 = read_texel(mytexture1, tex_coord); + float4 texpaint2 = read_texel(mytexture2, tex_coord); + float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); + + #ifdef _TRANSLUCENCY + if (!intersection.triangle_front_facing) { + float3 absorption = pow(max(texcolor, float3(0.001)), float3(intersection.distance * texpaint0.a)); + payload.color.rgb *= absorption; + } + #endif + + float3 tangent = float3(0, 0, 0); + float3 binormal = float3(0, 0, 0); + create_basis(n, tangent, binormal); + + texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); + texpaint1.g = -texpaint1.g; + n = float3x3(tangent, binormal, n) * texpaint1.rgb; + + uint bounce_seed = 0; + + float f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + + bool scatter = false; + + #ifdef _TRANSLUCENCY + if (f > texpaint0.a) { + float roughness = texpaint2.g; + float3 scatter_dir = cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + payload.ray_dir = normalize(mix(ray.direction, scatter_dir, roughness * roughness * 0.5)); + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + scatter = true; + } + #endif + + if (!scatter) { + #ifdef _TRANSLUCENCY + f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + #endif + + float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + + #ifdef _FRESNEL + float specular_chance = fresnel(n, ray.direction); + #else + const float specular_chance = 0.5; + #endif + + if (f < specular_chance) { + float3 specular_dir = reflect(ray.direction, n); + payload.ray_dir = mix(specular_dir, diffuse_dir, texpaint2.g * texpaint2.g); + float3 specular = surface_specular(texcolor, texpaint2.b); + payload.color.xyz *= specular; + + #ifdef _FRESNEL + payload.color.xyz /= specular_chance; + #endif + } + else { + payload.ray_dir = diffuse_dir; + payload.color.xyz *= surface_albedo(texcolor, texpaint2.b); + #ifdef _FRESNEL + payload.color.xyz /= 1.0 - specular_chance; + #endif + } + + #ifdef _FRESNEL + payload.color.xyz *= 0.5; + #endif + + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + + #ifdef _EMISSION + if (int(texpaint1.a * 255.0f) % 3 == 1) { // matid + payload.color.xyz *= 100.0f; + payload.color.a = -3.0; + } + #endif + + #ifdef _SUBSURFACE + if (int(texpaint1.a * 255.0f) % 3 == 2) { + float d = min(1.0 / min(intersection.distance * 2.0, 1.0) / 10.0, 0.5); + payload.color.xyz += payload.color.xyz * d; + if (f < 0.5) { + payload.ray_origin += ray.direction * f * 0.001; + } + } + #endif + } + + #ifdef _EMISSION + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; + } + #endif + + ray.origin = payload.ray_origin; + ray.direction = payload.ray_dir; + + #ifdef _ROULETTE + payload.color.rgb *= rr_factor; + #endif + } + } + + float3 color = render_target.read(tid).xyz; + accum = accum / SAMPLES; + + #ifdef _RENDER + float a = 1.0 / (constant_buffer.eye.w + 1); + float b = 1.0 - a; + color = color * b + accum * a; + render_target.write(float4(color, 1.0f), tid); + #else + if (constant_buffer.eye.w == 0) { + color = accum; + } + render_target.write(float4(mix(color, accum, 1.0 / 4.0), 1.0f), tid); + #endif +} diff --git a/base/shaders/raytrace/src/build_metal.sh b/base/shaders/raytrace/src/build_metal.sh index 0b6ef6f3..dd466509 100755 --- a/base/shaders/raytrace/src/build_metal.sh +++ b/base/shaders/raytrace/src/build_metal.sh @@ -1,5 +1,7 @@ -cp raytrace_brute_core.metal ../raytrace_brute_core.metal -cp raytrace_brute_full.metal ../raytrace_brute_full.metal +{ cat raytrace_brute.metal; } > ../raytrace_brute_core.metal +{ echo "#define _FULL"; cat raytrace_brute.metal; } > ../raytrace_brute_full.metal +{ echo "#define _MULTI"; cat raytrace_brute.metal; } > ../raytrace_brute_multi_core.metal +{ echo "#define _MULTI"; echo "#define _FULL"; cat raytrace_brute.metal; } > ../raytrace_brute_multi_full.metal cp raytrace_bake_ao.metal ../raytrace_bake_ao.metal cp raytrace_bake_light.metal ../raytrace_bake_light.metal cp raytrace_bake_bent.metal ../raytrace_bake_bent.metal diff --git a/base/shaders/raytrace/src/raytrace_brute_full.metal b/base/shaders/raytrace/src/raytrace_brute.metal similarity index 59% rename from base/shaders/raytrace/src/raytrace_brute_full.metal rename to base/shaders/raytrace/src/raytrace_brute.metal index ba0e0ff8..1e0de1fa 100644 --- a/base/shaders/raytrace/src/raytrace_brute_full.metal +++ b/base/shaders/raytrace/src/raytrace_brute.metal @@ -1,5 +1,3 @@ -#define _FULL - #ifdef _FULL #define _EMISSION #define _SUBSURFACE @@ -8,7 +6,7 @@ // #define _TRANSPARENCY // #define _FRESNEL #endif -#define _RENDER +// #define _RENDER using namespace metal; using namespace raytracing; @@ -20,6 +18,12 @@ struct Vertex { uint tex; }; +#ifdef _MULTI +typedef intersector intersector_t; +#else +typedef intersector intersector_t; +#endif + struct RayGenConstantBuffer { float4 eye; // xyz, frame float4x4 inv_vp; @@ -32,9 +36,9 @@ struct RayPayload { float3 ray_dir; }; -constant int SAMPLES = 2; // 64 +constant int SAMPLES = 8; #ifdef _TRANSLUCENCY -constant int DEPTH = 6; +constant int DEPTH = 16; #else constant int DEPTH = 3; // Opaque hits #endif @@ -102,7 +106,7 @@ float2 s16_to_f32(uint val) { return float2(a, b) / 32767.0f; } -float3 hit_world_position(ray ray, typename intersector::result_type intersection) { +float3 hit_world_position(ray ray, intersector_t::result_type intersection) { return ray.origin + ray.direction * intersection.distance; } @@ -138,19 +142,15 @@ float3 surface_specular(const float3 base_color, const float metalness) { return mix(float3(0.04, 0.04, 0.04), base_color, metalness); } -float3 env_brdf_approx(float3 specular, float roughness, float dotnv) { - const float4 c0 = float4(-1, -0.0275, -0.572, 0.022); - const float4 c1 = float4(1, 0.0425, 1.04, -0.04); - float4 r = roughness * c0 + c1; - float a004 = min(r.x * r.x, exp2(-9.28 * dotnv)) * r.x + r.y; - float2 ab = float2(-1.04, 1.04) * a004 + r.zw; - return specular * ab.x + ab.y; -} - float fresnel(float3 normal, float3 incident) { return mix(0.5, 1.0, pow(1.0 + dot(normal, incident), 5.0)); } +float4 read_texel(texture2d tex, float2 tex_coord) { + uint2 size = uint2(tex.get_width(), tex.get_height()); + return tex.read(uint2(fract(tex_coord) * float2(size)), 0); +} + kernel void raytracingKernel( uint2 tid [[thread_position_in_grid]], constant RayGenConstantBuffer &constant_buffer [[buffer(0)]], @@ -167,24 +167,32 @@ kernel void raytracingKernel( device void *indices [[buffer(2)]], device void *vertices [[buffer(3)]] ) { + uint2 dim = uint2(render_target.get_width(), render_target.get_height()); + if (tid.x >= dim.x || tid.y >= dim.y) { + return; + } + + int frame = int(constant_buffer.eye.w); uint thread_seed = 0; float3 accum = float3(0, 0, 0); for (int j = 0; j < SAMPLES; ++j) { + int sample_index = frame * SAMPLES + j; + // AA float2 xy = float2(tid) + float2(0.5f, 0.5f); - xy.x += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.x += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); thread_seed += 1; - xy.y += rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + xy.y += rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); - float2 screen_pos = xy / float2(render_target.get_width(), render_target.get_height()) * 2.0 - 1.0; + float2 screen_pos = xy / float2(dim) * 2.0 - 1.0; ray ray; ray.min_distance = 0.0001; - ray.max_distance = 10.0; + ray.max_distance = 100.0; generate_camera_ray(screen_pos, ray.origin, ray.direction, constant_buffer.eye.xyz, constant_buffer.inv_vp); RayPayload payload; - payload.color = float4(1, 1, 1, j); + payload.color = float4(1, 1, 1, sample_index); #ifdef _TRANSPARENCY int transparent_hits = 0; @@ -195,7 +203,7 @@ kernel void raytracingKernel( #ifdef _ROULETTE float rr_factor = 1.0; if (i >= rr_start) { - float f = rand(tid.x, tid.y, j, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + float f = rand(tid.x, tid.y, sample_index, thread_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); if (f <= rr_probability) { break; } @@ -203,91 +211,132 @@ kernel void raytracingKernel( } #endif - // #ifdef _SUBSURFACE - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE | RAY_FLAG_CULL_BACK_FACING_TRIANGLES, ~0, 0, 1, 0, ray, payload); - // #else - // TraceRay(scene, RAY_FLAG_FORCE_OPAQUE, ~0, 0, 1, 0, ray, payload); - // #endif - - intersector in; + intersector_t in; in.assume_geometry_type(geometry_type::triangle); in.force_opacity(forced_opacity::opaque); in.accept_any_intersection(false); - // in.set_triangle_cull_mode(triangle_cull_mode::none); - typename intersector::result_type intersection; + intersector_t::result_type intersection; intersection = in.intersect(ray, scene); + + // Miss if (intersection.type == intersection_type::none) { #ifdef _EMISSION - if (payload.color.a == -2.0) { - return; + if (payload.color.a == -3.0) { + accum += payload.color.rgb; + break; } #endif - float2 tex_coord = fract(equirect(ray.direction, constant_buffer.params.y)); - float3 texenv = mytexture_env.sample(linear_sampler, tex_coord).rgb * abs(constant_buffer.params.x); - payload.color = float4(payload.color.rgb * texenv.rgb, -1); + if (i == 0 && constant_buffer.params.x < 0.0) { // No envmap + payload.color.rgb = float3(0.0275, 0.0275, 0.0275); + } + else { + float2 tex_coord = equirect(ray.direction, constant_buffer.params.y); + float3 texenv = mytexture_env.sample(linear_sampler, tex_coord, level(0)).rgb * abs(constant_buffer.params.x); + payload.color.rgb *= texenv; + } + + accum += clamp(payload.color.rgb, 0.0, 8.0); + break; } - else { - device uint32_t *inda = (device uint32_t *)(indices); - uint3 indices_sample = uint3( - inda[intersection.primitive_id * 3], - inda[intersection.primitive_id * 3 + 1], - inda[intersection.primitive_id * 3 + 2] - ); - device Vertex *verta = (device Vertex *)(vertices); - float2 vertex_uvs[3] = { - s16_to_f32(verta[indices_sample[0]].tex), - s16_to_f32(verta[indices_sample[1]].tex), - s16_to_f32(verta[indices_sample[2]].tex) - }; - float2 barycentrics = intersection.triangle_barycentric_coord; - float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + device uint32_t *inda = (device uint32_t *)(indices); + uint base_index = intersection.primitive_id * 3; - uint2 size = uint2(mytexture0.get_width(), mytexture0.get_height()); - uint3 utex_coord = uint3(uint2((tex_coord - float2(uint2(tex_coord))) * float2(size)), 0); - float4 texpaint0 = mytexture0.read(utex_coord.xy, utex_coord.z); + #ifdef _MULTI + base_index += intersection.user_instance_id; + #endif - #ifdef _TRANSPARENCY - if (texpaint0.a <= 0.1) { - payload.ray_dir = ray.direction; - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; - payload.color.a = -2; - return; + uint3 indices_sample = uint3( + inda[base_index], + inda[base_index + 1], + inda[base_index + 2] + ); + + device Vertex *verta = (device Vertex *)(vertices); + float2 vertex_uvs[3] = { + s16_to_f32(verta[indices_sample[0]].tex), + s16_to_f32(verta[indices_sample[1]].tex), + s16_to_f32(verta[indices_sample[2]].tex) + }; + float2 barycentrics = intersection.triangle_barycentric_coord; + float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z; + + float3 hit = hit_world_position(ray, intersection); + float4 texpaint0 = read_texel(mytexture0, tex_coord); + + #ifdef _TRANSPARENCY + if (texpaint0.a <= 0.01) { + ray.origin = hit + ray.direction * 0.0001f; + if (transparent_hits < DEPTH_TRANSPARENT) { + payload.color.a = sample_index; + transparent_hits++; + i--; } - #endif + else { + payload.color.a = -2; + } + continue; + } + #endif - float3 vertex_normals[3] = { - float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), - float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), - float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) - }; - float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); + float3 vertex_normals[3] = { + float3(s16_to_f32(verta[indices_sample[0]].nor), s16_to_f32(verta[indices_sample[0]].poszw).y), + float3(s16_to_f32(verta[indices_sample[1]].nor), s16_to_f32(verta[indices_sample[1]].poszw).y), + float3(s16_to_f32(verta[indices_sample[2]].nor), s16_to_f32(verta[indices_sample[2]].poszw).y) + }; + float3 n = normalize(hit_attribute(vertex_normals, barycentrics)); - float4 texpaint1 = mytexture1.read(utex_coord.xy, utex_coord.z); - float4 texpaint2 = mytexture2.read(utex_coord.xy, utex_coord.z); - float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); + #ifdef _MULTI + float4x3 obj_to_world = intersection.object_to_world_transform; + n = normalize(float3x3(obj_to_world[0], obj_to_world[1], obj_to_world[2]) * n); + #endif - float3 tangent = float3(0, 0, 0); - float3 binormal = float3(0, 0, 0); - create_basis(n, tangent, binormal); + float4 texpaint1 = read_texel(mytexture1, tex_coord); + float4 texpaint2 = read_texel(mytexture2, tex_coord); + float3 texcolor = pow(texpaint0.rgb, float3(2.2, 2.2, 2.2)); - texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); - texpaint1.g = -texpaint1.g; - n = float3x3(tangent, binormal, n) * texpaint1.rgb; + #ifdef _TRANSLUCENCY + if (!intersection.triangle_front_facing) { + float3 absorption = pow(max(texcolor, float3(0.001)), float3(intersection.distance * texpaint0.a)); + payload.color.rgb *= absorption; + } + #endif - float f = rand(tid.x, tid.y, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - thread_seed += 1; + float3 tangent = float3(0, 0, 0); + float3 binormal = float3(0, 0, 0); + create_basis(n, tangent, binormal); + texpaint1.rgb = normalize(texpaint1.rgb * 2.0 - 1.0); + texpaint1.g = -texpaint1.g; + n = float3x3(tangent, binormal, n) * texpaint1.rgb; + + uint bounce_seed = 0; + + float f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; + + bool scatter = false; + + #ifdef _TRANSLUCENCY + if (f > texpaint0.a) { + float roughness = texpaint2.g; + float3 scatter_dir = cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + payload.ray_dir = normalize(mix(ray.direction, scatter_dir, roughness * roughness * 0.5)); + payload.ray_origin = hit + payload.ray_dir * 0.0001f; + scatter = true; + } + #endif + + if (!scatter) { #ifdef _TRANSLUCENCY - float3 diffuse_dir = texpaint0.a < f ? - cos_weighted_hemisphere_direction(tid, ray.direction, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank) : - cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); - #else - float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, thread_seed, constant_buffer.eye.w, mytexture_sobol, mytexture_scramble, mytexture_rank); + f = rand(tid.x, tid.y, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + bounce_seed += 1; #endif + float3 diffuse_dir = cos_weighted_hemisphere_direction(tid, n, payload.color.a, bounce_seed, frame, mytexture_sobol, mytexture_scramble, mytexture_rank); + #ifdef _FRESNEL float specular_chance = fresnel(n, ray.direction); #else @@ -295,12 +344,7 @@ kernel void raytracingKernel( #endif if (f < specular_chance) { - #ifdef _TRANSLUCENCY - float3 specular_dir = texpaint0.a < f * 2 ? ray.direction : reflect(ray.direction, n); - #else float3 specular_dir = reflect(ray.direction, n); - #endif - payload.ray_dir = mix(specular_dir, diffuse_dir, texpaint2.g * texpaint2.g); float3 specular = surface_specular(texcolor, texpaint2.b); payload.color.xyz *= specular; @@ -316,59 +360,44 @@ kernel void raytracingKernel( payload.color.xyz /= 1.0 - specular_chance; #endif } + #ifdef _FRESNEL payload.color.xyz *= 0.5; #endif - // float dotnv = abs(dot(n, -WorldRayDirection())); - // payload.ray_origin = hit_world_position() + n * mix(0.1f, 0.0001f, dotnv); - payload.ray_origin = hit_world_position(ray, intersection) + payload.ray_dir * 0.0001f; + payload.ray_origin = hit + payload.ray_dir * 0.0001f; #ifdef _EMISSION if (int(texpaint1.a * 255.0f) % 3 == 1) { // matid payload.color.xyz *= 100.0f; - payload.color.a = -2.0; + payload.color.a = -3.0; } #endif #ifdef _SUBSURFACE if (int(texpaint1.a * 255.0f) % 3 == 2) { - payload.ray_origin += ray.direction * f; + float d = min(1.0 / min(intersection.distance * 2.0, 1.0) / 10.0, 0.5); + payload.color.xyz += payload.color.xyz * d; + if (f < 0.5) { + payload.ray_origin += ray.direction * f * 0.001; + } } #endif } #ifdef _EMISSION - if (payload.color.a == -2) { + if (payload.color.a == -3.0) { accum += payload.color.rgb; break; } #endif - // Miss - if (payload.color.a < 0) { - #ifdef _TRANSPARENCY - if (payload.color.a == -2 && transparent_hits < DEPTH_TRANSPARENT) { - payload.color.a = j; - transparent_hits++; - i--; - } - #endif - - if (i == 0 && constant_buffer.params.x < 0) { // No envmap - payload.color.rgb = float3(0.0275, 0.0275, 0.0275); - } - - accum += clamp(payload.color.rgb, 0.0, 8.0); - break; - } + ray.origin = payload.ray_origin; + ray.direction = payload.ray_dir; #ifdef _ROULETTE payload.color.rgb *= rr_factor; #endif - - ray.origin = payload.ray_origin; - ray.direction = payload.ray_dir; } } @@ -384,6 +413,6 @@ kernel void raytracingKernel( if (constant_buffer.eye.w == 0) { color = accum; } - render_target.write(float4(mix(color, accum, 1.0 / 16.0), 1.0f), tid); + render_target.write(float4(mix(color, accum, 1.0 / 4.0), 1.0f), tid); #endif } diff --git a/base/sources/backends/metal_gpu.h b/base/sources/backends/metal_gpu.h index b6e428cd..c1eed64e 100644 --- a/base/sources/backends/metal_gpu.h +++ b/base/sources/backends/metal_gpu.h @@ -24,3 +24,5 @@ typedef struct { typedef struct { void *_accelerationStructure; } gpu_acceleration_structure_impl_t; + +#define GPU_RAYTRACE_MAX_OBJECTS 64 diff --git a/base/sources/backends/metal_gpu.m b/base/sources/backends/metal_gpu.m index 4835a386..c3eab382 100644 --- a/base/sources/backends/metal_gpu.m +++ b/base/sources/backends/metal_gpu.m @@ -704,13 +704,15 @@ static gpu_texture_t *_texenv; static gpu_texture_t *_texsobol; static gpu_texture_t *_texscramble; static gpu_texture_t *_texrank; -static gpu_buffer_t *vb[16]; -static gpu_buffer_t *vb_last[16]; -static gpu_buffer_t *ib[16]; +static gpu_buffer_t *vb[GPU_RAYTRACE_MAX_OBJECTS]; +static gpu_buffer_t *vb_last[GPU_RAYTRACE_MAX_OBJECTS]; +static gpu_buffer_t *ib[GPU_RAYTRACE_MAX_OBJECTS]; static int vb_count = 0; static int vb_count_last = 0; static inst_t instances[1024]; static int instances_count = 0; +static gpu_buffer_t *vb_full = NULL; +static gpu_buffer_t *ib_full = NULL; void gpu_raytrace_pipeline_init(gpu_raytrace_pipeline_t *pipeline, void *shader, int ray_shader_size, gpu_buffer_t *constant_buffer) { id device = get_metal_device(); @@ -774,7 +776,12 @@ id create_acceleration_sctructure(MTLAccelerationStruc void gpu_raytrace_acceleration_structure_init(gpu_acceleration_structure_t *accel) { vb_count = 0; instances_count = 0; - memset(vb_last, 0, sizeof(vb_last)); + if (gpu_raytrace_multi) { + memset(vb, 0, sizeof(vb)); + } + else { + memset(vb_last, 0, sizeof(vb_last)); + } } void gpu_raytrace_acceleration_structure_add(gpu_acceleration_structure_t *accel, gpu_buffer_t *_vb, gpu_buffer_t *_ib, mat4_t _transform) { @@ -787,12 +794,19 @@ void gpu_raytrace_acceleration_structure_add(gpu_acceleration_structure_t *accel } } if (vb_i == -1) { + if (vb_count >= GPU_RAYTRACE_MAX_OBJECTS) { + return; + } vb_i = vb_count; vb[vb_count] = _vb; ib[vb_count] = _ib; vb_count++; } + if (instances_count >= (int)(sizeof(instances) / sizeof(instances[0]))) { + return; + } + inst_t inst = {.i = vb_i, .m = _transform}; instances[instances_count] = inst; instances_count++; @@ -806,12 +820,14 @@ void _gpu_raytrace_acceleration_structure_destroy_bottom(gpu_acceleration_struct void _gpu_raytrace_acceleration_structure_destroy_top(gpu_acceleration_structure_t *accel) { _instance_accel = nil; + vb_full = NULL; + ib_full = NULL; } void gpu_raytrace_acceleration_structure_build(gpu_acceleration_structure_t *accel, gpu_buffer_t *_vb_full, gpu_buffer_t *_ib_full) { bool build_bottom = false; - for (int i = 0; i < 16; ++i) { + for (int i = 0; i < GPU_RAYTRACE_MAX_OBJECTS; ++i) { if (vb_last[i] != vb[i]) { build_bottom = true; } @@ -838,36 +854,59 @@ void gpu_raytrace_acceleration_structure_build(gpu_acceleration_structure_t *acc MTLResourceOptions options = MTLResourceStorageModeShared; - MTLAccelerationStructureTriangleGeometryDescriptor *descriptor = [MTLAccelerationStructureTriangleGeometryDescriptor descriptor]; - descriptor.indexType = MTLIndexTypeUInt32; - descriptor.indexBuffer = (__bridge id)ib[0]->impl.metal_buffer; - descriptor.vertexBuffer = (__bridge id)vb[0]->impl.metal_buffer; - descriptor.vertexStride = vb[0]->stride; - descriptor.triangleCount = ib[0]->count / 3; - descriptor.vertexFormat = MTLAttributeFormatShort4Normalized; + // Bottom level + if (build_bottom || _primitive_accels == nil) { + _primitive_accels = [[NSMutableArray alloc] init]; + for (int i = 0; i < vb_count; ++i) { + MTLAccelerationStructureTriangleGeometryDescriptor *descriptor = [MTLAccelerationStructureTriangleGeometryDescriptor descriptor]; + descriptor.indexType = MTLIndexTypeUInt32; + descriptor.indexBuffer = (__bridge id)ib[i]->impl.metal_buffer; + descriptor.vertexBuffer = (__bridge id)vb[i]->impl.metal_buffer; + descriptor.vertexStride = vb[i]->stride; + descriptor.triangleCount = ib[i]->count / 3; + descriptor.vertexFormat = MTLAttributeFormatShort4Normalized; - MTLPrimitiveAccelerationStructureDescriptor *accel_descriptor = [MTLPrimitiveAccelerationStructureDescriptor descriptor]; - accel_descriptor.geometryDescriptors = @[ descriptor ]; - id acceleration_structure = create_acceleration_sctructure(accel_descriptor); - _primitive_accels = [[NSMutableArray alloc] init]; - [_primitive_accels addObject:acceleration_structure]; + MTLPrimitiveAccelerationStructureDescriptor *accel_descriptor = [MTLPrimitiveAccelerationStructureDescriptor descriptor]; + accel_descriptor.geometryDescriptors = @[ descriptor ]; + [_primitive_accels addObject:create_acceleration_sctructure(accel_descriptor)]; + } + } - id instance_buffer = [device newBufferWithLength:sizeof(MTLAccelerationStructureInstanceDescriptor) * 1 options:options]; + // Top level + int instance_count = gpu_raytrace_multi ? instances_count : 1; - MTLAccelerationStructureInstanceDescriptor *instance_descriptors = (MTLAccelerationStructureInstanceDescriptor *)instance_buffer.contents; - instance_descriptors[0].accelerationStructureIndex = 0; - instance_descriptors[0].options = MTLAccelerationStructureInstanceOptionOpaque; - instance_descriptors[0].mask = 1; - instance_descriptors[0].transformationMatrix.columns[0] = MTLPackedFloat3Make(instances[0].m.m[0], instances[0].m.m[1], instances[0].m.m[2]); - instance_descriptors[0].transformationMatrix.columns[1] = MTLPackedFloat3Make(instances[0].m.m[4], instances[0].m.m[5], instances[0].m.m[6]); - instance_descriptors[0].transformationMatrix.columns[2] = MTLPackedFloat3Make(instances[0].m.m[8], instances[0].m.m[9], instances[0].m.m[10]); - instance_descriptors[0].transformationMatrix.columns[3] = MTLPackedFloat3Make(instances[0].m.m[12], instances[0].m.m[13], instances[0].m.m[14]); + id instance_buffer = [device newBufferWithLength:sizeof(MTLAccelerationStructureUserIDInstanceDescriptor) * instance_count options:options]; + MTLAccelerationStructureUserIDInstanceDescriptor *instance_descriptors = + (MTLAccelerationStructureUserIDInstanceDescriptor *)instance_buffer.contents; + + for (int i = 0; i < instance_count; ++i) { + float *m = instances[i].m.m; + + instance_descriptors[i].accelerationStructureIndex = instances[i].i; + instance_descriptors[i].options = MTLAccelerationStructureInstanceOptionOpaque; + instance_descriptors[i].mask = 1; + instance_descriptors[i].intersectionFunctionTableOffset = 0; + instance_descriptors[i].transformationMatrix.columns[0] = MTLPackedFloat3Make(m[0], m[1], m[2]); + instance_descriptors[i].transformationMatrix.columns[1] = MTLPackedFloat3Make(m[4], m[5], m[6]); + instance_descriptors[i].transformationMatrix.columns[2] = MTLPackedFloat3Make(m[8], m[9], m[10]); + instance_descriptors[i].transformationMatrix.columns[3] = MTLPackedFloat3Make(m[12], m[13], m[14]); + + uint32_t ib_off = 0; + for (int j = 0; j < instances[i].i; ++j) { + ib_off += ib[j]->count; + } + instance_descriptors[i].userID = ib_off; + } MTLInstanceAccelerationStructureDescriptor *inst_accel_descriptor = [MTLInstanceAccelerationStructureDescriptor descriptor]; inst_accel_descriptor.instancedAccelerationStructures = _primitive_accels; - inst_accel_descriptor.instanceCount = 1; + inst_accel_descriptor.instanceCount = instance_count; inst_accel_descriptor.instanceDescriptorBuffer = instance_buffer; + inst_accel_descriptor.instanceDescriptorType = MTLAccelerationStructureInstanceDescriptorTypeUserID; _instance_accel = create_acceleration_sctructure(inst_accel_descriptor); + + vb_full = gpu_raytrace_multi ? _vb_full : vb[0]; + ib_full = gpu_raytrace_multi ? _ib_full : ib[0]; } void gpu_raytrace_acceleration_structure_destroy(gpu_acceleration_structure_t *accel) {} @@ -899,6 +938,8 @@ void gpu_raytrace_dispatch_rays() { id device = get_metal_device(); if (!device.supportsRaytracing) return; + if (_instance_accel == nil || vb_full == NULL || ib_full == NULL) + return; dispatch_semaphore_wait(_semaphore, DISPATCH_TIME_FOREVER); id queue = get_metal_queue(); @@ -917,8 +958,8 @@ void gpu_raytrace_dispatch_rays() { id compute_encoder = [command_buffer computeCommandEncoder]; [compute_encoder setBuffer:(__bridge id)constant_buf->impl.metal_buffer offset:0 atIndex:0]; [compute_encoder setAccelerationStructure:_instance_accel atBufferIndex:1]; - [compute_encoder setBuffer:(__bridge id)ib[0]->impl.metal_buffer offset:0 atIndex:2]; - [compute_encoder setBuffer:(__bridge id)vb[0]->impl.metal_buffer offset:0 atIndex:3]; + [compute_encoder setBuffer:(__bridge id)ib_full->impl.metal_buffer offset:0 atIndex:2]; + [compute_encoder setBuffer:(__bridge id)vb_full->impl.metal_buffer offset:0 atIndex:3]; [compute_encoder setTexture:(__bridge id)output->impl._tex atIndex:0]; [compute_encoder setTexture:(__bridge id)_texpaint0->impl._tex atIndex:1]; [compute_encoder setTexture:(__bridge id)_texpaint1->impl._tex atIndex:2]; diff --git a/paint/sources/render/render_path_raytrace.c b/paint/sources/render/render_path_raytrace.c index 333588ca..9d3fa65e 100644 --- a/paint/sources/render/render_path_raytrace.c +++ b/paint/sources/render/render_path_raytrace.c @@ -76,13 +76,8 @@ void render_path_raytrace_commands(bool use_live_layer) { render_path_raytrace_f32a->buffer[1] = transform_world_y(ct); render_path_raytrace_f32a->buffer[2] = transform_world_z(ct); render_path_raytrace_f32a->buffer[3] = render_path_raytrace_frame; -#ifdef IRON_METAL - // render_path_raytrace_frame = (render_path_raytrace_frame % (16)) + 1; // _PAINT - render_path_raytrace_frame = render_path_raytrace_frame + 1; // _RENDER -#else - render_path_raytrace_frame = (render_path_raytrace_frame % 4) + 1; // _PAINT + render_path_raytrace_frame = (render_path_raytrace_frame % 4) + 1; // _PAINT // render_path_raytrace_frame = render_path_raytrace_frame + 1; // _RENDER -#endif render_path_raytrace_f32a->buffer[4] = render_path_raytrace_help_mat.m00; render_path_raytrace_f32a->buffer[5] = render_path_raytrace_help_mat.m01; render_path_raytrace_f32a->buffer[6] = render_path_raytrace_help_mat.m02; @@ -114,11 +109,7 @@ void render_path_raytrace_commands(bool use_live_layer) { _gpu_raytrace_dispatch_rays(framebuffer->_image, render_path_raytrace_f32a); if (g_context->ddirty == 1 || g_context->pdirty == 1) { -#ifdef IRON_METAL - g_context->rdirty = 128; -#else g_context->rdirty = 4; -#endif } g_context->ddirty--; g_context->pdirty--; @@ -193,14 +184,6 @@ void render_path_raytrace_draw(bool use_live_layer) { render_path_raytrace_frame = 0; } -#ifdef IRON_METAL - // Delay path tracing additional samples while painting - bool down = mouse_down("left") || pen_down("tip"); - if (context_in_3d_view() && down && !ui_menu_show) { - render_path_raytrace_frame = 0; - } -#endif - render_path_raytrace_commands(use_live_layer); render_path_set_target("buf", NULL, NULL, GPU_CLEAR_NONE, 0, 0.0); render_path_draw_meshes("overlay");