Skip to content

Commit

Permalink
MDL_renderer: Replaced double floating point immediate values for PI …
Browse files Browse the repository at this point in the history
…in texture_lookup.h with defines from vector_math.h.

Made the checks for valid idxCall values consistent.
Adjusted some comments.
Explained in README.md how to adjust 3rdparty.cmd cmake.exe call when not using CMake versions from cmake.org.
  • Loading branch information
droettger committed May 9, 2023
1 parent becec33 commit 37f2255
Show file tree
Hide file tree
Showing 7 changed files with 48 additions and 51 deletions.
8 changes: 4 additions & 4 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -210,7 +210,7 @@ Pre-requisites:

3rdparty library setup:
* From the Windows *Start Menu* (Windows' search bar might not find it!), open the *x64 Native Tools Command Prompt for VS2017* or *x64 Native Tools Command Prompt for VS2019* or *x64 Native Tools Command Prompt for VS2022*
* Change directory to the folder containing the `3rdparty.cmd`
* Change directory to the folder containing the `3rdparty.cmd`. This command assumes that CMake is installed in `C:\Program Files\cmake\bin\cmake.exe`. If you're using the CMake installed with Microsoft Visual Studio, please adjust `3rdparty.cmd` to the proper executable location or `cmake.exe` when it's inside the PATH environment variable, which is not the default when installing CMake versions from cmake.org.
* Execute the command `3rdparty.cmd`. This will automatically download GLFW 3.3, GLEW 2.1.0, and ASSIMP archives from sourceforge.com or github.com (see `3rdparty.cmake`) and unpack, compile and install them into the existing `3rdparty` folder in a few minutes.
* Close the *x64 Native Tools Command Prompt* after it finished.

Expand Down Expand Up @@ -261,7 +261,7 @@ Building the examples:

Adding the libraries and data (Yes, this could be done automatically but this is required only once.):
* Copy the x64 library DLLs: `cudart64_<toolkit_version>.dll` from the CUDA installation bin folder, and from the respective 3rdparty folders: `glew32.dll`, `DevIL.dll`, `ILU.dll`, `ILUT.dll`, `assimp-vc<compiler_version>-mt.dll` and `CuOmmBaking.dll` from the OptiX Toolkit into the build target folder with the executables (*bin/Release* or *bin/Debug*). (E.g. `cudart64_101.dll` from CUDA Toolkit 10.1 or cudart64_110.dll from the matching(!) CUDA Toolkit 11.x or 12.x version and `assimp-vc143-mt.dll` from the `3rdparty/assimp` folder when building with MSVS 2022.)
* IMPORTANT: Copy all files from the `data` folder into the build folder with the executables (`bin/Release` or `bin/Debug`). The executables search for the texture images relative to their module directory.
* **IMPORTANT**: Copy all files from the `data` folder into the build folder with the executables (`bin/Release` or `bin/Debug`). The executables search for their resources relative to their working directory.

**Linux**

Expand All @@ -287,13 +287,13 @@ Build the Examples:
* `OPTIX77_PATH=<path_to_optix_7.7.0> OPTIX_TOOLKIT_PATH=<path_to_optix_toolkit> MDL_SDK_PATH=<path_to_MDL_SDK> cmake ..`
* Similar for all other OptiX 7.x.0 SDKs by changing the minor version number accordingly. Some examples won't be built when using older OptiX SDK versions.
* `make`
* IMPORTANT: Copy all files from the `data` folder into the `bin` folder with the executables. The executables search for the texture images relative to their module directory.
* **IMPORTANT**: Copy all files from the `data` folder into the `bin` folder with the executables. The executables search for their resources relative to their working directory.

Instead of setting the temporary OPTIX77_PATH environment variable, you can also adjust the line `set(OPTIX77_PATH "~/NVIDIA-OptiX-SDK-7.7.0-linux64")` inside the `3rdparty/CMake/FindOptiX77.cmake` script to your local OptiX SDK 7.7.0 installation. Similar for the other OptiX 7.x.0 versions.

# Running

IMPORTANT: When running the examples from inside the debugger, make sure the working directory points to the folder with the executable because files are searched relative to that. In Visual Studio that is the same as `$(TargetDir)`. The default is `$(ProjectDir)` which will not work!
**IMPORTANT**: When running the examples from inside the debugger, make sure the working directory points to the folder with the executable because resources are searched relative to that. In Visual Studio that is the same as `$(TargetDir)`. The default is `$(ProjectDir)` which will not work!

Open a command prompt and change directory to the folder with the executables (same under Linux, just without the .exe suffix.)

Expand Down
3 changes: 1 addition & 2 deletions apps/MDL_renderer/shaders/config.h
Original file line number Diff line number Diff line change
Expand Up @@ -63,11 +63,10 @@

// The number of supported texture coordinate slots inside MDL shaders (state::texture_*(i)).
// This makes the Mdl_state bigger which will cost performance!
// It's also less convenient to use the TBN ortho-normal basis.
// hair_bsdf() requires two texture coordinates to communicate the intersection results
// and a per fiber texture coordinate which can be used to color each hair individually.
// That's the whole reason for this define.
// HACK The renderer's vertex attributes are not acctually managing two distinct texture coordinates.
// HACK The renderer's vertex attributes are not actually managing two distinct texture coordinates.
// The Mdl_state will simply duplicate the data of texture coordinate slot 0 to slot 1 everywhere except for the hair_bsdf().
#define NUM_TEXTURE_SPACES 2

Expand Down
12 changes: 6 additions & 6 deletions apps/MDL_renderer/shaders/hit.cu
Original file line number Diff line number Diff line change
Expand Up @@ -478,19 +478,19 @@ extern "C" __global__ void __closesthit__radiance()
if (isFrontFace) // Entered a volume.
{
float3 absorption = shaderConfiguration.absorption_coefficient;
if (0 < shaderConfiguration.idxCallVolumeAbsorptionCoefficient)
if (0 <= shaderConfiguration.idxCallVolumeAbsorptionCoefficient)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeAbsorptionCoefficient, &absorption, &state, &res_data, nullptr, material.arg_block);
}

