Merge pull request #5 from mortacious/devel

Update to OptiX 7.4

Author: mortacious

README.md

@@ -1,6 +1,6 @@
 # Python-OptiX
 
-Python wrapper for the OptiX 7.3 raytracing engine.
+Python wrapper for the OptiX 7 raytracing engine.
 
 Python-OptiX wraps the OptiX C++ API using Cython and provides a simplified 
 interface to the original C-like API using mainly the 
@@ -12,16 +12,18 @@ Only Linux is supported at the moment.
 
 ### OptiX Versions
 
-Python-OptiX currently only supports the most recent (7.3.0) release of OptiX
+Python-OptiX currently supports the OptiX releases 7.3.0 and 7.4.0
 
 ## Installation
 
 ### Dependencies
 
 Install a recent version of the [CUDA Toolkit](https://developer.nvidia.com/cuda-downloads)
-and the [OptiX 7.3.0 SDK](https://developer.nvidia.com/optix/downloads/7.3.0/linux64)
+and the [OptiX 7.4.0 SDK](https://developer.nvidia.com/optix/downloads/7.4.0/linux64-x86_64)
 
-Make sure the CUDA header files are installed as well
+Note: The older [OptiX 7.3.0 SDK](https://developer.nvidia.com/optix/downloads/7.4.0/linux64-x86_64) version is supported as well.
+
+Make sure the CUDA header files are installed as well.
 
 Add the locations of CUDA and OptiX to the system `PATH` variable if necessary.
 

examples/compile_with_tasks.py

@@ -0,0 +1,75 @@
+import concurrent.futures
+import optix as ox
+import argparse
+import logging
+import sys
+from concurrent.futures import ThreadPoolExecutor
+import time
+
+logging.basicConfig(stream=sys.stdout, level=logging.DEBUG)
+log = logging.getLogger()
+
+
+if __name__ == "__main__":
+    if ox.optix_version()[1] < 4:
+        raise NotImplementedError("Parallel tasks are not implemented in optix versions < 7.3.")
+
+    parser = argparse.ArgumentParser("Compile OptiX modules using parallel tasks")
+    parser.add_argument('file', nargs=1, help="The input file (.ptx or .cu) to compile")
+    parser.add_argument('-na', '--num-attributes', type=int, default=2, required=False,
+                        help="Number of attribute values (up to 8, default 2)")
+    parser.add_argument('-npv', '--num-payload-values', type=int, default=2, required=False,
+                        help=f"Number of payload values (up to {ox.PipelineCompileOptions.DEFAULT_MAX_PAYLOAD_VALUE_COUNT}, default 2)")
+    parser.add_argument('-npt', '--num-payload-types', type=int, default=1, required=False,
+                        help=f"Number of payload types (up to {ox.ModuleCompileOptions.DEFAULT_MAX_PAYLOAD_TYPE_COUNT}, default 1)")
+    parser.add_argument('-ni', '--num-iters', type=int, default=1, required=False,
+                        help="Number of iterations to compile. > 1 disables disk cache (default 1)")
+    parser.add_argument('-dt', '--disable-tasks', action='store_true', required=False,
+                        help="Disable compilation with tasks (default enabled)")
+    parser.add_argument('-nt', '--num-threads', type=int, default=1, required=False,
+                        help="Number of threads (default 1)")
+    parser.add_argument('-mt', '--max-num-tasks', type=int, default=2, required=False,
+                        help="Maximum number of additional tasks (default 2)")
+
+    args = parser.parse_args()
+
+    logger = ox.Logger(log)
+    ctx = ox.DeviceContext(validation_mode=True, log_callback_function=logger, log_callback_level=3)
+
+    if args.num_iters > 1:
+        ctx.cache_enabled = False
+
+    # compile the file content to ptx in case a .cu file is given
+    ptx = ox.Module.compile_cuda_ptx(args.file[0])
+
+    pipeline_options = ox.PipelineCompileOptions(num_payload_values=0,
+                                                 num_attribute_values=args.num_attributes)
+
+    payload_semantics = [ox.PayloadSemantics.DEFAULT] * args.num_payload_values
+    payload_types = [payload_semantics] * args.num_payload_types
+
+    compile_opts = ox.ModuleCompileOptions(payload_types=payload_types)
+
+    use_tasks = not args.disable_tasks
+
+    if use_tasks:
+        tic = time.time()
+        with ThreadPoolExecutor(max_workers=args.num_threads) as executor:
+            for i in range(args.num_iters):
+                module, task = ox.Module.create_as_task(ctx, ptx, module_compile_options=compile_opts, pipeline_compile_options=pipeline_options)
+                task_futures = {executor.submit(task.execute, args.max_num_tasks)}
+            while task_futures:
+                done, not_done = concurrent.futures.wait(task_futures, timeout=0.25, return_when=concurrent.futures.FIRST_COMPLETED)
+                for future in done:
+                    new_tasks = future.result()
+                    if len(new_tasks) > 0:
+                        task_futures.update({executor.submit(t.execute, args.max_num_tasks) for t in new_tasks})
+                    task_futures.remove(future)
+
+        # wait for the executor to finish here
+        print("Overall run time with tasks", time.time()-tic)
+    else:
+        tic = time.time()
+        for i in range(args.num_iters):
+            module = ox.Module(ctx, ptx, module_compile_options=compile_opts, pipeline_compile_options=pipeline_options)
+        print("Overall run time without tasks", time.time()-tic)

examples/dynamic_geometry.py

@@ -220,6 +220,7 @@ def display_subframe(output_buffer, gl_display, window):
                         framebuf_res_x, framebuf_res_y,
                         output_buffer.get_pbo() )
 
+
 def init_camera_state(state):
     camera = state.camera
     camera.eye = (0, 1, -20)
@@ -233,12 +234,14 @@ def init_camera_state(state):
     trackball.set_reference_frame([1,0,0], [0,0,1], [0,1,0])
     trackball.reinitialize_orientation_from_camera()
 
+
 def create_context(state):
     logger = ox.Logger(log)
     ctx = ox.DeviceContext(validation_mode=False, log_callback_function=logger, log_callback_level=4)
     ctx.cache_enabled = False
     state.ctx = ctx
 
+
 def generate_animated_vertices(out_vertices, animation_mode, time, width, height):
     threads_per_block = 128
     num_blocks = (width*height + threads_per_block - 1) // threads_per_block
@@ -251,12 +254,13 @@ def generate_animated_vertices(out_vertices, animation_mode, time, width, height
 def launch_generate_animated_vertices(state, animation_mode):
     generate_animated_vertices(state.d_temp_vertices, animation_mode, state.time, g_tessellation_resolution, g_tessellation_resolution)
 
+
 def update_mesh_accel(state):
     # first sphere is static
 
     # second sphere moves by updating its transform matrix
-    transform = state.ias_build_input.get_transform_view(1)
-    transform[1,-1] = np.sin(4*state.time)
+    transform = state.ias_build_input.view_instance_transform(1)
+    transform[1, -1] = np.sin(4*state.time)
 
     # third sphere deforms
     launch_generate_animated_vertices(state, AnimationMode.DEFORM)
@@ -270,13 +274,14 @@ def update_mesh_accel(state):
         state.last_exploding_sphere_rebuild_time = state.time
         state.exploding_gas = ox.AccelerationStructure(state.ctx, state.gas_build_input,
                 compact=True, allow_update=True, random_vertex_access=True)
-        state.ias_build_input.instances[3].update_traversable(state.exploding_gas)
+        state.ias_build_input[3].traversable = state.exploding_gas
         state.ias_build_input.update_instance(3)
     else:
         state.exploding_gas.update(state.gas_build_input)
 
     state.ias.update(state.ias_build_input)
 
+
 def build_vertex_generation_kernel(state):
     cuda_source = os.path.join(script_dir, 'cuda', 'dynamic_geometry_vertex_generation.cu')
     example_include_path = os.path.dirname(cuda_source)
@@ -289,6 +294,7 @@ def build_vertex_generation_kernel(state):
     state.generate_vertices_kernel = cp.RawKernel(code=code, backend='nvrtc',
             options=build_flags, name='generate_vertices')
 
+
 def build_mesh_accel(state):
     # Allocate temporary space for vertex generation.
     # The same memory space is reused for generating the deformed and exploding vertices before updates.
@@ -335,9 +341,10 @@ def create_module(state):
     else:
         exception_flags=ox.ExceptionFlags.NONE
 
+    print("Triangle value", ox.PrimitiveTypeFlags.TRIANGLE.value)
     pipeline_opts = ox.PipelineCompileOptions(
             uses_motion_blur=False,
-            uses_primitive_type_flags = ox.PrimitiveTypeFlags.TRIANGLE,
+            uses_primitive_type_flags =ox.PrimitiveTypeFlags.TRIANGLE,
             traversable_graph_flags=ox.TraversableGraphFlags.ALLOW_SINGLE_LEVEL_INSTANCING,
             exception_flags=exception_flags,
             num_payload_values=3,
@@ -347,7 +354,7 @@ def create_module(state):
     compile_opts = ox.ModuleCompileOptions(
             max_register_count=ox.ModuleCompileOptions.DEFAULT_MAX_REGISTER_COUNT,
             opt_level=ox.CompileOptimizationLevel.DEFAULT,
-            debug_level=ox.CompileDebugLevel.LINEINFO)
+            debug_level=ox.CompileDebugLevel.MODERATE)
 
     cuda_source = os.path.join(script_dir, 'cuda', 'dynamic_geometry.cu')
     state.module = ox.Module(state.ctx, cuda_source, compile_opts, pipeline_opts)
@@ -364,7 +371,7 @@ def create_pipeline(state):
     program_grps = [state.raygen_grp, state.miss_grp, state.hit_grp]
 
     link_opts = ox.PipelineLinkOptions(max_trace_depth=1,
-                                       debug_level=ox.CompileDebugLevel.LINEINFO)
+                                       debug_level=ox.CompileDebugLevel.MODERATE)
 
     pipeline = ox.Pipeline(state.ctx,
                            compile_options=state.pipeline_opts,

examples/dynamic_materials.py

@@ -194,7 +194,7 @@ def create_module(state):
     compile_opts = ox.ModuleCompileOptions(
             max_register_count=ox.ModuleCompileOptions.DEFAULT_MAX_REGISTER_COUNT,
             opt_level=ox.CompileOptimizationLevel.DEFAULT,
-            debug_level=ox.CompileDebugLevel.LINEINFO)
+            debug_level=ox.CompileDebugLevel.MODERATE)
 
     source = os.path.join(script_dir, 'cuda', 'dynamic_materials.cu')
     state.module = ox.Module(state.ctx, source, compile_opts, pipeline_opts)

examples/hello.py

@@ -1,20 +1,15 @@
-import os, sys, logging
-
 import optix as ox
 import cupy as cp
 import numpy as np
-
 from PIL import Image, ImageOps
-
-script_dir = os.path.dirname(os.path.abspath(__file__))
-
+import logging
+import sys
 logging.basicConfig(stream=sys.stdout, level=logging.DEBUG)
 log = logging.getLogger()
 
 def create_module(ctx, pipeline_opts):
-    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.LINEINFO)
-    source = os.path.join(script_dir, 'cuda', 'hello.cu')
-    module = ox.Module(ctx, source, compile_opts, pipeline_opts)
+    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.FULL, opt_level=ox.CompileOptimizationLevel.LEVEL_0)
+    module = ox.Module(ctx, 'cuda/hello.cu', compile_opts, pipeline_opts)
     return module
 
 

examples/spheres.py

@@ -1,13 +1,9 @@
-import os, sys, logging
-
+import optix as ox
 import cupy as cp
 import numpy as np
-import optix as ox
-
 from PIL import Image, ImageOps
-
-script_dir = os.path.dirname(os.path.abspath(__file__))
-
+import logging
+import sys
 logging.basicConfig(stream=sys.stdout, level=logging.DEBUG)
 log = logging.getLogger()
 img_size = (1024, 768)
@@ -26,9 +22,8 @@ def create_acceleration_structure(ctx, bboxes):
 
 
 def create_module(ctx, pipeline_opts):
-    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.LINEINFO)
-    source = os.path.join(script_dir, 'cuda', 'spheres.cu')
-    module = ox.Module(ctx, source, compile_opts, pipeline_opts)
+    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.FULL, opt_level=ox.CompileOptimizationLevel.LEVEL_0)
+    module = ox.Module(ctx, 'cuda/spheres.cu', compile_opts, pipeline_opts)
     return module
 
 

