Move closure pool to render state.

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

Author: curtisblack

src/testrender/cuda/optix_raytracer.cu

@@ -48,16 +48,18 @@ __device__ __constant__ RenderParams render_params;
 
 
 static inline __device__ void
-execute_shader(ShaderGlobalsType& sg, const int shader_id, char* closure_pool)
+execute_shader(ShaderGlobalsType& sg, const int shader_id, StackClosurePool& closure_pool)
 {
     if (shader_id < 0) {
         // TODO: should probably never get here ...
         return;
     }
 
-    // Pack the "closure pool" into one of the ShaderGlobals pointers
-    *(int*)&closure_pool[0] = 0;
-    sg.renderstate          = &closure_pool[0];
+    closure_pool.reset();
+    RenderState renderState;
+    // TODO: renderState.context = ...
+    renderState.closure_pool = &closure_pool;
+    sg.renderstate           = &renderState;
 
     // Pack the pointers to the options structs in a faux "context",
     // which is a rough stand-in for the host ShadingContext.

src/testrender/simpleraytracer.cpp

@@ -945,7 +945,7 @@ SimpleRaytracer::eval_background(const Dual2<Vec3>& dir, ShadingContext* ctx,
 #ifndef __CUDACC__
     shadingsys->execute(*ctx, *m_shaders[backgroundShaderID].surf, sg);
 #else
-    alignas(8) char closure_pool[256];
+    StackClosurePool closure_pool;
     execute_shader(sg, render_params.bg_id, closure_pool);
 #endif
     return process_background_closure((const ClosureColor*)sg.Ci);
@@ -957,8 +957,8 @@ SimpleRaytracer::subpixel_radiance(float x, float y, Sampler& sampler,
 {
 #ifdef __CUDACC__
     // Scratch space for the output closures
-    alignas(8) char closure_pool[256];
-    alignas(8) char light_closure_pool[256];
+    StackClosurePool closure_pool;
+    StackClosurePool light_closure_pool;
 #endif
 
     constexpr float inf = std::numeric_limits<float>::infinity();

src/testshade/cuda/optix_grid_renderer.cu

@@ -98,7 +98,7 @@ __raygen__()
     //       networks, so there should be (at least) some mechanism to issue a
     //       warning or error if the closure or param storage can possibly be
     //       exceeded.
-    alignas(8) char closure_pool[256];
+    StackClosurePool closure_pool;
     alignas(8) char params[256];
 
     OSL_CUDA::ShaderGlobals sg;
@@ -137,8 +137,11 @@ __raygen__()
     sg.object2common = reinterpret_cast<void*>(render_params.object2common);
 
     // Pack the "closure pool" into one of the ShaderGlobals pointers
-    *(int*)&closure_pool[0] = 0;
-    sg.renderstate          = &closure_pool[0];
+    closure_pool.reset();
+    RenderState renderState;
+    // TODO: renderState.context = ...
+    renderState.closure_pool = &closure_pool;
+    sg.renderstate           = &renderState;
 
     // Run the OSL group and init functions
     if (render_params.fused_callable)

src/testshade/render_state.h

@@ -10,7 +10,7 @@
 // All the the state free functions in rs_simplerend.cpp will need to do their job
 // NOTE:  Additional data is here that will be used by rs_simplerend.cpp in future PR's
 //        procedurally generating ShaderGlobals.
-struct RenderState {
+struct RenderContext {
     int xres;
     int yres;
     OSL::Matrix44 world_to_camera;
@@ -24,6 +24,35 @@ struct RenderState {
     void* journal_buffer;
 };
 
+class StackClosurePool {
+    alignas(8) char buffer[256];
+    void* ptr;
+
+public:
+    StackClosurePool() { reset(); }
+
+    void reset()
+    {
+      ptr = &buffer[0];
+      *(int*)ptr = 0;
+    }
+
+    void* allocate(size_t size, size_t alignment)
+    {
+        uintptr_t p = OIIO::round_to_multiple_of_pow2((uintptr_t)ptr,
+                                                      alignment);
+        ptr         = (void*)(p + size);
+        if (ptr <= &buffer[256])
+            return p;
+        return nullptr;
+    }
+}
+
+struct RenderState {
+    RenderContext* context;
+    StackClosurePool* closure_pool;
+};
+
 
 // Create constexpr hashes for all strings used by the free function renderer services.
 // NOTE:  Actually ustring's should also be instantiated in host code someplace as well

src/testshade/rs_simplerend.cpp

@@ -66,7 +66,7 @@ rs_get_inverse_matrix_space_time(OSL::OpaqueExecContextPtr ec,
     using OSL::Matrix44;
 
 
-    auto rs = OSL::get_rs<RenderState>(ec);
+    auto rs = OSL::get_rs<RenderState>(ec)->context;
     if (to == OSL::Hashes::camera || to == OSL::Hashes::screen
         || to == OSL::Hashes::NDC || to == RS::Hashes::raster) {
         Matrix44 M { rs->world_to_camera };
@@ -372,11 +372,8 @@ rs_trace_get(OSL::OpaqueExecContextPtr ec, OSL::ustringhash name,
 OSL_RSOP OSL_HOSTDEVICE void*
 rs_allocate_closure(OSL::OpaqueExecContextPtr ec, size_t size, size_t alignment)
 {
-    auto sg       = (OSL::ShaderGlobals*)ec;
-    uintptr_t ptr = OIIO::round_to_multiple_of_pow2((uintptr_t)sg->renderstate,
-                                                    alignment);
-    sg->renderstate = (void*)(ptr + size);
-    return (void*)ptr;
+    auto rs = OSL::get_rs<RenderState>(ec);
+    return rs->closure_pool->allocate(size, alignment);
 }
 #endif
 
@@ -503,7 +500,7 @@ rs_get_attribute(OSL::OpaqueExecContextPtr oec, OSL::ustringhash_pod object_,
     auto object              = OSL::ustringhash_from(object_);
     auto name                = OSL::ustringhash_from(name_);
     const OSL::TypeDesc type = OSL::TypeDesc_from(_type);
-    auto rs                  = OSL::get_rs<RenderState>(oec);
+    auto rs                  = OSL::get_rs<RenderState>(oec)->context;
 
     // The many branches in the code below handle the case where we don't know
     // the attribute name at compile time. In the case it is known, dead-code
@@ -648,7 +645,7 @@ rs_errorfmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash fmt_specification,
             int32_t arg_count, const OSL::EncodedType* argTypes,
             uint32_t argValuesSize, uint8_t* argValues)
 {
-    auto rs = OSL::get_rs<RenderState>(ec);
+    auto rs = OSL::get_rs<RenderState>(ec)->context;
 
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_errorfmt(OSL::get_thread_index(ec), OSL::get_shade_index(ec),
@@ -661,7 +658,7 @@ rs_warningfmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash fmt_specification,
               int32_t arg_count, const OSL::EncodedType* argTypes,
               uint32_t argValuesSize, uint8_t* argValues)
 {
-    auto rs = OSL::get_rs<RenderState>(ec);
+    auto rs = OSL::get_rs<RenderState>(ec)->context;
 
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_warningfmt(OSL::get_max_warnings_per_thread(ec),
@@ -676,7 +673,7 @@ rs_printfmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash fmt_specification,
             int32_t arg_count, const OSL::EncodedType* argTypes,
             uint32_t argValuesSize, uint8_t* argValues)
 {
-    auto rs = OSL::get_rs<RenderState>(ec);
+    auto rs = OSL::get_rs<RenderState>(ec)->context;
 
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_printfmt(OSL::get_thread_index(ec), OSL::get_shade_index(ec),
@@ -691,7 +688,7 @@ rs_filefmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash filename_hash,
            const OSL::EncodedType* argTypes, uint32_t argValuesSize,
            uint8_t* argValues)
 {
-    auto rs = OSL::get_rs<RenderState>(ec);
+    auto rs = OSL::get_rs<RenderState>(ec)->context;
 
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_filefmt(OSL::get_thread_index(ec), OSL::get_shade_index(ec),

src/testshade/simplerend.cpp

@@ -1049,7 +1049,7 @@ SimpleRenderer::add_output(string_view varname_, string_view filename,
 
 
 void
-SimpleRenderer::export_state(RenderState& state) const
+SimpleRenderer::export_state(RenderContext& state) const
 {
     state.xres   = m_xres;
     state.yres   = m_yres;
@@ -1073,7 +1073,7 @@ SimpleRenderer::errorfmt(OSL::ShaderGlobals* sg,
                          const EncodedType* arg_types, uint32_t arg_values_size,
                          uint8_t* argValues)
 {
-    RenderState* rs = reinterpret_cast<RenderState*>(sg->renderstate);
+    RenderContext* rs = reinterpret_cast<RenderState*>(sg->renderstate)->context;
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_errorfmt(OSL::get_thread_index(sg), OSL::get_shade_index(sg),
                        fmt_specification, arg_count, arg_types, arg_values_size,
@@ -1086,7 +1086,7 @@ SimpleRenderer::warningfmt(OSL::ShaderGlobals* sg,
                            int32_t arg_count, const EncodedType* arg_types,
                            uint32_t arg_values_size, uint8_t* argValues)
 {
-    RenderState* rs = reinterpret_cast<RenderState*>(sg->renderstate);
+    RenderContext* rs = reinterpret_cast<RenderState*>(sg->renderstate)->context;
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_warningfmt(OSL::get_max_warnings_per_thread(sg),
                          OSL::get_thread_index(sg), OSL::get_shade_index(sg),
@@ -1102,7 +1102,7 @@ SimpleRenderer::printfmt(OSL::ShaderGlobals* sg,
                          const EncodedType* arg_types, uint32_t arg_values_size,
                          uint8_t* argValues)
 {
-    RenderState* rs = reinterpret_cast<RenderState*>(sg->renderstate);
+    RenderContext* rs = reinterpret_cast<RenderState*>(sg->renderstate)->context;
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_printfmt(OSL::get_thread_index(sg), OSL::get_shade_index(sg),
                        fmt_specification, arg_count, arg_types, arg_values_size,
@@ -1115,7 +1115,7 @@ SimpleRenderer::filefmt(OSL::ShaderGlobals* sg, OSL::ustringhash filename_hash,
                         const EncodedType* arg_types, uint32_t arg_values_size,
                         uint8_t* argValues)
 {
-    RenderState* rs = reinterpret_cast<RenderState*>(sg->renderstate);
+    RenderContext* rs = reinterpret_cast<RenderState*>(sg->renderstate)->context;
     OSL::journal::Writer jw { rs->journal_buffer };
     jw.record_filefmt(OSL::get_thread_index(sg), OSL::get_shade_index(sg),
                       filename_hash, fmt_specification, arg_count, arg_types,

src/testshade/simplerend.h

@@ -145,7 +145,7 @@ class SimpleRenderer : public RendererServices {
     size_t noutputs() const { return m_outputbufs.size(); }
 
     virtual void init_shadingsys(ShadingSystem* ss) { shadingsys = ss; }
-    virtual void export_state(RenderState&) const;
+    virtual void export_state(RenderContext&) const;
     virtual void prepare_render() {}
     virtual void warmup() {}
     virtual void render(int /*xres*/, int /*yres*/) {}

src/testshade/testshade.cpp

@@ -945,21 +945,23 @@ setup_transformations(SimpleRenderer& rend, OSL::Matrix44& Mshad,
     rend.name_transform("myspace", Mmyspace);
 }
 
-// NOTE:  each host thread could end up with its own RenderState.
-//        Starting simple with a single instance for now
-static RenderState theRenderState;
+// A single render context shared by all render threads.
+static RenderContext theRenderState;
 
 
 // Set up the ShaderGlobals fields for pixel (x,y).
 static void
-setup_shaderglobals(ShaderGlobals& sg, ShadingSystem* shadingsys, int x, int y)
+setup_shaderglobals(ShaderGlobals& sg, ShadingSystem* shadingsys,
+                    RenderState& renderState, int x, int y)
 {
     // Just zero the whole thing out to start
     memset((char*)&sg, 0, sizeof(ShaderGlobals));
 
     // Any state data needed by SimpleRenderer or its free function equivalent
     // will need to be passed here the ShaderGlobals.
-    sg.renderstate = &theRenderState;
+    renderState.context = &theRenderState;
+    renderState.closure_pool = nullptr; // Use inbuilt closure pool.
+    sg.renderstate      = &renderState;
 
     // Set "shader" space to be Mshad.  In a real renderer, this may be
     // different for each shader group.
@@ -1182,7 +1184,8 @@ setup_output_images(SimpleRenderer* rend, ShadingSystem* shadingsys,
         ShadingContext* ctx             = shadingsys->get_context(thread_info);
         raytype_bit = shadingsys->raytype_bit(ustring(raytype_name));
         ShaderGlobals sg;
-        setup_shaderglobals(sg, shadingsys, 0, 0);
+        RenderState renderState;
+        setup_shaderglobals(sg, shadingsys, renderState, 0, 0);
 
 #if OSL_USE_BATCHED
         if (batched) {
@@ -1586,6 +1589,7 @@ shade_region(SimpleRenderer* rend, ShaderGroup* shadergroup, OIIO::ROI roi,
 
     // Set up shader globals and a little test grid of points to shade.
     ShaderGlobals shaderglobals;
+    RenderState renderState;
 
     raytype_bit = shadingsys->raytype_bit(ustring(raytype_name));
 
@@ -1606,7 +1610,7 @@ shade_region(SimpleRenderer* rend, ShaderGroup* shadergroup, OIIO::ROI roi,
             // set it up rigged to look like we're rendering a single
             // quadrilateral that exactly fills the viewport, and that
             // setup is done in the following function call:
-            setup_shaderglobals(shaderglobals, shadingsys, x, y);
+            setup_shaderglobals(shaderglobals, shadingsys, renderState, x, y);
 
             if (this_threads_index == uninitialized_thread_index) {
                 this_threads_index = next_thread_index.fetch_add(1u);

testsuite/example-cuda/cuda_grid_renderer.cu

@@ -74,7 +74,7 @@ shade(float3* Cout, int w, int h)
     //       networks, so there should be (at least) some mechanism to issue a
     //       warning or error if the closure or param storage can possibly be
     //       exceeded.
-    alignas(8) char closure_pool[256];
+    StackClosurePool closure_pool;
     alignas(8) char params[256];
 
     const float invw = 1.0 / w;
@@ -115,8 +115,11 @@ shade(float3* Cout, int w, int h)
     sg.flipHandedness = 0;
 
     // Pack the "closure pool" into one of the ShaderGlobals pointers
-    *(int*)&closure_pool[0] = 0;
-    sg.renderstate          = &closure_pool[0];
+    closure_pool.reset();
+    RenderState renderState;
+    // TODO: renderState.context = ...
+    renderState.closure_pool = &closure_pool;
+    sg.renderstate           = &renderState;
 
     // Run the shader
     __osl__init(&sg, params);