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