examples/triangle.py

@@ -1,13 +1,8 @@
-import os
-
+import optix as ox
 import cupy as cp
 import numpy as np
-import optix as ox
-
 from PIL import Image, ImageOps
 
-script_dir = os.path.dirname(os.path.abspath(__file__))
-
 img_size = (1024, 768)
 
 # use a regular function for logging
@@ -23,9 +18,8 @@ def create_acceleration_structure(ctx, vertices):
 
 
 def create_module(ctx, pipeline_opts):
-    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.LINEINFO)
-    source = os.path.join(script_dir, 'cuda', 'triangle.cu')
-    module = ox.Module(ctx, source, compile_opts, pipeline_opts)
+    compile_opts = ox.ModuleCompileOptions(debug_level=ox.CompileDebugLevel.FULL, opt_level=ox.CompileOptimizationLevel.LEVEL_0)
+    module = ox.Module(ctx, 'cuda/triangle.cu', compile_opts, pipeline_opts)
     return module
 
 

optix/__init__.py

@@ -1,6 +1,6 @@
-from .context import DeviceContext
+from .context import DeviceContext, optix_version
 from .build import *
-from .module import Module, ModuleCompileOptions, CompileOptimizationLevel, CompileDebugLevel
+from .module import Module, ModuleCompileOptions, CompileOptimizationLevel, CompileDebugLevel, PayloadSemantics, Task
 from .program_group import ProgramGroup
 from .struct import SbtRecord, LaunchParamsRecord
 from .shader_binding_table import ShaderBindingTable