OptiX basic visualization problem

Hi, I am new to using optix and have just been able to create my own program which generates an image of a 2D plane from a camera point of view. Now I would like to expand this project to do the following:

  1. Implement an “antenna” in the form of a source point above the plane.
  2. Shoot a ray from the antenna towards the plane and make the intersection a different color if it hits.
  3. Still be able to look through the cameras point of view and see the impact on the plane caused by the ray from the antenna.

My first thought how to approach this problem is to create a new rayGen program called something like “rayGenAntenna” to generate rays from the antenna, however I dont seem to be able to capture the color changes from the camera point of view. Is this a correct approach or should I try something different?

This is what my current output looks like, where the plane is centered in the origin and is 1x1 in size. The camera is positioned in (5, 0, 5) looking at the origin with FOV 45 degrees:

Here is a sketch of what I would like to get, basically the same as above but with a yellow dot from the antenna ray hitting the plane:

Here is my code so far if it gives a better understanding of what I have done:

Main application code

#include <optix.h>
#include <cuda_runtime.h>
#include <optix_stubs.h>
#include <optix_function_table_definition.h>
#include <iostream>
#include <../inc/exception.h>
#include <optix_types.h>
#include <sutil/sutil.h>

#include <optix_stack_size.h>

#include <fstream>
#include <sstream>
#include <string>

#include <GLFW/glfw3.h>
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "../inc/stb_image_write.h"

#include "CommonStructs.h"

// Initialize OptiX and create context
void initOptix(CUcontext& cuContext, CUstream& stream) {
    // Initialize CUDA
    CUdevice cuDevice;
    cuInit(0);
    cuDeviceGet(&cuDevice, 0); // Assumes device 0 is suitable
    cuCtxCreate(&cuContext, 0, cuDevice);

    // Create a CUDA stream
    cuStreamCreate(&stream, CU_STREAM_DEFAULT);

    // Initialize OptiX
    optixInit();
}

// Camera struct
//struct Camera {
//    float3 eye;         // Camera position
//    float3 lookAt;      // Look-at point
//    float3 up;          // Up vector
//    float fov;          // Field of view in degrees
//    float aspectRatio;  // Aspect ratio of the image
//};


// Vertex structure for a simple 2D plane
struct Vertex {
    float x, y, z; // Position
};

//Define vertices of a square plane centered at the origin
Vertex vertices[] = {
    {-0.5f, -0.5f, 0.0f}, // Bottom left
    {0.5f, -0.5f, 0.0f},  // Bottom right
    {0.5f, 0.5f, 0.0f},   // Top right
    {-0.5f, 0.5f, 0.0f}   // Top left
};


// Define indices for two triangles that make up the square
unsigned int indices[] = {
    0, 1, 2, // First triangle
    2, 3, 0  // Second triangle
};



// Load PTX files path
const char* ptxPathRayGen = "C:/Users/PC/Desktop/new_example_5/cuda/cuda_output/rayGen_v4.ptx";
const char* ptxPathhit = "C:/Users/PC/Desktop/new_example_5/cuda/cuda_output/hit_v4.ptx";
const char* ptxPathMiss = "C:/Users/PC/Desktop/new_example_5/cuda/cuda_output/miss_v4.ptx";




OptixTraversableHandle createAccelerationStructure(OptixDeviceContext context, CUdeviceptr& d_vertices, CUdeviceptr& d_indices, CUdeviceptr& d_tempBufferGas, CUdeviceptr& d_gasOutputBuffer) {
    // Assume 'vertices' and 'indices' are already defined as shown above

    size_t vertices_size = sizeof(vertices);
    size_t indices_size = sizeof(indices);

    // Allocate device memory for vertices and indices
    cuMemAlloc(&d_vertices, vertices_size);
    cuMemAlloc(&d_indices, indices_size);

    // Copy vertices and indices to device memory
    cuMemcpyHtoD(d_vertices, vertices, vertices_size);
    cuMemcpyHtoD(d_indices, indices, indices_size);

    // Define build input
    OptixBuildInput buildInput = {};
    buildInput.type = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;

    // Setup vertex buffer
    OptixBuildInputTriangleArray triangleArray = {};
    triangleArray.vertexFormat = OPTIX_VERTEX_FORMAT_FLOAT3;
    triangleArray.vertexStrideInBytes = sizeof(Vertex);
    triangleArray.numVertices = 4;
    triangleArray.vertexBuffers = &d_vertices;

    // Setup index buffer
    triangleArray.indexFormat = OPTIX_INDICES_FORMAT_UNSIGNED_INT3;
    triangleArray.indexStrideInBytes = sizeof(unsigned int) * 3;
    triangleArray.numIndexTriplets = 2;
    triangleArray.indexBuffer = d_indices;

    // Set the flags for each SBT record
    unsigned int inputFlags[1] = { OPTIX_GEOMETRY_FLAG_NONE };
    triangleArray.flags = inputFlags; // Correctly assign flags to the triangleArray
    triangleArray.numSbtRecords = 1; // Correctly assign the number of SBT records

    // Now, correctly link the fully configured triangleArray to the buildInput
    buildInput.triangleArray = triangleArray;

    // Specify build options
    OptixAccelBuildOptions accelOptions = {};
    accelOptions.buildFlags = OPTIX_BUILD_FLAG_NONE;
    accelOptions.operation = OPTIX_BUILD_OPERATION_BUILD;

    // Output buffers for acceleration structure
    OptixAccelBufferSizes gasBufferSizes;
    optixAccelComputeMemoryUsage(context, &accelOptions, &buildInput, 1, &gasBufferSizes);

    // Allocate memory for acceleration structure
    cuMemAlloc(&d_tempBufferGas, gasBufferSizes.tempSizeInBytes);
    cuMemAlloc(&d_gasOutputBuffer, gasBufferSizes.outputSizeInBytes);

    // Build acceleration structure
    OptixTraversableHandle gasHandle = 0;
    optixAccelBuild(context, 0, &accelOptions, &buildInput, 1,
        d_tempBufferGas, gasBufferSizes.tempSizeInBytes,
        d_gasOutputBuffer, gasBufferSizes.outputSizeInBytes,
        &gasHandle, nullptr, 0);

    // Return the handle to the created acceleration structure
    return gasHandle;

    // Do not forget to free the allocated buffers after they are no longer needed
}


