rtTrace only results in miss program invocation

Gaming and Visualization Technologies Visualization OptiX
1

Hi everybody!

I’m working on a project in which I’m trying to use Optix for ray tracing. I’ve based my code off of several samples, in particular the ones that use triangle meshes for the geometry, as well as the programming guide.

So, as a first step, I’m currently trying to generate rays from spherical area lights and to trace them until their first intersection with the scene geometry. In trying to achieve this I find myself running into a problem: the intersection, bounding box, any hit and closest hit programs are never invoked. The only programs that are actually executed are the ray generation and miss programs (I was under the impression that the intersection program is what determines whether a ray hits or misses geometry, so I’m not sure what’s going on here really).

The ray generation program currently generates rays by uniformly sampling the surface of a sphere light. All of the other programs are empty, save for rtPrintf calls to let me know whether or not the programs are actually invoked. Seeing as the intersection/bounding box programs are never called I’m thinking I’ve set them up wrongly. But, frankly, I’m not sure where to look anymore.

So my question to you: does anyone see anything glaringly wrong here (I’ll add the code below)? And, perhaps more importantly, do you have any tips as to how I can find out what’s wrong?

I’m using Optix 3.6.3 along with CUDA 6.0. I use a GeForce GTX 760 and pass “compute_20,sm_20” to nvcc.

The code listing below shows everything using the host API. Here, the Init() function is called once prior to the LoadMesh() function (which is also called once). The TracePhotons function is called every frame.

#include "OptixTracer.h"

#include <optixu/optixu_aabb_namespace.h>

#include "Scene.h"

OptixTracer::OptixTracer()
	: mContext()
	, mMaterial()
	, mGeometryGroup()
{
	mContext = optix::Context::create();
}

OptixTracer::~OptixTracer()
{
	mGeometryGroup->destroy();
	mMaterial->destroy();
	mContext->destroy();
}

bool OptixTracer::Init()
{
	// Check for compute capabilities
	int computeCapabilities[2];
	if(rtDeviceGetAttribute(0, RT_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY, 
		sizeof(computeCapabilities), &computeCapabilities))
	{
		return false;
	}

	// We'll use the default device for now (Specific selections can be made if necessary)
	int deviceId = 0;
	mContext->setDevices(&deviceId, &deviceId + 1);

	mContext->setRayTypeCount(1u);
	mContext->setEntryPointCount(1u);
	mContext->setStackSize(1200);
	mContext->setPrintEnabled(true);

	// Setup output buffer
	optix::Variable outputBuffer = mContext["outputBuffer"];
	optix::Buffer buffer = mContext->createBuffer(RT_BUFFER_INPUT_OUTPUT, 
		RT_FORMAT_FLOAT3, 1024);
	outputBuffer->set(buffer);

	// Ray Generation Program
	std::string photonTestPTXPath("assets/optix_programs/photon_test.cu.ptx");
	optix::Program rayGeneration = mContext->createProgramFromPTXFile(
		photonTestPTXPath, "PhotonEmission");
	mContext->setRayGenerationProgram(0u, rayGeneration);

	// Exception Program
	optix::Program exceptionProgram = mContext->createProgramFromPTXFile(
		photonTestPTXPath, "ExceptionProgram");
	mContext->setExceptionProgram(0u, exceptionProgram);
	mContext->setExceptionEnabled(RT_EXCEPTION_ALL, true);

	// Miss Program
	optix::Program missProgram = mContext->createProgramFromPTXFile(
		photonTestPTXPath, "MissProgram");
	mContext->setMissProgram(0u, missProgram);

	// Material Program
	std::string materialPTXPath("assets/optix_programs/default_material.cu.ptx");
	optix::Program closestHit = mContext->createProgramFromPTXFile(
		materialPTXPath, "ClosestHit");
	optix::Program anyHit = mContext->createProgramFromPTXFile(
		materialPTXPath, "AnyHit");

	mMaterial = mContext->createMaterial();
	mMaterial->setClosestHitProgram(0u, closestHit);
	mMaterial->setAnyHitProgram(0u, anyHit);
		
	return true;
}

