Skip to content

Commit 3525cd4

Browse files
jan-van-bergenjan-van-bergen
committed
Curvature is now negated when on the opposite side of the Triangle -- side of the Triangle is now decided based on the geometric normal
1 parent 41c1a6a commit 3525cd4

File tree

4 files changed

+78
-73
lines changed

4 files changed

+78
-73
lines changed

‎Src/CUDA/Pathtracer.cu‎

Lines changed: 38 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -618,15 +618,12 @@ __device__ void shade_material(int bounce, int sample_index, int buffer_size) {
618618

619619
normal = normalize(normal);
620620

621-
// Make sure the normal is always pointing outwards
622-
bool entering_material = dot(ray_direction, normal) < 0.0f;
623-
if (!entering_material) {
624-
normal = -normal;
625-
}
621+
float mesh_scale_inv = 1.0f / mesh_get_scale(hit.mesh_id);
626622

627623
// Load and propagate Ray Cone
628624
float cone_angle;
629625
float cone_width;
626+
float curvature = 0.0f;
630627
if (config.enable_mipmapping) {
631628
if (bounce == 0) {
632629
cone_angle = camera.pixel_spread_angle;
@@ -635,6 +632,30 @@ __device__ void shade_material(int bounce, int sample_index, int buffer_size) {
635632
cone_angle = material_buffer.buffer->cone_angle[index];
636633
cone_width = material_buffer.buffer->cone_width[index] + cone_angle * hit.t;
637634
}
635+
636+
// Calculate Triangle curvature here,
637+
// not yet needed (see below) but after this step the Triangle position edges are transformed into world space
638+
curvature = triangle_get_curvature(
639+
hit_triangle.position_edge_1,
640+
hit_triangle.position_edge_2,
641+
hit_triangle.normal_edge_1,
642+
hit_triangle.normal_edge_2
643+
) * mesh_scale_inv;
644+
}
645+
646+
matrix3x4_transform_direction(world, hit_triangle.position_edge_1);
647+
matrix3x4_transform_direction(world, hit_triangle.position_edge_2);
648+
649+
// Calculate geometric normal (in world space) to the Triangle
650+
float3 geometric_normal = cross(hit_triangle.position_edge_1, hit_triangle.position_edge_2);
651+
float triangle_double_area_inv = 1.0f / length(geometric_normal);
652+
geometric_normal *= triangle_double_area_inv; // Normalize
653+
654+
// Check which side of the Triangle we are on based on its geometric normal
655+
bool entering_material = dot(ray_direction, geometric_normal) < 0.0f;
656+
if (!entering_material) {
657+
normal = -normal;
658+
curvature = -curvature;
638659
}
639660

640661
// Construct TBN frame
@@ -643,6 +664,8 @@ __device__ void shade_material(int bounce, int sample_index, int buffer_size) {
643664

644665
float3 omega_i = world_to_local(-ray_direction, tangent, bitangent, normal);
645666

667+
if (omega_i.z <= 0.0f) return; // Below hemisphere, reject
668+
646669
// Initialize BSDF
647670
int material_id = mesh_get_material_id(hit.mesh_id);
648671

@@ -656,37 +679,21 @@ __device__ void shade_material(int bounce, int sample_index, int buffer_size) {
656679
bsdf.omega_i = omega_i;
657680
bsdf.init(bounce, entering_material, material_id);
658681

659-
// Calculate texture level of detail
660-
float mesh_scale = mesh_get_scale(hit.mesh_id);
661-
662-
// Calculate geometric normal to the triangle.
663-
// NOTE: geometric normal is kept in local space since ray_cone_get_texture_gradients() requires this,
664-
// it is only later transformed into world space
665-
float3 geometric_normal = cross(hit_triangle.position_edge_1, hit_triangle.position_edge_2);
666-
float triangle_area_inv = 1.0f / length(geometric_normal);
667-
geometric_normal *= triangle_area_inv; // Normalize
668-
669682
if (BSDF::HAS_ALBEDO) {
670683
TextureLOD lod;
671684

672685
if (config.enable_mipmapping && bsdf.has_texture()) {
673686
if (use_anisotropic_texture_sampling(bounce)) {
674-
float3 ellipse_axis_1, ellipse_axis_2;
675-
ray_cone_get_ellipse_axes(ray_direction, normal, cone_width, ellipse_axis_1, ellipse_axis_2);
676-
677-
ray_cone_get_texture_gradients(
678-
mesh_scale,
679-
geometric_normal,
680-
triangle_area_inv,
681-
hit_triangle.position_0, hit_triangle.position_edge_1, hit_triangle.position_edge_2,
682-
hit_triangle.tex_coord_0, hit_triangle.tex_coord_edge_1, hit_triangle.tex_coord_edge_2,
683-
hit_point_local, tex_coord,
684-
ellipse_axis_1, ellipse_axis_2,
685-
lod.aniso.gradient_1, lod.aniso.gradient_2
686-
);
687+
float3 ellipse_axis_1;
688+
float3 ellipse_axis_2;
689+
ray_cone_get_ellipse_axes(ray_direction, geometric_normal, cone_width, ellipse_axis_1, ellipse_axis_2);
690+
691+
lod.aniso.gradient_1 = ray_cone_ellipse_axis_to_gradient(hit_triangle, triangle_double_area_inv, geometric_normal, hit_point, tex_coord, ellipse_axis_1);
692+
lod.aniso.gradient_2 = ray_cone_ellipse_axis_to_gradient(hit_triangle, triangle_double_area_inv, geometric_normal, hit_point, tex_coord, ellipse_axis_2);
687693
} else {
688-
float lod_triangle = sqrtf(triangle_get_lod(mesh_scale, triangle_area_inv, hit_triangle.tex_coord_edge_1, hit_triangle.tex_coord_edge_2));
689-
float lod_ray_cone = ray_cone_get_lod(ray_direction, normal, cone_width);
694+
float lod_triangle = triangle_get_lod(triangle_double_area_inv, hit_triangle.tex_coord_edge_1, hit_triangle.tex_coord_edge_2);
695+
float lod_ray_cone = ray_cone_get_lod(ray_direction, geometric_normal, cone_width);
696+
690697
lod.iso.lod = log2f(lod_triangle * lod_ray_cone);
691698
}
692699
}
@@ -701,18 +708,8 @@ __device__ void shade_material(int bounce, int sample_index, int buffer_size) {
701708
aov_framebuffer_set(AOVType::POSITION, pixel_index, make_float4(hit_point));
702709
}
703710

704-
matrix3x4_transform_direction(world, geometric_normal);
705-
geometric_normal = normalize(geometric_normal);
706-
707-
// Calulate new Ray Cone angle based on Mesh curvature
711+
// Calulate new Ray Cone angle
708712
if (config.enable_mipmapping) {
709-
float curvature = triangle_get_curvature(
710-
hit_triangle.position_edge_1,
711-
hit_triangle.position_edge_2,
712-
hit_triangle.normal_edge_1,
713-
hit_triangle.normal_edge_2
714-
) / mesh_scale;
715-
716713
cone_angle -= 2.0f * curvature * fabsf(cone_width) / dot(normal, ray_direction); // Eq. 5 (Akenine-Möller 2021)
717714
}
718715

‎Src/CUDA/RayCone.h‎

Lines changed: 14 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
11
#pragma once
22
#include "Config.h"
3+
#include "Triangle.h"
34

45
union TextureLOD {
56
struct {
@@ -27,6 +28,7 @@ __device__ inline float3 sample_albedo(int bounce, const float3 & diffuse, int t
2728
}
2829
}
2930

31+
// Project the Ray Cone onto the Triangle and obtain two axes that describe the resulting ellipse (in world space)
3032
__device__ inline void ray_cone_get_ellipse_axes(
3133
const float3 & ray_direction,
3234
const float3 & geometric_normal,
@@ -41,36 +43,21 @@ __device__ inline void ray_cone_get_ellipse_axes(
4143
ellipse_axis_2 = cone_width / max(0.0001f, length(h_2 - dot(ray_direction, h_2) * ray_direction)) * h_2;
4244
}
4345

44-
__device__ inline void ray_cone_get_texture_gradients(
45-
float mesh_scale,
46-
const float3 & geometric_normal,
47-
float triangle_area_inv,
48-
const float3 & position_0,
49-
const float3 & position_edge_1,
50-
const float3 & position_edge_2,
51-
const float2 & tex_coord_0,
52-
const float2 & tex_coord_edge_1,
53-
const float2 & tex_coord_edge_2,
54-
const float3 & hit_point,
55-
const float2 & hit_tex_coord,
56-
const float3 & ellipse_axis_1,
57-
const float3 & ellipse_axis_2,
58-
float2 & gradient_1,
59-
float2 & gradient_2
46+
// Convert an ellipse axis into texture space
47+
__device__ inline float2 ray_cone_ellipse_axis_to_gradient(
48+
const TrianglePosNorTex & triangle,
49+
float triangle_double_area_inv,
50+
const float3 & geometric_normal,
51+
const float3 & hit_point,
52+
const float2 & hit_tex_coord,
53+
const float3 & ellipse_axis
6054
) {
61-
float3 e_p = hit_point + ellipse_axis_1 - position_0;
62-
63-
float inv_mesh_scale = 1.0f / mesh_scale;
64-
65-
float u_1 = dot(geometric_normal, cross(e_p, position_edge_2)) * triangle_area_inv;
66-
float v_1 = dot(geometric_normal, cross(position_edge_1, e_p)) * triangle_area_inv;
67-
gradient_1 = inv_mesh_scale * (barycentric(u_1, v_1, tex_coord_0, tex_coord_edge_1, tex_coord_edge_2) - hit_tex_coord);
55+
float3 e_p = hit_point + ellipse_axis - triangle.position_0;
6856

69-
e_p = hit_point + ellipse_axis_2 - position_0;
57+
float u = dot(geometric_normal, cross(e_p, triangle.position_edge_2)) * triangle_double_area_inv;
58+
float v = dot(geometric_normal, cross(triangle.position_edge_1, e_p)) * triangle_double_area_inv;
7059

71-
float u_2 = dot(geometric_normal, cross(e_p, position_edge_2)) * triangle_area_inv;
72-
float v_2 = dot(geometric_normal, cross(position_edge_1, e_p)) * triangle_area_inv;
73-
gradient_2 = inv_mesh_scale * (barycentric(u_2, v_2, tex_coord_0, tex_coord_edge_1, tex_coord_edge_2) - hit_tex_coord);
60+
return barycentric(u, v, triangle.tex_coord_0, triangle.tex_coord_edge_1, triangle.tex_coord_edge_2) - hit_tex_coord;
7461
}
7562

7663
__device__ inline float ray_cone_get_lod(const float3 & ray_direction, const float3 & geometric_normal, float cone_width) {

‎Src/CUDA/Raytracing/Triangle.h‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -103,17 +103,15 @@ __device__ inline TrianglePosNorTex triangle_get_positions_normals_and_tex_coord
103103

104104
// Triangle texture base LOD as described in "Texture Level of Detail Strategies for Real-Time Ray Tracing"
105105
__device__ inline float triangle_get_lod(
106-
float mesh_scale,
107-
float triangle_area_inv,
106+
float triangle_double_area_world_space_inv,
108107
const float2 & tex_coord_edge_1,
109108
const float2 & tex_coord_edge_2
110109
) {
111-
float t_a = fabsf(
110+
float triangle_double_area_texel_space = fabsf(
112111
tex_coord_edge_1.x * tex_coord_edge_2.y -
113112
tex_coord_edge_2.x * tex_coord_edge_1.y
114113
);
115-
116-
return t_a * triangle_area_inv / (mesh_scale * mesh_scale);
114+
return sqrtf(triangle_double_area_texel_space * triangle_double_area_world_space_inv);
117115
}
118116

119117
__device__ inline float triangle_get_curvature(

‎Src/Renderer/Triangle.h‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,8 @@
66
#include "Math/Vector2.h"
77
#include "Math/Vector3.h"
88

9+
#include "Util/Util.h"
10+
911
struct Triangle {
1012
Vector3 position_0;
1113
Vector3 position_1;
@@ -55,6 +57,27 @@ struct Triangle {
5557
if (normal_0_invalid) this->normal_0 = geometric_normal;
5658
if (normal_1_invalid) this->normal_1 = geometric_normal;
5759
if (normal_2_invalid) this->normal_2 = geometric_normal;
60+
} else {
61+
Vector3 geometric_normal = Vector3::normalize(Vector3::cross(
62+
position_1 - position_0,
63+
position_2 - position_0
64+
));
65+
bool all_normals_have_wrong_orientation =
66+
Vector3::dot(geometric_normal, normal_0) < 0.0f &&
67+
Vector3::dot(geometric_normal, normal_1) < 0.0f &&
68+
Vector3::dot(geometric_normal, normal_2) < 0.0f;
69+
bool some_normals_have_wrong_orientation =
70+
Vector3::dot(geometric_normal, normal_0) < 0.0f &&
71+
Vector3::dot(geometric_normal, normal_1) < 0.0f &&
72+
Vector3::dot(geometric_normal, normal_2) < 0.0f;
73+
74+
if (all_normals_have_wrong_orientation) {
75+
Util::swap(this->position_1, this->position_2);
76+
Util::swap(this->normal_1, this->normal_2);
77+
Util::swap(this->tex_coord_1, this->tex_coord_2);
78+
} else if (some_normals_have_wrong_orientation) {
79+
__debugbreak();
80+
}
5881
}
5982
}
6083

0 commit comments

Comments
 (0)