// Function to load the contents of a PTX file into a string
std::string loadPtx(const std::string& filepath) {
    std::ifstream file(filepath.c_str());
    if (!file) {
        std::cerr << "Failed to open PTX file: " << filepath << std::endl;
        return "";
    }

    std::stringstream buffer;
    buffer << file.rdbuf();
    return buffer.str();
}


int main() {
    // Step 1: Setup OptiX context (for OptiX 7+, prepare to use CUDA directly)
    CUcontext cuContext;
    CUstream stream;
    OptixDeviceContext optixContext = nullptr;

    // Initialize OptiX and CUDA context
    initOptix(cuContext, stream);

    CUdeviceptr d_vertices = 0;
    CUdeviceptr d_indices = 0;
    CUdeviceptr d_tempBufferGas = 0;
    CUdeviceptr d_gasOutputBuffer = 0;

    // Create an OptiX device context using the CUDA context
    OPTIX_CHECK(optixDeviceContextCreate(cuContext, nullptr, &optixContext));
    OPTIX_CHECK(optixDeviceContextSetLogCallback(optixContext, [](unsigned int level, const char* tag, const char* message, void* cbdata) {
        std::cerr << "[" << tag << "] " << message << "\n";
        }, nullptr, 4)); // LogLevel=4 for verbosity

    // Step 2: Define geometry and create acceleration structures
    OptixTraversableHandle gasHandle = createAccelerationStructure(optixContext, d_vertices, d_indices, d_tempBufferGas, d_gasOutputBuffer);


    // Module and Program Group Compile Options
    OptixModuleCompileOptions moduleCompileOptions = {};
    moduleCompileOptions.maxRegisterCount = OPTIX_COMPILE_DEFAULT_MAX_REGISTER_COUNT;
    moduleCompileOptions.optLevel = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
    moduleCompileOptions.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE;

    OptixPipelineCompileOptions pipelineCompileOptions = {};
    pipelineCompileOptions.usesMotionBlur = false;
    pipelineCompileOptions.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;
    pipelineCompileOptions.numPayloadValues = 3; // Adjust based on your needs<--------------------------------------- CHANGED from 2
    pipelineCompileOptions.numAttributeValues = 2; // Adjust based on your needs
    pipelineCompileOptions.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE; // or other flags as necessary
    pipelineCompileOptions.pipelineLaunchParamsVariableName = "params";

    // Step 3: Create shaders (ray generation, hit, miss)
    char log[2048]; // For logging
    size_t logSize = sizeof(log) - 1;
    log[0] = '\0'; // Initialize log with empty string


    // Load the PTX source code
    std::string rayGenPTX = loadPtx(ptxPathRayGen);
    OptixModule rayGenModule;
    OPTIX_CHECK(optixModuleCreate(
        optixContext,
        &moduleCompileOptions,
        &pipelineCompileOptions,
        rayGenPTX.c_str(), rayGenPTX.size(),
        log, &logSize,
        &rayGenModule
    ));

    std::string hitPTX = loadPtx(ptxPathhit);
    OptixModule hitModule;
    OPTIX_CHECK(optixModuleCreate(
        optixContext,
        &moduleCompileOptions,
        &pipelineCompileOptions,
        hitPTX.c_str(), hitPTX.size(),
        log, &logSize,
        &hitModule
    ));


    std::string missPTX = loadPtx(ptxPathMiss);
    OptixModule missModule;
    OPTIX_CHECK(optixModuleCreate(
        optixContext,
        &moduleCompileOptions,
        &pipelineCompileOptions,
        missPTX.c_str(), missPTX.size(),
        log, &logSize,
        &missModule
    ));


    // Define program group options - typically, this is left as default for basic usage
    OptixProgramGroupOptions programGroupOptions = {};


    // Create program group for ray generation shader
    OptixProgramGroupDesc raygenPGDesc = {};
    raygenPGDesc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
    raygenPGDesc.raygen.module = rayGenModule; // Ensure rayGenModule is correctly initialized
    raygenPGDesc.raygen.entryFunctionName = "__raygen__rg";

    OptixProgramGroup raygenProgramGroup;
    OPTIX_CHECK(optixProgramGroupCreate(
        optixContext,
        &raygenPGDesc,
        1, // Number of program group descriptions
        &programGroupOptions, // Corrected to use program group options
        log, &logSize, // Log buffer and its size
        &raygenProgramGroup // The created program group
    ));

    // Create program group for closest hit shader
    OptixProgramGroupDesc hitPGDesc = {};
    hitPGDesc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
    hitPGDesc.hitgroup.moduleCH = hitModule; // Use your hit module here
    hitPGDesc.hitgroup.entryFunctionNameCH = "__closesthit__ch"; // Entry point for your hit shader

    OptixProgramGroup hitProgramGroup;
    OPTIX_CHECK(optixProgramGroupCreate(
        optixContext,
        &hitPGDesc, // Use the hit program group descriptor
        1, // One program group description
        &programGroupOptions, // Assuming programGroupOptions is already defined
        log, &logSize, // Log buffer and size
        &hitProgramGroup // The created program group
    ));

    // Create program group for miss shader
    OptixProgramGroupDesc missPGDesc = {};
    missPGDesc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
    missPGDesc.miss.module = missModule; // Ensure missModule is correctly initialized
    missPGDesc.miss.entryFunctionName = "__miss__ms";

    OptixProgramGroup missProgramGroup;
    OPTIX_CHECK(optixProgramGroupCreate(
        optixContext,
        &missPGDesc,
        1, // Number of program group descriptions
        &programGroupOptions, // Program group options
        log, &logSize, // Log buffer and its size
        &missProgramGroup // The created program group
    ));



    // Step 4: Setup ray tracing pipeline
    OptixPipeline pipeline;

    OptixPipelineLinkOptions pipelineLinkOptions = {};
    pipelineLinkOptions.maxTraceDepth = 2;
    OptixProgramGroup programGroups[] = { raygenProgramGroup, hitProgramGroup, missProgramGroup };

    OPTIX_CHECK(optixPipelineCreate(
        optixContext,
        &pipelineCompileOptions,
        &pipelineLinkOptions,
        programGroups,
        sizeof(programGroups) / sizeof(programGroups[0]), // Number of program groups
        log, &logSize,
        &pipeline
    ));


    // STACK SIZES
    OptixStackSizes ssp = {}; // Whole pipeline.

    for (auto pg : programGroups)
    {
        OptixStackSizes ss;

#if (OPTIX_VERSION >= 70700)
        OPTIX_CHECK(optixProgramGroupGetStackSize(pg, &ss, pipeline));
#else
        OPTIX_CHECK(optixProgramGroupGetStackSize(pg, &ss));
#endif

        ssp.cssRG = std::max(ssp.cssRG, ss.cssRG);
        ssp.cssMS = std::max(ssp.cssMS, ss.cssMS);
        ssp.cssCH = std::max(ssp.cssCH, ss.cssCH);
        ssp.cssAH = std::max(ssp.cssAH, ss.cssAH);
        ssp.cssIS = std::max(ssp.cssIS, ss.cssIS);
        ssp.cssCC = std::max(ssp.cssCC, ss.cssCC);
        ssp.dssDC = std::max(ssp.dssDC, ss.dssDC);
    }

    // Temporaries
    unsigned int cssCCTree = ssp.cssCC; // Should be 0. No continuation callables in this pipeline. // maxCCDepth == 0
    unsigned int cssCHOrMSPlusCCTree = std::max(ssp.cssCH, ssp.cssMS) + cssCCTree;

    const unsigned int maxDCDepth = 0; // The __direct_callable__light_mesh calls other direct callables from MDL expressions.

    // Arguments

    unsigned int directCallableStackSizeFromTraversal = ssp.dssDC * maxDCDepth; // FromTraversal: DC is invoked from IS or AH.      // Possible stack size optimizations.
    unsigned int directCallableStackSizeFromState = ssp.dssDC * maxDCDepth; // FromState:     DC is invoked from RG, MS, or CH. // Possible stack size optimizations.
    unsigned int continuationStackSize = ssp.cssRG + cssCCTree + cssCHOrMSPlusCCTree * (std::max(1u, pipelineLinkOptions.maxTraceDepth) - 1u) +
        std::min(1u, pipelineLinkOptions.maxTraceDepth) * std::max(cssCHOrMSPlusCCTree, ssp.cssAH + ssp.cssIS);
    unsigned int maxTraversableGraphDepth = 1;

    OPTIX_CHECK(optixPipelineSetStackSize(pipeline, directCallableStackSizeFromTraversal, directCallableStackSizeFromState, continuationStackSize, maxTraversableGraphDepth));


    // Step 5: Render the scene

    struct __align__(OPTIX_SBT_RECORD_ALIGNMENT) RayGenSbtRecord {
        __align__(OPTIX_SBT_RECORD_HEADER_SIZE) char header[OPTIX_SBT_RECORD_HEADER_SIZE];
        // Add ray generation data here
        //OptixProgramGroup raygenProgram;
    };

    struct __align__(OPTIX_SBT_RECORD_ALIGNMENT) MissSbtRecord {
        __align__(OPTIX_SBT_RECORD_HEADER_SIZE) char header[OPTIX_SBT_RECORD_HEADER_SIZE];
        // Add miss data here
        //OptixProgramGroup missProgram;
    };

    struct __align__(OPTIX_SBT_RECORD_ALIGNMENT) HitGroupSbtRecord {
        __align__(OPTIX_SBT_RECORD_HEADER_SIZE) char header[OPTIX_SBT_RECORD_HEADER_SIZE];
        // Add hit group data here
        //OptixProgramGroup closestHitProgram;
    };

    RayGenSbtRecord rgSbt;
    MissSbtRecord msSbt;
    HitGroupSbtRecord hgSbt;

    optixSbtRecordPackHeader(raygenProgramGroup, &rgSbt);
    optixSbtRecordPackHeader(missProgramGroup, &msSbt);
    optixSbtRecordPackHeader(hitProgramGroup, &hgSbt);

    // Allocate and copy SBT records to device memory
    CUdeviceptr d_raygenRecords, d_missRecords, d_hitgroupRecords;
    size_t sbtSize = sizeof(RayGenSbtRecord); // Same for Miss and HitGroup if they don't have additional data
    cudaMalloc(reinterpret_cast<void**>(&d_raygenRecords), sbtSize);
    cudaMalloc(reinterpret_cast<void**>(&d_missRecords), sbtSize);
    cudaMalloc(reinterpret_cast<void**>(&d_hitgroupRecords), sbtSize);

    cudaMemcpy(reinterpret_cast<void*>(d_raygenRecords), &rgSbt, sbtSize, cudaMemcpyHostToDevice);
    cudaMemcpy(reinterpret_cast<void*>(d_missRecords), &msSbt, sbtSize, cudaMemcpyHostToDevice);
    cudaMemcpy(reinterpret_cast<void*>(d_hitgroupRecords), &hgSbt, sbtSize, cudaMemcpyHostToDevice);

    OptixShaderBindingTable sbt = {};
    sbt.raygenRecord = d_raygenRecords;
    sbt.missRecordBase = d_missRecords;
    sbt.missRecordStrideInBytes = sizeof(MissSbtRecord);
    sbt.missRecordCount = 1;
    sbt.hitgroupRecordBase = d_hitgroupRecords;
    sbt.hitgroupRecordStrideInBytes = sizeof(HitGroupSbtRecord);
    sbt.hitgroupRecordCount = 1;

    //struct LaunchParams {
    //    int width;
    //    int height;
    //    CUdeviceptr outputBuffer;
    //    Camera camera;
    //    OptixTraversableHandle handle;
    //    // Add other parameters as needed
    //};
    LaunchParams params;
    params.width = 800;
    params.height = 600;

    //Camera params
    params.camera.eye = make_float3(0.0f, 5.0f, 5.0f);
    params.camera.lookAt = make_float3(0.0f, 0.0f, 0.0f);
    params.camera.up = make_float3(0.0f, 1.0f, 0.0f);
    params.camera.fov = 45.0f;
    params.camera.aspectRatio = static_cast<float>(params.width) / static_cast<float>(params.height);
    params.handle = gasHandle;

    params.antenna.position = make_float3(0.0f, 0.0f, 2.0f);

    CUdeviceptr d_params;
    cudaMalloc(reinterpret_cast<void**>(&d_params), sizeof(LaunchParams));
    cudaMemcpy(reinterpret_cast<void*>(d_params), &params, sizeof(LaunchParams), cudaMemcpyHostToDevice);

    CUdeviceptr d_outputBuffer;
    cudaMalloc(reinterpret_cast<void**>(&d_outputBuffer), params.width * params.height * sizeof(float4)); // Assuming float4 per pixel
    params.outputBuffer = d_outputBuffer;
    // Update params on the device
    cudaMemcpy(reinterpret_cast<void*>(d_params), &params, sizeof(LaunchParams), cudaMemcpyHostToDevice);

    optixLaunch(
        pipeline,
        stream,
        d_params,
        sizeof(LaunchParams),
        &sbt,
        params.width,  // launch width
        params.height, // launch height
        1              // launch depth
    );
    cudaStreamSynchronize(stream); // Wait for completion

    float4* h_outputBuffer = new float4[params.width * params.height];
    cudaMemcpy(h_outputBuffer, reinterpret_cast<void*>(d_outputBuffer), params.width * params.height * sizeof(float4), cudaMemcpyDeviceToHost);

    //Now, `h_outputBuffer` contains your rendered image. You can save it to a file or display it.

    //This saves a png file
    //Convert color values from [0.0, 1.0] to [0, 255]
    unsigned char* uc_outputBuffer = new unsigned char[params.width * params.height * 4];
    for (int i = 0; i < params.width * params.height; ++i) {
        uc_outputBuffer[i * 4] = static_cast<unsigned char>(std::min(h_outputBuffer[i].x * 255.0f, 255.0f));
        uc_outputBuffer[i * 4 + 1] = static_cast<unsigned char>(std::min(h_outputBuffer[i].y * 255.0f, 255.0f));
        uc_outputBuffer[i * 4 + 2] = static_cast<unsigned char>(std::min(h_outputBuffer[i].z * 255.0f, 255.0f));
        uc_outputBuffer[i * 4 + 3] = 255; // Full opacity
    }

    // Save the image to a file
    const char* outputFile = "output.png";
    stbi_write_png(outputFile, params.width, params.height, 4, uc_outputBuffer, params.width * 4);

    std::cout << "Saved rendered image to " << outputFile << std::endl;


    // Step 6: Cleanup and free resources
    delete[] h_outputBuffer; // Free the original output buffer 
    delete[] uc_outputBuffer; // Don't forget to free the uchar buffer after writing the image

    cudaFree(reinterpret_cast<void*>(d_raygenRecords));
    cudaFree(reinterpret_cast<void*>(d_missRecords));
    cudaFree(reinterpret_cast<void*>(d_hitgroupRecords));
    cudaFree(reinterpret_cast<void*>(d_outputBuffer));
    cudaFree(reinterpret_cast<void*>(d_params));


    OPTIX_CHECK(optixProgramGroupDestroy(raygenProgramGroup));
    OPTIX_CHECK(optixProgramGroupDestroy(hitProgramGroup));
    OPTIX_CHECK(optixProgramGroupDestroy(missProgramGroup));
    OPTIX_CHECK(optixModuleDestroy(rayGenModule));
    OPTIX_CHECK(optixModuleDestroy(hitModule));
    OPTIX_CHECK(optixModuleDestroy(missModule));


    cuMemFree(d_vertices);
    cuMemFree(d_indices);
    cuMemFree(d_tempBufferGas);
    cuMemFree(d_gasOutputBuffer);

    OPTIX_CHECK(optixDeviceContextDestroy(optixContext));
    cuStreamDestroy(stream);
    cuCtxDestroy(cuContext);

    return 0;
}

