Skip to content

Commit ae990c3

Browse files
curtisblackcurtisblack
committed
cleanup
Signed-off-by: Curtis Black <curtis.w.black@gmail.com>
1 parent aef2a80 commit ae990c3

3 files changed

Lines changed: 36 additions & 41 deletions

File tree

‎src/testrender/cuda/rend_lib.cu‎

Lines changed: 16 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -39,14 +39,15 @@ __direct_callable__dummy_rend_lib()
3939
{
4040
}
4141

42-
__device__ void*
42+
43+
__device__ const void*
4344
osl_add_closure_closure(void* sg_, const void* a_, const void* b_)
4445
{
45-
a_ = __builtin_assume_aligned(a_, alignof(float));
46-
b_ = __builtin_assume_aligned(b_, alignof(float));
47-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
48-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
49-
const OSL::ClosureColor* b = (const OSL::ClosureColor*)b_;
46+
a_ = __builtin_assume_aligned(a_, alignof(float));
47+
b_ = __builtin_assume_aligned(b_, alignof(float));
48+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
49+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
50+
const OSL::ClosureColor* b = (const OSL::ClosureColor*)b_;
5051
if (a == NULL)
5152
return b;
5253
if (b == NULL)
@@ -63,15 +64,15 @@ osl_add_closure_closure(void* sg_, const void* a_, const void* b_)
6364
return add;
6465
}
6566

66-
__device__ void*
67+
__device__ const void*
6768
osl_mul_closure_color(void* sg_, const void* a_, const void* w_)
6869
{
6970
a_ = __builtin_assume_aligned(a_, alignof(float));
7071
w_ = __builtin_assume_aligned(w_, alignof(float));
7172

72-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
73-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
74-
const OSL::Color3* w = (const OSL::Color3*)w_;
73+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
74+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
75+
const OSL::Color3* w = (const OSL::Color3*)w_;
7576
if (a == NULL)
7677
return NULL;
7778
if (w->x == 0.0f && w->y == 0.0f && w->z == 0.0f)
@@ -90,13 +91,13 @@ osl_mul_closure_color(void* sg_, const void* a_, const void* w_)
9091
return mul;
9192
}
9293

93-
__device__ void*
94+
__device__ const void*
9495
osl_mul_closure_float(void* sg_, const void* a_, float w)
9596
{
9697
a_ = __builtin_assume_aligned(a_, alignof(float));
9798

98-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
99-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
99+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
100+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
100101
if (a == NULL)
101102
return NULL;
102103
if (w == 0.0f)
@@ -118,8 +119,8 @@ osl_mul_closure_float(void* sg_, const void* a_, float w)
118119
__device__ void*
119120
osl_allocate_closure_component(void* sg_, int id, int size)
120121
{
121-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
122-
auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
122+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
123+
auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
123124
// Allocate the component and the mul back to back
124125
const size_t needed = sizeof(OSL::ClosureComponent) + size;
125126
OSL::ClosureComponent* comp

‎src/testshade/render_state.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,14 +29,17 @@ class StackClosurePool {
2929
void* ptr;
3030

3131
public:
32+
OSL_HOSTDEVICE
3233
StackClosurePool() { reset(); }
3334

35+
OSL_HOSTDEVICE
3436
void reset()
3537
{
3638
ptr = &buffer[0];
3739
*(int*)ptr = 0;
3840
}
3941

42+
OSL_HOSTDEVICE
4043
void* allocate(size_t size, size_t alignment)
4144
{
4245
uintptr_t p = OIIO::round_to_multiple_of_pow2((uintptr_t)ptr,

‎testsuite/example-cuda/rend_lib.cu‎

Lines changed: 17 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -14,27 +14,18 @@ extern __device__ char* s_color_system;
1414
}
1515
OSL_NAMESPACE_END
1616

17-
// Taken from the SimplePool class
18-
__device__ static inline size_t
19-
alignment_offset_calc(void* ptr, size_t alignment)
20-
{
21-
uintptr_t ptrbits = reinterpret_cast<uintptr_t>(ptr);
22-
uintptr_t offset = ((ptrbits + alignment - 1) & -alignment) - ptrbits;
23-
return offset;
24-
}
25-
2617
// These functions are declared extern to prevent name mangling.
2718
extern "C" {
2819

2920

30-
__device__ void*
21+
__device__ const void*
3122
osl_add_closure_closure(void* sg_, const void* a_, const void* b_)
3223
{
33-
a_ = __builtin_assume_aligned(a_, alignof(float));
34-
b_ = __builtin_assume_aligned(b_, alignof(float));
35-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
36-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
37-
const OSL::ClosureColor* b = (const OSL::ClosureColor*)b_;
24+
a_ = __builtin_assume_aligned(a_, alignof(float));
25+
b_ = __builtin_assume_aligned(b_, alignof(float));
26+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
27+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
28+
const OSL::ClosureColor* b = (const OSL::ClosureColor*)b_;
3829
if (a == NULL)
3930
return b;
4031
if (b == NULL)
@@ -51,15 +42,15 @@ osl_add_closure_closure(void* sg_, const void* a_, const void* b_)
5142
return add;
5243
}
5344

54-
__device__ void*
45+
__device__ const void*
5546
osl_mul_closure_color(void* sg_, const void* a_, const void* w_)
5647
{
5748
a_ = __builtin_assume_aligned(a_, alignof(float));
5849
w_ = __builtin_assume_aligned(w_, alignof(float));
5950

60-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
61-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
62-
const OSL::Color3* w = (const OSL::Color3*)w_;
51+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
52+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
53+
const OSL::Color3* w = (const OSL::Color3*)w_;
6354
if (a == NULL)
6455
return NULL;
6556
if (w->x == 0.0f && w->y == 0.0f && w->z == 0.0f)
@@ -78,13 +69,13 @@ osl_mul_closure_color(void* sg_, const void* a_, const void* w_)
7869
return mul;
7970
}
8071

81-
__device__ void*
72+
__device__ const void*
8273
osl_mul_closure_float(void* sg_, const void* a_, float w)
8374
{
8475
a_ = __builtin_assume_aligned(a_, alignof(float));
8576

86-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
87-
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
77+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
78+
const OSL::ClosureColor* a = (const OSL::ClosureColor*)a_;
8879
if (a == NULL)
8980
return NULL;
9081
if (w == 0.0f)
@@ -106,8 +97,8 @@ osl_mul_closure_float(void* sg_, const void* a_, float w)
10697
__device__ void*
10798
osl_allocate_closure_component(void* sg_, int id, int size)
10899
{
109-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
110-
auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
100+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
101+
auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;
111102
// Allocate the component and the mul back to back
112103
const size_t needed = sizeof(OSL::ClosureComponent) + size;
113104
OSL::ClosureComponent* comp
@@ -126,8 +117,8 @@ osl_allocate_weighted_closure_component(void* sg_, int id, int size,
126117
{
127118
w_ = __builtin_assume_aligned(w_, alignof(float));
128119

129-
ShaderGlobals* sg = (ShaderGlobals*)sg_;
130-
const OSL::Color3* w = (const OSL::Color3*)w_;
120+
OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
121+
const OSL::Color3* w = (const OSL::Color3*)w_;
131122
if (w->x == 0.0f && w->y == 0.0f && w->z == 0.0f)
132123
return NULL;
133124
auto* closure_pool = ((RenderState*)sg->renderstate)->closure_pool;

0 commit comments

Comments
 (0)