Hi,
I’m new to raytracing and I am building a simple renderer for a university project. I created a single OptixModule with one Pipeline and was able to trace Radiance and Shadow Rays.
However now i wanted to create a second Module and Pipeline which Creates Photon rays from my light sources into the world.
At the top of my file I created:
OptixModule optixModule = nullptr;
OptixModule photonModule = nullptr;
OptixPipeline pipeline = nullptr;
OptixPipeline photonPipeline = nullptr;
OptixProgramGroup raygenPG = nullptr;
OptixProgramGroup missPG = nullptr;
OptixProgramGroup hitgroupPG = nullptr;
OptixProgramGroup shadowMissPG = nullptr;
OptixProgramGroup shadowHitgroupPG = nullptr;
OptixProgramGroup photonRaygenPG = nullptr;
OptixProgramGroup photonMissPG = nullptr;
OptixProgramGroup photonHitgroupPG = nullptr
And in my init function I do:
CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&p_d_params), sizeof(PhotonLaunchParams)));
CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&d_params), sizeof(LaunchParams)));
and call:
void initOptixPipeline() {
// 1. Compile & Link Options
OptixModuleCompileOptions moduleCompileOptions = {};
moduleCompileOptions.maxRegisterCount = OPTIX_COMPILE_DEFAULT_MAX_REGISTER_COUNT;
moduleCompileOptions.optLevel = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
moduleCompileOptions.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_DEFAULT;
OptixPipelineCompileOptions pipelineCompileOptions = {};
pipelineCompileOptions.usesMotionBlur = false;
pipelineCompileOptions.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
pipelineCompileOptions.numPayloadValues = 3; // p0, p1, p2 for rgb
pipelineCompileOptions.numAttributeValues = 2; // Barycentrics
pipelineCompileOptions.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
pipelineCompileOptions.pipelineLaunchParamsVariableName = "params";
OptixPipelineCompileOptions photonPipelineCompileOptions = {};
photonPipelineCompileOptions.usesMotionBlur = false;
photonPipelineCompileOptions.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
photonPipelineCompileOptions.numPayloadValues = 1; // just p0
photonPipelineCompileOptions.numAttributeValues = 2;
photonPipelineCompileOptions.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
photonPipelineCompileOptions.pipelineLaunchParamsVariableName = "pparams";
// 2. PTX to Module
std::string ptxCode = readPtxFile("out/build/x64-Debug/raytrace.ptx");
char log[2048];
size_t sizeof_log = sizeof(log);
OPTIX_CHECK(optixModuleCreate(
optixContext,
&moduleCompileOptions,
&pipelineCompileOptions,
ptxCode.c_str(),
ptxCode.size(),
log, &sizeof_log,
&optixModule
));
std::string photonPtx = readPtxFile("out/build/x64-Debug/photons.ptx");
OPTIX_CHECK(optixModuleCreate(
optixContext,
&moduleCompileOptions,
&photonPipelineCompileOptions,
photonPtx.c_str(),
photonPtx.size(),
log,
&sizeof_log,
&photonModule
));
// 3. Program Groups
OptixProgramGroupOptions pgOptions = {};
// Raygen PG
sizeof_log = sizeof(log);
OptixProgramGroupDesc raygenDesc = {};
raygenDesc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
raygenDesc.raygen.module = optixModule;
raygenDesc.raygen.entryFunctionName = "__raygen__render";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &raygenDesc, 1, &pgOptions, log, &sizeof_log, &raygenPG));
// Miss PG
sizeof_log = sizeof(log);
OptixProgramGroupDesc missDesc = {};
missDesc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
missDesc.miss.module = optixModule;
missDesc.miss.entryFunctionName = "__miss__radiance";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &missDesc, 1, &pgOptions, log, &sizeof_log, &missPG));
// Closest Hit
sizeof_log = sizeof(log);
OptixProgramGroupDesc hitDesc = {};
hitDesc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
hitDesc.hitgroup.moduleCH = optixModule;
hitDesc.hitgroup.entryFunctionNameCH = "__closesthit__radiance";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &hitDesc, 1, &pgOptions, log, &sizeof_log, &hitgroupPG));
sizeof_log = sizeof(log);
OptixProgramGroupDesc shadowMissDesc = {};
shadowMissDesc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
shadowMissDesc.miss.module = optixModule;
shadowMissDesc.miss.entryFunctionName = "__miss__shadow";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &shadowMissDesc, 1, &pgOptions, log, &sizeof_log, &shadowMissPG));
sizeof_log = sizeof(log);
OptixProgramGroupDesc shadowHitDesc = {};
shadowHitDesc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
shadowHitDesc.hitgroup.moduleAH = optixModule;
shadowHitDesc.hitgroup.entryFunctionNameAH = "__anyhit__shadow";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &shadowHitDesc, 1, &pgOptions, log, &sizeof_log, &shadowHitgroupPG));
// Photon Raygen PG
sizeof_log = sizeof(log);
OptixProgramGroupDesc photonRaygenDesc = {};
photonRaygenDesc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
photonRaygenDesc.raygen.module = photonModule;
photonRaygenDesc.raygen.entryFunctionName = "__raygen__photon";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &photonRaygenDesc, 1, &pgOptions, log, &sizeof_log, &photonRaygenPG));
// Photon Miss PG
sizeof_log = sizeof(log);
OptixProgramGroupDesc photonMissDesc = {};
photonMissDesc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
photonMissDesc.miss.module = photonModule;
photonMissDesc.miss.entryFunctionName = "__miss__photon";
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &photonMissDesc, 1, &pgOptions, log, &sizeof_log, &photonMissPG));
// Photon Closest Hit
sizeof_log = sizeof(log);
OptixProgramGroupDesc photonHitDesc = {};
photonHitDesc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
photonHitDesc.hitgroup.moduleCH = photonModule;
photonHitDesc.hitgroup.entryFunctionNameCH = "__closesthit__photon";
// Wir brauchen keinen Any-Hit oder Intersection-Shader für einfache Dreiecke
OPTIX_CHECK(optixProgramGroupCreate(optixContext, &photonHitDesc, 1, &pgOptions, log, &sizeof_log, &photonHitgroupPG));
// 4. Pipelines verknüpfen
sizeof_log = sizeof(log);
OptixProgramGroup programGroups[] = { raygenPG, missPG, shadowMissPG, hitgroupPG, shadowHitgroupPG };
OptixProgramGroup photonProgramGroups[] = { photonRaygenPG, photonMissPG, photonHitgroupPG };
OptixPipelineLinkOptions pipelineLinkOptions = {};
pipelineLinkOptions.maxTraceDepth = 3; // Wie oft ein Strahl reflektieren darf
OPTIX_CHECK(optixPipelineCreate(
optixContext,
&pipelineCompileOptions,
&pipelineLinkOptions,
programGroups, sizeof(programGroups) / sizeof(programGroups[0]),
log, &sizeof_log,
&pipeline
));
OPTIX_CHECK(optixPipelineCreate(
optixContext,
&photonPipelineCompileOptions,
&pipelineLinkOptions,
photonProgramGroups, sizeof(photonProgramGroups) / sizeof(photonProgramGroups[0]),
log, &sizeof_log,
&photonPipeline
));
OptixStackSizes stackSizes = {};
for (OptixProgramGroup programGroup : programGroups) {
OPTIX_CHECK(optixUtilAccumulateStackSizes(programGroup, &stackSizes, pipeline));
}
unsigned int directCallableStackSizeFromTraversal = 0;
unsigned int directCallableStackSizeFromState = 0;
unsigned int continuationStackSize = 0;
OPTIX_CHECK(optixUtilComputeStackSizes(
&stackSizes,
pipelineLinkOptions.maxTraceDepth,
0,
0,
&directCallableStackSizeFromTraversal,
&directCallableStackSizeFromState,
&continuationStackSize
));
OPTIX_CHECK(optixPipelineSetStackSize(
pipeline,
directCallableStackSizeFromTraversal,
directCallableStackSizeFromState,
continuationStackSize,
2
));
}
After that in my render loop I call
// Photon Pipeline Parameter
PhotonLaunchParams photonParams = {};
photonParams.traversable = iasHandle;
photonParams.lightCount = static_cast<int>(lights.size());
for (int i = 0; i < photonParams.lightCount; ++i) {
photonParams.lights[i] = lights[i];
}
photonParams.photonCount = 1000000; // Anzahl der Photonen, die wir abfeuern wollen
photonParams.lvData = reinterpret_cast<float*>(d_volumeData);
photonParams.lvDim = make_uint3(32, 32, 32);
photonParams.lvMin = make_float3(-18.0f, -7.0f, -18.0f);
photonParams.lvMax = make_float3(18.0f, 10.0f, 18.0f);
CUDA_CHECK(cudaMemcpy(reinterpret_cast<void*>(p_d_params), &photonParams, sizeof(PhotonLaunchParams), cudaMemcpyHostToDevice));
// Photon aufrufen
OPTIX_CHECK(optixLaunch(
photonPipeline,
0, // CUDA stream
p_d_params, // Pointer zu unseren gemappten LaunchParams
sizeof(PhotonLaunchParams),
&sbt, // Shader Binding Table
photonParams.photonCount, // Threads in X
photonParams.lightCount, // Threads in Y
1 // Threads in Z
));
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
before I call the OptixLaunch for my raytracing pipeline.
Somehow it gives me
cudaDeviceSynchronize() : an illegal memory access was encountered
right there and if I start it with the compute sanitizer it says this:
========= Program hit cudaErrorUnknown (error 999) due to “unknown error” on CUDA API call to cudaDeviceSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: cuProfilerStop [0x7ffed6c610d2] in nvcuda64.dll
========= Host Frame: _cudaGetProcAddress [0x7ffef4cb195f] in nvcudart_hybrid64.dll
========= Host Frame: MainApp::render in main.cpp:1342 [0x12eb2d] in cgintro_animation.exe
========= Host Frame: App::run in app.cpp:238 [0x139c70] in cgintro_animation.exe
========= Host Frame: main in main.cpp:1426 [0x525c6] in cgintro_animation.execudaDeviceSynchronize() : unknown error
So I can’t really figure out why this is happening and any attempts to fix it didn’t succeed or just created more errors.
I would really appreciate help :)