rayGen program

#include <optix.h>
#include "../sutil/vec_math.h"
#include <cuda_runtime.h>
#include "../src/CommonStructs.h"

extern "C" {
    __constant__ LaunchParams params;
}

extern "C" __global__ void __raygen__rg() {
    const uint3 idx = optixGetLaunchIndex();
    const float u = (static_cast<float>(idx.x) / static_cast<float>(params.width) - 0.5f) * 2.0f;
    const float v = (static_cast<float>(idx.y) / static_cast<float>(params.height) - 0.5f) * 2.0f;

    float3 direction = normalize(params.camera.lookAt - params.camera.eye);
    float3 right = normalize(cross(direction, params.camera.up));
    float3 up = cross(right, direction);

    float fovScale = tanf(params.camera.fov * 0.5f * M_PIf / 180.0f);
    float aspectRatio = params.camera.aspectRatio;

    float3 rayDirection = normalize(direction + fovScale * (u * aspectRatio * right + v * up));



    unsigned int p0, p1, p2; // Payloads to encode float3 color
    optixTrace(
        params.handle, 
        params.camera.eye, 
        rayDirection, 
        0.0f,    // Min intersection distance
        1e16f,   // Max intersection distance
        0.0f,    // rayTime -- for motion blur
        OptixVisibilityMask(1),
        OPTIX_RAY_FLAG_NONE,
        0,       // SBT offset
        1,       // SBT stride
        0,       // missSBTIndex
        p0, p1, p2); // Payloads
    
    float3 color = make_float3(
        __uint_as_float(p0),
        __uint_as_float(p1),
        __uint_as_float(p2));
    
    // Write the color to the output buffer
    float4* outputBuffer = reinterpret_cast<float4*>(params.outputBuffer);
    outputBuffer[idx.y * params.width + idx.x] = make_float4(color, 1.0f);

}

