Skip to content

Commit e84a8c1

Browse files
author
Felix Igelbrink
committed
adapted examples for api changes and fixed a few bugs
1 parent 49a9166 commit e84a8c1

8 files changed

Lines changed: 78 additions & 26 deletions

File tree

examples/compile_with_tasks.py

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,4 @@
11
import concurrent.futures
2-
3-
import direct.task.Task
4-
52
import optix as ox
63
import argparse
74
import logging

examples/dynamic_geometry.py

Lines changed: 13 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -220,6 +220,7 @@ def display_subframe(output_buffer, gl_display, window):
220220
framebuf_res_x, framebuf_res_y,
221221
output_buffer.get_pbo() )
222222

223+
223224
def init_camera_state(state):
224225
camera = state.camera
225226
camera.eye = (0, 1, -20)
@@ -233,12 +234,14 @@ def init_camera_state(state):
233234
trackball.set_reference_frame([1,0,0], [0,0,1], [0,1,0])
234235
trackball.reinitialize_orientation_from_camera()
235236

237+
236238
def create_context(state):
237239
logger = ox.Logger(log)
238240
ctx = ox.DeviceContext(validation_mode=False, log_callback_function=logger, log_callback_level=4)
239241
ctx.cache_enabled = False
240242
state.ctx = ctx
241243

244+
242245
def generate_animated_vertices(out_vertices, animation_mode, time, width, height):
243246
threads_per_block = 128
244247
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
251254
def launch_generate_animated_vertices(state, animation_mode):
252255
generate_animated_vertices(state.d_temp_vertices, animation_mode, state.time, g_tessellation_resolution, g_tessellation_resolution)
253256

257+
254258
def update_mesh_accel(state):
255259
# first sphere is static
256260

257261
# second sphere moves by updating its transform matrix
258-
transform = state.ias_build_input.get_transform_view(1)
259-
transform[1,-1] = np.sin(4*state.time)
262+
transform = state.ias_build_input.view_instance_transform(1)
263+
transform[1, -1] = np.sin(4*state.time)
260264

261265
# third sphere deforms
262266
launch_generate_animated_vertices(state, AnimationMode.DEFORM)
@@ -270,13 +274,14 @@ def update_mesh_accel(state):
270274
state.last_exploding_sphere_rebuild_time = state.time
271275
state.exploding_gas = ox.AccelerationStructure(state.ctx, state.gas_build_input,
272276
compact=True, allow_update=True, random_vertex_access=True)
273-
state.ias_build_input.instances[3].update_traversable(state.exploding_gas)
277+
state.ias_build_input[3].traversable = state.exploding_gas
274278
state.ias_build_input.update_instance(3)
275279
else:
276280
state.exploding_gas.update(state.gas_build_input)
277281

278282
state.ias.update(state.ias_build_input)
279283

284+
280285
def build_vertex_generation_kernel(state):
281286
cuda_source = os.path.join(script_dir, 'cuda', 'dynamic_geometry_vertex_generation.cu')
282287
example_include_path = os.path.dirname(cuda_source)
@@ -289,6 +294,7 @@ def build_vertex_generation_kernel(state):
289294
state.generate_vertices_kernel = cp.RawKernel(code=code, backend='nvrtc',
290295
options=build_flags, name='generate_vertices')
291296

297+
292298
def build_mesh_accel(state):
293299
# Allocate temporary space for vertex generation.
294300
# The same memory space is reused for generating the deformed and exploding vertices before updates.
@@ -335,9 +341,10 @@ def create_module(state):
335341
else:
336342
exception_flags=ox.ExceptionFlags.NONE
337343

