Fix the OptiX path in testrender and testshade (#1896)

The recent free function change (#1852) put the OptiX paths in testrender and testshade in a non-working state (issue #1894). The bitcode for the free functions wasn't plumbed through for OptiX. This PR rolls the free function definitions into the existing "rendlib" bitcode and PTX. This is somewhat of a stopgap fix; eventually we'll want to handle the free function more explicitly.

Fixes #1894 

Signed-off-by: Tim Grant <[email protected]>

Author: tgrant-nv

src/include/OSL/oslconfig.h.in

@@ -210,7 +210,7 @@ static_assert(sizeof(TypeDesc_pod) == sizeof(TypeDesc),
               "TypeDesc size differs from its POD counterpart");
 
 /// Convenience function to convert to a TypeDesc.
-inline TypeDesc
+OSL_HOSTDEVICE inline TypeDesc
 TypeDesc_from(TypeDesc_pod type)
 {
     return OSL::bitcast<OSL::TypeDesc>(type);

src/testrender/CMakeLists.txt

@@ -21,6 +21,7 @@ if (OSL_USE_OPTIX)
 
     set (testrender_rend_lib_srcs
         cuda/rend_lib.cu
+        ../testshade/rs_simplerend.cpp
         )
 
     # We need to make sure that the PTX files are regenerated whenever these

src/testrender/cuda/rend_lib.cu

@@ -364,35 +364,6 @@ make_float3(const float4& a)
 
 
 
-// FIXME:
-// clang++ 9.0 seems to have trouble with tex2d<float4>() look-ups,
-// so we'll declare this external and implement texture lookups in
-// CUDA files compiled by nvcc (optix_grid_renderer.cu and
-// optix_raytrace.cu).
-// (clang++ 9.0 error 'undefined __nv_tex_surf_handler')
-extern __device__ float4
-osl_tex2DLookup(void* handle, float s, float t);
-
-__device__ int
-osl_texture(void* sg_, OSL::ustringhash_pod name, void* handle, void* opt_,
-            float s, float t, float dsdx, float dtdx, float dsdy, float dtdy,
-            int chans, void* result, void* dresultdx, void* dresultdy,
-            void* alpha, void* dalphadx, void* dalphady,
-            void* ustringhash_errormessage)
-{
-    if (!handle)
-        return 0;
-    // cudaTextureObject_t texID = cudaTextureObject_t(handle);
-    float4 fromTexture = osl_tex2DLookup(handle, s, t);
-    // see note above
-    // float4 fromTexture = tex2D<float4>(texID, s, t);
-    *((float3*)result) = make_float3(fromTexture.x, fromTexture.y,
-                                     fromTexture.z);
-    return 1;
-}
-
-
-
 __device__ int
 osl_range_check_err(int indexvalue, int length, OSL::ustringhash_pod symname,
                     void* sg, OSL::ustringhash_pod sourcefile, int sourceline,

src/testshade/CMakeLists.txt

@@ -21,6 +21,7 @@ if (OSL_USE_OPTIX)
 
     set (testshade_rend_lib_srcs
         ../testrender/cuda/rend_lib.cu
+        rs_simplerend.cpp
         )
 
     set ( testshade_cuda_headers

src/testshade/cuda/optix_grid_renderer.cu

@@ -168,8 +168,6 @@ __raygen__()
     output_buffer[pixel] = { f_output[1], f_output[2], f_output[3] };
 }
 
-
-
 // Because clang++ 9.0 seems to have trouble with some of the texturing "intrinsics"
 // let's do the texture look-ups in this file.
 extern "C" __device__ float4

src/testshade/rs_simplerend.cpp

@@ -2,14 +2,16 @@
 // SPDX-License-Identifier: BSD-3-Clause
 // https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
 
-#ifndef OSL_HOST_RS_BITCODE
+#if !defined(__CUDACC__) && !defined(OSL_HOST_RS_BITCODE)
 #    error OSL_HOST_RS_BITCODE must be defined by your build system.
 #endif
 
 #include <OpenImageIO/fmath.h>
 
-#include <OSL/fmt_util.h>
-#include <OSL/journal.h>
+#ifndef __CUDACC__
+#    include <OSL/fmt_util.h>
+#    include <OSL/journal.h>
+#endif
 #include <OSL/rendererservices.h>
 #include <OSL/rs_free_function.h>
 
@@ -181,6 +183,14 @@ rs_transform_points(OSL::OpaqueExecContextPtr /*ec*/, OSL::ustringhash /*from*/,
     return false;
 }
 
+#ifdef __CUDACC__
+// This texture lookup function needs to be compiled by NVCC because clang
+// doesn't know how to handle CUDA texture intrinsics. This function must be
+// defined in the CUDA source for testshade and testrender.
+extern "C" __device__ float4
+osl_tex2DLookup(void* handle, float s, float t);
+#endif
+
 OSL_RSOP OSL_HOSTDEVICE bool
 rs_texture(OSL::OpaqueExecContextPtr ec, OSL::ustringhash filename,
            OSL::TextureSystem::TextureHandle* texture_handle,
@@ -196,7 +206,12 @@ rs_texture(OSL::OpaqueExecContextPtr ec, OSL::ustringhash filename,
                                  nchannels, result, dresultds, dresultdt,
                                  errormessage);
 #else
-    return false;
+    if (!texture_handle)
+        return false;
+    const float4 fromTexture = osl_tex2DLookup((void*)texture_handle, s, t);
+    *((float3*)result)       = make_float3(fromTexture.x, fromTexture.y,
+                                           fromTexture.z);
+    return true;
 #endif
 }
 
@@ -613,6 +628,7 @@ rs_get_interpolated_test(void* val)
     return true;
 }
 
+#ifndef __CUDACC__
 OSL_RSOP OSL_HOSTDEVICE void
 rs_errorfmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash fmt_specification,
             int32_t arg_count, const OSL::EncodedType* argTypes,
@@ -668,3 +684,4 @@ rs_filefmt(OSL::OpaqueExecContextPtr ec, OSL::ustringhash filename_hash,
                       filename_hash, fmt_specification, arg_count, argTypes,
                       argValuesSize, argValues);
 }
+#endif  // #ifndef __CUDACC__