metal: raytrace fixes

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<triangle_data, instancing, world_space_data> intersector_t;
+#else
+typedef intersector<triangle_data, instancing> 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<triangle_data, instancing>::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<float, access::read> 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,37 +211,47 @@ 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<triangle_data, instancing> 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<triangle_data, instancing>::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;
+			}
+
 			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[intersection.primitive_id * 3],
-					inda[intersection.primitive_id * 3 + 1],
-					inda[intersection.primitive_id * 3 + 2]
+				inda[base_index],
+				inda[base_index + 1],
+				inda[base_index + 2]
 			);
 
 			device Vertex *verta = (device Vertex *)(vertices);
@@ -243,16 +263,21 @@ kernel void raytracingKernel(
 			float2 barycentrics = intersection.triangle_barycentric_coord;
 			float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
 
-				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);
+			float3 hit = hit_world_position(ray, intersection);
+			float4 texpaint0 = read_texel(mytexture0, tex_coord);
 
 			#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;
+			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;
-					return;
+				}
+				continue;
 			}
 			#endif
 
@@ -263,10 +288,22 @@ kernel void raytracingKernel(
 			};
 			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);
+			#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);
@@ -275,17 +312,31 @@ kernel void raytracingKernel(
 			texpaint1.g = -texpaint1.g;
 			n = float3x3(tangent, binormal, n) * texpaint1.rgb;
 
-				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;
+			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
-				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);
+			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
@@ -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
 }

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<triangle_data, instancing, world_space_data> intersector_t;
+#else
+typedef intersector<triangle_data, instancing> 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<triangle_data, instancing>::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<float, access::read> 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,37 +212,47 @@ 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<triangle_data, instancing> 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<triangle_data, instancing>::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;
+			}
+
 			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[intersection.primitive_id * 3],
-					inda[intersection.primitive_id * 3 + 1],
-					inda[intersection.primitive_id * 3 + 2]
+				inda[base_index],
+				inda[base_index + 1],
+				inda[base_index + 2]
 			);
 
 			device Vertex *verta = (device Vertex *)(vertices);
@@ -245,16 +264,21 @@ kernel void raytracingKernel(
 			float2 barycentrics = intersection.triangle_barycentric_coord;
 			float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
 
-				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);
+			float3 hit = hit_world_position(ray, intersection);
+			float4 texpaint0 = read_texel(mytexture0, tex_coord);
 
 			#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;
+			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;
-					return;
+				}
+				continue;
 			}
 			#endif
 
@@ -265,10 +289,22 @@ kernel void raytracingKernel(
 			};
 			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);
+			#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);
@@ -277,17 +313,31 @@ kernel void raytracingKernel(
 			texpaint1.g = -texpaint1.g;
 			n = float3x3(tangent, binormal, n) * texpaint1.rgb;
 
-				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;
+			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
-				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);
+			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
@@ -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
 }

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<triangle_data, instancing, world_space_data> intersector_t;
+#else
+typedef intersector<triangle_data, instancing> 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<triangle_data, instancing>::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<float, access::read> 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,37 +212,47 @@ 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<triangle_data, instancing> 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<triangle_data, instancing>::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;
+			}
+
 			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[intersection.primitive_id * 3],
-					inda[intersection.primitive_id * 3 + 1],
-					inda[intersection.primitive_id * 3 + 2]
+				inda[base_index],
+				inda[base_index + 1],
+				inda[base_index + 2]
 			);
 
 			device Vertex *verta = (device Vertex *)(vertices);
@@ -243,16 +264,21 @@ kernel void raytracingKernel(
 			float2 barycentrics = intersection.triangle_barycentric_coord;
 			float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
 
-				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);
+			float3 hit = hit_world_position(ray, intersection);
+			float4 texpaint0 = read_texel(mytexture0, tex_coord);
 
 			#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;
+			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;
-					return;
+				}
+				continue;
 			}
 			#endif
 
@@ -263,10 +289,22 @@ kernel void raytracingKernel(
 			};
 			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);
