trending/armorpaint
1
0
Fork 0
mirror of https://github.com/armory3d/armorpaint.git synced 2026-10-02 02:54:34 +08:00
Code Issues Packages Projects Releases Wiki Activity

base: update raytrace shaders

Browse Source
This commit is contained in:
luboslenco
2026-09-17 21:15:14 +02:00
parent feff1fd87d
commit 678c1fd4d5
6 changed files with 103 additions and 88 deletions
Download Patch File Download Diff File Expand all files Collapse all files
base/shaders/raytrace/raytrace_bake_ao.metal
+1 -3
View File
@@ -76,9 +76,7 @@ kernel void raytracingKernel(
texture2d<float, access::read> mytexture_sobol [[texture(5)]],
texture2d<float, access::read> mytexture_scramble [[texture(6)]],
texture2d<float, access::read> mytexture_rank [[texture(7)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
instance_acceleration_structure scene [[buffer(1)]]
) {
uint seed = 0;
base/shaders/raytrace/raytrace_bake_bent.metal
+1 -3
View File
@@ -76,9 +76,7 @@ kernel void raytracingKernel(
texture2d<float, access::read> mytexture_sobol [[texture(5)]],
texture2d<float, access::read> mytexture_scramble [[texture(6)]],
texture2d<float, access::read> mytexture_rank [[texture(7)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
instance_acceleration_structure scene [[buffer(1)]]
) {
uint seed = 0;
base/shaders/raytrace/raytrace_bake_light.metal
+26 -19
View File
@@ -9,6 +9,19 @@ struct Vertex {
uint tex;
};
struct Instance {
constant uint *vertex_buffer;
constant uint *index_buffer;
uint stride; // Vertex size in bytes
uint geometry; // Index into the geometry textures
};
struct GeometryTextures {
texture2d<float, access::read> texpaint0;
texture2d<float, access::read> texpaint1;
texture2d<float, access::read> texpaint2;
};
struct RayGenConstantBuffer {
float4 v0; // frame, strength, radius, offset
float4 v1;
@@ -103,8 +116,8 @@ kernel void raytracingKernel(
texture2d<float, access::read> mytexture_scramble [[texture(6)]],
texture2d<float, access::read> mytexture_rank [[texture(7)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
constant Instance *instances [[buffer(2)]],
constant GeometryTextures *geometry_textures [[buffer(3)]]
) {
uint seed = 0;
@@ -143,31 +156,25 @@ kernel void raytracingKernel(
payload.color = float4(texenv.rgb, -1);
}
else {
device uint32_t *inda = (device uint32_t *)(indices);
constant Instance &inst = instances[intersection.user_instance_id];
uint3 indices_sample = uint3(
inda[intersection.primitive_id * 3],
inda[intersection.primitive_id * 3 + 1],
inda[intersection.primitive_id * 3 + 2]
inst.index_buffer[intersection.primitive_id * 3],
inst.index_buffer[intersection.primitive_id * 3 + 1],
inst.index_buffer[intersection.primitive_id * 3 + 2]
);
device Vertex *verta = (device Vertex *)(vertices);
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)
};
// Texture coordinates are the last 4 bytes of the base vertex layout
float2 barycentrics = intersection.triangle_barycentric_coord;
float3 n = normalize(hit_attribute(vertex_normals, barycentrics));
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)
s16_to_f32(inst.vertex_buffer[(indices_sample[0] * inst.stride + 12) / 4]),
s16_to_f32(inst.vertex_buffer[(indices_sample[1] * inst.stride + 12) / 4]),
s16_to_f32(inst.vertex_buffer[(indices_sample[2] * inst.stride + 12) / 4])
};
float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics);
uint2 size = uint2(mytexture2.get_width(), mytexture2.get_height());
float3 texpaint2 = pow(mytexture2.read(uint2(tex_coord * float2(size)), 0).rgb, 2.2); // layer base
texture2d<float, access::read> hit_texture2 = geometry_textures[inst.geometry].texpaint2;
uint2 size = uint2(hit_texture2.get_width(), hit_texture2.get_height());
float3 texpaint2 = pow(hit_texture2.read(uint2(tex_coord * float2(size)), 0).rgb, 2.2); // layer base
payload.color.rgb = texpaint2.rgb;
}
base/shaders/raytrace/raytrace_bake_thick.metal
+1 -3
View File
@@ -76,9 +76,7 @@ kernel void raytracingKernel(
texture2d<float, access::read> mytexture_sobol [[texture(5)]],
texture2d<float, access::read> mytexture_scramble [[texture(6)]],
texture2d<float, access::read> mytexture_rank [[texture(7)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
instance_acceleration_structure scene [[buffer(1)]]
) {
uint seed = 0;
base/shaders/raytrace/raytrace_brute_core.metal
+37 -30
View File
@@ -22,11 +22,20 @@ struct Vertex {
uint tex;
};
#ifdef _MULTI
struct Instance {
constant uint *vertex_buffer;
constant uint *index_buffer;
uint stride; // Vertex size in bytes
uint geometry; // Index into the geometry textures
};
struct GeometryTextures {
texture2d<float, access::read> texpaint0;
texture2d<float, access::read> texpaint1;
texture2d<float, access::read> texpaint2;
};
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
@@ -377,6 +386,16 @@ bool occluded(float3 origin, float3 dir, instance_acceleration_structure scene)
}
#endif
Vertex load_vertex(constant uint *vb, uint offset) {
uint base = offset / 4;
Vertex v;
v.posxy = vb[base];
v.poszw = vb[base + 1];
v.nor = vb[base + 2];
v.tex = vb[base + 3];
return v;
}
uint2 texel_coord(float2 tex_coord, float2 tex_size) {
return uint2(fract(tex_coord) * tex_size);
}
@@ -385,9 +404,6 @@ 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)]],
@@ -397,8 +413,8 @@ kernel void raytracingKernel(
#endif
sampler linear_sampler [[sampler(0)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
constant Instance *instances [[buffer(2)]],
constant GeometryTextures *geometry_textures [[buffer(3)]]
) {
uint2 dim = uint2(render_target.get_width(), render_target.get_height());
if (tid.x >= dim.x || tid.y >= dim.y) {
@@ -409,7 +425,6 @@ kernel void raytracingKernel(
uint2 scramble_cache = init_scramble(tid, frame, mytexture_scramble);
uint4 rank_cache = init_rank(tid, frame, mytexture_rank);
float2 tex_size = float2(mytexture0.get_width(), mytexture0.get_height());
float3 accum = float3(0, 0, 0);
for (int j = 0; j < SAMPLES; ++j) {
@@ -474,23 +489,17 @@ kernel void raytracingKernel(
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
constant Instance &inst = instances[intersection.user_instance_id];
uint3 indices_sample = uint3(
inda[base_index],
inda[base_index + 1],
inda[base_index + 2]
inst.index_buffer[base_index],
inst.index_buffer[base_index + 1],
inst.index_buffer[base_index + 2]
);
device Vertex *verta = (device Vertex *)(vertices);
Vertex a0 = verta[indices_sample[0]];
Vertex a1 = verta[indices_sample[1]];
Vertex a2 = verta[indices_sample[2]];
Vertex a0 = load_vertex(inst.vertex_buffer, indices_sample[0] * inst.stride);
Vertex a1 = load_vertex(inst.vertex_buffer, indices_sample[1] * inst.stride);
Vertex a2 = load_vertex(inst.vertex_buffer, indices_sample[2] * inst.stride);
float2 vertex_uvs[3] = {
s16_to_f32(a0.tex),
@@ -501,8 +510,10 @@ kernel void raytracingKernel(
float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
float3 hit = hit_world_position(ray, intersection);
constant GeometryTextures &textures = geometry_textures[inst.geometry];
float2 tex_size = float2(textures.texpaint0.get_width(), textures.texpaint0.get_height());
uint2 texel = texel_coord(tex_coord, tex_size);
float4 texpaint0 = mytexture0.read(texel, 0);
float4 texpaint0 = textures.texpaint0.read(texel, 0);
float2 zw0 = s16_to_f32(a0.poszw);
float2 zw1 = s16_to_f32(a1.poszw);
@@ -525,12 +536,10 @@ kernel void raytracingKernel(
float3 n_object = n;
#ifdef _MULTI
float4x3 o2w = intersection.object_to_world_transform;
float3x3 obj_to_world = float3x3(o2w[0], o2w[1], o2w[2]);
n = normalize(obj_to_world * n);
ng = normalize(obj_to_world * ng);
#endif
bool back_face = dot(ng, ray.direction) > 0.0;
if (back_face) {
@@ -546,7 +555,7 @@ kernel void raytracingKernel(
float4 texpaint1 = float4(0.0);
if (normal_map) {
texpaint1 = mytexture1.read(texel, 0);
texpaint1 = textures.texpaint1.read(texel, 0);
}
float3 texcolor = srgb_to_linear(texpaint0.rgb);
@@ -564,7 +573,7 @@ kernel void raytracingKernel(
}
#endif
float4 texpaint2 = mytexture2.read(texel, 0);
float4 texpaint2 = textures.texpaint2.read(texel, 0);
float f = rand(tid.x, tid.y, sample_index, dim_base + DIM_SELECT, frame, scramble_cache, rank_cache, mytexture_sobol, mytexture_rank);
@@ -592,12 +601,10 @@ kernel void raytracingKernel(
create_uv_basis(vertex_positions[0], vertex_positions[1], vertex_positions[2],
vertex_uvs[0], vertex_uvs[1], vertex_uvs[2], n_object, tangent, binormal);
#ifdef _MULTI
tangent = obj_to_world * tangent;
binormal = obj_to_world * binormal;
tangent = normalize(tangent - n * dot(n, tangent));
binormal = normalize(binormal - n * dot(n, binormal) - tangent * dot(tangent, binormal));
#endif
if (back_face) {
binormal = -binormal;
base/shaders/raytrace/raytrace_brute_full.metal
+37 -30
View File
@@ -23,11 +23,20 @@ struct Vertex {
uint tex;
};
#ifdef _MULTI
struct Instance {
constant uint *vertex_buffer;
constant uint *index_buffer;
uint stride; // Vertex size in bytes
uint geometry; // Index into the geometry textures
};
struct GeometryTextures {
texture2d<float, access::read> texpaint0;
texture2d<float, access::read> texpaint1;
texture2d<float, access::read> texpaint2;
};
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
@@ -378,6 +387,16 @@ bool occluded(float3 origin, float3 dir, instance_acceleration_structure scene)
}
#endif
Vertex load_vertex(constant uint *vb, uint offset) {
uint base = offset / 4;
Vertex v;
v.posxy = vb[base];
v.poszw = vb[base + 1];
v.nor = vb[base + 2];
v.tex = vb[base + 3];
return v;
}
uint2 texel_coord(float2 tex_coord, float2 tex_size) {
return uint2(fract(tex_coord) * tex_size);
}
@@ -386,9 +405,6 @@ 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)]],
@@ -398,8 +414,8 @@ kernel void raytracingKernel(
#endif
sampler linear_sampler [[sampler(0)]],
instance_acceleration_structure scene [[buffer(1)]],
device void *indices [[buffer(2)]],
device void *vertices [[buffer(3)]]
constant Instance *instances [[buffer(2)]],
constant GeometryTextures *geometry_textures [[buffer(3)]]
) {
uint2 dim = uint2(render_target.get_width(), render_target.get_height());
if (tid.x >= dim.x || tid.y >= dim.y) {
@@ -410,7 +426,6 @@ kernel void raytracingKernel(
uint2 scramble_cache = init_scramble(tid, frame, mytexture_scramble);
uint4 rank_cache = init_rank(tid, frame, mytexture_rank);
float2 tex_size = float2(mytexture0.get_width(), mytexture0.get_height());
float3 accum = float3(0, 0, 0);
for (int j = 0; j < SAMPLES; ++j) {
@@ -475,23 +490,17 @@ kernel void raytracingKernel(
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
constant Instance &inst = instances[intersection.user_instance_id];
uint3 indices_sample = uint3(
inda[base_index],
inda[base_index + 1],
inda[base_index + 2]
inst.index_buffer[base_index],
inst.index_buffer[base_index + 1],
inst.index_buffer[base_index + 2]
);
device Vertex *verta = (device Vertex *)(vertices);
Vertex a0 = verta[indices_sample[0]];
Vertex a1 = verta[indices_sample[1]];
Vertex a2 = verta[indices_sample[2]];
Vertex a0 = load_vertex(inst.vertex_buffer, indices_sample[0] * inst.stride);
Vertex a1 = load_vertex(inst.vertex_buffer, indices_sample[1] * inst.stride);
Vertex a2 = load_vertex(inst.vertex_buffer, indices_sample[2] * inst.stride);
float2 vertex_uvs[3] = {
s16_to_f32(a0.tex),
@@ -502,8 +511,10 @@ kernel void raytracingKernel(
float2 tex_coord = hit_attribute2d(vertex_uvs, barycentrics) * constant_buffer.params.z;
float3 hit = hit_world_position(ray, intersection);
constant GeometryTextures &textures = geometry_textures[inst.geometry];
float2 tex_size = float2(textures.texpaint0.get_width(), textures.texpaint0.get_height());
uint2 texel = texel_coord(tex_coord, tex_size);
float4 texpaint0 = mytexture0.read(texel, 0);
float4 texpaint0 = textures.texpaint0.read(texel, 0);
float2 zw0 = s16_to_f32(a0.poszw);
float2 zw1 = s16_to_f32(a1.poszw);
@@ -526,12 +537,10 @@ kernel void raytracingKernel(
float3 n_object = n;
#ifdef _MULTI
float4x3 o2w = intersection.object_to_world_transform;
float3x3 obj_to_world = float3x3(o2w[0], o2w[1], o2w[2]);
n = normalize(obj_to_world * n);
ng = normalize(obj_to_world * ng);
#endif
bool back_face = dot(ng, ray.direction) > 0.0;
if (back_face) {
@@ -547,7 +556,7 @@ kernel void raytracingKernel(
float4 texpaint1 = float4(0.0);
if (normal_map) {
texpaint1 = mytexture1.read(texel, 0);
texpaint1 = textures.texpaint1.read(texel, 0);
}
float3 texcolor = srgb_to_linear(texpaint0.rgb);
@@ -565,7 +574,7 @@ kernel void raytracingKernel(
}
#endif
float4 texpaint2 = mytexture2.read(texel, 0);
float4 texpaint2 = textures.texpaint2.read(texel, 0);
float f = rand(tid.x, tid.y, sample_index, dim_base + DIM_SELECT, frame, scramble_cache, rank_cache, mytexture_sobol, mytexture_rank);
@@ -593,12 +602,10 @@ kernel void raytracingKernel(
create_uv_basis(vertex_positions[0], vertex_positions[1], vertex_positions[2],
vertex_uvs[0], vertex_uvs[1], vertex_uvs[2], n_object, tangent, binormal);
#ifdef _MULTI
tangent = obj_to_world * tangent;
binormal = obj_to_world * binormal;
tangent = normalize(tangent - n * dot(n, tangent));
binormal = normalize(binormal - n * dot(n, binormal) - tangent * dot(tangent, binormal));
#endif
if (back_face) {
binormal = -binormal;