eic-opticks/CSGOptiX/IAS_Builder.cc

0001 #include <cassert>
0002 #include <iostream>
0003 #include <iomanip>
0004 #include <map>
0005 #include <optix.h>
0006 #include <optix_stubs.h>
0007
0008 #include <cuda_runtime.h>
0009 #include "scuda.h"    // roundUp
0010
0011 #include "CUDA_CHECK.h"
0012 #include "OPTIX_CHECK.h"
0013
0014 #include "Ctx.h"
0015 #include "Properties.h"
0016
0017 #include "GAS.h"
0018 #include "IAS.h"
0019 #include "IAS_Builder.h"
0020 #include "SBT.h"
0021
0022
0023 /**
0024 IAS_Builder::Build_OLD
0025 -----------------------
0026
0027 NB this is no longer standardly used, being replaced by WITH_SOPTIX_ACCEL approach, see SBT::createIAS
0028
0029 Boilerplate turning the vector of OptixInstance into an IAS.
0030
0031
0032 **/
0033
0034 void IAS_Builder::Build_OLD(IAS& ias, const std::vector<OptixInstance>& instances)
0035 {
0036     unsigned numInstances = instances.size() ;
0037
0038     unsigned numBytes = sizeof( OptixInstance )*numInstances ;
0039
0040
0041     CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &ias.d_instances ), numBytes ) );
0042     CUDA_CHECK( cudaMemcpy(
0043                 reinterpret_cast<void*>( ias.d_instances ),
0044                 instances.data(),
0045                 numBytes,
0046                 cudaMemcpyHostToDevice
0047                 ) );
0048
0049
0050     OptixBuildInput buildInput = {};
0051
0052     buildInput.type = OPTIX_BUILD_INPUT_TYPE_INSTANCES;
0053     buildInput.instanceArray.instances = ias.d_instances ;
0054     buildInput.instanceArray.numInstances = numInstances ;
0055
0056
0057     OptixAccelBuildOptions accel_options = {};
0058     accel_options.buildFlags =
0059         OPTIX_BUILD_FLAG_PREFER_FAST_TRACE |
0060         OPTIX_BUILD_FLAG_ALLOW_COMPACTION ;
0061     accel_options.operation  = OPTIX_BUILD_OPERATION_BUILD;
0062
0063     OptixAccelBufferSizes as_buffer_sizes;
0064
0065     OPTIX_CHECK( optixAccelComputeMemoryUsage( Ctx::context, &accel_options, &buildInput, 1, &as_buffer_sizes ) );
0066     CUdeviceptr d_temp_buffer_as;
0067     CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_temp_buffer_as ), as_buffer_sizes.tempSizeInBytes ) );
0068
0069     // non-compacted output
0070     CUdeviceptr d_buffer_temp_output_as_and_compacted_size;
0071     size_t      compactedSizeOffset = roundUp<size_t>( as_buffer_sizes.outputSizeInBytes, 8ull );
0072     CUDA_CHECK( cudaMalloc(
0073                 reinterpret_cast<void**>( &d_buffer_temp_output_as_and_compacted_size ),
0074                 compactedSizeOffset + 8
0075                 ) );
0076
0077     OptixAccelEmitDesc emitProperty = {};
0078     emitProperty.type               = OPTIX_PROPERTY_TYPE_COMPACTED_SIZE;
0079     emitProperty.result             = ( CUdeviceptr )( (char*)d_buffer_temp_output_as_and_compacted_size + compactedSizeOffset );
0080
0081
0082     OPTIX_CHECK( optixAccelBuild( Ctx::context,
0083                                   0,                  // CUDA stream
0084                                   &accel_options,
0085                                   &buildInput,
0086                                   1,                  // num build inputs
0087                                   d_temp_buffer_as,
0088                                   as_buffer_sizes.tempSizeInBytes,
0089                                   d_buffer_temp_output_as_and_compacted_size,
0090                                   as_buffer_sizes.outputSizeInBytes,
0091                                   &ias.handle,
0092                                   &emitProperty,      // emitted property list
0093                                   1                   // num emitted properties
0094                                   ) );
0095
0096     CUDA_CHECK( cudaFree( (void*)d_temp_buffer_as ) );
0097
0098     size_t compacted_as_size;
0099     CUDA_CHECK( cudaMemcpy( &compacted_as_size, (void*)emitProperty.result, sizeof(size_t), cudaMemcpyDeviceToHost ) );
0100
0101
0102     if( compacted_as_size < as_buffer_sizes.outputSizeInBytes )
0103     {
0104         CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &ias.d_buffer ), compacted_as_size ) );
0105
0106         // use ias.handle as input and output
0107         OPTIX_CHECK( optixAccelCompact( Ctx::context, 0, ias.handle, ias.d_buffer, compacted_as_size, &ias.handle ) );
0108
0109         CUDA_CHECK( cudaFree( (void*)d_buffer_temp_output_as_and_compacted_size ) );
0110
0111         std::cerr
0112             << "(compacted is smaller) "
0113             << " compacted_as_size : " << compacted_as_size
0114             << " as_buffer_sizes.outputSizeInBytes : " << as_buffer_sizes.outputSizeInBytes
0115             << "\n"
0116             ;
0117
0118     }
0119     else
0120     {
0121         ias.d_buffer = d_buffer_temp_output_as_and_compacted_size;
0122
0123         std::cerr
0124             << "(compacted not smaller) "
0125             << " compacted_as_size : " << compacted_as_size
0126             << " as_buffer_sizes.outputSizeInBytes : " << as_buffer_sizes.outputSizeInBytes
0127             << "\n"
0128             ;
0129     }
0130 }
0131
0132