From 3ab9c8f1099eec4c72f9b5d16f38a7d18c5e50f9 Mon Sep 17 00:00:00 2001 From: Enigmatisms <984041003@qq.com> Date: Mon, 27 May 2024 01:43:31 +0800 Subject: [PATCH] WIP: area source bug, emitter eval_le not used. --- scene/meshes/cupt/light.obj | 10 +++--- scene/xml/cbox-point.xml | 4 +-- src/core/emitter.cuh | 37 ++++++++++--------- src/core/scene.cuh | 69 +++++++++++++++++++++++------------- src/core/virtual_funcs.cuh | 6 ++-- src/renderer/path_tracer.cuh | 8 +++-- 6 files changed, 80 insertions(+), 54 deletions(-) diff --git a/scene/meshes/cupt/light.obj b/scene/meshes/cupt/light.obj index 5aaae1d..0bbb0c8 100644 --- a/scene/meshes/cupt/light.obj +++ b/scene/meshes/cupt/light.obj @@ -1,11 +1,11 @@ # Blender 4.0.2 # www.blender.org o Plane -v -0.250000 0.800000 0.250000 -v 0.250000 0.800000 0.250000 -v -0.250000 0.800000 -0.250000 -v 0.250000 0.800000 -0.250000 -vn -0.0000 -1.0000 -0.0000 +v -0.250000 0.250000 0.9900000 +v 0.250000 0.250000 0.9900000 +v -0.250000 -0.250000 0.9900000 +v 0.250000 -0.250000 0.9900000 +vn -0.0000 -0.0000 -1.0000 vt 0.000000 1.000000 vt 1.000000 0.000000 vt 0.000000 0.000000 diff --git a/scene/xml/cbox-point.xml b/scene/xml/cbox-point.xml index e2af73b..c4d92c5 100644 --- a/scene/xml/cbox-point.xml +++ b/scene/xml/cbox-point.xml @@ -57,13 +57,13 @@ - + diff --git a/src/core/emitter.cuh b/src/core/emitter.cuh index 79bf86c..1401e57 100644 --- a/src/core/emitter.cuh +++ b/src/core/emitter.cuh @@ -13,16 +13,11 @@ CPT_CPU_GPU_INLINE float distance_attenuate(Vec3&& diff) { return min(1.f / max(diff.length2(), 1e-5f), 1.f); } -enum EmitterBinding: uint8_t { - POINT = 0x00, - TRIANGLE = 0x01, - SPHERE = 0x02 -}; - class Emitter { protected: Vec3 Le; - EmitterBinding obj_ref_id; // index pointing to object, -1 means the emitter is delta_pos + int obj_ref_id; + bool is_sphere; // whether the emitter binds to a sphere public: /** * Sample a point on the emitter (useful for non-infinitesimal emitters) @@ -30,8 +25,8 @@ public: CPT_CPU_GPU Emitter() {} CONDITION_TEMPLATE(VecType, Vec3) - CPT_CPU_GPU Emitter(VecType&& le, EmitterBinding obj_ref_id = POINT): - Le(std::forward(le)), obj_ref_id(obj_ref_id) {} + CPT_CPU_GPU Emitter(VecType&& le, int obj_ref = -1, bool is_sphere = false): + Le(std::forward(le)), obj_ref_id(obj_ref), is_sphere(is_sphere) {} CPT_GPU_INLINE virtual Vec3 sample(const Vec3& hit_pos, Vec3& le, float& pdf, Vec2&&, ConstPrimPtr, ConstPrimPtr, int) const { pdf = 1; @@ -44,12 +39,16 @@ public: } CPT_GPU_INLINE bool non_delta() const noexcept { - return this->obj_ref_id; + return this->is_sphere; } CPT_GPU_INLINE Vec3 get_le() const noexcept { return Le; } + + CPT_CPU_GPU int get_obj_ref() const noexcept { + return obj_ref_id * (obj_ref_id >= 0); + } }; class PointSource: public Emitter { @@ -80,8 +79,8 @@ public: CPT_CPU_GPU AreaSource() {} CONDITION_TEMPLATE(VType, Vec3) - CPT_CPU_GPU AreaSource(VType&& le, EmitterBinding obj_ref_id = TRIANGLE): - Emitter(std::forward(le), obj_ref_id) {} + CPT_CPU_GPU AreaSource(VType&& le, int obj_ref, bool is_sphere = false): + Emitter(std::forward(le), obj_ref, is_sphere) {} CPT_GPU_INLINE Vec3 sample( const Vec3& hit_pos, Vec3& le, float& pdf, @@ -90,27 +89,31 @@ public: float sample_sum = uv.x() + uv.y(); uv = select(uv, -uv + 1.f, sample_sum < 1.f); float diff_x = 1.f - uv.x(), diff_y = 1.f - uv.y(); - Vec3 sampled = uv.x() * prims->y[sampled_index] + uv.y() * prims->z[sampled_index] - (uv.x() + uv.y()) * prims->x[sampled_index]; + Vec3 sampled = uv.x() * prims->y[sampled_index] + uv.y() * prims->z[sampled_index] - (uv.x() + uv.y() - 1) * prims->x[sampled_index]; Vec3 normal = (norms->x[sampled_index] * diff_x * diff_y + \ norms->y[sampled_index] * uv.x() * diff_y + \ norms->z[sampled_index] * uv.y() * diff_x).normalized(); - Vec3 sphere_normal = sample_uniform_sphere(std::move(uv), sample_sum); + // if (sampled_index == 10 || sampled_index == 11) + // printf("Normal (%d): %f, %f, %f\n", sampled_index, normal.x(), normal.y(), normal.z()); + Vec3 sphere_normal = sample_uniform_sphere(select(uv, -uv + 1.f, sample_sum < 1.f), sample_sum); sampled = select( sampled, sphere_normal * prims->y[sampled_index].x() + prims->x[sampled_index], - obj_ref_id == TRIANGLE + is_sphere == false ); - normal = select(normal, sphere_normal, obj_ref_id == TRIANGLE); + normal = select(normal, sphere_normal, is_sphere == false); // normal needs special calculation sphere_normal = hit_pos - sampled; pdf *= sphere_normal.length2(); sphere_normal.normalize(); sample_sum = normal.dot(sphere_normal); // dot_light + // printf("Sampled sum: %f\n", sample_sum); pdf *= float(sample_sum > 0) / sample_sum; le = Le * float(sample_sum > 0); return sampled; } CPT_GPU virtual Vec3 eval_le(const Vec3* const inci_dir, const Vec3* const normal) const override { - return select(Le, Vec3(0, 0, 0), inci_dir->dot(*normal) < 0); + printf("valid: %d\n", int(inci_dir->dot(*normal) < 0)); + return select(this->Le, Vec3(0, 0, 0), inci_dir->dot(*normal) < 0); } }; \ No newline at end of file diff --git a/src/core/scene.cuh b/src/core/scene.cuh index 314f047..e132e4b 100644 --- a/src/core/scene.cuh +++ b/src/core/scene.cuh @@ -140,16 +140,27 @@ void parseBSDF(const tinyxml2::XMLElement* bsdf_elem, std::unordered_map& emitter_map +) { + int idx = 0; + while (emitter_elem) { + std::string id = emitter_elem->Attribute("id"); + emitter_map[id] = idx++; + emitter_elem = emitter_elem->NextSiblingElement("emitter"); + } +} + void parseEmitter( const tinyxml2::XMLElement* emitter_elem, - std::unordered_map& emitter_map, + std::unordered_map& emitter_obj_map, // key emitter name, value object_id std::vector obj_ref_names, Emitter** emitters, int index ) { std::string type = emitter_elem->Attribute("type"); std::string id = emitter_elem->Attribute("id"); - emitter_map[id] = index; obj_ref_names.push_back(id); Vec3 emission(0, 0, 0), scaler(0, 0, 0); @@ -193,11 +204,8 @@ void parseEmitter( std::cerr << "Bound primitive is not specified for area source '" << id << "', name: "<< element->Attribute("name") << std::endl; throw std::runtime_error("Bound primitive is not specified for area source"); } - EmitterBinding bound_type = TRIANGLE; - if (element->Attribute("value") == "sphere") { - bound_type = SPHERE; - } - create_area_source<<<1, 1>>>(emitters[index], emission * scaler, bound_type); + bool spherical_bound = element->Attribute("value") == "sphere"; + create_area_source<<<1, 1>>>(emitters[index], emission * scaler, emitter_obj_map[id], spherical_bound); } } @@ -216,9 +224,10 @@ void parseSphereShape( const tinyxml2::XMLElement* shapeElement, const std::unordered_map& bsdf_map, const std::unordered_map& emitter_map, + std::unordered_map& emitter_obj_map, std::vector& objects, std::array& verts_list, std::array& norms_list, std::array& uvs_list, - int& prim_offset, std::string folder_prefix + int& prim_offset, std::string folder_prefix, int index ) { int bsdf_id = -1, emitter_id = 0; @@ -231,6 +240,7 @@ void parseSphereShape( bsdf_id = get_map_id(bsdf_map, id); } else if (type == "emitter") { emitter_id = get_map_id(emitter_map, id); + emitter_obj_map[id] = index; } element = element->NextSiblingElement("ref"); } @@ -265,9 +275,10 @@ void parseObjShape( const tinyxml2::XMLElement* shapeElement, const std::unordered_map& bsdf_map, const std::unordered_map& emitter_map, + std::unordered_map& emitter_obj_map, std::vector& objects, std::array& verts_list, std::array& norms_list, std::array& uvs_list, - int& prim_offset, std::string folder_prefix + int& prim_offset, std::string folder_prefix, int index ) { std::string filename, name; int bsdf_id = -1, emitter_id = 0; @@ -289,6 +300,7 @@ void parseObjShape( bsdf_id = get_map_id(bsdf_map, id); } else if (type == "emitter") { emitter_id = get_map_id(emitter_map, id); + emitter_obj_map[id] = index; } element = element->NextSiblingElement("ref"); } @@ -388,7 +400,7 @@ public: *emitter_elem = scene_elem->FirstChildElement("emitter"), *sensor_elem = scene_elem->FirstChildElement("sensor"), *ptr = nullptr; - std::unordered_map bsdf_map, emitter_map; + std::unordered_map bsdf_map, emitter_map, emitter_obj_map; std::vector emitter_names; emitter_names.reserve(9); emitter_map.reserve(9); @@ -406,17 +418,9 @@ public: } CUDA_CHECK_RETURN(cudaDeviceSynchronize()); - // ------------------------- (2) parse all emitters -------------------------- - ptr = emitter_elem; - for (; ptr != nullptr; ++ num_emitters) - ptr = ptr->NextSiblingElement("emitter"); - CUDA_CHECK_RETURN(cudaMalloc(&emitters, sizeof(Emitter*) * (num_emitters + 1))); - create_abstract_source<<<1, 1>>>(emitters[0]); - for (int i = 1; i <= num_emitters; i++) { - parseEmitter(emitter_elem, emitter_map, emitter_names, emitters, i); - emitter_elem = emitter_elem->NextSiblingElement("emitter"); - } - CUDA_CHECK_RETURN(cudaDeviceSynchronize()); + + // ------------------------- (2) parse emitter names ------------------------- + parseEmitterNames(emitter_elem, emitter_map); // ------------------------- (3) parse all objects ------------------------- ptr = shape_elem; @@ -436,19 +440,34 @@ public: for (int i = 0; i < num_objects; i++) { std::string type = shape_elem->Attribute("type"); if (type == "obj") - parseObjShape(shape_elem, bsdf_map, emitter_map, objects, verts_list, norms_list, uvs_list, prim_offset, folder_prefix); + parseObjShape(shape_elem, bsdf_map, emitter_map, emitter_obj_map, objects, + verts_list, norms_list, uvs_list, prim_offset, folder_prefix, i); else if (type == "sphere") - parseSphereShape(shape_elem, bsdf_map, emitter_map, objects, verts_list, norms_list, uvs_list, prim_offset, folder_prefix); + parseSphereShape(shape_elem, bsdf_map, emitter_map, emitter_obj_map, objects, + verts_list, norms_list, uvs_list, prim_offset, folder_prefix, i); sphere_objs[i] = type == "sphere"; shape_elem = shape_elem->NextSiblingElement("shape"); } num_prims = prim_offset; - // ------------------------- (4) parse camera & scene config ------------------------- + + // ------------------------- (4) parse all emitters -------------------------- + ptr = emitter_elem; + for (; ptr != nullptr; ++ num_emitters) + ptr = ptr->NextSiblingElement("emitter"); + CUDA_CHECK_RETURN(cudaMalloc(&emitters, sizeof(Emitter*) * (num_emitters + 1))); + create_abstract_source<<<1, 1>>>(emitters[0]); + for (int i = 1; i <= num_emitters; i++) { + parseEmitter(emitter_elem, emitter_obj_map, emitter_names, emitters, i); + emitter_elem = emitter_elem->NextSiblingElement("emitter"); + } + CUDA_CHECK_RETURN(cudaDeviceSynchronize()); + + // ------------------------- (5) parse camera & scene config ------------------------- cam = DeviceCamera::from_xml(sensor_elem); config = RenderingConfig::from_xml(sensor_elem); - // ------------------------- (5) initialize shapes ------------------------- + // ------------------------- (6) initialize shapes ------------------------- shapes.resize(num_prims); prim_offset = 0; for (int obj_id = 0; obj_id < num_objects; obj_id ++) { diff --git a/src/core/virtual_funcs.cuh b/src/core/virtual_funcs.cuh index 27fc47a..1c84b58 100644 --- a/src/core/virtual_funcs.cuh +++ b/src/core/virtual_funcs.cuh @@ -26,15 +26,15 @@ __global__ void create_point_source(Emitter* &dst, Vec3 le, Vec3 pos) { } } -__global__ void create_area_source(Emitter* &dst, Vec3 le, EmitterBinding obj_ref) { +__global__ void create_area_source(Emitter* &dst, Vec3 le, int obj_ref, bool is_sphere) { if (threadIdx.x == 0 && blockIdx.x == 0) { - dst = new AreaSource(std::move(le), obj_ref); + dst = new AreaSource(std::move(le), obj_ref, is_sphere); } } __global__ void create_abstract_source(Emitter* &dst) { if (threadIdx.x == 0 && blockIdx.x == 0) { - dst = new Emitter(Vec3(0, 0, 0), POINT); + dst = new Emitter(Vec3(0, 0, 0)); } } diff --git a/src/renderer/path_tracer.cuh b/src/renderer/path_tracer.cuh index 6d14a17..f7a7065 100644 --- a/src/renderer/path_tracer.cuh +++ b/src/renderer/path_tracer.cuh @@ -62,7 +62,7 @@ CPT_GPU Emitter* sample_emitter(Sampler& sampler, float& pdf, int num, int no_sa pdf = 1.f / float(num); // when no_sample == 0 (means, we do not intersect any emitter) or num > 1 (there are more than 1 emitters) // the sample will be valid - return c_emitter[emit_id * (no_sample == 0 || num > 1)]; + return c_emitter[emit_id * uint32_t(no_sample == 0 || num > 1)]; } /** @@ -169,7 +169,11 @@ __global__ static void render_pt_kernel( Emitter* emitter = sample_emitter(sampler, direct_pdf, num_emitter, emitter_id); // (3) sample a point on the emitter (we avoid sampling the hit emitter) - emitter_id = objects[object_id].sample_emitter_primitive(sampler.discrete1D(), direct_pdf); + // printf("Num of emitters: %d, %x, %x, %x, %d\n", num_emitter, int(c_emitter[0]), int(c_emitter[1]), int(emitter), int(hit_emitter)); + + emitter_id = objects[emitter->get_obj_ref()].sample_emitter_primitive(sampler.discrete1D(), direct_pdf); + // if (objects[object_id].emitter_id) + // printf("Object id: %d, sampled id: %d, emitter: %d\n", object_id, emitter_id, objects[object_id].emitter_id); Ray shadow_ray(ray.o + ray.d * min_dist, Vec3(0, 0, 0)); // use ray.o to avoid creating another shadow_int variable shadow_ray.d = emitter->sample(shadow_ray.o, ray.o, direct_pdf, sampler.next2D(), verts, norms, emitter_id) - shadow_ray.o;