Need help understanding why optixLaunch is failing

Gaming and Visualization Technologies Visualization OptiX
1

I have been experimenting with the optixCustomPrimitive example from the optix SDK 8.0.0. My goal is to load a volume into memory and to render it using optix passing the bounding box of the volume as a custom primitive. To achieve this I have tinkered with the code and read through the optix documentation but my application always fails on optixLaunch with an “invalid value” error. I am having difficulty debugging this error as it is unclear what the reason is. I am attaching the optix calls and configuration that I have used and my shader.cu file. I am sorry if this is a very foolish question but are c++ namespaces not allowed as part of the launchParams since the params are placed in a extern C block in the shaders.cu?

// sbt and params structs


#include <optix.h>
#include <cuda_runtime.h>
#include "common.h"
#include <stdint.h>

namespace core {
	namespace optix {
		struct Light {
			float3 position;
			float3 color;
			float intensity;
			char pad[4];
		};
		struct Volume {
			unsigned char* data;
			float3 position;
			int3 dimensions;
			float3 spacing;
			float scale;
			common::VolumeTypes type;
			char pad[12];
		};
		struct Camera {
			float3 position, view, up;
			int32_t imageHeight, imageWidth;
			float fov;
			float IMAGE_PLANE_DISTANCE;
			char pad[12];
		};

		struct OptixParams {
			unsigned char* canvas;
			OptixTraversableHandle handle;
			Light* lights;
			Camera activeCamera;
			Volume volume;
		};
		struct RayGenSbtRecord { 
			__align__(OPTIX_SBT_RECORD_ALIGNMENT)
				char header[OPTIX_SBT_RECORD_HEADER_SIZE];
		};
		struct MissSbtRecord { 
			__align__(OPTIX_SBT_RECORD_ALIGNMENT)
				char header[OPTIX_SBT_RECORD_HEADER_SIZE];
		};
		struct HitGroupSbtRecord { 
			__align__(OPTIX_SBT_RECORD_ALIGNMENT)
				char header[OPTIX_SBT_RECORD_HEADER_SIZE];
		};
	};
};

//----------------------------------------------------------------
// shaders.cu

#include <optix.h>
#include <cuda_runtime.h>
#include "optix-params.h"

extern "C" {
	__constant__ core::optix::OptixParams params;
}

extern "C"
__global__ void __raygen__generateTraceRays() {
}

extern "C"
__global__ void __miss__sampleEnvironmentMap() {
}

extern "C"
__global__ void __closesthit__renderVolume() {
}

//------------------------------------------------------------------
// Launch setup

// init

m_cudaContext = 0;
optixInit();
cudaFree(0);
optixDeviceContextCreate(m_cudaContext, nullptr, &m_optixContext);

// module create 

module_compile_options.optLevel = OPTIX_COMPILE_OPTIMIZATION_LEVEL_0;
module_compile_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_FULL;
pipeline_compile_options.usesMotionBlur = false;
pipeline_compile_options.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;
pipeline_compile_options.numPayloadValues = 2;
pipeline_compile_options.numAttributeValues = 2;
pipeline_compile_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;  // TODO: should be OPTIX_EXCEPTION_FLAG_STACK_OVERFLOW;
pipeline_compile_options.pipelineLaunchParamsVariableName = "params";
optixModuleCreate(
	m_optixContext,
	&module_compile_options,
	&pipeline_compile_options,
	moduleOptixIR.c_str(),
	moduleOptixIR.size(),
	logString, &logStringSize,
	&module
);

// raygen program

raygen_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
raygen_prog_group_desc.raygen.module = module;
raygen_prog_group_desc.raygen.entryFunctionName = "__raygen__generateTraceRays";

optixProgramGroupCreate(
	m_optixContext,
	&raygen_prog_group_desc,
	1,   // num program groups
	&program_group_options,
	logString, &logStringSize,
	&m_raygen_prog_group
);

// miss program

miss_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
miss_prog_group_desc.miss.module = module;
miss_prog_group_desc.miss.entryFunctionName = "__miss__sampleEnvironmentMap";

optixProgramGroupCreate(
	m_optixContext,
	&miss_prog_group_desc,
	1,   // num program groups
	&program_group_options,
	logString, &logStringSize,
	&m_miss_prog_group
);

// hitgroup program

hitgroup_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
hitgroup_prog_group_desc.hitgroup.moduleCH = module;
hitgroup_prog_group_desc.hitgroup.entryFunctionNameCH = "__closesthit__renderVolume";
hitgroup_prog_group_desc.hitgroup.moduleAH = nullptr;
hitgroup_prog_group_desc.hitgroup.entryFunctionNameAH = nullptr;
hitgroup_prog_group_desc.hitgroup.moduleIS = nullptr;//intersectionModule;
hitgroup_prog_group_desc.hitgroup.entryFunctionNameIS = nullptr;

optixProgramGroupCreate(
	m_optixContext,
	&hitgroup_prog_group_desc,
	1,   // num program groups
	&program_group_options,
	logString, &logStringSize,
	&m_hitgroup_prog_group
);

const uint32_t    max_trace_depth = 1;
OptixProgramGroup program_groups[] = { m_raygen_prog_group, m_miss_prog_group, m_hitgroup_prog_group };

// pipeline create 

OptixPipelineLinkOptions pipeline_link_options = {};
pipeline_link_options.maxTraceDepth = max_trace_depth;

optixPipelineCreate(
	m_optixContext,
	&pipeline_compile_options,
	&pipeline_link_options,
	program_groups,
	sizeof(program_groups) / sizeof(program_groups[0]),
	logString, &logStringSize,
	&m_pipeline
);

// set stack size 