float3 scattering = shaderConfiguration.scattering_coefficient;
if (0 < shaderConfiguration.idxCallVolumeScatteringCoefficient)
if (0 <= shaderConfiguration.idxCallVolumeScatteringCoefficient)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeScatteringCoefficient, &scattering, &state, &res_data, nullptr, material.arg_block);
}

float bias = shaderConfiguration.directional_bias;
if (0 < shaderConfiguration.idxCallVolumeDirectionalBias)
if (0 <= shaderConfiguration.idxCallVolumeDirectionalBias)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeDirectionalBias, &bias, &state, &res_data, nullptr, material.arg_block);
}
Expand Down Expand Up @@ -791,19 +791,19 @@ extern "C" __global__ void __closesthit__radiance_no_emission()
if (isFrontFace) // Entered a volume.
{
float3 absorption = shaderConfiguration.absorption_coefficient;
if (0 < shaderConfiguration.idxCallVolumeAbsorptionCoefficient)
if (0 <= shaderConfiguration.idxCallVolumeAbsorptionCoefficient)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeAbsorptionCoefficient, &absorption, &state, &res_data, nullptr, material.arg_block);
}

float3 scattering = shaderConfiguration.scattering_coefficient;
if (0 < shaderConfiguration.idxCallVolumeScatteringCoefficient)
if (0 <= shaderConfiguration.idxCallVolumeScatteringCoefficient)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeScatteringCoefficient, &scattering, &state, &res_data, nullptr, material.arg_block);
}

