diff --git a/demos/optical_bench/src/Optix Files/glass_iterative_camera.cu b/demos/optical_bench/src/Optix Files/glass_iterative_camera.cu
index 44bbcaa9cfa725a94d8aea1de6e1da46349a84ef..7522a7d81a3360f13b06c3321a27ec50446ec492 100644
--- a/demos/optical_bench/src/Optix Files/glass_iterative_camera.cu
+++ b/demos/optical_bench/src/Optix Files/glass_iterative_camera.cu
@@ -101,7 +101,6 @@ RT_PROGRAM void closest_hit_radiance()
const float3 fhp = rtTransformPoint(RT_OBJECT_TO_WORLD, front_hit_point);
prd_radiance.origin = fhp;
prd_radiance.direction = w_in;
- prd_radiance.reflectance *= reflection_color*transmittance;
}else{
prd_radiance.done = true;
}
@@ -110,8 +109,7 @@ RT_PROGRAM void closest_hit_radiance()
const float3 w_in = w_t;
const float3 bhp = rtTransformPoint(RT_OBJECT_TO_WORLD, back_hit_point);
prd_radiance.origin = bhp;
- prd_radiance.direction = w_in;
- prd_radiance.reflectance *= refraction_color*transmittance;
+ prd_radiance.direction = w_in;
}
// Note: we do not trace the ray for the next bounce here, we just set it up for
diff --git a/demos/optical_bench/src/Optix Files/laser_caster.cu b/demos/optical_bench/src/Optix Files/laser_caster.cu
index 11316aa2f4521202ffbaa4853be1f80a4b42f81c..4eb8721443614d178d23ee8229e62a6a923d0ae7 100644
--- a/demos/optical_bench/src/Optix Files/laser_caster.cu
+++ b/demos/optical_bench/src/Optix Files/laser_caster.cu
@@ -6,8 +6,8 @@
using namespace optix;
-rtDeclareVariable(uint2, launch_index, rtLaunchIndex, );
-rtDeclareVariable(uint2, launch_dim, rtLaunchDim, );
+rtDeclareVariable(uint3, launch_index, rtLaunchIndex, );
+rtDeclareVariable(uint3, launch_dim, rtLaunchDim, );
rtDeclareVariable(rtObject, top_object, , );
rtDeclareVariable(float, scene_epsilon, , );
rtDeclareVariable(int, max_depth_laser, , );
@@ -23,51 +23,50 @@ rtDeclareVariable(float3, laser_up, , );
rtBuffer<int, 2> laserIndex;
rtBuffer<float3, 1> result_laser;
+
RT_PROGRAM void laser_caster()
{
const float laserSize = 0.04;
- const float laserBeamWidth = 0.005;
-
- unsigned int seed = tea<16>(launch_dim.x*launch_index.y+launch_index.x, random_frame_seed);
+ float2 d = make_float2(launch_index.x, launch_index.y) / make_float2(launch_dim.x, launch_dim.y) - make_float2(0.5f, 0.5f);
+ float3 ray_origin = laser_origin + laser_right * laserSize * d.x + laser_up * laserSize * d.y;
+
+ const float laserBeamWidth = 0.025;
+ unsigned int seed = tea<16>(launch_dim.x*launch_index.y+launch_index.x*launch_index.z, random_frame_seed);
float2 random = make_float2(rnd(seed), rnd(seed));
- float r = sqrt(-2*log(random.x));
+ float r = sqrtf(-2*log(random.x));
float theta = 2*3.141592654f*random.y;
- random = clamp(make_float2(r*cos(theta), r*sin(theta)), -4.5f, 4.5f) * laserBeamWidth/4.5f;
+ random = clamp(make_float2(r*cosf(theta), r*sinf(theta)), -4.5f, 4.5f) * laserBeamWidth/4.5f;
+ if(launch_index.z != 0) ray_origin += laser_right * random.x + laser_up * random.y;
- float2 d = make_float2(launch_index) / make_float2(launch_dim) - make_float2(0.5f, 0.5f);
- float3 ray_origin = laser_origin + laser_right * laserSize * d.x + laser_up * laserSize * d.y + laser_right * random.x + laser_up * random.y;
- float3 ray_direction = laser_forward;
-
PerRayData_radiance_iterative prd;
+ optix::Ray ray(ray_origin, laser_forward, /*ray type*/ 1, scene_epsilon );
prd.depth = 0;
- prd.seed = launch_index.y + launch_index.x;
prd.done = false;
prd.hit_lens = false; //track if the ray ever hit the lens
- prd.radiance = make_float3(1) * (1 - length(random)/laserBeamWidth);
+ prd.power = (launch_index.z > 0) * 1; //No power for launch index 0
// next ray to be traced
prd.origin = ray_origin;
- prd.direction = ray_direction;
+ prd.direction = laser_forward;
unsigned int startIndex = (launch_index.y * 50 + launch_index.x) * max_depth_laser * 2;
+ bool cast_ray = laserIndex[make_uint2(launch_index)] < 0;
for(int i = 0; i < max_depth_laser * 2; i += 2){
//Determine if this launch_index, depth or last ray should trigger new cast
- if(laserIndex[launch_index] < 0 || prd.done || prd.depth >= max_depth_laser){ // just write rest of data as "invalid"
- result_laser[startIndex + i] = make_float3(0,-1,0);
- result_laser[startIndex + i + 1] = make_float3(0,-1,0);
+ if(cast_ray || prd.done || prd.depth >= max_depth_laser){ // just write rest of data as "invalid"
+ if(launch_index.z == 0) result_laser[startIndex + i] = make_float3(0,-1,0);
+ if(launch_index.z == 0) result_laser[startIndex + i + 1] = make_float3(0,-1,0);
continue;
}
- result_laser[startIndex + i] = ray_origin;
+ if(launch_index.z == 0) result_laser[startIndex + i] = prd.origin;
- optix::Ray ray(ray_origin, ray_direction, /*ray type*/ 1, scene_epsilon );
+ ray.origin = prd.origin;
+ ray.direction = prd.direction;
rtTrace(top_object, ray, prd);
-
// Update ray data for the next path segment
prd.depth++;
- ray_origin = prd.origin;
- ray_direction = prd.direction;
- result_laser[startIndex + i + 1] = ray_origin;
- }
+ if(launch_index.z == 0) result_laser[startIndex + i + 1] = prd.origin;
+ }
}
\ No newline at end of file
diff --git a/demos/optical_bench/src/Optix Files/laser_target_material.cu b/demos/optical_bench/src/Optix Files/laser_target_material.cu
index 38ba778b1da13a2bb1066334dcb8c239cbd06024..669e5d1f872147d2bdb5b5ce5185c219c7394001 100644
--- a/demos/optical_bench/src/Optix Files/laser_target_material.cu
+++ b/demos/optical_bench/src/Optix Files/laser_target_material.cu
@@ -27,12 +27,27 @@ rtDeclareVariable(PerRayData_radiance_iterative, prd_radiance_it, rtPayload, );
rtDeclareVariable(float2, texture_coord, attribute texture_coord, );
rtDeclareVariable(int, writeable_surface, attribute writeable_surface, );
rtDeclareVariable(float, hit_depth, rtIntersectionDistance, );
+
+rtDeclareVariable(float, max_power, , );
rtDeclareVariable(optix::Ray, ray, rtCurrentRay, );
-rtBuffer<float3, 2> targetBuffer;
+rtBuffer<float, 1> targetBufferMax;
+rtBuffer<float, 2> targetBuffer;
rtDeclareVariable(float2, targetBufferDim, , );
rtDeclareVariable(int, targetBufferWrite, , );
+//thanks stackoverflow!
+__device__ static float atomicMax(float* address, float val){
+ int* address_as_i = (int*) address;
+ int old = *address_as_i, assumed;
+ do {
+ assumed = old;
+ old = ::atomicCAS(address_as_i, assumed,
+ __float_as_int(::fmaxf(val, __int_as_float(assumed))));
+ } while (assumed != old);
+ return __int_as_float(old);
+}
+
RT_PROGRAM void any_hit_radiance(){
rtIgnoreIntersection();
}
@@ -42,5 +57,9 @@ RT_PROGRAM void closest_hit_iterative()
float3 hit_point = ray.origin + hit_depth * ray.direction;
prd_radiance_it.origin = hit_point;
prd_radiance_it.done = true;
- if(targetBufferWrite && writeable_surface) targetBuffer[make_uint2(texture_coord * targetBufferDim)] = prd_radiance_it.radiance;
+
+ if(targetBufferWrite && writeable_surface){
+ atomicAdd(&targetBuffer[make_uint2(texture_coord * targetBufferDim)], prd_radiance_it.power);
+ atomicMax(&targetBufferMax[0], targetBuffer[make_uint2(texture_coord * targetBufferDim)]);
+ }
}
diff --git a/demos/optical_bench/src/Optix Files/normal_shader_new.cu b/demos/optical_bench/src/Optix Files/normal_shader_new.cu
index 68da742d718f875d88cc381e9911bf4b9b6cba6e..6c5906e8af21a826bbc2e51c87e12ccc6b983eb0 100644
--- a/demos/optical_bench/src/Optix Files/normal_shader_new.cu
+++ b/demos/optical_bench/src/Optix Files/normal_shader_new.cu
@@ -38,6 +38,6 @@ rtDeclareVariable(PerRayData_radiance_iterative, prd_radiance, rtPayload, );
RT_PROGRAM void closest_hit_radiance()
{
- prd_radiance.radiance = normalize(rtTransformNormal(RT_OBJECT_TO_WORLD, shading_normal))*0.5f + 0.5f;
+ prd_radiance.power = normalize(rtTransformNormal(RT_OBJECT_TO_WORLD, shading_normal)).x * 0.5f + 0.5f;
prd_radiance.done = true;
}
diff --git a/demos/optical_bench/src/Optix Files/prd.h b/demos/optical_bench/src/Optix Files/prd.h
index 8c818de9237e5c3205805d98be693b003c88e8b9..9a3b97f68f118c6a0b26797d8f0ed2ecfa35a646 100644
--- a/demos/optical_bench/src/Optix Files/prd.h
+++ b/demos/optical_bench/src/Optix Files/prd.h
@@ -30,33 +30,26 @@
#include <optixu/optixu_vector_types.h>
-struct PerRayData_radiance_iterative
-{
- int depth;
- unsigned int seed;
-
- // shading state
- bool done;
- float3 reflectance;
- float3 radiance;
- float3 origin;
- float3 direction;
-
- bool hit_lens;
+struct PerRayData_radiance_iterative{
+ float3 origin;
+ float3 direction;
+ float power;
+ int depth;
+
+ bool done;
+ bool hit_lens;
};
-struct PerRayData_radiance
-{
+struct PerRayData_radiance{
float3 result;
float hit_depth;
float importance;
int depth;
bool miss;
- bool hit_lens;
+ bool hit_lens;
};
-struct PerRayData_shadow
-{
+struct PerRayData_shadow{
float3 attenuation;
};
diff --git a/demos/optical_bench/src/laser_menu.cpp b/demos/optical_bench/src/laser_menu.cpp
index e050af0a1d47abb531ae073ba0042f1a0e07b87f..b40b7a6d2d02a659d923971616474b75c7c53441 100644
--- a/demos/optical_bench/src/laser_menu.cpp
+++ b/demos/optical_bench/src/laser_menu.cpp
@@ -106,7 +106,7 @@ LaserMenu::LaserMenu(phx::Scene* scene, phx::Engine* engine) {
clear_target_->GetTransform()->SetLocalTranslation(glm::vec3(0));
clear_target_->SetHoldable(false);
clear_target_->SetOnPress(
- [this, rayPass](phx::Transform*, glm::mat4) { rayPass->clearTarget(); });
+ [this, rayPass](phx::Transform*, glm::mat4) { rayPass->clearTargetBuffers(); });
}
LaserMenu::~LaserMenu() {
diff --git a/demos/optical_bench/src/optix_context_manager.cpp b/demos/optical_bench/src/optix_context_manager.cpp
index 2195007fd7a09b660a2e05cdc4d517a46e745545..6e350088caf41d499de4e045c3869824dcb7f747 100644
--- a/demos/optical_bench/src/optix_context_manager.cpp
+++ b/demos/optical_bench/src/optix_context_manager.cpp
@@ -39,7 +39,7 @@ OptixContextManager::OptixContextManager() {
// context_->setPrintEnabled(true);
// context_->setPrintBufferSize(1024);
- context_->setStackSize(5000);
+ context_->setStackSize(5000l);
context_["max_depth"]->setInt(10);
context_["scene_epsilon"]->setFloat(1.e-4f);
diff --git a/demos/optical_bench/src/ray_pass.cpp b/demos/optical_bench/src/ray_pass.cpp
index 55e98214041a1ef851f86ddfb3d22ca8e5d0c1f9..cc1e7a6fc68b3e3daa72be1496674205c9789669 100644
--- a/demos/optical_bench/src/ray_pass.cpp
+++ b/demos/optical_bench/src/ray_pass.cpp
@@ -26,6 +26,7 @@
#include <stdlib.h>
#include <cassert>
+#include <cstdlib>
#include <iostream>
#include <memory>
#include <string>
@@ -70,19 +71,21 @@ RayPass::~RayPass() {
// delete openGL resources
glDeleteBuffers(1, &bufferOpenGLLaserVertices);
glDeleteBuffers(1, &bufferOpenGLLaserColors);
- glDeleteBuffers(1, &targetBufferOpenGL);
glDeleteVertexArrays(1, &vaoOpenGL);
bufferLaserOptix->destroy();
bufferLaserIndexOptix->destroy();
targetBuffer->destroy();
+ targetMaxBuffer->destroy();
targetGroup->destroy();
targetGeometry->destroy();
targetTransform->destroy();
targetMaterial->destroy();
targetAcc->destroy();
+ free(textureDataBuffer);
+
delete rod_;
}
@@ -210,16 +213,14 @@ void RayPass::launchLaser() {
optixContextManager_->getContext()["random_frame_seed"]->setUint(
random_seed++); // totally "random" (only used as seed)
try {
- optixContextManager_->getContext()->launch(laser_entry, 50, 50);
+ optixContextManager_->getContext()->launch(laser_entry, 50, 50,
+ laser_traces_per_frame);
} catch (optix::Exception e) {
phx::error(e.getErrorString().c_str());
}
- // copy buffer from optix
- glBindBuffer(GL_PIXEL_UNPACK_BUFFER, targetBufferOpenGL);
- target_texture->set_sub_image(0, 0, 0, target_res, target_res, GL_RGB,
- GL_FLOAT, nullptr);
- glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+ // copy retrieved intersection data to texture
+ copyAndScaleTargetData();
}
void RayPass::createLaser() {
@@ -280,16 +281,15 @@ void RayPass::createTarget() {
}
target_texture = target_screen_->GetMaterial()->GetAmbientTexture();
- glGenBuffers(1, &targetBufferOpenGL);
- glBindBuffer(GL_ARRAY_BUFFER, targetBufferOpenGL);
- glBufferData(GL_ARRAY_BUFFER, target_res * target_res * sizeof(float) * 3l,
- nullptr, GL_STREAM_DRAW);
+ // used to expand the data in RAM and map back to GPU
+ textureDataBuffer =
+ (float*)std::malloc(target_res * target_res * sizeof(float) * 3l);
- targetBuffer =
- context->createBufferFromGLBO(RT_BUFFER_INPUT_OUTPUT, targetBufferOpenGL);
- targetBuffer->setFormat(RT_FORMAT_FLOAT3);
- targetBuffer->setSize(target_res, target_res);
- clearTarget();
+ targetBuffer = context->createBuffer(RT_BUFFER_INPUT_OUTPUT, RT_FORMAT_FLOAT,
+ target_res, target_res);
+ targetMaxBuffer =
+ context->createBuffer(RT_BUFFER_INPUT_OUTPUT, RT_FORMAT_FLOAT, 1);
+ clearTargetBuffers();
// Material
std::string ptx =
@@ -329,6 +329,7 @@ void RayPass::createTarget() {
targetGeometry["targetBufferWrite"]->setInt(1);
targetGeometry["targetBuffer"]->setBuffer(targetBuffer);
+ targetGeometry["targetBufferMax"]->setBuffer(targetMaxBuffer);
targetGeometry["targetBufferDim"]->setFloat((float)target_res,
(float)target_res);
@@ -357,16 +358,42 @@ void RayPass::createTarget() {
targetSupportRod = new SupportRod(engine_->GetScene().get(), selector);
}
-void RayPass::clearTarget() {
+void RayPass::copyAndScaleTargetData() {
+ // fetch max
+ float max =
+ *(static_cast<float*>(targetMaxBuffer->map(0, RT_BUFFER_MAP_READ)));
+ targetMaxBuffer->unmap();
+ if (max == 0) max = 1;
+
+ // expand & scale buffer
+ float curr_val = 0;
+ float* buffer_data =
+ static_cast<float*>(targetBuffer->map(0, RT_BUFFER_MAP_READ));
+ for (unsigned long i = 0; i < target_res * target_res; i++) {
+ curr_val = buffer_data[i] / max;
+ textureDataBuffer[i * 3] = curr_val;
+ textureDataBuffer[i * 3 + 1] = curr_val;
+ textureDataBuffer[i * 3 + 2] = curr_val;
+ }
+ targetBuffer->unmap();
+
+ // copy to texture
+ target_texture->set_sub_image(0, 0, 0, target_res, target_res, GL_RGB,
+ GL_FLOAT, textureDataBuffer);
+}
+
+void RayPass::clearTargetBuffers() {
// null initially
float* buffer_data =
static_cast<float*>(targetBuffer->map(0, RT_BUFFER_MAP_WRITE_DISCARD));
- for (unsigned long i = 0; i < target_res * target_res * 3; i = i + 3) {
+ for (unsigned long i = 0; i < target_res * target_res; i++) {
buffer_data[i] = 0.0f;
- buffer_data[i + 1] = 0.0f;
- buffer_data[i + 2] = 0.0f;
}
targetBuffer->unmap();
+
+ *(static_cast<float*>(targetMaxBuffer->map(0, RT_BUFFER_MAP_WRITE_DISCARD))) =
+ 0.0f;
+ targetMaxBuffer->unmap();
}
void RayPass::moveLaser(glm::mat4 transform) {
diff --git a/demos/optical_bench/src/ray_pass.hpp b/demos/optical_bench/src/ray_pass.hpp
index ff15e6ddc7c0722b9e754e8b2389a66852607250..d3069b24a7fa587e5bc080a6b4e3bcb9512c2b4f 100644
--- a/demos/optical_bench/src/ray_pass.hpp
+++ b/demos/optical_bench/src/ray_pass.hpp
@@ -81,7 +81,7 @@ class RayPass : public RenderPass {
std::string GetCurrentPatternName();
- void clearTarget();
+ void clearTargetBuffers();
private:
const unsigned int laser_max_depth = 20;
@@ -113,23 +113,31 @@ class RayPass : public RenderPass {
std::vector<phx::ResourcePointer<phx::Image>> patterns_;
unsigned int current_pattern_index_;
- // The resource image needs to be adjusted to set the right texture size
+ // target buffer
+ // The resource image needs to be adjusted to set the right texture resolution
const unsigned long target_res = 512l;
- gl::texture_2d* target_texture;
- GLuint targetBufferOpenGL = 0;
- phx::MaterialHandle* target_screen_;
- void createLaser();
- void createTarget();
optix::Buffer targetBuffer;
+ optix::Buffer targetMaxBuffer;
+
+ // target object
+ void createTarget();
optix::GeometryGroup targetGroup;
optix::GeometryInstance targetGeometry;
optix::Transform targetTransform;
optix::Material targetMaterial;
optix::Acceleration targetAcc;
+ gl::texture_2d* target_texture;
+ phx::MaterialHandle* target_screen_;
phx::Entity* targetEntity;
SupportRod* targetSupportRod;
+ const unsigned int laser_traces_per_frame = 300u;
+
+ // used to copy and expand data
+ float* textureDataBuffer = nullptr;
+ void copyAndScaleTargetData();
// laser object
+ void createLaser();
phx::Entity* laser_object_;
SupportRod* rod_;
};