Trimesh.ray equivalent in Optix

Gaming and Visualization Technologies Visualization OptiX
1

Hello,

I’m new starting to wok with Optix.

I’m trying to reproduce the same output of as

mesh_rot = trimesh.Trimesh(vertices=origins_rot, faces=mesh_for_shadow.faces,process=False)

-- >shadowed = mesh_rot.ray.intersects_any( ray_origins=origins, ray_directions=dirs)

this is my code:

#include "optix_shadow.h"
#include <cuda_runtime.h>
#include <cuda.h>

#include <optix.h>
#include <optix_stubs.h>
#include <optix_function_table_definition.h>
#include <vector>
#include <iostream>
#include <fstream>
#include <cstring>

// ─────────────────────────────────────────────────────────────
// ERROR CHECKING
// ─────────────────────────────────────────────────────────────

#define OPTIX_CHECK(call)                                             \
    do {                                                              \
        OptixResult res = call;                                       \
        if (res != OPTIX_SUCCESS) {                                   \
            std::cerr << "OptiX Error: " << res << std::endl;         \
            exit(1);                                                  \
        }                                                             \
    } while(0)


#define CUDA_CHECK(call)                                      \
    do {                                                      \
        cudaError_t err = call;                               \
        if (err != cudaSuccess) {                             \
            std::cerr << "CUDA error: "                        \
                      << cudaGetErrorString(err)              \
                      << " at " << __FILE__ << ":" << __LINE__\
                      << std::endl;                           \
            exit(1);                                          \
        }                                                     \
    } while(0)
// ─────────────────────────────────────────────────────────────
// GLOBAL STATE
// ─────────────────────────────────────────────────────────────

static OptixDeviceContext context = nullptr;
static OptixTraversableHandle gas_handle = 0;
static CUdeviceptr d_gas_output = 0;

static OptixPipeline pipeline = nullptr;
static OptixModule module = nullptr;

static OptixProgramGroup raygen_prog_group = nullptr;
static OptixProgramGroup miss_prog_group = nullptr;
static OptixProgramGroup hitgroup_prog_group = nullptr;

static OptixShaderBindingTable sbt = {};

static CUstream stream = 0;

static Params params;
static CUdeviceptr d_params;

// ─────────────────────────────────────────────────────────────
// MODULE
// ─────────────────────────────────────────────────────────────

void create_module(const char* ptx_code)
{
    OptixModuleCompileOptions mopt = {};
    mopt.maxRegisterCount = OPTIX_COMPILE_DEFAULT_MAX_REGISTER_COUNT;
    mopt.optLevel = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
    mopt.debugLevel =  OPTIX_COMPILE_DEBUG_LEVEL_MODERATE;

    OptixPipelineCompileOptions popt = {};
    popt.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;
    popt.numPayloadValues = 1;
    popt.numAttributeValues = 2;   // ⚠️ REQUIRED for triangles
    popt.pipelineLaunchParamsVariableName = "params";

    char log[2048];
    size_t log_size = sizeof(log);

    OPTIX_CHECK(
        optixModuleCreate(
            context,
            &mopt,
            &popt,
            ptx_code,
            strlen(ptx_code),
            log,
            &log_size,
            &module
        )
    );

    if (log_size > 1)
        std::cout << "MODULE LOG:\n" << log << std::endl;
}

// ─────────────────────────────────────────────────────────────
// PROGRAM GROUPS
// ─────────────────────────────────────────────────────────────