float bias = shaderConfiguration.directional_bias;
if (0 < shaderConfiguration.idxCallVolumeDirectionalBias)
if (0 <= shaderConfiguration.idxCallVolumeDirectionalBias)
{
optixDirectCall<void>(shaderConfiguration.idxCallVolumeDirectionalBias, &bias, &state, &res_data, nullptr, material.arg_block);
}
Expand Down
10 changes: 6 additions & 4 deletions apps/MDL_renderer/shaders/miss.cu
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,7 @@ extern "C" __global__ void __miss__env_constant()
if (sysData.directLighting)
{
// If the last surface intersection was diffuse or glossy which was directly lit with multiple importance sampling,
// then calculate light emission with multiple importance sampling as well.
// then calculate implicit light emission with multiple importance sampling as well.
const float weightMIS = (thePrd->eventType & (mi::neuraylib::BSDF_EVENT_DIFFUSE | mi::neuraylib::BSDF_EVENT_GLOSSY))
? balanceHeuristic(thePrd->pdf, 0.25f * M_1_PIf)
: 1.0f;
Expand Down Expand Up @@ -123,18 +123,20 @@ extern "C" __global__ void __miss__env_sphere()
// The environment light is always in the first element.
const LightDefinition& light = sysData.lightDefinitions[0];

const float3 R = transformVector(light.oriInv, thePrd->wi); // Transform the ray.direction from world space to light object space.
// Transform the ray.direction from world space to light object space.
const float3 R = transformVector(light.oriInv, thePrd->wi);

// All lights shine down the positive z-axis in this renderer.
const float u = (atan2f(-R.x, R.z) + M_PIf) * 0.5f * M_1_PIf;
const float v = acosf(-R.y) * M_1_PIf; // Texture is with origin at lower left, v == 0.0f is south pole.
// Texture is with origin at lower left, v == 0.0f is south pole.
const float v = acosf(-R.y) * M_1_PIf;

float3 emission = make_float3(tex2D<float4>(light.textureEmission, u, v));

if (sysData.directLighting)
{
// If the last surface intersection was a diffuse event which was directly lit with multiple importance sampling,
// then calculate light emission with multiple importance sampling for this implicit light hit as well.
// then calculate implicit light emission with multiple importance sampling as well.
if (thePrd->eventType & (mi::neuraylib::BSDF_EVENT_DIFFUSE | mi::neuraylib::BSDF_EVENT_GLOSSY))
{
// For simplicity we pretend that we perfectly importance-sampled the actual texture-filtered environment map
Expand Down
2 changes: 1 addition & 1 deletion apps/MDL_renderer/shaders/raygeneration.cu
Original file line number Diff line number Diff line change
Expand Up @@ -165,7 +165,7 @@ __forceinline__ __device__ float3 integrator(PerRayData& prd)
}
}

// Note that the primary rays (or volume scattering miss cases) wouldn't normally offset the ray t_min by sysSceneEpsilon. Keep it simple here.
// Note that the primary rays and volume scattering miss cases do not offset the ray t_min by sysSceneEpsilon.
optixTrace(sysData.topObject,
prd.pos, prd.wi, // origin, direction
epsilon, prd.distance, 0.0f, // tmin, tmax, time
Expand Down
2 changes: 1 addition & 1 deletion apps/MDL_renderer/shaders/system_data.h
Original file line number Diff line number Diff line change
Expand Up @@ -94,8 +94,8 @@ struct SystemData
int typeLens; // Camera type.

int numCameras;
int numMaterials;
int numLights;
int numMaterials;

int directLighting;
};
Expand Down
62 changes: 29 additions & 33 deletions apps/MDL_renderer/shaders/texture_lookup.h
Original file line number Diff line number Diff line change
Expand Up @@ -36,18 +36,14 @@
#include <cuda.h>
#include <cuda_runtime.h>

#include "vector_math.h"
#include "texture_handler.h"

#include <mi/neuraylib/target_code_types.h>

// PERF Disabled to not slow down the texure lookup functions.
//#define USE_SMOOTHERSTEP_FILTER

#ifndef M_PI
#define M_PI 3.14159265358979323846
#endif

#define M_ONE_OVER_PI 0.318309886183790671538

typedef mi::neuraylib::tct_deriv_float tct_deriv_float;
typedef mi::neuraylib::tct_deriv_float2 tct_deriv_float2;
Expand Down Expand Up @@ -681,10 +677,10 @@ extern "C" __device__ float df_light_profile_evaluate(
float u = (theta_phi[0] - lp.theta_phi_start.x) * lp.theta_phi_inv_delta.x * lp.inv_angular_resolution.x;

// converting input phi from -pi..pi to 0..2pi
float phi = (theta_phi[1] > 0.0f) ? theta_phi[1] : (float(2.0 * M_PI) + theta_phi[1]);
float phi = (theta_phi[1] > 0.0f) ? theta_phi[1] : 2.0f * M_PIf + theta_phi[1];

// floorf wraps phi range into 0..2pi
phi = phi - lp.theta_phi_start.y - floorf((phi - lp.theta_phi_start.y) * float(0.5 / M_PI)) * float(2.0 * M_PI);
phi = phi - lp.theta_phi_start.y - floorf((phi - lp.theta_phi_start.y) * 0.5f * M_1_PIf) * (2.0f * M_PIf);

// (phi < 0.0f) is no problem, this is handle by the (black) border
// since it implies lp.theta_phi_start.y > 0 (and we really have "no data" below that)
Expand Down Expand Up @@ -783,13 +779,13 @@ extern "C" __device__ void df_light_profile_sample(
result[1] = start.y + (float(idx_phi) + xi0) * delta.y;

// align phi
if (result[1] > float(2.0 * M_PI))
if (result[1] > 2.0f * M_PIf)
{
result[1] -= float(2.0 * M_PI); // wrap
result[1] -= 2.0f * M_PIf; // wrap
}
if (result[1] > float(1.0 * M_PI))
if (result[1] > M_PIf)
{
result[1] = float(-2.0 * M_PI) + result[1]; // to [-pi, pi]
result[1] = -2.0f * M_PIf + result[1]; // to [-pi, pi]
}

// compute pdf
Expand Down Expand Up @@ -822,10 +818,10 @@ extern "C" __device__ float df_light_profile_pdf(
const int idx_theta = int(theta * lp.theta_phi_inv_delta.x);

// converting input phi from -pi..pi to 0..2pi
float phi = (theta_phi[1] > 0.0f) ? theta_phi[1] : (float(2.0 * M_PI) + theta_phi[1]);
float phi = (theta_phi[1] > 0.0f) ? theta_phi[1] : (2.0f * M_PIf + theta_phi[1]);

// floorf wraps phi range into 0..2pi
phi = phi - lp.theta_phi_start.y - floorf((phi - lp.theta_phi_start.y) * float(0.5 / M_PI)) * float(2.0 * M_PI);
phi = phi - lp.theta_phi_start.y - floorf((phi - lp.theta_phi_start.y) * (0.5f * M_1_PIf)) * (2.0f * M_PIf);

// (phi < 0.0f) is no problem, this is handle by the (black) border
// since it implies lp.theta_phi_start.y > 0 (and we really have "no data" below that)
Expand Down Expand Up @@ -933,16 +929,16 @@ __forceinline__ __device__ float3 bsdf_compute_uvw(const float theta_phi_in[2],
float u = theta_phi_out[1] - theta_phi_in[1];
if (u < 0.0)
{
u += float(2.0 * M_PI);
u += 2.0f * M_PIf;
}
if (u > float(1.0 * M_PI))
if (u > M_PIf)
{
u = float(2.0 * M_PI) - u;
u = 2.0f * M_PIf - u;
}
u *= M_ONE_OVER_PI;
u *= M_1_PIf;

const float v = theta_phi_out[0] * float(2.0 / M_PI);
const float w = theta_phi_in[0] * float(2.0 / M_PI);
const float v = theta_phi_out[0] * M_2_PIf;
const float w = theta_phi_in[0] * M_2_PIf;

return make_float3(u, v, w);
}
Expand Down Expand Up @@ -1035,7 +1031,7 @@ extern "C" __device__ void df_bsdf_measurement_sample(
const float* sample_data = bm.sample_data[part_index];

// compute the theta_in index (flipping input and output, BSDFs are symmetric)
unsigned int idx_theta_in = (unsigned int)(theta_phi_out[0] * M_ONE_OVER_PI * 2.0f * float(res.x));
unsigned int idx_theta_in = (unsigned int)(theta_phi_out[0] * M_2_PIf * float(res.x));
idx_theta_in = min(idx_theta_in, res.x - 1);

// sample theta_out
Expand Down Expand Up @@ -1082,8 +1078,8 @@ extern "C" __device__ void df_bsdf_measurement_sample(
//-------------------------------------------
const float2 inv_res = bm.inv_angular_resolution[part_index];

const float s_theta = float(0.5 * M_PI) * inv_res.x;
const float s_phi = float(1.0 * M_PI) * inv_res.y;
const float s_theta = M_PI_2f * inv_res.x;
const float s_phi = M_PIf * inv_res.y;

const float cos_theta_0 = cosf(float(idx_theta_out) * s_theta);
const float cos_theta_1 = cosf(float(idx_theta_out + 1u) * s_theta);
Expand All @@ -1094,18 +1090,18 @@ extern "C" __device__ void df_bsdf_measurement_sample(

if (flip)
{
result[1] = float(2.0 * M_PI) - result[1]; // phi \in [0, 2pi]
result[1] = 2.0f * M_PIf - result[1]; // phi \in [0, 2pi]
}

// align phi
result[1] += (theta_phi_out[1] > 0) ? theta_phi_out[1] : (float(2.0 * M_PI) + theta_phi_out[1]);
if (result[1] > float(2.0 * M_PI))
result[1] += (theta_phi_out[1] > 0) ? theta_phi_out[1] : (2.0f * M_PIf + theta_phi_out[1]);
if (result[1] > 2.0f * M_PIf)
{
result[1] -= float(2.0 * M_PI);
result[1] -= 2.0f * M_PIf;
}
if (result[1] > float(1.0 * M_PI))
if (result[1] > M_PIf)
{
result[1] = float(-2.0 * M_PI) + result[1]; // to [-pi, pi]
result[1] = -2.0f * M_PIf + result[1]; // to [-pi, pi]
}

// compute pdf
Expand Down Expand Up @@ -1141,8 +1137,8 @@ extern "C" __device__ float df_bsdf_measurement_pdf(

// compute indices in the CDF data
float3 uvw = bsdf_compute_uvw(theta_phi_in, theta_phi_out); // phi_delta, theta_out, theta_in
unsigned int idx_theta_in = (unsigned int)(theta_phi_in[0] * M_ONE_OVER_PI * 2.0f * float(res.x));
unsigned int idx_theta_out = (unsigned int)(theta_phi_out[0] * M_ONE_OVER_PI * 2.0f * float(res.x));
unsigned int idx_theta_in = (unsigned int)(theta_phi_in[0] * M_2_PIf * float(res.x));
unsigned int idx_theta_out = (unsigned int)(theta_phi_out[0] * M_2_PIf * float(res.x));
unsigned int idx_phi_out = (unsigned int)(uvw.x * float(res.y));

idx_theta_in = min(idx_theta_in, res.x - 1);
Expand Down Expand Up @@ -1173,8 +1169,8 @@ extern "C" __device__ float df_bsdf_measurement_pdf(
// compute probability to select a position in the sphere patch
float2 inv_res = bm.inv_angular_resolution[part_index];

const float s_theta = float(0.5 * M_PI) * inv_res.x;
const float s_phi = float(1.0 * M_PI) * inv_res.y;
const float s_theta = M_PI_2f * inv_res.x;
const float s_phi = M_PIf * inv_res.y;

const float cos_theta_0 = cosf(float(idx_theta_out) * s_theta);
const float cos_theta_1 = cosf(float(idx_theta_out + 1u) * s_theta);
Expand All @@ -1200,7 +1196,7 @@ __forceinline__ __device__ void df_bsdf_measurement_albedo(
}

const uint2 res = bm.angular_resolution[part_index];
unsigned int idx_theta = (unsigned int)(theta_phi[0] * float(2.0 / M_PI) * float(res.x));
unsigned int idx_theta = (unsigned int)(theta_phi[0] * M_2_PIf * float(res.x));

idx_theta = min(idx_theta, res.x - 1u);
result[0] = bm.albedo_data[part_index][idx_theta];
Expand Down

0 comments on commit 37f2255

Please sign in to comment.