Hit program

#include <optix.h>
#include "../src/CommonStructs.h"

extern "C" {
__constant__ LaunchParams params; // Even if not used directly, it needs to be declared to fix the error.
}


extern "C" __global__ void __closesthit__ch() {
    const float3 hitColor = make_float3(1.0f, 0.0f, 0.0f); // Example: Set hit color to red.
    optixSetPayload_0(__float_as_uint(hitColor.x));
    optixSetPayload_1(__float_as_uint(hitColor.y));
    optixSetPayload_2(__float_as_uint(hitColor.z));
}

Miss program

#include <optix.h>
#include "../src/CommonStructs.h"

extern "C" {
__constant__ LaunchParams params; // This assumes you have a similar struct in your host code.
}

__constant__ float3 bgColor = {0.5f, 0.1f, 1.0f}; // Background color


extern "C" __global__ void __miss__ms() { 
    const float3 hitColor = make_float3(1.0f, 0.0f, 0.0f); // Example: Set hit color to red.
    optixSetPayload_0(__float_as_uint(bgColor.x));
    optixSetPayload_1(__float_as_uint(bgColor.y));
    optixSetPayload_2(__float_as_uint(bgColor.z));
}

CommonStructs.h

#ifndef COMMON_STRUCTS_H
#define COMMON_STRUCTS_H

