api: Ustringhash phase 3 (#1732)

Move OSL library functions and remaining execution time interface from ustring to ustringhash. Intent is all code for cpu or device just use ustringhash only. This allows us to remove STRING_PARAM, ustringrep, and OSL_USTRINGREP_IS_HASH. All interfaces and library code should explicitly use ustringhash or ustring, or for library functions llvm can call ustringhash_pod or ustring_pod Compile time strings used in the OSL Library are represented in strdecls.h and will be available as ustring variables in the Strings namespace as well as available as unstringhash constexpr variables in the Hashes namespace. Compile time strings used by the Renderer should adopt a similiar strategy as exemplified by TestShade which keeps its strings in rs_strdecls.h and made available as ustring variables in the RS::Strings namespace as well as available as unstringhash constexpr variables in the RS::Hashes namespace. Added helper utility functions to generate ustring from ustringhash_pods Constant string arrays: Previously, array copys required copying data via address of symbol on host. To enable working on device, replaced memcpy with array allocation and loading of values. Constant arrays are now present as global arrays in final llvm module. On the Batched side compatability/integration: Batched OSL Library functions continue to use ustrings (not ustringhashes to reduce scope/risk). Batched OSL library functions that call their scalar counterparts were adjusted to pass arguments as ustringhash_pods and store results as ustrings STRING_PARAMS, device_string.h, and string_hash.h: Built-in strings on the CUDA side are now accessed via Hashes namespace, instead of STRING_PARAMS and are evaluated as ustringhashes. Similarly, DeviceString, StringParams, HDSTR are obsolete, and device_string.h is eliminated. Hashes.h, which only contains OIIO's strhash and pvtstrlen functions (both formerly in string_hash.h), replaces all inclusions of device_string.h. Namespace Hashes in hashes.h contains pre-compiled ustringhashes representing strings from strdecls.h. Texture, dictionary, closure: Helper implementation functions are used by both scalar and batched versions. However, since scalar uses ustringhashes and the batched version uses ustrings added a flag to inform functions on treating TypeDesc::STRING has either a ustring or a ustringhash testing: Added testing of bitcode rendereservices for unoptimized OSL shader. array-copy tests regression: Added new test to verify copying of arrays of unequal lengths Signed-off-by: Steena Monteiro <[email protected]> * Deleted liboslexec/string_hash.h Signed-off-by: Steena Monteiro <[email protected]> * Modified rs_get_attribute to use OpaqueExecutionContext and ustringhash. Changed example programs--testshade, testrenderer, and osltoy--to use device friendly constexpr (RS|OSL)::Hashes and (RS|OSL)::Strings variables with rs_strdecls.h vs constructing ustrings on the fly. Signed-off-by: Steena Monteiro <[email protected]> * Fixed clang format issues; added uint64_t casts for constant64; removed unused variable; added array-copy to testsuite Signed-off-by: Steena Monteiro <[email protected]> * Fixed clang format issues in llvm_instance.cpp Signed-off-by: Steena Monteiro <[email protected]> * Added fix in llvm_assign_impl for assignment of non-constant scalar to an array Signed-off-by: Steena Monteiro <[email protected]> * Fixed clang-format in simplerend.cpp. Signed-off-by: Steena Monteiro <[email protected]> * Addressed review comments -- (1) Added ll.constant(ustring().cstr()) in llvm_gen and batched_llvm_gen and (2) Added comments for boolean arguments that indicate ustrings are treated as hashes Signed-off-by: Steena Monteiro <[email protected]> * Addressed comments on Optix build failures Signed-off-by: Steena Monteiro <[email protected]> * Resolved merge conflicts Signed-off-by: Steena Monteiro <[email protected]> * Fixed clang-format Signed-off-by: Steena Monteiro <[email protected]> * Fixed split-reg for batched Signed-off-by: Steena Monteiro <[email protected]> * Replaced llvm_load_stringhash with llvm_const_hash; Changed format string to ustringhash_pod in osl_printf, osl_fprintf, osl_warning, osl_error stubs; Fixed ShaderGlobals' position, u, and v in optix_grid_renderer.cu. Signed-off-by: Steena Monteiro <[email protected]> --------- Signed-off-by: Steena Monteiro <[email protected]> Signed-off-by: Larry Gritz <[email protected]> Co-authored-by: Larry Gritz <[email protected]>
AcademySoftwareFoundation · Dec 21, 2023 · 12530f1 · 12530f1
1 parent 253c410
commit 12530f1
Show file tree

Hide file tree

Showing 80 changed files with 2,046 additions and 1,797 deletions.
diff --git a/src/cmake/testing.cmake b/src/cmake/testing.cmake
@@ -133,6 +133,17 @@ macro ( TESTSUITE )
             add_one_testsuite ("${_testname}" "${_testsrcdir}"
                                ENV TESTSHADE_OPT=0 )
         endif ()
+        # Run the same test again using bitcode based free function renderer 
+        # services keeping unoptimized, unless it matches a few patterns that
+        # we don't test unoptimized (or has an OPTIMIZEONLY marker file).
+        if (NOT _testname MATCHES "optix"
+            AND NOT EXISTS "${_testsrcdir}/NOSCALAR"
+            AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION"
+            AND NOT EXISTS "${_testsrcdir}/OPTIMIZEONLY"
+            AND NOT EXISTS "${_testsrcdir}/NORSBITCODE")
+            add_one_testsuite ("${_testname}.rs_bitcode" "${_testsrcdir}"
+                               ENV TESTSHADE_OPT=0 TESTSHADE_RS_BITCODE=1)
+        endif ()
         # Run the same test again with aggressive -O2 runtime
         # optimization, triggered by setting TESTSHADE_OPT env variable.
         # Skip OptiX-only tests and those with a NOOPTIMIZE marker file.
@@ -143,7 +154,8 @@ macro ( TESTSUITE )
             add_one_testsuite ("${_testname}.opt" "${_testsrcdir}"
                                ENV TESTSHADE_OPT=2 )
         endif ()
-        # Run the same test again with aggressive -O2 runtime
+        # Run the same test again using bitcode based free function renderer 
+        # services keeping aggressive -O2 runtime
         # optimization, triggered by setting TESTSHADE_OPT env variable.
         # Skip OptiX-only tests and those with a NOOPTIMIZE marker file.
         if (NOT _testname MATCHES "optix"
@@ -218,6 +230,13 @@ macro ( TESTSUITE )
         endif ()
 
         # if there is an RS_BITCODE marker file in the directory.
+        if ((EXISTS "${_testsrcdir}/RS_BITCODE" OR test_all_rs_bitcode)
+            AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION"
+            AND NOT EXISTS "${_testsrcdir}/RS_BITCODE_REGRESSION"
+            AND NOT EXISTS "${_testsrcdir}/OPTIMIZEONLY")
+            add_one_testsuite ("${_testname}.rsbitcode" "${_testsrcdir}"
+                                ENV TESTSHADE_OPT=0 TESTSHADE_RS_BITCODE=1 )
+        endif ()
         if ((EXISTS "${_testsrcdir}/RS_BITCODE" OR test_all_rs_bitcode)
             AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION"
             AND NOT EXISTS "${_testsrcdir}/RS_BITCODE_REGRESSION"
@@ -244,7 +263,7 @@ macro (osl_add_all_tests)
     # special installed tests.
     TESTSUITE ( aastep allowconnect-err andor-reg and-or-not-synonyms
                 arithmetic area-reg arithmetic-reg
-                array array-reg array-copy-reg array-derivs array-range
+                array array-reg array-copy array-copy-reg array-derivs array-range
                 array-aassign array-assign-reg array-length-reg
                 bitwise-and-reg bitwise-or-reg bitwise-shl-reg  bitwise-shr-reg bitwise-xor-reg
                 blackbody blackbody-reg blendmath breakcont breakcont-reg
@@ -356,8 +375,8 @@ macro (osl_add_all_tests)
                 struct-nested struct-nested-assign struct-nested-deep
                 ternary
                 testshade-expr
-		test-fmt-arrays test-fmt-fileprint
-		test-fmt-cxpf  test-fmt-noise test-fmt-matrixcolor 
+                test-fmt-arrays test-fmt-fileprint
+                test-fmt-cxpf  test-fmt-noise test-fmt-matrixcolor 
                 test-fmt-stpf test-fmt-errorwarning test-fmt-errorwarning-repeats
                 texture-alpha texture-alpha-derivs
                 texture-blur texture-connected-options

diff --git a/src/include/OSL/device_string.h b/src/include/OSL/device_string.h
diff --git a/src/include/OSL/hashes.h b/src/include/OSL/hashes.h
@@ -0,0 +1,67 @@
+// Copyright Contributors to the Open Shading Language project.
+// SPDX-License-Identifier: BSD-3-Clause
+// https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
+
+#pragma once
+
+#include <OSL/oslconfig.h>
+#include <OpenImageIO/detail/farmhash.h>
+#include <OpenImageIO/oiioversion.h>
+
+// USAGE NOTES:
+//
+// To define a "standard" DeviceString, add a STRDECL to <OSL/strdecls.h>
+// specifying the string literal and the name to use for the variable.
+
+
+OSL_NAMESPACE_ENTER
+
+namespace pvt {
+
+OIIO_CONSTEXPR14 inline size_t
+pvtstrlen(const char* s)
+{
+    if (s == nullptr)
+        return 0;
+    size_t len = 0;
+    while (s[len] != 0)
+        len++;
+    return len;
+}
+
+}  // namespace pvt
+
+// The string_view(const char *) is only constexpr for c++17
+// which would prevent OIIO::Strutil::strhash from being
+// constexpr for c++14.
+// workaround by using local version here with a private
+// constexpr strlen
+OIIO_CONSTEXPR14 inline size_t
+strhash(const char* s)
+{
+    size_t len = pvt::pvtstrlen(s);
+    return OIIO::Strutil::strhash(OIIO::string_view(s, len));
+}
+
+// Template to ensure the hash is evaluated at compile time.
+template<size_t V> static constexpr size_t HashConstEval = V;
+#define OSL_HASHIFY(unquoted_string) \
+    HashConstEval<OSL::strhash(__OSL_STRINGIFY(unquoted_string))>
+
+namespace {  // Scope Hashes variables to just this translation unit
+namespace Hashes {
+#ifdef __CUDA_ARCH__  // TODO: restrict to CUDA version < 11.4, otherwise the contexpr should work
+#    define STRDECL(str, var_name) \
+        __device__ const OSL::ustringhash var_name(OSL::strhash(str));
+#else
+#    define STRDECL(str, var_name) \
+        constexpr OSL::ustringhash var_name(OSL::strhash(str));
+#endif
+#include <OSL/strdecls.h>
+#undef STRDECL
+};  // namespace Hashes
+}  // unnamed namespace
+
+
+
+OSL_NAMESPACE_EXIT
diff --git a/src/include/OSL/llvm_util.h b/src/include/OSL/llvm_util.h
@@ -27,6 +27,7 @@ class DISubroutineType;
 class ExecutionEngine;
 class Function;
 class FunctionType;
+class GlobalVariable;
 class SectionMemoryManager;
 class JITEventListener;
 class Linker;
@@ -546,6 +547,7 @@ class OSLEXECPUBLIC LLVM_Util {
     llvm::Type* type_matrix() const { return m_llvm_type_matrix; }
     llvm::Type* type_typedesc() const { return m_llvm_type_longlong; }
     llvm::Type* type_ustring() { return m_llvm_type_ustring; }
+    llvm::Type* type_real_ustring() { return m_llvm_type_real_ustring; }
     llvm::PointerType* type_void_ptr() const { return m_llvm_type_void_ptr; }
     llvm::PointerType* type_ustring_ptr() const
     {
@@ -718,6 +720,10 @@ class OSLEXECPUBLIC LLVM_Util {
     llvm::Value* constant(ustring s);
     llvm::Value* constant(string_view s) { return constant(ustring(s)); }
 
+    llvm::Constant* constant_array(cspan<llvm::Constant*> constants);
+    llvm::GlobalVariable* create_global_constant(llvm::Constant* initializer,
+                                                 const std::string& llname = {});
+
     llvm::Value* llvm_mask_to_native(llvm::Value* llvm_mask);
     llvm::Value* native_to_llvm_mask(llvm::Value* native_mask);
 
@@ -1087,7 +1093,9 @@ class OSLEXECPUBLIC LLVM_Util {
     llvm::Type* m_llvm_type_void;
     llvm::Type* m_llvm_type_triple;
     llvm::Type* m_llvm_type_matrix;
-    llvm::Type* m_llvm_type_ustring;
+    llvm::Type* m_llvm_type_real_ustring;  // True const char *
+    llvm::Type*
+        m_llvm_type_ustring;  // UStringRep can change between int64(hash) and a const char *
     llvm::PointerType* m_llvm_type_void_ptr;
     llvm::PointerType* m_llvm_type_char_ptr;
     llvm::PointerType* m_llvm_type_bool_ptr;