#pragma warning(disable : 4996)
void OptixTracer::LoadMesh(const std::shared_ptr<Mesh> mesh,
	const MaterialCache* materialCache)
{
	assert(materialCache != nullptr);

	// Load index/vertex data
	unsigned int numIndices = mesh->mIndices.size();
	unsigned int numVertices = mesh->mVertices.size();

	optix::Buffer iBuffer = mContext->createBuffer(
		RT_BUFFER_INPUT, RT_FORMAT_UNSIGNED_INT, numIndices);
	unsigned int* iBufferData = static_cast<unsigned int*>(iBuffer->map());
	std::copy(mesh->mIndices.begin(), mesh->mIndices.end(), iBufferData);

	optix::Buffer vBuffer = mContext->createBuffer(
		RT_BUFFER_INPUT, RT_FORMAT_USER, numVertices);
	vBuffer->setElementSize(sizeof(Vertex));
	Vertex* vBufferData = static_cast<Vertex*>(vBuffer->map());
	std::copy(mesh->mVertices.begin(), mesh->mVertices.end(), vBufferData);

	// Update/compute bounds
	optix::Aabb aabb;
	for(unsigned int i = 0; i < numVertices; ++i)
	{
		glm::vec3 position = vBufferData[i].mPosition;
		aabb.include(optix::make_float3(position.x, position.y, position.z));
	}

	// Unmap buffers
	iBuffer->unmap();
	vBuffer->unmap();

	// Create geometry instances
	std::string triangleMeshPTXPath("assets/optix_programs/triangle_mesh.cu.ptx");
	optix::Program intersectionProgram = mContext->createProgramFromPTXFile(
		triangleMeshPTXPath, "MeshIntersect");
	optix::Program boundingBoxProgram = mContext->createProgramFromPTXFile(
		triangleMeshPTXPath, "MeshBounds");

	std::vector<optix::GeometryInstance> instances;
	for(size_t i = 0; i < mesh->mSubmeshes.size(); ++i)
	{
		Submesh submesh = mesh->mSubmeshes[i];
		unsigned int numTriangles = submesh.mNumIndices / 3;
		assert(submesh.mNumIndices % 3 == 0);

		optix::Geometry geometry = mContext->createGeometry();
		geometry->setPrimitiveCount(numTriangles);
		geometry->setIntersectionProgram(intersectionProgram);
		geometry->setBoundingBoxProgram(boundingBoxProgram);
		geometry["indexBuffer"]->setBuffer(iBuffer);
		geometry["vertexBuffer"]->setBuffer(vBuffer);
		geometry["indexOffset"]->setUint(submesh.mOffset);

		optix::GeometryInstance instance = mContext->createGeometryInstance();
		instance->setGeometry(geometry);
		instance->setMaterialCount(1u);
		instance->setMaterial(0u, mMaterial);

		LoadMaterialParameters(instance, materialCache->GetMaterial(submesh.mMaterialID));
		instances.push_back(instance);
	}

	// Set up geometry group
	mGeometryGroup = mContext->createGeometryGroup();
	mGeometryGroup->setChildCount(static_cast<unsigned int>(instances.size()));
	
	optix::Acceleration acceleration = mContext->createAcceleration("Trbvh", "Bvh");
	mGeometryGroup->setAcceleration(acceleration);
	acceleration->markDirty();

	for(size_t i = 0; i < instances.size(); ++i)
	{
		mGeometryGroup->setChild(i, instances[i]);
	}

	mContext["topObject"]->set(mGeometryGroup);

	mContext->validate();
	mContext->compile();
}

void OptixTracer::LoadMaterialParameters(
	optix::GeometryInstance instance, const Material* material)
{
	assert(material != nullptr);

	instance["Kd"]->set3fv(&material->mKd[0]);
	instance["Ks"]->set3fv(&material->mKs[0]);
	instance["Ni"]->setFloat(material->mNi);
	instance["Ns"]->setFloat(material->mNs);
}