void create_program_groups()
{
    OptixProgramGroupOptions opts = {};

    // ── Raygen ─────────────────────────────
    OptixProgramGroupDesc raygen_desc = {};
    raygen_desc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
    raygen_desc.raygen.module = module;
    raygen_desc.raygen.entryFunctionName = "__raygen__rg";

    OPTIX_CHECK(optixProgramGroupCreate(
        context, &raygen_desc, 1, &opts, nullptr, nullptr, &raygen_prog_group));

    // ── Miss ───────────────────────────────
    OptixProgramGroupDesc miss_desc = {};
    miss_desc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
    miss_desc.miss.module = module;
    miss_desc.miss.entryFunctionName = "__miss__ms";

    OPTIX_CHECK(optixProgramGroupCreate(
        context, &miss_desc, 1, &opts, nullptr, nullptr, &miss_prog_group));

    // ── Hitgroup ───────────────────────────
    OptixProgramGroupDesc hit_desc = {};
    hit_desc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
    
    //closest hit
    hit_desc.hitgroup.moduleCH = nullptr;
    hit_desc.hitgroup.entryFunctionNameCH = nullptr;

    //any hit
    hit_desc.hitgroup.moduleAH = module;
    hit_desc.hitgroup.entryFunctionNameAH = "__anyhit__ah";  // must be non-null to enable any-hit shader
    
    //intersection
    hit_desc.hitgroup.moduleIS = nullptr;
    hit_desc.hitgroup.entryFunctionNameIS = nullptr;

    std::cerr << "Creating hitgroup program group with: ";

    if (hit_desc.hitgroup.entryFunctionNameCH)
        std::cerr << hit_desc.hitgroup.entryFunctionNameCH;
    else
        std::cerr << "NULL";
        
    std::cerr << " (CH) and ";
        
    if (hit_desc.hitgroup.entryFunctionNameAH)
        std::cerr << hit_desc.hitgroup.entryFunctionNameAH;
    else
        std::cerr << "NULL";
        
    std::cerr << " (AH)" << std::endl;
    OPTIX_CHECK(optixProgramGroupCreate(
        context, &hit_desc, 1, &opts, nullptr, nullptr, &hitgroup_prog_group));
}

// ─────────────────────────────────────────────────────────────
// PIPELINE
// ─────────────────────────────────────────────────────────────

void create_pipeline()
{
    OptixPipelineCompileOptions popt = {};
    popt.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;
    popt.numPayloadValues = 1;
    popt.numAttributeValues = 2;
    popt.pipelineLaunchParamsVariableName = "params";

    OptixPipelineLinkOptions link = {};
    link.maxTraceDepth = 1;

    OptixProgramGroup groups[] = {
        raygen_prog_group,
        miss_prog_group,
        hitgroup_prog_group
    };

    OPTIX_CHECK(optixPipelineCreate(
        context,
        &popt,
        &link,
        groups,
        3,
        nullptr,
        nullptr,
        &pipeline));

    // CRITICAL
    OPTIX_CHECK(optixPipelineSetStackSize(
        pipeline,
        2048,
        2048,
        2048,
        1
    ));
}

// ─────────────────────────────────────────────────────────────
// SBT
// ─────────────────────────────────────────────────────────────

struct RayGenRecord { char header[OPTIX_SBT_RECORD_HEADER_SIZE]; };
struct MissRecord   { char header[OPTIX_SBT_RECORD_HEADER_SIZE]; };
struct HitRecord    { char header[OPTIX_SBT_RECORD_HEADER_SIZE]; };

void create_sbt()
{
    // Raygen
    RayGenRecord rg;
    optixSbtRecordPackHeader(raygen_prog_group, &rg);

    cudaMalloc((void**)&sbt.raygenRecord, sizeof(rg));
    cudaMemcpy((void*)sbt.raygenRecord, &rg, sizeof(rg), cudaMemcpyHostToDevice);

    // Miss
    MissRecord ms;
    optixSbtRecordPackHeader(miss_prog_group, &ms);

    cudaMalloc((void**)&sbt.missRecordBase, sizeof(ms));
    cudaMemcpy((void*)sbt.missRecordBase, &ms, sizeof(ms), cudaMemcpyHostToDevice);

    sbt.missRecordStrideInBytes = sizeof(MissRecord);
    sbt.missRecordCount = 1;

    // Hit
    HitRecord hg;
    optixSbtRecordPackHeader(hitgroup_prog_group, &hg);

    cudaMalloc((void**)&sbt.hitgroupRecordBase, sizeof(hg));
    cudaMemcpy((void*)sbt.hitgroupRecordBase, &hg, sizeof(hg), cudaMemcpyHostToDevice);

    sbt.hitgroupRecordStrideInBytes = sizeof(HitRecord);
    sbt.hitgroupRecordCount = 1;
}

// ─────────────────────────────────────────────────────────────
// INIT
// ─────────────────────────────────────────────────────────────

