Updating GAS with custom primitives

[tintingai]

Hello,
I have an “invalid value” error when updating my GAS with custom primitives.
I already have an Optix program which is loading triangles and throws rays on these triangles and works pretty well. This is a simple program where the GAS is generated at the initialization and there is a loop after where the triangles are moving a bit and the GAS is updated, for N frames, and then I am freeing everything and the program stop.

Now I would like to change my triangles by custom primitives.
In the GAS construction, I changed those lines:

state.triangle_input.type = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;
  state.triangle_input.triangleArray.vertexFormat = OPTIX_VERTEX_FORMAT_FLOAT3;
  state.triangle_input.triangleArray.numVertices = static_cast<unsigned int>(nverts);
  state.triangle_input.triangleArray.vertexBuffers = &state.d_temp_vertices;
  state.triangle_input.triangleArray.flags = &state.triangle_flags;
  state.triangle_input.triangleArray.numSbtRecords = 1;
  state.triangle_input.triangleArray.indexFormat = OPTIX_INDICES_FORMAT_UNSIGNED_INT3;
  state.triangle_input.triangleArray.numIndexTriplets = static_cast<unsigned int>(ntris);
  state.triangle_input.triangleArray.indexBuffer = reinterpret_cast<CUdeviceptr>(devTriangles);

with those lines:

state.triangle_input.type = OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES;
  state.triangle_input.customPrimitiveArray.aabbBuffers = reinterpret_cast<CUdeviceptr*>(&d_aabb_array);
  state.triangle_input.customPrimitiveArray.numPrimitives = 1;
  uint32_t aabb_input_flags[1] = {OPTIX_GEOMETRY_FLAG_NONE};
  state.triangle_input.customPrimitiveArray.flags = aabb_input_flags;
  state.triangle_input.customPrimitiveArray.numSbtRecords = 1;

with the same code following for construction:

  OptixAccelBuildOptions accel_options = {};
  accel_options.buildFlags = OPTIX_BUILD_FLAG_ALLOW_UPDATE;
  accel_options.operation = OPTIX_BUILD_OPERATION_BUILD;

  OptixAccelBufferSizes gas_buffer_sizes;
  OPTIX_CHECK( optixAccelComputeMemoryUsage(state.context, &accel_options, &state.triangle_input, 1, &gas_buffer_sizes) );
  state.temp_buffer_size = gas_buffer_sizes.tempSizeInBytes;

  CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>(&state.d_temp_buffer), gas_buffer_sizes.tempSizeInBytes) );

  // non-compact output
  CUdeviceptr d_buffer_temp_output_gas_and_compacted_size;
  size_t compactedSizeOffset = roundUp<size_t>(gas_buffer_sizes.outputSizeInBytes, 8ull);
  CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>(&d_buffer_temp_output_gas_and_compacted_size), compactedSizeOffset + 8) );

  OptixAccelEmitDesc emitProperty = {};
  emitProperty.type = OPTIX_PROPERTY_TYPE_AABBS;
  emitProperty.result = (CUdeviceptr)((char *)d_buffer_temp_output_gas_and_compacted_size + compactedSizeOffset);

  OPTIX_CHECK( optixAccelBuild(
        state.context,
        0,
        &accel_options,
        &state.triangle_input,
        1,
        state.d_temp_buffer,
        gas_buffer_sizes.tempSizeInBytes,
        d_buffer_temp_output_gas_and_compacted_size,
        gas_buffer_sizes.outputSizeInBytes,
        &state.gas_handle,
        &emitProperty, 1)
  );
  state.d_gas_output_buffer = d_buffer_temp_output_gas_and_compacted_size;
  state.gas_output_buffer_size = gas_buffer_sizes.outputSizeInBytes;

And, because the documentation do not seem to mention a special case about custom primitives for dynamic updates, I kept exactly the same code for updating my GAS:

void updateASFromDevice(GASstate &state) {
    OptixAccelBuildOptions gas_accel_options = {};
    gas_accel_options.buildFlags = OPTIX_BUILD_FLAG_ALLOW_UPDATE;
    gas_accel_options.operation = OPTIX_BUILD_OPERATION_UPDATE;
    OPTIX_CHECK(optixAccelBuild(
        state.context,
        0,
        &gas_accel_options,
        &state.triangle_input,
        1,
        state.d_temp_buffer,
        state.temp_buffer_size,
        state.d_gas_output_buffer,
        state.gas_output_buffer_size,
        &state.gas_handle,
        nullptr,
        0)
    ); 
}