struct Antenna {
    float3 position; // Position of the antenna in 3D space
    // Additional properties like direction, emission pattern, etc., can be added here
};

struct Camera {
    float3 eye;         // Camera position
    float3 lookAt;      // Look-at point
    float3 up;          // Up vector
    float fov;          // Field of view in degrees
    float aspectRatio;  // Aspect ratio of the image
};

struct LaunchParams {
    int width;
    int height;
    CUdeviceptr outputBuffer;
    Camera camera;
    OptixTraversableHandle handle;
    Antenna antenna;
    // Add other parameters as needed
};

#endif // COMMON_STRUCTS_H

There are two ways to approach that problem of finding intersections of rays from that antenna to the plane.

1.) When you want to do that inside your current implementation which renders the above scene screenshot, then you would shoot rays from your camera into the scene, and inside the closest hit program assigned to your plane, you would shoot a visibility test (“shadow”) ray from that surface hit point to the antenna world position.
Means the tmax stops at that distance.

If nothing is blocking the visibility, then there is a direct connection between these two points and you could evaluate any required information, like distance etc. If the antenna has a specific distribution function, you would be able to calculate things like signal strength etc. from some additional antenna orientation and that connecting ray direction and distance.

Doing it like this will only care about direct connections of hit surface points to the antenna position.

Just think of the antenna as point light and all these calculations as direct lighting.
Think of it as spot light if there is some specific distribution function.

Example code doing exactly that direct lighting can be found in all my examples.
This would, for example, be the direct lighting calculation of a diffuse BRDF:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo12/shaders/brdf_diffuse.cu#L180
and if you assume that your antenna position is a point light, then this explicit light sampling routine would be used to fill the appropriate light sample fields of a singular point light:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo12/shaders/light_sample.cu#L301
Since that is not part of the scene geometry it cannot be hit implicitly, so only direct lighting makes this work.
There is no need to implement that sampling as direct callable program when you have only that one type of antenna.

2.) If you wanted to implement the antenna distribution function (let’s say it’s spherical) and capture the hits arriving on the plane you would need to represent the plane as some output buffer, think of having a 2D texture of a discrete resolution (e.g. 1024 x 1024) mapped onto your plane.
You could of course shoot random rays from the antenna according to the distribution function into the world and check if you hit the plane geometry, and then use atomics to add each hit to the respective texture cell, because that is a scatter algorithm (multiple rays could hit the same cell) but that doesn’t only sound inefficient, it’s really bad.
Instead you would again generate rays starting on the plane, let’s say in the center of each texel and shoot them all to the antenna position and see if there is a direct connection (visibility test succeeded) and whatever other information you need to calculate. This is a gather algorithm because each output buffer cell maps to specific rays.
Now you would have a buffer or texture with the data and could assign that to hit surface points on your plane inside the renderer which displays the above scene.
Means this would require two different optixLaunch calls with completely different raygen programs but the same geometry.

When shooting rays from a surface always take care to prevent self-intersections of the ray with the geometry you started on. That’s usually done by offsetting the ray origin or tmin value a little from the start point on the surface. Same on the other side of the ray (tmax) when using geometry lights (see the comments in my direct
lighting example code.)

So in summary: Prefer gather over scatter algorithms. Handle these problems similar to direct lighting.

