-
Notifications
You must be signed in to change notification settings - Fork 7
Standardize code for HLS #42
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
Changes from 8 commits
ee32c11
3979043
1893a47
1741496
8d3c225
6ceb4ce
60c883a
8a73547
4644777
70c34f9
b92d9c2
635037d
5e96817
092b972
8afb9cb
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -25,24 +25,22 @@ class constant_medium { | |
| , neg_inv_density { -1 / d } | ||
| , phase_function { isotropic_material { a } } {} | ||
|
|
||
| bool hit(auto& ctx, const ray& r, real_t min, real_t max, hit_record& rec, | ||
| bool hit(const ray& r, real_t min, real_t max, hit_record& rec, | ||
| material_t& hit_material_type) const { | ||
| auto& rng = ctx.rng; | ||
| hit_material_type = phase_function; | ||
| material_t temp_material_type; | ||
| hit_record rec1, rec2; | ||
| if (!dev_visit( | ||
| [&](auto&& arg) { | ||
| return arg.hit(ctx, r, -infinity, infinity, rec1, | ||
|
lforg37 marked this conversation as resolved.
|
||
| return arg.hit(r, -infinity, infinity, rec1, | ||
| temp_material_type); | ||
| }, | ||
| boundary)) { | ||
| return false; | ||
| } | ||
|
|
||
| if (!dev_visit( | ||
| [&](auto&& arg) { | ||
| return arg.hit(ctx, r, rec1.t + 0.0001f, infinity, rec2, | ||
| return arg.hit(r, rec1.t + 0.0001f, infinity, rec2, | ||
| temp_material_type); | ||
| }, | ||
| boundary)) { | ||
|
|
@@ -62,7 +60,8 @@ class constant_medium { | |
| /// Distance between the two hitpoints affect of probability | ||
| /// of the ray hitting a smoke particle | ||
| const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; | ||
| const auto hit_distance = neg_inv_density * sycl::log(rng.float_t()); | ||
| auto rng = LocalPseudoRNG { toseed(r.direction()) }; | ||
|
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Ah, I realize you have removed the context passed everywhere and replaced basically the random generator with some local computation.
Contributor
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. As getting a value from the RNG update its internal state, having a shared RNG creates a very long chain of read after write dependency that prevent the HLS compiler to parallelize the otherwise independent computation steps.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Perhaps we could have this context with some special HLS decorations to ignore some dependencies.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Otherwise, perhaps something like this pseudo HLS-SYCL code: auto rng(auto local_seed = std::source_location::current()) {
return [=] {
static auto state = local_seed.line();
#pragma HLS dependence variable=state inter false
state = crunch(state);
return mix(state);
};
}or auto rng(auto local_seed = std::source_location::current()) {
return [state = local_seed.line()] mutable {
#pragma HLS dependence variable=state inter false
state = crunch(state);
return mix(state);
};
}Perhaps we need extension intel/llvm#3746 to have the big picture working.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Above is basically a good idea, if I understand the intention, to separate the state into multiple independent RNGs. I'm not sure what you're trying to achieve with the false dependence pragma, however.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. @stephenneuendorffer yes this is the idea to distribute the state here in a lazy way to even initialize the seed with something depending from the source location (here with the line number, we could also use the file name, etc.). #pragma HLS dependence variable=state inter falseas a way to have an |
||
| const auto hit_distance = neg_inv_density * sycl::log(rng.real()); | ||
|
|
||
| /// With lower density, hit_distance has higher probabilty | ||
| /// of being greater than distance_inside_boundary | ||
|
|
||
| Original file line number | Diff line number | Diff line change | ||||||||||||||||||||||||||||||||||
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| @@ -0,0 +1,23 @@ | ||||||||||||||||||||||||||||||||||||
| #ifndef HIT_RECORD_HPP | ||||||||||||||||||||||||||||||||||||
| #define HIT_RECORD_HPP | ||||||||||||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||||||||||||
| #include "rtweekend.hpp" | ||||||||||||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||||||||||||
| class hit_record { | ||||||||||||||||||||||||||||||||||||
| public: | ||||||||||||||||||||||||||||||||||||
| real_t t; // | ||||||||||||||||||||||||||||||||||||
| point p; // hit point | ||||||||||||||||||||||||||||||||||||
| vec normal; // normal at hit point | ||||||||||||||||||||||||||||||||||||
| bool front_face; // to check if hit point is on the outer surface | ||||||||||||||||||||||||||||||||||||
| /*local coordinates for rectangles | ||||||||||||||||||||||||||||||||||||
| and mercator coordintes for spheres */ | ||||||||||||||||||||||||||||||||||||
| real_t u; | ||||||||||||||||||||||||||||||||||||
| real_t v; | ||||||||||||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||||||||||||
| // To set if the hit point is on the front face | ||||||||||||||||||||||||||||||||||||
|
Comment on lines
+10
to
+17
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. While you are here, add more Doxygen
Suggested change
|
||||||||||||||||||||||||||||||||||||
| void set_face_normal(const ray& r, const vec& outward_normal) { | ||||||||||||||||||||||||||||||||||||
| front_face = dot(r.direction(), outward_normal) < 0; | ||||||||||||||||||||||||||||||||||||
| normal = front_face ? outward_normal : vec {} - outward_normal; | ||||||||||||||||||||||||||||||||||||
| } | ||||||||||||||||||||||||||||||||||||
| }; | ||||||||||||||||||||||||||||||||||||
| #endif | ||||||||||||||||||||||||||||||||||||
|
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Missing end-of-line.
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Curious... An IDE configuration bug? |
||||||||||||||||||||||||||||||||||||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I have the feeling that if we remove the
std::monostatefrom the material it will always create aspherehere even when there is no intersection.