344+
print("Triangle value", ox.PrimitiveTypeFlags.TRIANGLE.value)
338345
pipeline_opts = ox.PipelineCompileOptions(
339346
uses_motion_blur=False,
340-
uses_primitive_type_flags = ox.PrimitiveTypeFlags.TRIANGLE,
347+
uses_primitive_type_flags =ox.PrimitiveTypeFlags.TRIANGLE,
341348
traversable_graph_flags=ox.TraversableGraphFlags.ALLOW_SINGLE_LEVEL_INSTANCING,
342349
exception_flags=exception_flags,
343350
num_payload_values=3,
@@ -347,7 +354,7 @@ def create_module(state):
347354
compile_opts = ox.ModuleCompileOptions(
348355
max_register_count=ox.ModuleCompileOptions.DEFAULT_MAX_REGISTER_COUNT,
349356
opt_level=ox.CompileOptimizationLevel.DEFAULT,
350-
debug_level=ox.CompileDebugLevel.LINEINFO)
357+
debug_level=ox.CompileDebugLevel.MODERATE)
351358

352359
cuda_source = os.path.join(script_dir, 'cuda', 'dynamic_geometry.cu')
353360
state.module = ox.Module(state.ctx, cuda_source, compile_opts, pipeline_opts)
@@ -364,7 +371,7 @@ def create_pipeline(state):
364371
program_grps = [state.raygen_grp, state.miss_grp, state.hit_grp]
365372

366373
link_opts = ox.PipelineLinkOptions(max_trace_depth=1,
367-
debug_level=ox.CompileDebugLevel.LINEINFO)
374+
debug_level=ox.CompileDebugLevel.MODERATE)
368375

369376
pipeline = ox.Pipeline(state.ctx,
370377
compile_options=state.pipeline_opts,

examples/dynamic_materials.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -194,7 +194,7 @@ def create_module(state):
194194
compile_opts = ox.ModuleCompileOptions(
195195
max_register_count=ox.ModuleCompileOptions.DEFAULT_MAX_REGISTER_COUNT,
196196
opt_level=ox.CompileOptimizationLevel.DEFAULT,
197-
debug_level=ox.CompileDebugLevel.LINEINFO)
197+
debug_level=ox.CompileDebugLevel.MODERATE)
198198

199199
source = os.path.join(script_dir, 'cuda', 'dynamic_materials.cu')
200200
state.module = ox.Module(state.ctx, source, compile_opts, pipeline_opts)

optix/build.pxd

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -303,12 +303,12 @@ cdef class BuildInputCurveArray(BuildInputArray):
303303

304304
cdef class Instance(OptixObject):
305305
cdef OptixInstance instance
306-
cdef AccelerationStructure traversable
306+
cdef AccelerationStructure _traversable
307307

308308

309309
cdef class BuildInputInstanceArray(BuildInputArray):
310310
cdef OptixBuildInputInstanceArray build_input
311-
cdef public object instances
311+
cdef object instances
312312
cdef object _d_instances
313313

314314

optix/build.pyx

Lines changed: 59 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -459,11 +459,9 @@ cdef class Instance(OptixObject):
459459
visibility_mask=None):
460460
if transform is None:
461461
transform = np.eye(3, 4, dtype=np.float32)
462-
transform = np.ascontiguousarray(np.asarray(transform, dtype=np.float32).reshape(3,4))
463-
cdef float[:, ::1] c_transform = transform
464-
memcpy(&self.instance.transform, &c_transform[0, 0], sizeof(float) * 12)
462+
self.transform = transform
465463
self.traversable = traversable
466-
self.instance.traversableHandle = self.traversable.handle
464+
467465
self.instance.instanceId = instance_id
468466
self.instance.flags = flags.value
469467
self.instance.sbtOffset = sbt_offset
@@ -474,8 +472,14 @@ cdef class Instance(OptixObject):
474472
raise ValueError(f"Too many entries in visibility mask. Got {visibility_mask.bit_length()} but supported are only {max_visibility_mask_bits}")
475473
self.instance.visibilityMask = visibility_mask
476474

477-
def update_traversable(self, AccelerationStructure traversable):
478-
self.traversable = traversable
475+
@property
476+
def traversable(self):
477+
return self._traversable
478+
479+
@traversable.setter
480+
def traversable(self, AccelerationStructure traversable):
481+
self._traversable = traversable
482+
# update the handle as well
479483
self.instance.traversableHandle = self.traversable.handle
480484

481485
def __deepcopy__(self, memodict={}):
@@ -488,6 +492,17 @@ cdef class Instance(OptixObject):
488492

489493
return result
490494

