Update optix code to use closure pool.

Signed-off-by: Curtis Black <curtis.w.black@gmail.com>

Author: curtisblack

src/testrender/cuda/rend_lib.cu

@@ -30,17 +30,6 @@ extern __device__ CUdeviceptr xform_buffer;
 }  // namespace pvt
 OSL_NAMESPACE_END
 
-
-// Taken from the SimplePool class
-__device__ static inline size_t
-alignment_offset_calc(void* ptr, size_t alignment)
-{
-    uintptr_t ptrbits = reinterpret_cast<uintptr_t>(ptr);
-    uintptr_t offset  = ((ptrbits + alignment - 1) & -alignment) - ptrbits;
-    return offset;
-}
-
-
 // These functions are declared extern to prevent name mangling.
 extern "C" {
 
@@ -50,200 +39,120 @@ __direct_callable__dummy_rend_lib()
 {
 }
 
-
-__device__ void*
-closure_component_allot(void* pool, int id, size_t prim_size,
-                        const OSL::Color3& w)
-{
-    ((OSL::ClosureComponent*)pool)->id = id;
-    ((OSL::ClosureComponent*)pool)->w  = w;
-
-    size_t needed = (sizeof(OSL::ClosureComponent) + prim_size
-                     + (alignof(OSL::ClosureComponent) - 1))
-                    & ~(alignof(OSL::ClosureComponent) - 1);
-    char* char_ptr = (char*)pool;
-
-    return (void*)&char_ptr[needed];
-}
-
-
-__device__ void*
-closure_mul_allot(void* pool, const OSL::Color3& w, OSL::ClosureColor* c)
-{
-    ((OSL::ClosureMul*)pool)->id      = OSL::ClosureColor::MUL;
-    ((OSL::ClosureMul*)pool)->weight  = w;
-    ((OSL::ClosureMul*)pool)->closure = c;
-
-    size_t needed = (sizeof(OSL::ClosureMul)
-                     + (alignof(OSL::ClosureComponent) - 1))
-                    & ~(alignof(OSL::ClosureComponent) - 1);
-    char* char_ptr = (char*)pool;
-
-    return &char_ptr[needed];
-}
-
-
-__device__ void*
-closure_mul_float_allot(void* pool, const float& w, OSL::ClosureColor* c)
-{
-    ((OSL::ClosureMul*)pool)->id       = OSL::ClosureColor::MUL;
-    ((OSL::ClosureMul*)pool)->weight.x = w;
-    ((OSL::ClosureMul*)pool)->weight.y = w;
-    ((OSL::ClosureMul*)pool)->weight.z = w;
-    ((OSL::ClosureMul*)pool)->closure  = c;
-
-    size_t needed = (sizeof(OSL::ClosureMul)
-                     + (alignof(OSL::ClosureComponent) - 1))
-                    & ~(alignof(OSL::ClosureComponent) - 1);
-    char* char_ptr = (char*)pool;
-
-    return &char_ptr[needed];
-}
-
-
 __device__ void*
-closure_add_allot(void* pool, OSL::ClosureColor* a, OSL::ClosureColor* b)
+osl_add_closure_closure(void* sg_, const void* a_, const void* b_)
 {
-    ((OSL::ClosureAdd*)pool)->id       = OSL::ClosureColor::ADD;
-    ((OSL::ClosureAdd*)pool)->closureA = a;
-    ((OSL::ClosureAdd*)pool)->closureB = b;
-
-    size_t needed = (sizeof(OSL::ClosureAdd)
-                     + (alignof(OSL::ClosureComponent) - 1))
-                    & ~(alignof(OSL::ClosureComponent) - 1);
-    char* char_ptr = (char*)pool;
-
-    return &char_ptr[needed];
+    a_                         = __builtin_assume_aligned(a_, alignof(float));
+    b_                         = __builtin_assume_aligned(b_, alignof(float));
+    ShaderGlobals* sg          = (ShaderGlobals*)sg_;
+    const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
+    const OSL::ClosureColor* b = (const OSL::ClosureColor*)b_;
+    if (a == NULL)
+        return b;
+    if (b == NULL)
+        return a;
+    auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
+    OSL::ClosureAdd* add
+        = (OSL::ClosureAdd*)closure_pool->allocate(sizeof(OSL::ClosureAdd),
+                                                   alignof(OSL::ClosureAdd));
+    if (add) {
+        add->id       = OSL::ClosureColor::ADD;
+        add->closureA = a;
+        add->closureB = b;
+    }
+    return add;
 }
 
-
 __device__ void*
-osl_allocate_closure_component(void* sg_, int id, int size)
+osl_mul_closure_color(void* sg_, const void* a_, const void* w_)
 {
-    OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
-
-    OSL::Color3 w = OSL::Color3(1, 1, 1);
-    // Fix up the alignment
-    void* ret = ((char*)sg_ptr->renderstate)
-                + alignment_offset_calc(sg_ptr->renderstate,
-                                        alignof(OSL::ClosureComponent));
-
-    size = max(4, size);
-
-    sg_ptr->renderstate = closure_component_allot(ret, id, size, w);
+    a_ = __builtin_assume_aligned(a_, alignof(float));
+    w_ = __builtin_assume_aligned(w_, alignof(float));
 
-    return ret;
-}
-
-
-
-__device__ void*
-osl_allocate_weighted_closure_component(void* sg_, int id, int size,
-                                        const void* w)
-{
-    OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
-
-    const OSL::Color3* wc
-        = (const OSL::Color3*)__builtin_assume_aligned(w, alignof(float));
-
-    if (wc->x == 0.0f && wc->y == 0.0f && wc->z == 0.0f) {
+    ShaderGlobals* sg          = (ShaderGlobals*)sg_;
+    const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
+    const OSL::Color3* w       = (const OSL::Color3*)w_;
+    if (a == NULL)
+        return NULL;
+    if (w->x == 0.0f && w->y == 0.0f && w->z == 0.0f)
         return NULL;
+    if (w->x == 1.0f && w->y == 1.0f && w->z == 1.0f)
+        return a;
+    auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
+    OSL::ClosureMul* mul
+        = (OSL::ClosureMul*)closure_pool->allocate(sizeof(OSL::ClosureMul),
+                                                   alignof(OSL::ClosureMul));
+    if (mul) {
+        mul->id      = OSL::ClosureColor::MUL;
+        mul->weight  = *w;
+        mul->closure = a;
     }
-
-    size = max(4, size);
-
-    // Fix up the alignment
-    void* ret = ((char*)sg_ptr->renderstate)
-                + alignment_offset_calc(sg_ptr->renderstate,
-                                        alignof(OSL::ClosureComponent));
-    sg_ptr->renderstate = closure_component_allot(ret, id, size, *wc);
-
-    return ret;
+    return mul;
 }
 
-
-
 __device__ void*
-osl_mul_closure_color(void* sg_, void* a, const void* w)
+osl_mul_closure_float(void* sg_, const void* a_, float w)
 {
-    OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
-    const OSL::Color3* wc
-        = (const OSL::Color3*)__builtin_assume_aligned(w, alignof(float));
+    a_ = __builtin_assume_aligned(a_, alignof(float));
 
-    if (a == NULL) {
+    ShaderGlobals* sg          = (ShaderGlobals*)sg_;
+    const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
+    if (a == NULL)
         return NULL;
-    }
-
-    if (wc->x == 0.0f && wc->y == 0.0f && wc->z == 0.0f) {
+    if (w == 0.0f)
         return NULL;
-    }
-
-    if (wc->x == 1.0f && wc->y == 1.0f && wc->z == 1.0f) {
+    if (w == 1.0f)
         return a;
+    auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
+    OSL::ClosureMul* mul
+        = (OSL::ClosureMul*)closure_pool->allocate(sizeof(OSL::ClosureMul),
+                                                   alignof(OSL::ClosureMul));
+    if (mul) {
+        mul->id      = OSL::ClosureColor::MUL;
+        mul->weight  = OSL::Color3(w);
+        mul->closure = a;
     }
-
-    // Fix up the alignment
-    void* ret = ((char*)sg_ptr->renderstate)
-                + alignment_offset_calc(sg_ptr->renderstate,
-                                        alignof(OSL::ClosureComponent));
-    sg_ptr->renderstate = closure_mul_allot(ret, *wc, (OSL::ClosureColor*)a);
-
-    return ret;
+    return mul;
 }
 
-
-
 __device__ void*
-osl_mul_closure_float(void* sg_, void* a, float w)
+osl_allocate_closure_component(void* sg_, int id, int size)
 {
-    a = __builtin_assume_aligned(a, alignof(float));
-
-    OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
-
-    if (a == NULL || w == 0.0f) {
-        return NULL;
-    }
-
-    if (w == 1.0f) {
-        return a;
+    ShaderGlobals* sg  = (ShaderGlobals*)sg_;
+    auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
+    // Allocate the component and the mul back to back
+    const size_t needed = sizeof(OSL::ClosureComponent) + size;
+    OSL::ClosureComponent* comp
+        = (OSL::ClosureComponent*)
+              closure_pool->allocate(needed, alignof(OSL::ClosureComponent));
+    if (comp) {
+        comp->id = id;
+        comp->w  = OSL::Color3(1.0f);
     }
-
-    // Fix up the alignment
-    void* ret = ((char*)sg_ptr->renderstate)
-                + alignment_offset_calc(sg_ptr->renderstate,
-                                        alignof(OSL::ClosureComponent));
-    sg_ptr->renderstate = closure_mul_float_allot(ret, w,
-                                                  (OSL::ClosureColor*)a);
-
-    return ret;
+    return comp;
 }
 
-
-
 __device__ void*
-osl_add_closure_closure(void* sg_, void* a, void* b)
+osl_allocate_weighted_closure_component(void* sg_, int id, int size,
+                                        const void* w_)
 {
-    a = __builtin_assume_aligned(a, alignof(float));
-    b = __builtin_assume_aligned(b, alignof(float));
-
-    OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
-
-    if (a == NULL) {
-        return b;
-    }
+    w_ = __builtin_assume_aligned(w_, alignof(float));
 
-    if (b == NULL) {
-        return a;
+    ShaderGlobals* sg    = (ShaderGlobals*)sg_;
+    const OSL::Color3* w = (const OSL::Color3*)w_;
+    if (w->x == 0.0f && w->y == 0.0f && w->z == 0.0f)
+        return NULL;
+    auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
+    // Allocate the component and the mul back to back
+    const size_t needed = sizeof(OSL::ClosureComponent) + size;
+    OSL::ClosureComponent* comp
+        = (OSL::ClosureComponent*)
+              closure_pool->allocate(needed, alignof(OSL::ClosureComponent));
+    if (comp) {
+        comp->id = id;
+        comp->w  = *w;
     }
-
-    // Fix up the alignment
-    void* ret = ((char*)sg_ptr->renderstate)
-                + alignment_offset_calc(sg_ptr->renderstate,
-                                        alignof(OSL::ClosureComponent));
-    sg_ptr->renderstate = closure_add_allot(ret, (OSL::ClosureColor*)a,
-                                            (OSL::ClosureColor*)b);
-
-    return ret;
+    return comp;
 }
 
 

src/testshade/render_state.h

@@ -33,8 +33,8 @@ class StackClosurePool {
 
     void reset()
     {
-      ptr = &buffer[0];
-      *(int*)ptr = 0;
+        ptr        = &buffer[0];
+        *(int*)ptr = 0;
     }
 
     void* allocate(size_t size, size_t alignment)
@@ -43,7 +43,7 @@ class StackClosurePool {
                                                       alignment);
         ptr         = (void*)(p + size);
         if (ptr <= &buffer[256])
-            return p;
+            return (void*)p;
         return nullptr;
     }
 }