1 // Copyright Contributors to the Open Shading Language project. 2 // SPDX-License-Identifier: BSD-3-Clause 3 // https://github.com/AcademySoftwareFoundation/OpenShadingLanguage 4 5 #pragma once 6 7 #include <OSL/oslconfig.h> 8 9 #ifdef __CUDA_ARCH__ 10 #include <optix.h> 11 #endif 12 13 // USAGE NOTES: 14 // 15 // To define a "standard" DeviceString, add a STRDECL to <OSL/strdecls.h> 16 // specifying the string literal and the name to use for the variable. 17 18 19 OSL_NAMESPACE_ENTER 20 21 22 // Strings are stored in a block of CUDA device memory. String variables hold a 23 // pointer to the start of the char array for each string. Each canonical string 24 // has a unique entry in the table, so two strings can be tested for equality by 25 // comparing their addresses. 26 // 27 // As a convenience, the ustring hash and the length of the string are also 28 // stored in the table, in the 16 bytes preceding the characters. 29 30 struct DeviceString { 31 const char* m_chars; 32 hashDeviceString33 OSL_HOSTDEVICE uint64_t hash() const 34 { 35 return *(uint64_t*)(m_chars - sizeof(uint64_t) - sizeof(uint64_t)); 36 } 37 lengthDeviceString38 OSL_HOSTDEVICE uint64_t length() const 39 { 40 return *(uint64_t*)(m_chars - sizeof(uint64_t)); 41 } 42 c_strDeviceString43 OSL_HOSTDEVICE const char* c_str() const 44 { 45 return m_chars; 46 } 47 48 OSL_HOSTDEVICE bool operator== (const DeviceString& other) const 49 { 50 return m_chars == other.m_chars; 51 } 52 53 OSL_HOSTDEVICE bool operator!= (const DeviceString& other) const 54 { 55 return m_chars != other.m_chars; 56 } 57 }; 58 59 60 // Choose the right cast for string parameters depending on the target. The 61 // macro is the same as the USTR macro defined in oslexec_pvt.h when compiling 62 // for the host. 63 #ifndef __CUDA_ARCH__ 64 # define HDSTR(cstr) (*((ustring *)&cstr)) 65 #else 66 # define HDSTR(cstr) (*(OSL::DeviceString*)&cstr) 67 #endif 68 69 70 // When compiling shadeops C++ sources for CUDA devices, we need to use 71 // DeviceString instead of ustring for some input parameters, so we use this 72 // typedef to select the correct type depending on the target. 73 #ifndef __CUDA_ARCH__ 74 typedef ustring StringParam; 75 #else 76 typedef DeviceString StringParam; 77 #endif 78 79 80 #ifdef __CUDA_ARCH__ 81 namespace DeviceStrings { 82 #define STRDECL(str,var_name) \ 83 extern __device__ OSL_NAMESPACE::DeviceString var_name; 84 #include <OSL/strdecls.h> 85 #undef STRDECL 86 } 87 #endif 88 89 90 OSL_NAMESPACE_EXIT 91 92 93 #ifdef __CUDA_ARCH__ 94 namespace StringParams = OSL_NAMESPACE::DeviceStrings; 95 #else 96 namespace StringParams = OSL_NAMESPACE::Strings; 97 #endif 98