Yes, all distance (depth) values should have been float types.
Yes, you need to set the device pointers of all your input or output buffers inside the launch parameter structure to be able to access them inside the OptiX device code.
Here are some other code changes which would make things a little simpler. See comments marked with // NV.
struct
{
vec3f vertical;
vec3f direction;
vec3f* cameraLocationsBuffer;
uint32_t* depthBuffer; // NV must be float*
vec3i size; // NV Redundant when optixGetLaunchDimensions matches these values.
} cameras;
extern "C" __global__ void __raygen__renderTrajectoryElevations()
{
const uint3 idx = optixGetLaunchIndex();
const uint3 dim = optixGetLaunchDimensions();
const uint32_t camIndex = idx.z * dim.y * dim.x +
idx.y * dim.x +
idx.x; // NV Correct, linear index of 3D launch.
// depth payload
// NV If you initialize this for the miss case with -1.0f, then you wouldn't need a miss program.
// (OptixProgramDescription needs to be set to nullptr inside the OptixProgramGroupDesc.)
uint32_t u0 = __float_as_uint(-1.0f);
// NV I would recommend to use CUDA vector types (here float3) inside device code, because that is what optixTrace arguments expect.
vec3f camPosition = optixLaunchParams.cameras.cameraLocationsBuffer[camIndex];
// NV I'm assuming this is only debug code and should later render with some projection per camera?
vec3f rayDir = optixLaunchParams.cameras.direction;
optixTrace(optixLaunchParams.traversable,
camPosition, // NV This expects a float3
rayDir,
0.0f, // t_min
1e3f, // t_max // NV Make this as small as possible to capture all your results.
0.0f, // rayTime
OptixVisibilityMask(255),
OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE,
SURFACE_RAY_TYPE, // SBT Offset
RAY_TYPE_COUNT, // SBT Stride
SURFACE_RAY_TYPE, // miss SBT Index // NV There wouldn't need to be a miss program when all that does is returning -1.0f inside the payload register.
u0);
// get depth from payload
const float depth = __uint_as_float(u0); // NV, Yes, positive when hit. A better name would be "distance".
// and write to frame buffer ...
optixLaunchParams.cameras.depthBuffer[camIndex] = depth;
}
extern "C" __global__ void __closesthit__radiance()
{
// get depth
const float t = optixGetRayTmax(); // NV Correct. It's the intersection "distance".
optixSetPayload_0(__float_as_uint(t));
}
if (launchParams.trajectory_cameras.size.x == 0)
return;
launchParamsBuffer.upload(&launchParams, 1);
OPTIX_CHECK(optixLaunch(pipeline,
stream,
launchParamsBuffer.d_pointer(),
launchParamsBuffer.sizeInBytes,
&sbt,
launchParams.cameras.size.x, // NV CAMERA_XRES
launchParams.cameras.size.y, // NV CAMERA_YRES
launchParams.cameras.size.z)); // NV NUM_CAMERAS
CUDA_SYNC_CHECK();
Model* model = loadOBJ("../models/myscene_smooth.obj");
SampleRenderer sample(model);
// should I use BOOST here? vector<vector<double>>(4, vector<double>(5)));
// NV This can be done much simpler with a linear array.
// NV I'm assuming the final goal is to rendere a number of camera intersection distance images with some projection in a single launch.
// NV If all images have the same dimension, using a 3D launch is actually working well,
// because the internal warps will use a 2D layout on the individual 2D slices of the 3D volume, so memory locality is good.
//std::vector<std::vector<std::vector<vec3f>>> camLocations(2, std::vector<std::vector<vec3f>>(2, std::vector<vec3f>(2)));
//
//const int K = camLocations.size();
//const int M = camLocations[0].size();
//const int N = camLocations[0][0].size();
//
//// add some random camera locations to test with
//for (int k = 0; k < K; k++) {
// for (int m = 0; m < M; m++) {
// for (int n = 0; n < N; n++) {
// camLocations[k][m][n] = vec3f(-23.0f * float(k + 1), 4.6f * float(m + 1), 30.7f * float(n + 1));
// }
// }
//}
// NV Assuming there are some defines or function parameters defining the camera x- and y-resolution and the number of cameras.
// NV Using your K, M, N values, define as size_t to match following size calculations. Could also use ull suffix
#define CAMERA_XRES size_t(2)
#define CAMERA_YRES size_t(2)
#define NUM_CAMERAS size_t(2)
const size_t numPoints = CAMERA_XRES * CAMERA_YRES * NUM_CAMERAS;
std::vector<vec3f> camLocations(numPoints);
vec3f* p = camLocations.data(); // NV Local running pointer.
// NV Could also use explicit dynamic allocations, but that wouldn't work with those API wrappers alloc_and_upload().
// vec3f* camLocations = new vec3f[CAMERA_XRES * CAMERA_YRES * NUM_CAMERAS];
// vec3f* p = camLocations;
// add some random camera locations to test with
// NV This exactly matches the linear index calculation inside the OptiX device code.
for (size_t z = 0; z < NUM_CAMERAS; ++z)
{
for (size_t y = 0; y < CAMERA_YRES; ++y)
{
for (size_t x = 0; x < CAMERA_XRES; ++x)
{
*p++ = vec3f(-23.0f * float(x + 1), // NV Not the same order because you used k for x, etc.
4.6f * float(y + 1),
30.7f * float(z + 1));
}
}
}
// get number of cameras
// NV I'm assuming this is only debug code because this is the number of rays and later each camera should shoot CAMERA_XRES * CAMERA_YRES rays?
// Moved up to use for the host size allocation as well.
// const int numPoints = K * M * N;
// all cameras look down(at surface)
// NV This is also hopefully debug code as well, because if all rays go straight down onto an elevation map aligned with the xy-plane,
// you wouldn't need a raytracer to determine these results.
// NV Also note that having a direction collinear with the camera up-axis will not work if that should define a projection later.
const vec3f camDirection = (0.0f, 0.0f, -1.0f); // NV Careful with immediate double values. Always append "f" for floats.
const vec3f camUp = (0.0f, 0.0f, 1.0f); // NV Try is this can use the constructor directly in these: const vec3f camUp(0.0f, 0.0f, 1.0f);
sample.launchParams.cameras.direction = camDirection;
sample.launchParams.cameras.vertical = camUp;
// allocate memory
//sample.cameraLocationsBuffer.resize(numPoints * sizeof(vec3f));
sample.cameraLocationsBuffer.alloc_and_upload(camLocations);
// delete [] camLocations; // NV when using new[] above.
sample.resize_elevationDepthBuffer(vec3i(CAMERA_XRES, CAMERA_YRES, NUM_CAMERAS));
sample.render_trajectory_elevations(); // fails on this call
I’m assuming the final goal will be to renderer a number of 2D distance images from different camera positions and actual camera projections. So you wouldn’t actually define all ray origins and directions but instead a camera and a 2D resolution?
Because when all ray directions are straight down as set in const vec3f camDirection = (0.0f, 0.0f, -1.0f); then you do not need a ray tracer at all to determine the distance on an elevation map with no caves. That could be determined with a texture map lookup instead.
Also please note that the optixLaunch dimension is limited to 2^30.
Please read these related threads:
https://forums.developer.nvidia.com/t/rendering-for-multiple-predefined-camera-locations/295480/2
https://forums.developer.nvidia.com/t/3d-optixlaunch-to-accommodate-multiple-viewpoints/160421/2