I added the intersection shader in the program group. My shaders are all empty except for the raygen shader and the intersection shader, which is a simple print regarding which ray interact with which AABB

The program is working really well if I am not updating the GAS and rebuilding it each frame, but when I want to update the GAS instead of rebuilding it, I have this error message during the execution:

/home/hbec/RTX-CUDA-toy/./src/rtx_functions.h:472 Optix Error: 'Invalid value'

Line 472 referring to the optixAccelBuild function in my updateASFromDevice function (see above).

I don’t understand what else do I have to change. Is there a specific case for custom primitive for the update of GAS ? And also do you know how to have more precision on the error message ?

Thank you in advance for your interest in my problem, I am at your disposal for any further information,
tintingai.

[dhart]

Hi @tintingai,

Have you tried using OptiX validation mode? I don’t know if that will clarify the error, but always good to check. See OptixDeviceContextOptions.validationMode.

Are you updating anything aside from custom primitive AABB locations? I can’t tell from the code if anything in the BVH is changing. The rule for BVH UPDATE is that everything must remain the same except for positions.

“When updating an existing acceleration structure, only the device pointers and/or their buffer content may be changed. You cannot change the number of build inputs, the build input types, build flags, traversable handles for instances (for an instance-AS), or the number of vertices, indices, AABBs, instances, SBT records or motion keys. Changes to any of these things may result in undefined behavior, including GPU faults.”

https://raytracing-docs.nvidia.com/optix7/guide/index.html#acceleration_structures#dynamic-updates

–
David.

[tintingai]

Hello dhart,

Thank you for your rapid answer.
Indeed, my logCallbackLevel was only to 1, this is why I didn’t get any precision. I set it to 4 and here is the complete error message:

Optix Log[2][ERROR]: 'Invalid value (8546) for "buildInputs[0].customPrimitiveArray.flags[0]"'
/home/hbec/RTX-CUDA-toy/./src/rtx_functions.h:474 Optix Error: 'Optix Log[2][ERROR]: 'Invalid value (8546) for "buildInputs[0].customPrimitiveArray.flags[0]"'
Invalid value'

It is about these lines:

uint32_t aabb_input_flags[1] = {OPTIX_GEOMETRY_FLAG_NONE};
state.triangle_input.customPrimitiveArray.flags = aabb_input_flags;

Apparently the flag OPTIX_GEOMETRY_FLAG_NONE is an invalid value, but I don’t understand why. I tried by curiosity the other flags and even leaving it empty, but the same error appears.

Concerning the change which could appear between the construction and the update, I removed all the part which update the positions, and the error remain.

Here is my main loop:

// construction of the aabb
OptixAabb* aabb_array = new OptixAabb[n];
for(int i=0; i<n; i++){
    float3 pos = points[i];
    aabb_array[i].minX = pos.x - 1;
    aabb_array[i].minY = pos.y - 1;
    aabb_array[i].minZ = pos.z - 1;
    aabb_array[i].maxX = pos.x + 1;
    aabb_array[i].maxY = pos.y + 1;
    aabb_array[i].maxZ = pos.z + 1;
}
OptixAabb* d_aabb_array;
CUDA_CHECK( cudaMalloc(&d_aabb_array, sizeof(OptixAabb) * n) );
CUDA_CHECK( cudaMemcpy(d_aabb_array, aabb_array, sizeof(OptixAabb) * n, cudaMemcpyHostToDevice) );

// Build Acceleration Structure
printf("%sBuild AS on GPU......................", AC_MAGENTA); fflush(stdout);
timer.restart();
buildASFromDeviceDataAABB(state, n, d_aabb_array);    //Custom primitive version
//buildASFromDeviceData(state, 3*n, n, devVertices, devTriangles);    //Triangle version
cudaDeviceSynchronize();
timer.stop();
printf("done: %f ms%s\n", timer.get_elapsed_ms(), AC_RESET);

// enter parameters
Params params, *device_params;
params.handle = state.gas_handle;
CUDA_CHECK(cudaMalloc(&device_params, sizeof(Params)));
CUDA_CHECK(cudaMemcpy(device_params, &params, sizeof(Params), cudaMemcpyHostToDevice));