Have a look at these related posts:
https://forums.developer.nvidia.com/t/electromagnetic-wave-simulation-using-optix/221893
https://forums.developer.nvidia.com/t/sphere-intersection-with-ray-distance-dependent-radius/60405/6
https://forums.developer.nvidia.com/t/closest-approach-of-ray-to-a-point/231750

1 Like

Hello droettger, thank you for a quick reply!

My apologies, I think that I described my problem a bit poorly or I do not fully understand. Using your solution, visibility checking by shooting “shadow rays” from the plane’s intersections with the rays from the camera towards the antenna, seeing if anything is in the way, lets me color the entire plane the same color as everything is visible to the antenna. However, what I really want to do is the other way around.

Something like this:

  • Shoot a ray (say one ray to begin with for simplicity) from the antenna towards the plane.
  • Check if it hit the plane.
  • If yes, make only that intersection of the plane and antenna ray a certain color (yellow in this case).
  • With all of the above satisfied I want to look at this from the camera point of view.

The end-goal of this project is to simulate a radar but I’m taking one step at a time. Using “shadow rays” from the plane to the antenna tells me if the antenna is visible or not from a certain point, which I don’t think is relevant in my case?

Did I understand you correctly? If so, do you think a solution of this kind would be applicable and are there any examples or tips on how to implement it?

Best regards,
P

I have understood what you wanted. The issue is that the probability of doing this ray-ray connection is not working how you think.

The issue with shooting a single ray with a fixed direction from the antenna to the plane is that this is effectively a Dirac distribution.
When doing that there is exactly one and only one possible connection from the antenna to the plane to the camera, and the probability to hit that exactly when shooting rays from either start point is zero.
Same probability to hit an infinitely small point in space randomly. (That’s why point lights and directional lights are singular lights.)

Means when you shoot a ray from the camera which hits the plane and you shoot a single ray from the antenna, there is a zero chance to connect these two exactly.

What that would require is some epsilon around the hit point (which in light transport terms goes into the realm of photon mapping) or an explicit connection from the plane position hit with the antenna ray to the camera plane for which in a pinhole camera (also an infinitely small point) there is exactly one ray which establishes that.

1.) Version which is similar to final gathering inside a photon mapper:

So what you would need to do to get this working the way you intended, would be to cast a single ray inside the ray generation program from the antenna into your scene.

Inside your per ray data (payload registers), add a flag which indicates that this is the ray from the antenna.
Then inside the closest hit program of the plane, check that flag and if it’s enabled, only store the hit position of the surface point on the plane into your per ray payload (let’s call this “photon world position”).

Now back inside the ray generation, disable that flag, put that photon world position into your per ray payload, and then shoot the primary rays from your camera to render your standard image.
Inside the plane’s closest hit program now each hit would need to check if it’s in an epsilon environment of that photon world position and if yes, color the result accordingly.

That epsilon environment is required because it’s again a zero probability to hit an infinitely small point with a ray which is not constructed to actually hit that, which the primary camera rays aren’t.

For a single antenna ray this doesn’t make too much sense.
The issue with this approach is that the antenna ray is shot for all launch indices which is a waste of time.
It would be better to do two optixLaunch calls instead again with some flag indicating what mode you want.
One optixLaunch for the single antenna ray with launch dimension 1x1, which returns a single float4 hit result and distance with the world coordinate on the plane (in the xyz-components) and distance (in the w-component, positive means hit!) in a single float4-sized output buffer, or a miss indicated by a negative distance.
Then copy that float4 from that output buffer into your device launch parameters (can use device-to-device copy) and then launch your standard rendering which does that epsilon environment check on the primary ray hit from camera to plane, only if the photon world position is valid, means its distance in that float4 is positive.
Same thing as before, just only calculating the antenna ray once.

2.) Version which explicitly connects the antenna ray with the camera plane.

Inside the ray generation program, render your standard image in one launch.
Then launch the antenna ray, again you’ll get at most a single photon world position on the plane, then project that 3D world position onto the camera plane by manually constructing the line from photon world position to the camera position, and determine the pixel that hits on your camera plane and set exactly that pixel inside the full resolution output image to your desired color (yellow). Doing it this way doesn’t need atomics as long as this is a single antenna ray.
This would potentially require visibility checks again if there is some other geometry between the plane and the camera.

The end-goal of this project is to simulate a radar but I’m taking one step at a time.

Typical XY Problem. Please ignore the above and take a step back and describe what the real goal is you need to solve.

What kind of radar? The sweeping things resulting in a 2D image like on ships or 3D point cloud results for car navigation?

Either would require to shoot rays from the antenna and gather hit points.
The visualization of those is a completely separate process and you might not even need raytracing for that part. This is definitely not something you do in a single launch.

Synthetic aperture radar (SAR) and inverse synthetic aperture radar (ISAR), so the sweeping things resulting in a 2D image like on ships.

I understand that this visualization part may not be entirely necessary for the final solution, however it helps me better understand what I am doing in OptiX. Perhaps there is an easier alternate way of doing it?

The real goal, without breaking it into too small parts would be the following:

Main goal: Create synthetic SAR- and ISAR images.

Sub goals:

  • Have some sort of geometry for target and antenna (2D plane is enough for the target).
  • Shoot rays from antenna towards target.
  • Be able to move antenna or target.
  • Rays bounce off target, some return to the antenna, save relevant physical information in payload. (Repeat from step 2).
  • Perform calculations on payload data.
  • Construct radar image from results.

With that in mind, would you recommend starting in a different end, or keep going with one of the approaches you explain above?

/P

If you want to experiment with that single antenna ray some more, all the methods I described would work for that specific case.

The easiest way which would generate some image with your yellow point in the current setup while shooting rays from the camera onto the plane would be to handle the antenna ray like a very focused spot light.

