Building curve input in optix 7.2

I’m trying to modify the optixTriangle SDK example to render a rectangle by this cod but it does not work. I feel like it’s something I do wrong in building the curve input.

    // accel handling
    OptixTraversableHandle gas_handle;
    CUdeviceptr            d_gas_output_buffer;
        const int NUM_KEYS = 1, degree = 1;  

        // Use default options for simplicity.  In a real use case we would want to
        // enable compaction, etc
        OptixAccelBuildOptions accel_options = {};
        accel_options.buildFlags = OPTIX_BUILD_FLAG_NONE;
        accel_options.operation  = OPTIX_BUILD_OPERATION_BUILD;
        accel_options.motionOptions.numKeys   = NUM_KEYS;
        // rectangle build input
        std::vector<float3> vertices;
        std::vector<float>  widths;
        vertices.push_back( make_float3( 0.0f, 0.0f, 0.0f ) );
        vertices.push_back( make_float3( 0.5f, 0.0f, 0.0f ) );
        vertices.push_back( make_float3( 0.5f, 0.5f, 0.0f ) );
        vertices.push_back( make_float3( 0.0f, 0.5f, 0.0f ) );
        vertices.push_back( make_float3( 0.0f, 0.0f, 0.0f ) );

        int nPrimitives = vertices.size() - 1;

        const size_t vertices_size = sizeof( float3 ) * vertices.size();
        CUdeviceptr  d_vertices    = 0;
        CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_vertices ), vertices_size ) );
        CUDA_CHECK( cudaMemcpy( reinterpret_cast<void*>( d_vertices ),, vertices_size, cudaMemcpyHostToDevice ) );

        const size_t widthsSize = sizeof( float ) * widths.size();
        CUdeviceptr  d_widths   = 0;
        CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_widths ), widthsSize ) );
        CUDA_CHECK( cudaMemcpy( reinterpret_cast<void*>( d_widths ),, widthsSize, cudaMemcpyHostToDevice ) );
        CUdeviceptr vertexBufferPointers[NUM_KEYS];
        CUdeviceptr widthBufferPointers[NUM_KEYS];
        for( int i = 0; i < NUM_KEYS; ++i ) {
            vertexBufferPointers[i] = d_vertices + i * (degree + 1) * sizeof(float3);
            widthBufferPointers[i] = d_widths + i * (degree + 1) * sizeof(float);

        // Curve build intput: with a single segment the index array
        // contains index of first vertex.
        const std::array<int, 4> segmentIndices     = {0,1,2,3};
        const size_t             segmentIndicesSize = sizeof( int ) * segmentIndices.size();
        CUdeviceptr              d_segementIndices  = 0;
        CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_segementIndices ), segmentIndicesSize ) );
        CUDA_CHECK( cudaMemcpy( reinterpret_cast<void*>( d_segementIndices ),,
                                segmentIndicesSize, cudaMemcpyHostToDevice ) );

        OptixBuildInput curve_input = {};

        curve_input.type = OPTIX_BUILD_INPUT_TYPE_CURVES;
        curve_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR;

        curve_input.curveArray.numPrimitives        = nPrimitives;
        curve_input.curveArray.vertexBuffers        = vertexBufferPointers;
        curve_input.curveArray.numVertices          = static_cast<uint32_t>( vertices.size() );
        curve_input.curveArray.vertexStrideInBytes  = sizeof( float3 );
        curve_input.curveArray.widthBuffers         = widthBufferPointers;
        curve_input.curveArray.widthStrideInBytes   = sizeof( float );
        curve_input.curveArray.normalBuffers        = 0;
        curve_input.curveArray.normalStrideInBytes  = 0;
        curve_input.curveArray.indexBuffer          = d_segementIndices;
        curve_input.curveArray.indexStrideInBytes   = sizeof( int );
        curve_input.curveArray.flag                 = OPTIX_GEOMETRY_FLAG_NONE;
        curve_input.curveArray.primitiveIndexOffset = 0;

        OptixAccelBufferSizes gas_buffer_sizes;
        OPTIX_CHECK( optixAccelComputeMemoryUsage(
                    1, // Number of build inputs
                    ) );
        CUdeviceptr d_temp_buffer_gas;
        CUDA_CHECK( cudaMalloc(
                    reinterpret_cast<void**>( &d_temp_buffer_gas ),
                    ) );
        CUDA_CHECK( cudaMalloc(
                    reinterpret_cast<void**>( &d_gas_output_buffer ),
                    ) );

        OPTIX_CHECK( optixAccelBuild(
                    0,                  // CUDA stream
                    1,                  // num build inputs
                    nullptr,            // emitted property list
                    0                   // num emitted properties
                    ) );

        // We can now free the scratch space buffer used during build and the vertex
        // inputs, since they are not needed by our trivial shading method
        CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_temp_buffer_gas ) ) );
        CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_vertices        ) ) );

Copying the answer from your OptiX-Help email thread:

Please have a look into the online OptiX Programming Guide:
When searching for curve inside the top right search bar, you’ll find all topics related to built-in curve primitives in OptiX.

In contrast to OptiX’ built-in triangle primitives, the built-in curve primitives need to have one of the pre-defined intersection programs for curve primitives set in the shader binding table hit record. Otherwise there is no way to reach the any-hit or closest-hit programs assigned to built-in curve primitives.
This chapter explains the necessary call to optixBuiltinISModuleGet() to get one of the three built-in curve intersection programs matching the curve type you want to use:

You can also not simply use the implementation for curves because triangles return two barycentric coordinates as attributes per triangle hit while curve intersection programs only return a single attribute value per curve hit. Means the closest hit programs for triangles and curves need to distinguish between the hit primitive types to correctly handle the vertex attribute calculations. That can either be done with different programs on the hit records per primtive type, or when you want to share the same hit programs, by analyzing the optixGetHitKind() results to determine which final vertex attribute calculation code path needs to be taken.
Explained here:

The OptiX SDK example program optixCurves and optixHair shows how to put this all together.

Note that there has just been an OptiX SDK 7.5.0 released and the online manuals are matching that newest version.
It’s always recommended to use the newest available OptiX SDK versions because of bug fixes and additional features.
Please read the OptiX release notes linked below each OptiX version’s download button for required display driver versions and changes to the individual OptiX SDK releases.

Thank you for your response. I have just downloaded optix 7.5 and compiled the sdk examples. When I ran the optixCurve example, I got the error

OPTIX_ERROR_UNSUPPORTED_ABI_VERSION in optixInit(), optixCurve.cpp line 188.

I have V100 GPU, driver version 510.47.03, cuda version 11.6.

Any idea?

OptiX 7.5.0 requires drivers from the release 515 branch.

Again, please always read the OptiX release notes linked below each OptiX version’s download button for required display driver versions and changes to the individual OptiX SDK releases.

the optix website still says driver R495

Yeah, that’s wrong. I’ll inform the web team to fix that. The OptiX Release Notes document has it right.