495+
@property
496+
def transform(self):
497+
cdef float [:] transform_view = self.instance.transform
498+
return np.asarray(transform_view).reshape(3,4)
499+
500+
@transform.setter
501+
def transform(self, tf):
502+
transform = np.ascontiguousarray(np.asarray(tf, dtype=np.float32).reshape(3,4))
503+
cdef float[:, ::1] c_transform = transform
504+
memcpy(&self.instance.transform, &c_transform[0, 0], sizeof(float) * 12)
505+
491506

492507
cdef class BuildInputInstanceArray(BuildInputArray):
493508
"""
@@ -521,13 +536,48 @@ cdef class BuildInputInstanceArray(BuildInputArray):
521536

522537
cdef size_t num_elements(self):
523538
return self.build_input.numInstances
524-
539+
540+
def __getitem__(self, index):
541+
return self.instances[index]
542+
543+
def __setitem__(self, index, instance):
544+
if not isinstance(instance, Instance):
545+
raise TypeError("Only instance objects.")
546+
self.instances[index] = instance
547+
self.update_instance(index)
548+
525549
def update_instance(self, index):
526-
src_ptr = <size_t>&((<Instance>(self.instances[index])).instance)
550+
"""
551+
Update the instance at index in gpu memory from the instances list in host memory.
552+
553+
Parameters
554+
----------
555+
index: int
556+
The index to update
557+
"""
558+
# update the value in the cuda buffer
559+
src_ptr = <size_t>&(<Instance>(self.instances[index])).instance
527560
dst_ptr = self._d_instances.ptr + index*sizeof(OptixInstance)
528561
cp.cuda.runtime.memcpy(dst_ptr, src_ptr, sizeof(OptixInstance), cp.cuda.runtime.memcpyHostToDevice)
529562

530-
def get_transform_view(self, index):
563+
# TODO: still thinking of a better way to acomplish the transform access in an OO way.
564+
def view_instance_transform(self, index):
565+
"""
566+
Obtain a view of the transform parameter at index in gpu memory as a cupy array for direct modification
567+
568+
Parameters
569+
----------
570+
index: int
571+
The index of the transform this view should point to
572+
573+
Returns
574+
-------
575+
transform_view: cp.ndarray of shape (3, 4)
576+
A view of the transform matrix
577+
578+
"""
579+
if index < 0 or index >=len(self.instances):
580+
raise IndexError(f"Invalid index {index} for list of length {len(self.instances)}.")
531581
device_ptr = cp.cuda.MemoryPointer(mem=self._d_instances.mem, offset=<int>index*sizeof(OptixInstance))
532582
return cp.ndarray(shape=(3,4), dtype=np.float32, memptr=device_ptr)
533583

optix/context.pxd

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -59,4 +59,4 @@ cdef class DeviceContext(OptixObject):
5959

6060

6161
cdef class OptixContextObject(OptixObject):
62-
cdef DeviceContext context
62+
cdef public DeviceContext context

optix/context.pyx

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,9 +11,6 @@ OPTIX_VERSION = _OPTIX_VERSION
1111
def optix_version():
1212
return _OPTIX_VERSION_MAJOR, _OPTIX_VERSION_MINOR, _OPTIX_VERSION_MICRO
1313

14-
15-
16-
1714
cdef class _LogWrapper:
1815
def __init__(self, log_function):
1916
self.log_function = log_function
@@ -262,3 +259,4 @@ cdef class OptixContextObject(OptixObject):
262259
"""
263260
def __init__(self, DeviceContext context):
264261
self.context = context
262+

optix/pipeline.pyx

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,7 @@ IF _OPTIX_VERSION > 70300: # switch to new instance flags
5757
ROUND_CUBIC_BSPLINE = OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CUBIC_BSPLINE,
5858
ROUND_LINEAR = OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_LINEAR,
5959
ROUND_CATMULLROM = OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CATMULLROM,
60-
TRIANGLE = OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE
60+
TRIANGLE = <unsigned int>OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE # fixes negative number error
6161
ELSE:
6262
class CompileDebugLevel(IntEnum):
6363
"""

0 commit comments

Comments
 (0)