optixTriangle sample: how to share a variable among threads and bring it back to the main?

[fatecasino]

I am trying to understand a very basic concept of how to collect data from cuda kernel back to the main function.

My case scenario is based on the optixTriangle sample and I would like simply to add +1.0 to a double SUM every time a ray misses the triangle.

What modifications should I do to the optixTriangle sample in order to collect the SUM and print it in the main function? What is the proper way to achieve this?
Thanks!

[fatecasino]

Any instructions, example or documentation related to this question?

[droettger]

I am trying to understand a very basic concept of how to collect data from cuda kernel back to the main function.

Please read everything with the word payload inside the OptiX Programming Guide and OptiX API Reference documents.

For that you put the word payload into the search bar at the upper right and then click on all search results and read the respective chapters.

Then look at how other OptiX applications manage payload registers on the optixTrace call and set and get these payload registers inside the other program domains reached by that ray.

My case scenario is based on the optixTriangle sample and I would like simply to add +1.0 to a double SUM every time a ray misses the triangle.

Please avoid double variables at all costs! These are potentially slow.

If you really try to count things by incrementing a floating point or double precision variable by 1.0, that is a really bad idea.
Just use an unsigned int and increment it, which also matches the payload register type. That would give you 4 Gig values if that is sufficient.

Also mind that if you reach the miss program, the ray ends there, so each ray will only reach that once at maximum.
So if you really want to add something to a per ray payload inside the miss program, that would look like this (expanded to multiple lines for clarity):

extern "C" __global__ void __miss__ms()
{
  unsigned int payload0 = optixGetPayload_0();
  ++payload0;
  optixSetPayload_0(payload0);
}

Of course the optixTriangle example is writing the red color into the first of the three payload registers it uses so if you want to retain that functionality, you would need to add another payload register and that means changing the OptixPipelineCompileOptions numPayloadValues inside the host code.

It’s unclear what exactly you want to accumulate, given that each hit or miss is just a boolean condition for each ray. Something like ambient occlusion would accumulate such individual visibility results for a number of rays, but the increment of the sum would happen inside the caller depending on the resulting payload values (0 for hit (shadow), 1 for miss (light)).