void OptixTracer::TracePhotons(const Scene& scene)
{
	const unsigned int numSamples = 100;
	mContext["numSamples"]->setUint(numSamples);

	const std::vector<SphereLight>& lights = scene.mSphereLights;
	for(size_t i = 0; i < lights.size(); ++i)
	{
		const SphereLight& light = lights[i];
		glm::vec3 lightPosition = light.mPosition;

		mContext["lightPosition"]->setFloat(
			lightPosition.x, lightPosition.y, lightPosition.z);
		mContext["lightRadius"]->setFloat(light.mRadius);

		mContext->launch(0u, numSamples);

		/*optix::Buffer buffer = mContext["outputBuffer"]->getBuffer();
		optix::float3* data = static_cast<optix::float3*>(buffer->map());

		buffer->unmap();*/
	}
}

I’ll also add the device code below (I’ve left out the GeneratePositionOnSphere() function for brevity, don’t worry: it’s there :))

#include <optix_world.h>

using namespace optix;

const __device__ float PI = 3.141592654f;
const __device__ float T_MIN = 0.05f;

rtDeclareVariable(uint,		launchIndex, rtLaunchIndex, );
rtDeclareVariable(rtObject, topObject, , );

rtDeclareVariable(uint,		numSamples, , );
rtDeclareVariable(float3,	lightPosition, , );
rtDeclareVariable(float,	lightRadius, , );

struct Payload
{
	float3 result;
};

RT_PROGRAM void PhotonEmission()
{
	//rtPrintf("%u\n", launchIndex);

	float3 sphereSample = GeneratePositionOnSphere();
	float3 rayOrigin = lightPosition + sphereSample;
	float3 rayDirection = normalize(sphereSample);

	optix::Ray ray = optix::make_Ray(rayOrigin, rayDirection, 0, 
		T_MIN, RT_DEFAULT_MAX);

	Payload payload;

	rtTrace(topObject, ray, payload);
}

RT_PROGRAM void ExceptionProgram()
{
    rtPrintf("ExceptionProgram\n");
    rtPrintExceptionDetails();
}

RT_PROGRAM void MissProgram()
{
	rtPrintf("MissProgram\n");
}

#include <optix_world.h>

using namespace optix;

RT_PROGRAM void MeshIntersect(int triangleIndex)
{
	rtPrintf("MeshIntersect\n");
}

RT_PROGRAM void MeshBounds(int triangleIndex, float result[6])
{
	rtPrintf("MeshBounds\n");
}

#include <optix_world.h>

using namespace optix;

RT_PROGRAM void ClosestHit()
{
	rtPrintf("ClosestHit\n");
}

RT_PROGRAM void AnyHit()
{
	rtPrintf("AnyHit\n");
}

Hopefully I’ve provided enough info, and someone can point me in the right direction. If anyone wants to know more, let me know!

Thanks!

2

“All of the other programs are empty, save for rtPrintf calls to let me know whether or not the programs are actually invoked.”
That is the problem. The acceleration structure building didn’t happen and there is no traversal possible without (maybe for NoAccel, but you don’t want that on triangle meshes). You must implement the bounding box and intersection programs to be able to get any invocations of your anyhit and closesthit programs.

Other tips:

  • optix::Buffer buffer = mContext->createBuffer(RT_BUFFER_INPUT_OUTPUT, RT_FORMAT_FLOAT3, 1024);
    Try to avoid float3 output buffers if you want to be fast. Use float4 instead. See the OptiX multi-GPU tips in the Programming Guide.
    Do not use an INPUT_OUTPUT Buffer if you only use it for output on the device. Use RT_BUFFER_OUTPUT instead.
  • You can move the iBuffer->unmap() from line 112 to line 96. You’re not using it in the AABB calculation.
  • Always try the Bvh builder first. Then add a benchmark and try the others.
  • const device float PI = 3.141592654f;
    There is no need to define Pi. OptiX math headers contain M_PIf among others.
3

I realize that, however, I figured that the intersection/bounding box programs would be invoked. This was not the case, which seemed odd to me. I have implemented the intersection, bounding box, and closest hit programs now and they do seem to be working correctly now when using Bvh builder.

Thanks!