Skip to content

Commit b112bb5

Browse files
curtisblackcurtisblack
committed
restore cuda code
Signed-off-by: Curtis Black <curtis.w.black@gmail.com>
1 parent 00e3d5e commit b112bb5

3 files changed

Lines changed: 398 additions & 1 deletion

File tree

‎src/include/OSL/rs_free_function.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -317,7 +317,8 @@ rs_trace_get(OSL::OpaqueExecContextPtr oec, OSL::ustringhash name,
317317
OSL::TypeDesc type, void* val, bool derivatives);
318318

319319
/// Allocates memory for a closure color. May return null if no memory could
320-
/// be allocated.
320+
/// be allocated. It is the renderers responsibility to clean up these
321+
/// allocations after a shader is run and the closures have been processed.
321322
OSL_RSOP OSL_HOSTDEVICE void*
322323
rs_allocate_closure(OSL::OpaqueExecContextPtr oec, size_t size,
323324
size_t alignment);

‎src/testrender/cuda/rend_lib.cu‎

Lines changed: 208 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,17 @@ extern __device__ CUdeviceptr xform_buffer;
3030
} // namespace pvt
3131
OSL_NAMESPACE_END
3232

33+
34+
// Taken from the SimplePool class
35+
__device__ static inline size_t
36+
alignment_offset_calc(void* ptr, size_t alignment)
37+
{
38+
uintptr_t ptrbits = reinterpret_cast<uintptr_t>(ptr);
39+
uintptr_t offset = ((ptrbits + alignment - 1) & -alignment) - ptrbits;
40+
return offset;
41+
}
42+
43+
3344
// These functions are declared extern to prevent name mangling.
3445
extern "C" {
3546

@@ -39,6 +50,203 @@ __direct_callable__dummy_rend_lib()
3950
{
4051
}
4152

53+
54+
__device__ void*
55+
closure_component_allot(void* pool, int id, size_t prim_size,
56+
const OSL::Color3& w)
57+
{
58+
((OSL::ClosureComponent*)pool)->id = id;
59+
((OSL::ClosureComponent*)pool)->w = w;
60+
61+
size_t needed = (sizeof(OSL::ClosureComponent) + prim_size
62+
+ (alignof(OSL::ClosureComponent) - 1))
63+
& ~(alignof(OSL::ClosureComponent) - 1);
64+
char* char_ptr = (char*)pool;
65+
66+
return (void*)&char_ptr[needed];
67+
}
68+
69+
70+
__device__ void*
71+
closure_mul_allot(void* pool, const OSL::Color3& w, OSL::ClosureColor* c)
72+
{
73+
((OSL::ClosureMul*)pool)->id = OSL::ClosureColor::MUL;
74+
((OSL::ClosureMul*)pool)->weight = w;
75+
((OSL::ClosureMul*)pool)->closure = c;
76+
77+
size_t needed = (sizeof(OSL::ClosureMul)
78+
+ (alignof(OSL::ClosureComponent) - 1))
79+
& ~(alignof(OSL::ClosureComponent) - 1);
80+
char* char_ptr = (char*)pool;
81+
82+
return &char_ptr[needed];
83+
}
84+
85+
86+
__device__ void*
87+
closure_mul_float_allot(void* pool, const float& w, OSL::ClosureColor* c)
88+
{
89+
((OSL::ClosureMul*)pool)->id = OSL::ClosureColor::MUL;
90+
((OSL::ClosureMul*)pool)->weight.x = w;
91+
((OSL::ClosureMul*)pool)->weight.y = w;
92+
((OSL::ClosureMul*)pool)->weight.z = w;
93+
((OSL::ClosureMul*)pool)->closure = c;
94+
95+
size_t needed = (sizeof(OSL::ClosureMul)
96+
+ (alignof(OSL::ClosureComponent) - 1))
97+
& ~(alignof(OSL::ClosureComponent) - 1);
98+
char* char_ptr = (char*)pool;
99+
100+
return &char_ptr[needed];
101+
}
102+
103+
104+
__device__ void*
105+
closure_add_allot(void* pool, OSL::ClosureColor* a, OSL::ClosureColor* b)
106+
{
107+
((OSL::ClosureAdd*)pool)->id = OSL::ClosureColor::ADD;
108+
((OSL::ClosureAdd*)pool)->closureA = a;
109+
((OSL::ClosureAdd*)pool)->closureB = b;
110+
111+
size_t needed = (sizeof(OSL::ClosureAdd)
112+
+ (alignof(OSL::ClosureComponent) - 1))
113+
& ~(alignof(OSL::ClosureComponent) - 1);
114+
char* char_ptr = (char*)pool;
115+
116+
return &char_ptr[needed];
117+
}
118+
119+
120+
__device__ void*
121+
osl_allocate_closure_component(void* sg_, int id, int size)
122+
{
123+
OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
124+
125+
OSL::Color3 w = OSL::Color3(1, 1, 1);
126+
// Fix up the alignment
127+
void* ret = ((char*)sg_ptr->renderstate)
128+
+ alignment_offset_calc(sg_ptr->renderstate,
129+
alignof(OSL::ClosureComponent));
130+
131+
size = max(4, size);
132+
133+
sg_ptr->renderstate = closure_component_allot(ret, id, size, w);
134+
135+
return ret;
136+
}
137+
138+
139+
140+
__device__ void*
141+
osl_allocate_weighted_closure_component(void* sg_, int id, int size,
142+
const void* w)
143+
{
144+
OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
145+
146+
const OSL::Color3* wc
147+
= (const OSL::Color3*)__builtin_assume_aligned(w, alignof(float));
148+
149+
if (wc->x == 0.0f && wc->y == 0.0f && wc->z == 0.0f) {
150+
return NULL;
151+
}
152+
153+
size = max(4, size);
154+
155+
// Fix up the alignment
156+
void* ret = ((char*)sg_ptr->renderstate)
157+
+ alignment_offset_calc(sg_ptr->renderstate,
158+
alignof(OSL::ClosureComponent));
159+
sg_ptr->renderstate = closure_component_allot(ret, id, size, *wc);
160+
161+
return ret;
162+
}
163+
164+
165+
166+
__device__ void*
167+
osl_mul_closure_color(void* sg_, void* a, const void* w)
168+
{
169+
OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
170+
const OSL::Color3* wc
171+
= (const OSL::Color3*)__builtin_assume_aligned(w, alignof(float));
172+
173+
if (a == NULL) {
174+
return NULL;
175+
}
176+
177+
if (wc->x == 0.0f && wc->y == 0.0f && wc->z == 0.0f) {
178+
return NULL;
179+
}
180+
181+
if (wc->x == 1.0f && wc->y == 1.0f && wc->z == 1.0f) {
182+
return a;
183+
}
184+
185+
// Fix up the alignment
186+
void* ret = ((char*)sg_ptr->renderstate)
187+
+ alignment_offset_calc(sg_ptr->renderstate,
188+
alignof(OSL::ClosureComponent));
189+
sg_ptr->renderstate = closure_mul_allot(ret, *wc, (OSL::ClosureColor*)a);
190+
191+
return ret;
192+
}
193+
194+
195+
196+
__device__ void*
197+
osl_mul_closure_float(void* sg_, void* a, float w)
198+
{
199+
a = __builtin_assume_aligned(a, alignof(float));
200+
201+
OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
202+
203+
if (a == NULL || w == 0.0f) {
204+
return NULL;
205+
}
206+
207+
if (w == 1.0f) {
208+
return a;
209+
}
210+
211+
// Fix up the alignment
212+
void* ret = ((char*)sg_ptr->renderstate)
213+
+ alignment_offset_calc(sg_ptr->renderstate,
214+
alignof(OSL::ClosureComponent));
215+
sg_ptr->renderstate = closure_mul_float_allot(ret, w,
216+
(OSL::ClosureColor*)a);
217+
218+
return ret;
219+
}
220+
221+
222+
223+
__device__ void*
224+
osl_add_closure_closure(void* sg_, void* a, void* b)
225+
{
226+
a = __builtin_assume_aligned(a, alignof(float));
227+
b = __builtin_assume_aligned(b, alignof(float));
228+
229+
OSL_CUDA::ShaderGlobals* sg_ptr = (OSL_CUDA::ShaderGlobals*)sg_;
230+
231+
if (a == NULL) {
232+
return b;
233+
}
234+
235+
if (b == NULL) {
236+
return a;
237+
}
238+
239+
// Fix up the alignment
240+
void* ret = ((char*)sg_ptr->renderstate)
241+
+ alignment_offset_calc(sg_ptr->renderstate,
242+
alignof(OSL::ClosureComponent));
243+
sg_ptr->renderstate = closure_add_allot(ret, (OSL::ClosureColor*)a,
244+
(OSL::ClosureColor*)b);
245+
246+
return ret;
247+
}
248+
249+
42250
#define IS_STRING(type) (type.basetype == OSL::TypeDesc::STRING)
43251
#define IS_PTR(type) (type.basetype == OSL::TypeDesc::PTR)
44252
#define IS_COLOR(type) (type.vecsemantics == OSL::TypeDesc::COLOR)

0 commit comments

Comments
 (0)