void init_optix(const float* vertices, int num_vertices,
                const int32_t* faces, int num_faces)
{
    std::cerr << "Initializing OptiX...\n";

    CUDA_CHECK(cudaFree(0));  // init CUDA
    OPTIX_CHECK(optixInit());

    OptixDeviceContextOptions options = {};
    options.logCallbackLevel = 4;

    OPTIX_CHECK(optixDeviceContextCreate(0, &options, &context));

    std::cerr << "Context created.\n";

    // =========================
    // Upload geometry
    // =========================
    CUdeviceptr d_vertices = 0;
    CUdeviceptr d_indices  = 0;

    size_t v_bytes = num_vertices * 3 * sizeof(float);
    size_t i_bytes = num_faces * 3 * sizeof(uint32_t);

    std::cerr << "Allocating vertices: " << v_bytes << " bytes\n";
    CUDA_CHECK(cudaMalloc((void**)&d_vertices, v_bytes));

    std::cerr << "Copying vertices\n";
    CUDA_CHECK(cudaMemcpy((void*)d_vertices, vertices,
                          v_bytes, cudaMemcpyHostToDevice));

    std::cerr << "Allocating indices: " << i_bytes << " bytes\n";
    CUDA_CHECK(cudaMalloc((void**)&d_indices, i_bytes));

    std::cerr << "Copying indices\n";
    CUDA_CHECK(cudaMemcpy((void*)d_indices, faces,
                          i_bytes, cudaMemcpyHostToDevice));

    CUDA_CHECK(cudaDeviceSynchronize());

    std::cerr << "Geometry uploaded.\n";

    // =========================
    // Build GAS
    // =========================
    uint32_t triangle_input_flags[1] = { OPTIX_GEOMETRY_FLAG_NONE };
    OptixBuildInput build_input = {};
    build_input.type = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;

    CUdeviceptr vb[] = { d_vertices };

    build_input.triangleArray.vertexBuffers = vb;
    build_input.triangleArray.numVertices = num_vertices;
    build_input.triangleArray.vertexFormat = OPTIX_VERTEX_FORMAT_FLOAT3;
    build_input.triangleArray.vertexStrideInBytes = sizeof(float) * 3;

    build_input.triangleArray.indexBuffer = d_indices;
    build_input.triangleArray.numIndexTriplets = num_faces;
    build_input.triangleArray.indexFormat = OPTIX_INDICES_FORMAT_UNSIGNED_INT3;
    build_input.triangleArray.indexStrideInBytes = sizeof(uint32_t) * 3;

    build_input.triangleArray.flags = triangle_input_flags;
    build_input.triangleArray.numSbtRecords = 1;

    OptixAccelBuildOptions opts = {};
    opts.operation = OPTIX_BUILD_OPERATION_BUILD;
    
    OptixAccelBufferSizes sizes;

    OPTIX_CHECK(optixAccelComputeMemoryUsage(
        context, &opts, &build_input, 1, &sizes));

    std::cerr << "Temp size: " << sizes.tempSizeInBytes << "\n";
    std::cerr << "Output size: " << sizes.outputSizeInBytes << "\n";

    CUdeviceptr d_temp = 0;
    CUDA_CHECK(cudaMalloc((void**)&d_temp, sizes.tempSizeInBytes));

    CUDA_CHECK(cudaMalloc((void**)&d_gas_output, sizes.outputSizeInBytes));

    OPTIX_CHECK(optixAccelBuild(
        context, 0,
        &opts,
        &build_input, 1,
        d_temp, sizes.tempSizeInBytes,
        d_gas_output, sizes.outputSizeInBytes,
        &gas_handle,
        nullptr, 0
    ));

    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaFree((void*)d_temp));

    std::cerr << "GAS built successfully.\n";

    // =========================
    // Stream + params
    // =========================
    CUDA_CHECK(cudaStreamCreate(&stream));
    CUDA_CHECK(cudaMalloc((void**)&d_params, sizeof(Params)));

    // =========================
    // Load PTX
    // =========================
    std::ifstream file("build/optix_shadow.ptx");
    if (!file) {
        std::cerr << "Failed to open PTX file\n";
        exit(1);
    }

    std::string ptx((std::istreambuf_iterator<char>(file)),
                     std::istreambuf_iterator<char>());

    std::cerr << "Creating module...\n";
    create_module(ptx.c_str());

    std::cerr << "Creating program groups...\n";
    create_program_groups();

    std::cerr << "Creating pipeline...\n";
    create_pipeline();

    std::cerr << "Creating SBT...\n";
    create_sbt();

    std::cerr << "OptiX fully initialized.\n";
}