OptixStackSizes stack_sizes = {};
for (auto& prog_group : program_groups)
{
	util::OPTIX_CHECK(
		optixUtilAccumulateStackSizes(prog_group, &stack_sizes, m_pipeline),
		log, std::string(__FILE__) + " : " + std::to_string(__LINE__) + " " + "optixUtilAccumulateStackSizes"
	);
}

uint32_t direct_callable_stack_size_from_traversal;
uint32_t direct_callable_stack_size_from_state;
uint32_t continuation_stack_size;

optixUtilComputeStackSizes(
	&stack_sizes, max_trace_depth,
	0,  // maxCCDepth
	0,  // maxDCDEpth
	&direct_callable_stack_size_from_traversal,
	&direct_callable_stack_size_from_state, &continuation_stack_size
);
			
			
optixPipelineSetStackSize(
	m_pipeline, direct_callable_stack_size_from_traversal,
	direct_callable_stack_size_from_state, continuation_stack_size,
	1  // maxTraversableDepth
);

// SBT setup

CUdeviceptr  raygen_record;
const size_t raygen_record_size = sizeof(RayGenSbtRecord);

cudaMalloc(reinterpret_cast<void**>(&raygen_record), raygen_record_size),

RayGenSbtRecord rg_sbt;

optixSbtRecordPackHeader(m_raygen_prog_group, &rg_sbt),


cudaMemcpy(
	reinterpret_cast<void*>(raygen_record),
	&rg_sbt,
	raygen_record_size,
	cudaMemcpyHostToDevice
);

CUdeviceptr miss_record;
size_t      miss_record_size = sizeof(MissSbtRecord);
		
cudaMalloc(reinterpret_cast<void**>(&miss_record), miss_record_size),
		
MissSbtRecord ms_sbt;

optixSbtRecordPackHeader(m_miss_prog_group, &ms_sbt),
cudaMemcpy(
	reinterpret_cast<void*>(miss_record),
	&ms_sbt,
	miss_record_size,
	cudaMemcpyHostToDevice
);

CUdeviceptr hitgroup_record;
size_t      hitgroup_record_size = sizeof(HitGroupSbtRecord);

cudaMalloc(reinterpret_cast<void**>(&hitgroup_record), hitgroup_record_size),
			
HitGroupSbtRecord hg_sbt;
		
optixSbtRecordPackHeader(m_hitgroup_prog_group, &hg_sbt),
		
cudaMemcpy(
	reinterpret_cast<void*>(hitgroup_record),
	&hg_sbt,
	hitgroup_record_size,
	cudaMemcpyHostToDevice
);


m_sbt.raygenRecord = raygen_record;
m_sbt.missRecordBase = miss_record;
m_sbt.missRecordStrideInBytes = sizeof(MissSbtRecord);
m_sbt.missRecordCount = 1;
m_sbt.hitgroupRecordBase = hitgroup_record;
m_sbt.hitgroupRecordStrideInBytes = sizeof(HitGroupSbtRecord);
m_sbt.hitgroupRecordCount = 1;

// optix launch

optixLaunch(
	m_pipeline, 0, m_d_params, sizeof(OptixParams), &m_sbt,
	m_params.activeCamera.imageWidth, m_params.activeCamera.imageHeight, /*depth=*/1
);
2

The first thing which looks different is the order of CUDA and OptiX initializations.
All OptiX SDK examples initialize CUDA first:

    // Initialize CUDA
    CUDA_CHECK( cudaFree( 0 ) );

    OptixDeviceContext context;
    CUcontext          cu_ctx = 0;  // zero means take the current context
    OPTIX_CHECK( optixInit() );

Did you enable the OptiX validation mode while debugging your issue to maybe get more information from OptiX?
https://raytracing-docs.nvidia.com/optix8/guide/index.html#context#validation-mode

I would recommend starting with working code from one of the simple OptiX SDK examples and then change it to your needs.

If that doesn’t solve it, please provide a complete and minimal reproducer project instead of some source code excerpts.

Namespaces for the OptiX launch parameters should work. The OptiX SDK example optixWhitted is doing that as well:

extern "C" {
__constant__ whitted::LaunchParams params;
}
3

I would recommend starting with working code from one of the simple OptiX SDK examples and then change it to your needs.

I have used the optixCustomPrimitive example from the SDK as an example for this code. I am attaching a minimal example of the error here.

example.zip (3.2 KB)

4

After adding this:

#include <algorithm>
#include <iostream>
#include <mutex>

static std::mutex g_mutexLogger;

static void callbackLogger(unsigned int level, const char* tag, const char* message, void* cbdata)
{
  std::lock_guard<std::mutex> lock(g_mutexLogger);

  std::cerr << tag  << " (" << level << "): " << ((message) ? message : "(no message)") << '\n';
}

and this

  OptixDeviceContextOptions options = {};

  options.logCallbackFunction = &callbackLogger;
  options.logCallbackData     = nullptr;
  options.logCallbackLevel    = 4;
  options.validationMode      = OPTIX_DEVICE_CONTEXT_VALIDATION_MODE_ALL;

the OptiX validation output says:

ERROR (2): "sbt->exceptionRecord" points to a memory area which is not correctly aligned
OPTIX_ERROR_INVALID_VALUE : Invalid value

which is fixed by default-initializing the sbt variable:

	OptixShaderBindingTable sbt = {};

Please always default initialize all OptiX structures like this throughout all your code.
Sometimes a new OptiX SDK version adds fields to some API structures and usually the default value is 0 and that initialization makes sure that is always the case when switching OptiX SDKs, or like in this case.

I would also not handle OptiX IR code as std::string. It contains null bytes.
See this thread for pure binary data handling example code: https://forums.developer.nvidia.com/t/embedding-optix-ir/273199/2