for(int ki = 0; ki<steps; ++ki){
        // launch the ray
        printf("\t%sOptiX Launch [%-15s].......", AC_BOLDCYAN, algStr[alg]); fflush(stdout);
        timer.restart();
        OPTIX_CHECK(optixLaunch(state.pipeline, 0, reinterpret_cast<CUdeviceptr>(device_params), sizeof(Params), &state.sbt, 1, 1, 1));
        CUDA_CHECK(cudaDeviceSynchronize());
        timer.stop();
        printf("done: %f ms%s\n", timer.get_elapsed_ms(), AC_RESET);

        // update AS from device data
        printf("\t%sUpdating AS..........................", AC_YELLOW); fflush(stdout);
        timer.restart();
        updateASFromDevice(state);
        CUDA_CHECK(cudaDeviceSynchronize());
        timer.stop();
        printf("done: %f ms%s\n", timer.get_elapsed_ms(), AC_RESET);
        printf("\n");
    }

Thank you again for helping me,
tintingai

[dhart]

Notice the error says the value being passed is 8546, while the value of OPTIX_GEOMETRY_FLAG_NONE = 0.

I would suggest looking for uninitialized and/or overwritten memory. You could stop your program right before calling optixAccelBuild(), and verify that your build input flags really do contain this erroneous value. If so, then walk backward in the code until you find out who is setting that value.

–
David.

[droettger]

Note that decimal 8546 is OPTIX_BUILD_OPERATION_UPDATE = 0x2162.

You’re not actually using OPTIX_BUILD_FLAG_ALLOW_COMPACTION but copied code from somewhere calculating the compacted size and then changed it incorrectly.

The bug is that you’re querying the top-level AABB of that acceleration structure which is size OptixAabb * numMotionSteps (in your case 6 floats) and not the compacted size while you only allocated room for the compacted size (size_t) value:

  CUdeviceptr d_buffer_temp_output_gas_and_compacted_size;
  size_t compactedSizeOffset = roundUp<size_t>(gas_buffer_sizes.outputSizeInBytes, 8ull);
  CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>(&d_buffer_temp_output_gas_and_compacted_size), compactedSizeOffset + 8) ); // BUG: Not enough room for the AABB property queried below.

  OptixAccelEmitDesc emitProperty = {};
  emitProperty.type = OPTIX_PROPERTY_TYPE_AABBS; // BUG: Should be OPTIX_PROPERTY_TYPE_COMPACTED_SIZE
  emitProperty.result = (CUdeviceptr)((char *)d_buffer_temp_output_gas_and_compacted_size + compactedSizeOffset);

[tintingai]

Thank you for all your answers, we figure it out where was the error !

We found the problem for the Invalid value (8546) for "buildInputs[0].customPrimitiveArray.flags[0]" error, it was because uint32_t aabb_input_flags[1] = {OPTIX_GEOMETRY_FLAG_NONE}; was initialized inside the function so when the function finished, the pointer was pointing to nowhere.

But after, there was another error concerning “illegal memory access error” which occurred starting to the first GAS update (I put the error message for maybe helping future persons having the same error):

OptiX Launch ....................
        Updating AS..........................
/home/hbec/RTX-CUDA-toy/./src/rtx.h:120 CUDA Error: 'an illegal memory access was encountered'

        OptiX Launch ....................
Optix Log[2][ERROR]: 'Failed to synchronize with given stream (CUDA error string: an illegal memory access was encountered, CUDA error code: 700)
Validation mode found given stream in erroneous state'
/home/hbec/RTX-CUDA-toy/./src/rtx.h:94 Optix Error: 'Optix Log[2][ERROR]: 'Failed to synchronize with given stream (CUDA error string: an illegal memory access was encountered, CUDA error code: 700)
Validation mode found given stream in erroneous state'
Error during validation mode run'
/home/hbec/RTX-CUDA-toy/./src/rtx.h:95 CUDA Error: 'an illegal memory access was encountered'
[...]

The error were coming from this line:
state.triangle_input.customPrimitiveArray.aabbBuffers = reinterpret_cast<CUdeviceptr*>(&d_aabb_array);
because d_aabb_array in the function was a parameter of type OptixAabb* of the GAS building function, so its pointer, after the GAS building function finished, was pointing to nowhere !
So we converted and stored the value of d_aabb_array from OptixAabb* to CUdeviceptr in our GASstate structure, and we gave instead the pointer of this global value in the aabbBuffers:

state.d_temp_vertices = reinterpret_cast<CUdeviceptr>(d_aabb_array);
state.triangle_input.customPrimitiveArray.aabbBuffers = &state.d_temp_vertices;

Now the GAS update with custom primitives works perfectly, thank you again for your support, we hope this post will help in the future.
tintingai