// ─────────────────────────────────────────────────────────────
// QUERY
// ─────────────────────────────────────────────────────────────

std::vector<bool> shadow_rays(const float* origins,
                              const float* dirs,
                              int N)
{
    CUdeviceptr d_origins, d_dirs, d_results;

    cudaMalloc((void**)&d_origins, N*sizeof(float3));
    cudaMalloc((void**)&d_dirs,    N*sizeof(float3));
    cudaMalloc((void**)&d_results, N*sizeof(int));

    cudaMemcpy((void*)d_origins, origins, N*sizeof(float3), cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_dirs, dirs, N*sizeof(float3), cudaMemcpyHostToDevice);

    params.origins = (float3*)d_origins;
    params.dirs    = (float3*)d_dirs;
    params.results = (int*)d_results;
    params.N       = N;
    params.handle  = gas_handle;

    cudaMemcpy((void*)d_params, &params, sizeof(Params), cudaMemcpyHostToDevice);

    OPTIX_CHECK(optixLaunch(
        pipeline,
        stream,
        d_params,
        sizeof(Params),
        &sbt,
        N, 1, 1
    ));

    cudaDeviceSynchronize();

    std::vector<int> tmp(N);
    cudaMemcpy(tmp.data(), (void*)d_results, N*sizeof(int), cudaMemcpyDeviceToHost);

    std::vector<bool> out(N);
    for (int i = 0; i < N; ++i)
        out[i] = tmp[i] != 0;

    cudaFree((void*)d_origins);
    cudaFree((void*)d_dirs);
    cudaFree((void*)d_results);

    return out;
}

#include <optix.h>
#include <optix_device.h>
#include <cuda_runtime.h>
#include <cuda.h>

struct Params {
    float3* origins;
    float3* dirs;
    int* results;
    int N;
    OptixTraversableHandle handle;
};

extern "C" {
__constant__ Params params;
}

// ─────────────────────────────────────────────
// RAYGEN
// ─────────────────────────────────────────────

extern "C" __global__ void __raygen__rg()
{
    const int idx = optixGetLaunchIndex().x;
    if (idx >= params.N) return;

    float3 origin = params.origins[idx];
    float3 dir    = params.dirs[idx];

    unsigned int payload = 0;  // default = no hit

    optixTrace(
        params.handle,
        origin,
        dir,
        1e-3f,
        1e16f,
        0.0f,
        OptixVisibilityMask(255),
        OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT,
        0, 1, 0,
        payload
    );

    params.results[idx] = (payload != 0);
}

extern "C" __global__ void __anyhit__ah()
{
    optixSetPayload_0(1);   // mark hit
    optixTerminateRay();    // stop immediately
}
// ─────────────────────────────────────────────
// MISS
// ─────────────────────────────────────────────

extern "C" __global__ void __miss__ms()
{
    // payload stays 0 → no hit
}

// ─────────────────────────────────────────────
// HIT
// ─────────────────────────────────────────────

extern "C" __global__ void __closesthit__ch()
{
    optixSetPayload_0(1);  // mark hit
}
2

Hi @medamimessaoud,

Your code looks like a shadow feeler test. Did you compile it ? Does it work?

I’ve slightly modified your code, I use:

  • Windows
  • OptiX 9+
  • VS 2022
  • CUDA 13.1

please change CMakeLists.txt as needed. The modified demo just prints the result to the console.

For learning and mastering OptiX there are a few good repos, I recommend these:

  1. The OptiX SDK samples that you should have from the installer (note that a headers-only repo is GitHub - NVIDIA/optix-dev: OptiX SDK headers, everything needed to build & run OptiX applications. SDK samples not included. · GitHub).
  2. The Advanced Samples aka optix_apps: GitHub - NVIDIA/OptiX_Apps: Advanced Samples for the NVIDIA OptiX 7 Ray Tracing SDK · GitHub

and of course the forums.

By the way, the fastest way to check for shadows (when cutout opacity or other features aren’t needed), is explained in droettger’s answer here, with code snippet: Anyhit program as shadow ray with optix 7.2 - Visualization / OptiX - NVIDIA Developer Forums

optix_shadows.zip (6.8 KB)