Decoupling the launch dimension from the rendering resolution just means that you define the resolution of your output buffer inside the global system parameters and use that instead of the launch dimensions to calculate the primary rays.
Yes, if you have different ray generation programs you always need two launches.
1.) This example would use only one ray generation program and would still need two launches with different dimensions but has only one pipeline and SBT. (You can do with one as well, see below.)
struct SystemData
{
// 8-byte alignment
OptixTraversableHandle topObject;
// The accumulated linear color space output buffer.
// This is always sized to the resolution, not always matching the launch dimension.
// Using a CUdeviceptr here to allow for different buffer formats without too many casts.
CUdeviceptr outputBuffer;
int2 resolution; // The resolution of the full image in outputBuffer. Independent from the launch dimensions for some rendering strategies.
CUdeviceptr pickingBuffer; // Buffer which will receive all information of the primary hit when pickingEnabled != 0.
float2 pickingFragment; // Pixel center coordinate on the full resolution image to shoot a primary picking ray for with the current projection.
// 4 byte alignment
int pickingEnabled; // Flag indicating that the pickingPixel should be used to calculate the primary ray. Launch at 1x1 size in that case.
};
extern "C" __constant__ SystemData sysData;
extern "C" __global__ void __raygen__eye_path()
{
...
const uint2 theLaunchDim = make_uint2(optixGetLaunchDimensions());
const uint2 theLaunchIndex = make_uint2(optixGetLaunchIndex());
// Decoupling the launch dimension from the screen resolution will allow for partial rendering algorithms.
const float2 screen = make_float2(sysData.resolution); // Note, not using theLaunchDim here!
// E.g. assume theLaunchDim == sysData.resolution for rendering and theLaunchDim == (1,1) for picking.
const float2 fragment = (sysData.pickingEnabled) ? sysData.pickingFragment : make_float2(theLaunchIndex) + 0.5f; // E.g. not progressive, just shooting rays at pixel centers.
const float2 ndc = (fragment / screen) * 2.0f - 1.0f; // Normalized device coordinates in range [-1, 1].
// Assume sysData.camera contains the usual pinhole camera setup.
float3 origin = sysData.camera.P;
float3 direction = normalize(sysData.camera.U * ndc.x + sysData.camera.V * ndc.y + sysData.camera.W);
...
// Now you would initialize all your per ray data with the necessary information to distinguish rendering from picking.
// The closesthit programs would be responsible to return the desired picking information on the resp. per ray data fields and end the path if picking is enabled.
// Then you can write the results to sysData.pickingBuffer here.
2.) Since you said you need a continuous stream of picking results, it’s really simple to combine picking and rendering into one launch as well. Means one of the launch indices would return picking results on every render.
For that you would just need to have an uint2 pickingPixel; in the SystemData and then indicate on the per ray data (PRD) for the primary ray where theLaunchIndex == sysData.pickingPixel that this ray should also fill in the PRD field with the hit IDs.
That needs a depth field on the PRD to identify primary rays (just like above as well).
This is what I’m using.
#include "config.h"
#include <optix.h>
#include "system_data.h"
extern "C" __constant__ SystemData sysData;
extern "C" __global__ void __exception__all()
{
//const uint3 theLaunchDim = optixGetLaunchDimensions();
const uint3 theLaunchIndex = optixGetLaunchIndex();
const int theExceptionCode = optixGetExceptionCode();
printf("Exception %d at (%u, %u)\n", theExceptionCode, theLaunchIndex.x, theLaunchIndex.y);
// DAR FIXME This only works for render strategies where the launch dimension matches the outputBuffer resolution.
//float4* buffer = reinterpret_cast<float4*>(sysData.outputBuffer);
//const unsigned int index = theLaunchDim.x * theLaunchIndex.y + theLaunchIndex.x;
//buffer[index] = make_float4(1000000.0f, 0.0f, 1000000.0f, 1.0f); // super magenta
}