Means for each primary ray hit point on the plane,
1.) calculate the normalized direction vector L from antenna position to that hit point,
2.) calculate the cosine between that vector L and your normalized ray direction from the antenna N (the “spot light direction”, the “light normal”) which is simply dot(L, N).
3.) If that cosine is exactly 1.0f, the vectors L and N are identical but as explained, the probability for that is zero because that is a Dirac distribution on the antenna and that requires explicitly constructed vectors to connect.
That is why you need to check if the cosine is bigger than some value slightly smaller than 1.0f. (That difference to 1.0f is effectively a spot light spread half-angle.)
So check something like this:
float3 color = (dot(L, N) > 0.99f) ? colorYellow : colorPlane;
The smaller that spread threshold value, the bigger your yellow spot on the plane.

For the actual radar simulation it makes most sense to do one optixLaunch shooting the rays from the antenna and implement all necessary simulation results with your scene.
You need to define first what the results of your simulation need to be and then architect the programs around that.

The visualization of the results is some completely different thing. If that just needs to be some 2D image from above, you could just render the hit coordinates as points with some rasterizer in OpenGL, Vulkan, or DirectX (whatever you prefer).

Ray tracing these results would actually be more complicated. You cannot ray trace points. That would need to be some geometry you can actually hit, and even when rendering them as small spheres, then there are again discrete sampling theorem issues which would need to be considered to hit all of them independently of their distance.

I see. My understanding is that using a camera implementation is going to make it very difficult to simulate radar, thus using one of the rasterizers would be a better solution for visualization, if necessary.

I have removed the camera implementation and attempted to do the same thing as before without visualizing it with the camera - Shoot a ray from the antenna towards the plane and make it reflect off the surface.

So far, I have managed to rewrite my programs to handle the first part of generating a ray from the antenna towards the plane. I’ve also added information to the payload to see where the ray hits the plane, which gives correct results. As for the reflection part I am a bit confused, whether I should somehow generate a new ray in the reflected direction with another optixTrace() call or if there is a different easier solution? Is it possible to add payload information from the previous ray to the new reflected ray to keep the “history”?

Console output when running current version:

[DISKCACHE] Cache hit for key: 

[COMPILER]
[DISKCACHE] Cache miss for key: 

[COMPILER] Info: Pipeline parameter "params" size is 88 bytes

[COMPILER] Function properties for __closesthit__ch_ptID_0xf978d3bbb466b7b2
        register count                  :   112
        direct stack size (bytes)       :     0
        direct spills (bytes)           :     0
        continuation stack size (bytes) :     0
        continuation spills (bytes)     :     0

[DISKCACHE] Inserted module in cache with key: 
[COMPILER] Info: Module Statistics
        payload values        :          4
        attribute values      :          0
        Pipeline configuration:  (default)
Info: Properties for entry function "__closesthit__ch"
        semantic type                :             CLOSESTHIT
        trace call(s)                :                      0
        continuation callable call(s):                      0
        direct callable call(s)      :                      0
        basic block(s)               :                      1
        instruction(s)               :                     30
Info: Compiled Module Summary
        non-entry function(s):     0
        basic block(s)       :     0
        instruction(s)       :     0

[DISKCACHE] Cache hit for key: 

[COMPILER]
[COMPILER] Info: Pipeline statistics
        module(s)                            :     3
        entry function(s)                    :     3
        trace call(s)                        :     1
        continuation callable call(s)        :     0
        direct callable call(s)              :     0
        basic block(s) in entry functions    :     5
        instruction(s) in entry functions    :    90
        non-entry function(s)                :     0
        basic block(s) in non-entry functions:     0
        instruction(s) in non-entry functions:     0
        debug information                    :    no

Distance for trip: 2.000000 meters, Time taken: 6.666667e-09 seconds
Hitpoint: x = 0.000000 , y = 0.000000, z = 0.000000
[DISKCACHE] Closed database: "C:\Users\PC\AppData\Local\NVIDIA\OptixCache\optix7cache.db"
[DISKCACHE]     Cache data size: "33.7 MiB"

C:\Users\PC\Desktop\new_example_5\build\bin\Debug\new_example_5.exe (process 23832) exited with code 0.

Here are my programs if it is to any help:

rayGen:

#include <optix.h>
#include "../sutil/vec_math.h"
#include <cuda_runtime.h>
#include "../src/CommonStructs.h"

extern "C" {
    __constant__ LaunchParams params;
}

extern "C" __global__ void __raygen__rg() {
    const uint3 idx = optixGetLaunchIndex();
    float3 rayOrigin = params.antenna.position;
    float3 rayDirection = normalize(make_float3(0, 0, -1)); // Assuming the ray is downwards

    // Initialize payload
    //unsigned int payload = 0; // Will store the encoded distance on hit
    unsigned int p0, p1, p2, p3 = 0;


    optixTrace(
        params.handle, 
        rayOrigin, 
        rayDirection, 
        0.01f,    // Min intersection distance, set to a small positive number to avoid self-intersection
        1e20f,    // Max intersection distance
        0.0f,     // rayTime -- for motion blur
        OptixVisibilityMask(1),
        OPTIX_RAY_FLAG_NONE,
        0,        // SBT offset
        1,        // SBT stride
        0,        // missSBTIndex
        p0, p1, p2, p3); // Payload used to return the distance

    // Decode the distance from the payload
    float distance = __uint_as_float(p0);
    float hitpoint_x = __uint_as_float(p1);
    float hitpoint_y = __uint_as_float(p2);
    float hitpoint_z = __uint_as_float(p3);

    // Now you have the distance to the hit point (and back). You can use this to calculate the time
    // taken for the radar pulse to return. For actual radar, this would be distance / (speed of light).
    // Note: This is a simplified model and assumes the speed of light in a vacuum.

    if (distance > 0.0f) {
        // A hit occurred, now simulate the "bounce" or whatever the radar pulse does upon reflection
        // If calculating time taken for the pulse to return:
        float timeTaken = distance / 3e8f; // Considering the speed of light in m/s and round trip
        printf("Distance for round trip: %f meters, Time taken: %e seconds\n", distance, timeTaken);
        printf("Hitpoint: x = %f , y = %f, z = %f\n", hitpoint_x, hitpoint_y, hitpoint_z);
    }
}