+			#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);
@@ -275,17 +313,31 @@ kernel void raytracingKernel(
 			texpaint1.g = -texpaint1.g;
 			n = float3x3(tangent, binormal, n) * texpaint1.rgb;
 
-				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;
+			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
-				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);
+			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
@@ -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
 }

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<triangle_data, instancing, world_space_data> intersector_t;
+#else
+typedef intersector<triangle_data, instancing> 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<float, access::read> sobol, texture2d<float, access::read> scramble, texture2d<float, access::read> 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<float, access::read> sobol, texture2d<float, access::read> scramble, texture2d<float, access::read> 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<float, access::read> 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<float, access::read_write> render_target [[texture(0)]],
+	texture2d<float, access::read> mytexture0 [[texture(1)]],
+	texture2d<float, access::read> mytexture1 [[texture(2)]],
+	texture2d<float, access::read> mytexture2 [[texture(3)]],
+	texture2d<float, access::sample> mytexture_env [[texture(4)]],
+	texture2d<float, access::read> mytexture_sobol [[texture(5)]],
+	texture2d<float, access::read> mytexture_scramble [[texture(6)]],
+	texture2d<float, access::read> 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
+}

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

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<triangle_data, instancing, world_space_data> intersector_t;
+#else
+typedef intersector<triangle_data, instancing> 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<triangle_data, instancing>::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<float, access::read> 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,37 +211,47 @@ 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<triangle_data, instancing> 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<triangle_data, instancing>::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;
+			}
+
 			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[intersection.primitive_id * 3],
-					inda[intersection.primitive_id * 3 + 1],
-					inda[intersection.primitive_id * 3 + 2]
+				inda[base_index],
+				inda[base_index + 1],
+				inda[base_index + 2]
 			);
 
 			device Vertex *verta = (device Vertex *)(vertices);
@@ -245,16 +263,21 @@ kernel void raytracingKernel(
 			float2 barycentrics = intersection.triangle_barycentric_coord;
 			float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
 
-				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);
+			float3 hit = hit_world_position(ray, intersection);
+			float4 texpaint0 = read_texel(mytexture0, tex_coord);
 
 			#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;
+			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;
-					return;
+				}
+				continue;
 			}
 			#endif
 
@@ -265,10 +288,22 @@ kernel void raytracingKernel(
 			};
 			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);
+			#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);
@@ -277,17 +312,31 @@ kernel void raytracingKernel(
 			texpaint1.g = -texpaint1.g;
 			n = float3x3(tangent, binormal, n) * texpaint1.rgb;
 
-				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;
+			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
-				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);
+			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
@@ -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
 }

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

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<MTLDevice> device = get_metal_device();
@@ -774,7 +776,12 @@ id<MTLAccelerationStructure> create_acceleration_sctructure(MTLAccelerationStruc
 void gpu_raytrace_acceleration_structure_init(gpu_acceleration_structure_t *accel) {
 	vb_count        = 0;
 	instances_count = 0;
+	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;
 
+	// 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<MTLBuffer>)ib[0]->impl.metal_buffer;
-	descriptor.vertexBuffer                                        = (__bridge id<MTLBuffer>)vb[0]->impl.metal_buffer;
-	descriptor.vertexStride                                        = vb[0]->stride;
-	descriptor.triangleCount                                       = ib[0]->count / 3;
+			descriptor.indexBuffer                                         = (__bridge id<MTLBuffer>)ib[i]->impl.metal_buffer;
+			descriptor.vertexBuffer                                        = (__bridge id<MTLBuffer>)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<MTLAccelerationStructure> acceleration_structure           = create_acceleration_sctructure(accel_descriptor);
-	_primitive_accels                                             = [[NSMutableArray alloc] init];
-	[_primitive_accels addObject:acceleration_structure];
+			[_primitive_accels addObject:create_acceleration_sctructure(accel_descriptor)];
+		}
+	}
 
-	id<MTLBuffer> 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<MTLBuffer> 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<MTLDevice> 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<MTLCommandQueue>          queue          = get_metal_queue();
@@ -917,8 +958,8 @@ void gpu_raytrace_dispatch_rays() {
 	id<MTLComputeCommandEncoder> compute_encoder = [command_buffer computeCommandEncoder];
 	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)constant_buf->impl.metal_buffer offset:0 atIndex:0];
 	[compute_encoder setAccelerationStructure:_instance_accel atBufferIndex:1];
-	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)ib[0]->impl.metal_buffer offset:0 atIndex:2];
-	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)vb[0]->impl.metal_buffer offset:0 atIndex:3];
+	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)ib_full->impl.metal_buffer offset:0 atIndex:2];
+	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)vb_full->impl.metal_buffer offset:0 atIndex:3];
 	[compute_encoder setTexture:(__bridge id<MTLTexture>)output->impl._tex atIndex:0];
 	[compute_encoder setTexture:(__bridge id<MTLTexture>)_texpaint0->impl._tex atIndex:1];
 	[compute_encoder setTexture:(__bridge id<MTLTexture>)_texpaint1->impl._tex atIndex:2];

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 + 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");