hit:

#include <optix.h>
#include "../src/CommonStructs.h"
#include "../sutil/vec_math.h"

extern "C" {
__constant__ LaunchParams params; // Even if not used directly, it needs to be declared to fix the error.
}

// In __closesthit__ch:
extern "C" __global__ void __closesthit__ch() {
    const float3 normal = make_float3(0.0f, 1.0f, 0.0f);
    const float3 rayDir = optixGetWorldRayDirection();
    float3 hitPoint = optixGetWorldRayOrigin() + rayDir * optixGetRayTmax();

    // Reflect the ray direction around the normal
    float3 reflectDir = reflect(normalize(rayDir), normal);


    float3 antennaPos = params.antenna.position;
    float distance = length(hitPoint - antennaPos); // This ensures we're getting the spatial distance

    // Pass back the distance as the payload
    optixSetPayload_0(__float_as_uint(distance));
    optixSetPayload_1(__float_as_uint(hitPoint.x));
    optixSetPayload_2(__float_as_uint(hitPoint.y));
    optixSetPayload_3(__float_as_uint(hitPoint.z));
}

miss: (At the moment I dont think this does anything…)

#include <optix.h>
#include "../src/CommonStructs.h"

extern "C" {
__constant__ LaunchParams params; // This assumes you have a similar struct in your host code.
}

__constant__ float3 bgColor = {0.5f, 0.1f, 1.0f}; // Background color


extern "C" __global__ void __miss__ms() { 
    const float3 hitColor = make_float3(1.0f, 0.0f, 0.0f); 
    optixSetPayload_0(__float_as_uint(bgColor.x));
    optixSetPayload_1(__float_as_uint(bgColor.y));
    optixSetPayload_2(__float_as_uint(bgColor.z));
}

There are different mechanism to implement such reflection behavior.

Since the reflection direction is depending on the incoming direction of the previous ray and the material properties, it makes most sense to calculate the so called continuation ray’s origin and direction inside the closest hit program.

The continuation ray origin is obviously the hit point. The continuation ray direction depends on how the hit object should reflect such incoming rays. That goes from perfect mirror (specular) reflection (again a Dirac distribution and only exactly one possible continuation ray direction exists) to completely diffuse reflections (Lambert distribution, all directions above the hemisphere above the material’s normal have the same probability) to anywhere in between which are glossy reflections. (You would eventually need to consider transmission effects as well.)

Let’s start with the perfect mirror reflection you’ve used inside your code already. (The formula for that is super simple and also doesn’t care to which side the normal points. The continuation ray is always on the side the incoming ray hit.)

Now if you return your reflectDir to the caller as well by storing it into the per ray payload by adding another three registers, you can shoot the continuation ray inside the ray generation in a loop until nothing is hit anymore (miss program reached, or a maximum number of bounces has been reached. That limit is required to not inadvertently program an endless loop when a ray is bouncing between object infinitely.
(That “maximum path length” variable is the “depth” variable in my path tracer examples which indexes the path segments. depth == 0 is the primary ray.)

If the reflection of the antenna ray on the hit object is not perfectly specular, there would be different continuation ray directions possible from the exact same hit point.
Inside a progressive iterative Monte Carlo algorithm, that is usually done by picking two uniform random numbers in the range [0.0f, 1.0f) in each iteration, to sample the material reflection distribution function (when it’s not specular). So for each pair of uniform random numbers you would get a different reflection ray direction.

If your antenna would be a geometry inside the scene, some of these reflected rays would hit that randomly and you could detect because you would know what geometry was hit. The smaller and father away the antenna is, the less probable this gets.
If the antenna is just a point with no geometry, there would obviously be no chance to hit that with any ray and you would need to explicitly construct potential connections from hit points to that and determine if the reflection ray could come even near that antenna, which is effectively the same calculation as the focused spot light, just this time from the hit point and with the material distribution function which is the determining factor if such a constructed ray is even possible. (For purely diffuse reflections the answer is always yes. For specular the answer is no, except for the only case where the plane normal is the negative antenna ray direction.)

So there is no need to shoot the continuation ray inside the closesthit program. That would be a recursive raytracing algorithm which is more expensive and needs a lot more stack space depending on the recursion depth you allow.

What I described above is effectively an iterative path tracer because that follows exactly one continuation ray on each hit event only without any recursion.
If there is no need to bounce antenna rays between different object inside the scene multiple times, you could also shoot many continuation rays according the material properties inside the ray generation program, means sampling only one level of continuation ray directions there. For that you would need to know the material properties of plane hit point from the first antenna ray, so incoming ray direction, normal, roughness etc. would need to be stored on the ray payload inside the closest hit program to have them accessible inside the caller, your ray generation program.

Code comments:

  • If the __miss__ms program is used for the same ray type as the __closesthit__ch() program, than it’s incorrect because it writes to p0, but that should be left alone for miss events to keep the distance == 0.0f.
  • In your program setup, with rayDirection == (0, 0, -1) and plane normal = (0, 1, 0) the reflect() function will never produce a continuation ray which reaches the antenna.
  • Also they are perpendicular to each other (dot(rayDirection, normal) == 0.0f) which means that plane cannot be hit by the antenna ray at all.
  • For performance reasons, you might not want to normalize vectors which are already length 1.0f.
  • It’s also not necessary to initialize your payload register p1, p2, p3 because these are not read unless the distance inside the p0 register decides if the other payload values are read.