Date   

Re: Multithreading error

Larry Gritz
 

Are you saying you've fully figured it out?


On Sep 23, 2021, at 11:46 AM, Mark Bolstad <the.render.dude@...> wrote:

And that smell would be weird. I have two/N distinct contexts, but the way I was using them I would use an SC from one thread in a different one leading to the corruption.

Thanks,
Mark


--
Larry Gritz





Re: Multithreading error

Mark Bolstad
 

And that smell would be weird. I have two/N distinct contexts, but the way I was using them I would use an SC from one thread in a different one leading to the corruption.

Thanks,
Mark


Re: Multithreading error

Larry Gritz
 

It sure smells like you have two threads simultaneously messing with the same closure pool, which should never happen.

Are you EXTRA SURE that the threads are each using their own distinct ShadingContext? Like, I would print the addresses of the SC before you kick off an execute call to be sure the two threads are using different ones.


On Sep 23, 2021, at 10:17 AM, Mark Bolstad <the.render.dude@...> wrote:

  * frame #0: 0x00007fff7270733a libsystem_kernel.dylib`__pthread_kill + 10
    frame #1: 0x00007fff727c3e60 libsystem_pthread.dylib`pthread_kill + 430
    frame #2: 0x00007fff7268e808 libsystem_c.dylib`abort + 120
    frame #3: 0x0000000109225f81 liboslexec.1.11.dylib`OSL_v1_11::pvt::SimplePool<20480>::alloc(this=0x000000010f604410, size=56, alignment=16) at oslexec_pvt.h:1319:9
    frame #4: 0x0000000109225ad1 liboslexec.1.11.dylib`OSL_v1_11::ShadingContext::closure_component_allot(this=0x000000010f6042c0, id=0, prim_size=40, w=0x00007ffeefbeee40) at oslexec_pvt.h:1614:70
    frame #5: 0x0000000109225a87 liboslexec.1.11.dylib`::osl_allocate_closure_component(sg=0x00007ffeefbef138, id=0, size=40) at opclosure.cpp:54:25
    frame #6: 0x000000010f1e505e
    frame #7: 0x00000001090eeeb2 liboslexec.1.11.dylib`OSL_v1_11::ShadingContext::execute(this=0x000000010f6042c0, sgroup=0x000000010f70b0e0, ssg=0x00007ffeefbef138, run=true) at context.cpp:174:13
    frame #8: 0x00000001090628ae liboslexec.1.11.dylib`OSL_v1_11::pvt::ShadingSystemImpl::execute(this=0x000000011280da00, ctx=0x000000010f6042c0, group=0x000000010f70b0e0, ssg=0x00007ffeefbef138, run=true) at shadingsys.cpp:2783:24
    frame #9: 0x0000000109062824 liboslexec.1.11.dylib`OSL_v1_11::ShadingSystem::execute(this=0x000000010f510748, ctx=0x000000010f6042c0, group=0x000000010f70b0e0, globals=0x00007ffeefbef138, run=true) at shadingsys.cpp:262:20
    frame #10: 0x00000001002e1c36 libgraphics.dylib`papillon::execute_surface_shader(context=0x000000010f6042c0, shader=0x000000010f70b0e0, ctx=0x00007ffeefbef2f0) at renderer_services.cpp:50:25
...

--
Larry Gritz





Re: Multithreading error

Mark Bolstad
 

  * frame #0: 0x00007fff7270733a libsystem_kernel.dylib`__pthread_kill + 10
    frame #1: 0x00007fff727c3e60 libsystem_pthread.dylib`pthread_kill + 430
    frame #2: 0x00007fff7268e808 libsystem_c.dylib`abort + 120
    frame #3: 0x0000000109225f81 liboslexec.1.11.dylib`OSL_v1_11::pvt::SimplePool<20480>::alloc(this=0x000000010f604410, size=56, alignment=16) at oslexec_pvt.h:1319:9
    frame #4: 0x0000000109225ad1 liboslexec.1.11.dylib`OSL_v1_11::ShadingContext::closure_component_allot(this=0x000000010f6042c0, id=0, prim_size=40, w=0x00007ffeefbeee40) at oslexec_pvt.h:1614:70
    frame #5: 0x0000000109225a87 liboslexec.1.11.dylib`::osl_allocate_closure_component(sg=0x00007ffeefbef138, id=0, size=40) at opclosure.cpp:54:25
    frame #6: 0x000000010f1e505e
    frame #7: 0x00000001090eeeb2 liboslexec.1.11.dylib`OSL_v1_11::ShadingContext::execute(this=0x000000010f6042c0, sgroup=0x000000010f70b0e0, ssg=0x00007ffeefbef138, run=true) at context.cpp:174:13
    frame #8: 0x00000001090628ae liboslexec.1.11.dylib`OSL_v1_11::pvt::ShadingSystemImpl::execute(this=0x000000011280da00, ctx=0x000000010f6042c0, group=0x000000010f70b0e0, ssg=0x00007ffeefbef138, run=true) at shadingsys.cpp:2783:24
    frame #9: 0x0000000109062824 liboslexec.1.11.dylib`OSL_v1_11::ShadingSystem::execute(this=0x000000010f510748, ctx=0x000000010f6042c0, group=0x000000010f70b0e0, globals=0x00007ffeefbef138, run=true) at shadingsys.cpp:262:20
    frame #10: 0x00000001002e1c36 libgraphics.dylib`papillon::execute_surface_shader(context=0x000000010f6042c0, shader=0x000000010f70b0e0, ctx=0x00007ffeefbef2f0) at renderer_services.cpp:50:25
...


Re: Multithreading error

Larry Gritz
 

With a debug build, can you get a stack trace at the point that the assertion is hit?

That looks like an assertion in SimplePool::alloc(). But I'm not sure what could go wrong. Maybe something that contains a SimplePool that should be per thread is being shared?
 

On Sep 23, 2021, at 9:38 AM, Mark Bolstad <the.render.dude@...> wrote:

Hoping for some insight. I’ve bootstrapped OSL into a custom renderer and for the most part my simple scene is working on a single thread( the shader is three layers, marble->checkerboard (Cb)-> diffuse (base_color)).

As soon as I enable a 2nd thread, I get an assertion violation in oslexec_pvt.h:1319, “reinterpret_cast<uintptr_t>(ptr) % alignment == 0” failed.

AFAIK I’m creating the thread info and shading context per thread correctly, but obviously I’m doing something wrong.

This is with 11.15 and 11.13. Any ideas as to where I should poke to find the error?

Mark


--
Larry Gritz





Multithreading error

Mark Bolstad
 

Hoping for some insight. I’ve bootstrapped OSL into a custom renderer and for the most part my simple scene is working on a single thread( the shader is three layers, marble->checkerboard (Cb)-> diffuse (base_color)).

As soon as I enable a 2nd thread, I get an assertion violation in oslexec_pvt.h:1319, “reinterpret_cast<uintptr_t>(ptr) % alignment == 0” failed.

AFAIK I’m creating the thread info and shading context per thread correctly, but obviously I’m doing something wrong.

This is with 11.15 and 11.13. Any ideas as to where I should poke to find the error?

Mark


Upcoming Events #cal-summary

osl-dev@lists.aswf.io Calendar <osl-dev@...>
 

Open Shading Language discussion list Upcoming Events

OSL TSC meeting ( every other week )

When:
Thursday, September 30, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event


OSL TSC meeting ( every other week )

When:
Thursday, October 14, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event


Now: OSL TSC meeting ( every other week ) - 09/16/2021 #cal-notice

osl-dev@lists.aswf.io Calendar <noreply@...>
 

OSL TSC meeting ( every other week )

When:
09/16/2021
2:00pm to 3:00pm
(UTC-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

View Event

Description:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR


Event: OSL TSC meeting ( every other week ) - 09/16/2021 #cal-reminder

osl-dev@lists.aswf.io Calendar <noreply@...>
 

Reminder: OSL TSC meeting ( every other week )

When:
09/16/2021
2:00pm to 3:00pm
(UTC-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

View Event

Description:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR


Event: OSL TSC meeting ( every other week ) - 09/16/2021 #cal-reminder

osl-dev@lists.aswf.io Calendar <noreply@...>
 

Reminder: OSL TSC meeting ( every other week )

When:
09/16/2021
2:00pm to 3:00pm
(UTC-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

View Event

Description:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR


Re: getattribute() call on GPU

Nicolas Guiard
 

Thanks Larry!


Le mar. 14 sept. 2021 à 09:00, Larry Gritz <lg@...> a écrit :
And I've merged this PR into master now, so hopefully this will be fixed.

-- lg


On Sep 13, 2021, at 6:14 AM, Nicolas Guiard <nicolas@...> wrote:

Hi,

I've been able to make it work by changing the prototype of the osl_get_attribute callback and passing the TypeDesc casted into a long long instead of passing its address.
This is exactly what is already done in osl_bind_interpolated_param.


Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.



On Wed, Sep 1, 2021 at 11:43 PM Nicolas Guiard via lists.aswf.io <nicolas=isotropix.com@...> wrote:
Hi,

We are currently trying to make getattribute() calls work on GPU with OptiX 7.3 and OSL 1.11.9.0 but we are facing an issue.
We've implemented the osl_get_attribute() callback in cuda and it's called properly. But as soon as we try to access the content of the attribute type, we get an illegal memory access.

Here's the implementation of osl_get_attribute for our test:
__device__
int osl_get_attribute(void *sg_, int dest_derivs, void *obj_name_, void *attr_name_, int array_lookup, int index, const void *attr_type, void *attr_dest)
{
    float *f = reinterpret_cast<float *>(attr_dest);
    if (attr_type == nullptr) {
        f[0] = 0.5f;
        f[1] = 0.5f;
        f[2] = 0.5f;
        return 1;
    }
    const OSL::TypeDesc& type = *(const OSL::TypeDesc *)attr_type;
    f[0] = type.basetype;
    f[1] = type.aggregate;
    f[2] = type.vecsemantics;

    return 1;
}

If we remove the lines in bold red and replace them by the following ones, it works just fine:
    f[0] = 1.0f;
    f[1] = 0.5f;
    f[2] = 0.0f;

In our test, we have a simple shader doing this:
surface attrtest()
{
    color c;
    getattribute("foo", c);
    Ci = c * emission();
}

And here's the generated PTX code:
//
// Generated by LLVM NVPTX Back-End
//
.version 5.0
.target sm_35
.address_size 64
// .globl __direct_callable__group_unnamed_group_3_3_init
.extern .func  (.param .b32 func_retval0) osl_get_attribute
(
.param .b64 osl_get_attribute_param_0,
.param .b32 osl_get_attribute_param_1,
.param .b64 osl_get_attribute_param_2,
.param .b64 osl_get_attribute_param_3,
.param .b32 osl_get_attribute_param_4,
.param .b32 osl_get_attribute_param_5,
.param .b64 osl_get_attribute_param_6,
.param .b64 osl_get_attribute_param_7
)
;
.extern .func  (.param .b64 func_retval0) osl_allocate_weighted_closure_component
(
.param .b64 osl_allocate_weighted_closure_component_param_0,
.param .b32 osl_allocate_weighted_closure_component_param_1,
.param .b32 osl_allocate_weighted_closure_component_param_2,
.param .b64 osl_allocate_weighted_closure_component_param_3
)
;
.extern .global .align 8 .u64 ds_555c6f602f9383e3_0003;
.visible .func __direct_callable__group_unnamed_group_3_3_init(
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_0,
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_1
)
{

ret;
}
// .globl __direct_callable__unnamed_group_3_build___project_scene_attrtest_4
.visible .func __direct_callable__unnamed_group_3_build___project_scene_attrtest_4(
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0,
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_1
)
{
.local .align 8 .b8 __local_depot1[16];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<2>;
.reg .b16 %rs<2>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;
mov.u64 %SPL, __local_depot1;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd2, [__direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0];
ld.global.u64 %rd3, [ds_555c6f602f9383e3_0003];
mov.u64 %rd4, 0;
mov.u32 %r1, 0;
mov.u64 %rd5, 140457662007856;
add.u64 %rd6, %SP, 0;
{ // callseq 0, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r1;
.param .b64 param2;
st.param.b64 [param2+0], %rd4;
.param .b64 param3;
st.param.b64 [param3+0], %rd3;
.param .b32 param4;
st.param.b32 [param4+0], %r1;
.param .b32 param5;
st.param.b32 [param5+0], %r2;
.param .b64 param6;
st.param.b64 [param6+0], %rd5;
.param .b64 param7;
st.param.b64 [param7+0], %rd6;
.param .b32 retval0;
call.uni (retval0),
osl_get_attribute,
(
param0,
param1,
param2,
param3,
param4,
param5,
param6,
param7
);
ld.param.b32 %r3, [retval0+0];
} // callseq 0
mov.u32 %r5, 1;
{ // callseq 1, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r5;
.param .b32 param2;
st.param.b32 [param2+0], %r5;
.param .b64 param3;
st.param.b64 [param3+0], %rd6;
.param .b64 retval0;
call.uni (retval0),
osl_allocate_weighted_closure_component,
(
param0,
param1,
param2,
param3
);
ld.param.b64 %rd7, [retval0+0];
} // callseq 1
setp.eq.s64 %p1, %rd7, 0;
@%p1 bra LBB1_2;
mov.u16 %rs1, 0;
st.u8 [%rd7+16], %rs1;
LBB1_2:
st.u64 [%rd2+272], %rd7;
ret;
}

As we can see in bold red in the generated code, the address of the type is a constant value and this value changes at each execution. Is there something we are not configuring properly or is it a bug?

Thanks in advance for any help on this!

Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.




--
Larry Gritz





Re: getattribute() call on GPU

Larry Gritz
 

And I've merged this PR into master now, so hopefully this will be fixed.

-- lg


On Sep 13, 2021, at 6:14 AM, Nicolas Guiard <nicolas@...> wrote:

Hi,

I've been able to make it work by changing the prototype of the osl_get_attribute callback and passing the TypeDesc casted into a long long instead of passing its address.
This is exactly what is already done in osl_bind_interpolated_param.


Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.



On Wed, Sep 1, 2021 at 11:43 PM Nicolas Guiard via lists.aswf.io <nicolas=isotropix.com@...> wrote:
Hi,

We are currently trying to make getattribute() calls work on GPU with OptiX 7.3 and OSL 1.11.9.0 but we are facing an issue.
We've implemented the osl_get_attribute() callback in cuda and it's called properly. But as soon as we try to access the content of the attribute type, we get an illegal memory access.

Here's the implementation of osl_get_attribute for our test:
__device__
int osl_get_attribute(void *sg_, int dest_derivs, void *obj_name_, void *attr_name_, int array_lookup, int index, const void *attr_type, void *attr_dest)
{
    float *f = reinterpret_cast<float *>(attr_dest);
    if (attr_type == nullptr) {
        f[0] = 0.5f;
        f[1] = 0.5f;
        f[2] = 0.5f;
        return 1;
    }
    const OSL::TypeDesc& type = *(const OSL::TypeDesc *)attr_type;
    f[0] = type.basetype;
    f[1] = type.aggregate;
    f[2] = type.vecsemantics;

    return 1;
}

If we remove the lines in bold red and replace them by the following ones, it works just fine:
    f[0] = 1.0f;
    f[1] = 0.5f;
    f[2] = 0.0f;

In our test, we have a simple shader doing this:
surface attrtest()
{
    color c;
    getattribute("foo", c);
    Ci = c * emission();
}

And here's the generated PTX code:
//
// Generated by LLVM NVPTX Back-End
//
.version 5.0
.target sm_35
.address_size 64
// .globl __direct_callable__group_unnamed_group_3_3_init
.extern .func  (.param .b32 func_retval0) osl_get_attribute
(
.param .b64 osl_get_attribute_param_0,
.param .b32 osl_get_attribute_param_1,
.param .b64 osl_get_attribute_param_2,
.param .b64 osl_get_attribute_param_3,
.param .b32 osl_get_attribute_param_4,
.param .b32 osl_get_attribute_param_5,
.param .b64 osl_get_attribute_param_6,
.param .b64 osl_get_attribute_param_7
)
;
.extern .func  (.param .b64 func_retval0) osl_allocate_weighted_closure_component
(
.param .b64 osl_allocate_weighted_closure_component_param_0,
.param .b32 osl_allocate_weighted_closure_component_param_1,
.param .b32 osl_allocate_weighted_closure_component_param_2,
.param .b64 osl_allocate_weighted_closure_component_param_3
)
;
.extern .global .align 8 .u64 ds_555c6f602f9383e3_0003;
.visible .func __direct_callable__group_unnamed_group_3_3_init(
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_0,
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_1
)
{

ret;
}
// .globl __direct_callable__unnamed_group_3_build___project_scene_attrtest_4
.visible .func __direct_callable__unnamed_group_3_build___project_scene_attrtest_4(
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0,
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_1
)
{
.local .align 8 .b8 __local_depot1[16];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<2>;
.reg .b16 %rs<2>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;
mov.u64 %SPL, __local_depot1;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd2, [__direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0];
ld.global.u64 %rd3, [ds_555c6f602f9383e3_0003];
mov.u64 %rd4, 0;
mov.u32 %r1, 0;
mov.u64 %rd5, 140457662007856;
add.u64 %rd6, %SP, 0;
{ // callseq 0, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r1;
.param .b64 param2;
st.param.b64 [param2+0], %rd4;
.param .b64 param3;
st.param.b64 [param3+0], %rd3;
.param .b32 param4;
st.param.b32 [param4+0], %r1;
.param .b32 param5;
st.param.b32 [param5+0], %r2;
.param .b64 param6;
st.param.b64 [param6+0], %rd5;
.param .b64 param7;
st.param.b64 [param7+0], %rd6;
.param .b32 retval0;
call.uni (retval0),
osl_get_attribute,
(
param0,
param1,
param2,
param3,
param4,
param5,
param6,
param7
);
ld.param.b32 %r3, [retval0+0];
} // callseq 0
mov.u32 %r5, 1;
{ // callseq 1, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r5;
.param .b32 param2;
st.param.b32 [param2+0], %r5;
.param .b64 param3;
st.param.b64 [param3+0], %rd6;
.param .b64 retval0;
call.uni (retval0),
osl_allocate_weighted_closure_component,
(
param0,
param1,
param2,
param3
);
ld.param.b64 %rd7, [retval0+0];
} // callseq 1
setp.eq.s64 %p1, %rd7, 0;
@%p1 bra LBB1_2;
mov.u16 %rs1, 0;
st.u8 [%rd7+16], %rs1;
LBB1_2:
st.u64 [%rd2+272], %rd7;
ret;
}

As we can see in bold red in the generated code, the address of the type is a constant value and this value changes at each execution. Is there something we are not configuring properly or is it a bug?

Thanks in advance for any help on this!

Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.




--
Larry Gritz





Event: OSL TSC meeting ( every other week ) - 09/16/2021 #cal-reminder

osl-dev@lists.aswf.io Calendar <noreply@...>
 

Reminder: OSL TSC meeting ( every other week )

When:
09/16/2021
2:00pm to 3:00pm
(UTC-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

View Event

Description:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR


Re: getattribute() call on GPU

Nicolas Guiard
 

Hi,

I've been able to make it work by changing the prototype of the osl_get_attribute callback and passing the TypeDesc casted into a long long instead of passing its address.
This is exactly what is already done in osl_bind_interpolated_param.


Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.



On Wed, Sep 1, 2021 at 11:43 PM Nicolas Guiard via lists.aswf.io <nicolas=isotropix.com@...> wrote:
Hi,

We are currently trying to make getattribute() calls work on GPU with OptiX 7.3 and OSL 1.11.9.0 but we are facing an issue.
We've implemented the osl_get_attribute() callback in cuda and it's called properly. But as soon as we try to access the content of the attribute type, we get an illegal memory access.

Here's the implementation of osl_get_attribute for our test:
__device__
int osl_get_attribute(void *sg_, int dest_derivs, void *obj_name_, void *attr_name_, int array_lookup, int index, const void *attr_type, void *attr_dest)
{
    float *f = reinterpret_cast<float *>(attr_dest);
    if (attr_type == nullptr) {
        f[0] = 0.5f;
        f[1] = 0.5f;
        f[2] = 0.5f;
        return 1;
    }
    const OSL::TypeDesc& type = *(const OSL::TypeDesc *)attr_type;
    f[0] = type.basetype;
    f[1] = type.aggregate;
    f[2] = type.vecsemantics;

    return 1;
}

If we remove the lines in bold red and replace them by the following ones, it works just fine:
    f[0] = 1.0f;
    f[1] = 0.5f;
    f[2] = 0.0f;

In our test, we have a simple shader doing this:
surface attrtest()
{
    color c;
    getattribute("foo", c);
    Ci = c * emission();
}

And here's the generated PTX code:
//
// Generated by LLVM NVPTX Back-End
//
.version 5.0
.target sm_35
.address_size 64
// .globl __direct_callable__group_unnamed_group_3_3_init
.extern .func  (.param .b32 func_retval0) osl_get_attribute
(
.param .b64 osl_get_attribute_param_0,
.param .b32 osl_get_attribute_param_1,
.param .b64 osl_get_attribute_param_2,
.param .b64 osl_get_attribute_param_3,
.param .b32 osl_get_attribute_param_4,
.param .b32 osl_get_attribute_param_5,
.param .b64 osl_get_attribute_param_6,
.param .b64 osl_get_attribute_param_7
)
;
.extern .func  (.param .b64 func_retval0) osl_allocate_weighted_closure_component
(
.param .b64 osl_allocate_weighted_closure_component_param_0,
.param .b32 osl_allocate_weighted_closure_component_param_1,
.param .b32 osl_allocate_weighted_closure_component_param_2,
.param .b64 osl_allocate_weighted_closure_component_param_3
)
;
.extern .global .align 8 .u64 ds_555c6f602f9383e3_0003;
.visible .func __direct_callable__group_unnamed_group_3_3_init(
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_0,
.param .b64 __direct_callable__group_unnamed_group_3_3_init_param_1
)
{

ret;
}
// .globl __direct_callable__unnamed_group_3_build___project_scene_attrtest_4
.visible .func __direct_callable__unnamed_group_3_build___project_scene_attrtest_4(
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0,
.param .b64 __direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_1
)
{
.local .align 8 .b8 __local_depot1[16];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<2>;
.reg .b16 %rs<2>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;
mov.u64 %SPL, __local_depot1;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd2, [__direct_callable__unnamed_group_3_build___project_scene_attrtest_4_param_0];
ld.global.u64 %rd3, [ds_555c6f602f9383e3_0003];
mov.u64 %rd4, 0;
mov.u32 %r1, 0;
mov.u64 %rd5, 140457662007856;
add.u64 %rd6, %SP, 0;
{ // callseq 0, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r1;
.param .b64 param2;
st.param.b64 [param2+0], %rd4;
.param .b64 param3;
st.param.b64 [param3+0], %rd3;
.param .b32 param4;
st.param.b32 [param4+0], %r1;
.param .b32 param5;
st.param.b32 [param5+0], %r2;
.param .b64 param6;
st.param.b64 [param6+0], %rd5;
.param .b64 param7;
st.param.b64 [param7+0], %rd6;
.param .b32 retval0;
call.uni (retval0),
osl_get_attribute,
(
param0,
param1,
param2,
param3,
param4,
param5,
param6,
param7
);
ld.param.b32 %r3, [retval0+0];
} // callseq 0
mov.u32 %r5, 1;
{ // callseq 1, 0
.reg .b32 temp_param_reg;
.param .b64 param0;
st.param.b64 [param0+0], %rd2;
.param .b32 param1;
st.param.b32 [param1+0], %r5;
.param .b32 param2;
st.param.b32 [param2+0], %r5;
.param .b64 param3;
st.param.b64 [param3+0], %rd6;
.param .b64 retval0;
call.uni (retval0),
osl_allocate_weighted_closure_component,
(
param0,
param1,
param2,
param3
);
ld.param.b64 %rd7, [retval0+0];
} // callseq 1
setp.eq.s64 %p1, %rd7, 0;
@%p1 bra LBB1_2;
mov.u16 %rs1, 0;
st.u8 [%rd7+16], %rs1;
LBB1_2:
st.u64 [%rd2+272], %rd7;
ret;
}

As we can see in bold red in the generated code, the address of the type is a constant value and this value changes at each execution. Is there something we are not configuring properly or is it a bug?

Thanks in advance for any help on this!

Cheers

Face
Nicolas Guiard
Head of R&D
Phone+33 972 452 846
Isotropix
LinkedIn YouTube Facebook Twitter

This email and any files transmitted with it are confidential and intended solely for the use of the individual or entity to whom they are addressed.


Upcoming Events #cal-summary

osl-dev@lists.aswf.io Calendar <osl-dev@...>
 

Open Shading Language discussion list Upcoming Events

OSL TSC meeting ( every other week )

When:
Thursday, September 16, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event


OSL TSC meeting ( every other week )

When:
Thursday, September 30, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event


Event: OSL TSC meeting ( every other week ) - 09/16/2021 #cal-reminder

osl-dev@lists.aswf.io Calendar <noreply@...>
 

Reminder: OSL TSC meeting ( every other week )

When:
09/16/2021
2:00pm to 3:00pm
(UTC-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

View Event

Description:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR


Re: Problem to compile OSL (v1.11.14.2) on ubuntu

Patrick Hodoul
 
Edited

I found a solution by using the ASWF CI build image for OSL.
Refer here for details.


OSL shader language question

Patrick Hodoul
 

Hi,

OCIOv2 now generates an OSL shader program. But I face some irritating syntax changes to accommodate OSL when generating the shader code. To simplify the generation of the code, I would like to write something like:
vector4 res;
color4 color = res;


The idea is to perform some computations with the vector4 'res' variable and another color4 variable representing the input color, and finally store it in the color4 variable (representing the output color) which is passed to the next step in the computation pipeline. Note that the input and output color variables are in fact the same. The example above oversimplifies the code to better highlight the need. I use vector4 by default because of the 'uniform' access to the values i.e. res[x] or res.r, etc and it's close to the GLSL/HLSL syntax.

To make this work, the code recreates the color4 like below. But it will simplify the code to write "color = res;"

color = color4(vector(res.x, res.y, res.z), res.w);
or
color.rgb = color(res.x, res.y, res.z);
color.a = res.w;

Following the really helpful 'operator overloading' idea (that I use a lot) I was wondering if there is a way to implement "color = res;" and also "res = color;"

--
Patrick


Problem to compile OSL (v1.11.14.2) on ubuntu

Patrick Hodoul
 

Hi,

For the OCIOv2 project, I developed a unit test framework to compile and excecute an OSL shader code.
And, I now need to run the OSL unit tests on a Linux platform (I did all my tests on macOS for now).

First of all, I was not able to find the package in apt-get (using the apt-cache search  command). Does the package exists?

Second, I tried to compile OSL on Linux in a docker as it was quite simple on macOS.
Below are my tries to compile OSL:
  • ubuntu 18.04
    • tons of packages to install or manually compile
    • issue with OIIO
      • The installed OIIO is too old i.e. 1.7
      • The OIIO script from OSL fails
      • Manual compilation of OIIO RB-2.0 fails because install pybind11 is too old
      • I stopped my investigations :-(
  • ubuntu 20.04
    • tons of packages to install
    • installed OIIO detected :-)
    • issue with Clang+LLVM
      • the Clang+LLVM script from OSL fails
        • the script is built to run on CI only.
      • apt-get install llvm clang
        • clang is 10.0.0, gcc version 9.3.0
      • cmake -DCMAKE_CXX_STANDARD=14 -DUSE_PYTHON=OFF ../. fails.  :-(
        • /home/devel/osl_1/src/liboslcomp/oslcomp.cpp:24:10: fatal error: clang/Basic/TargetInfo.h: No such file or directory


I'm open to any suggestions, existing docker images, etc.

Thanks,
Patrick


Upcoming Events #cal-summary

osl-dev@lists.aswf.io Calendar <osl-dev@...>
 

Open Shading Language discussion list Upcoming Events

OSL TSC meeting ( every other week )

When:
Thursday, September 16, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event


OSL TSC meeting ( every other week )

When:
Thursday, September 30, 2021, 2:00pm to 3:00pm
(GMT-07:00) America/Los Angeles

Where:
https://zoom.us/j/100511909

Organizer: Chris Kulla ckulla@...

Details:

Every other week meeting of the OSL TSC.

Meeting Agenda / Notes: https://docs.google.com/document/d/1yf0bG6eoE2EvKZBNZX3nskdTvu99ADTDTNOknCDJd1I/

Confirm this meeting invite is still valid by finding the meeting at https://lists.aswf.io/calendar.

Join Zoom Meeting https://zoom.us/j/100511909

Meeting ID: 100 511 909

One tap mobile +16465588656,,100511909# US (New York) +13126266799,,100511909# US (Chicago)

Dial by your location +1 646 558 8656 US (New York) +1 312 626 6799 US (Chicago) +1 669 900 6833 US (San Jose) +1 253 215 8782 US +1 301 715 8592 US +1 346 248 7799 US (Houston) 877 369 0926 US Toll-free 855 880 1246 US Toll-free +1 587 328 1099 Canada +1 647 374 4685 Canada +1 647 558 0588 Canada +1 778 907 2071 Canada +1 438 809 7799 Canada 855 703 8985 Canada Toll-free Meeting ID: 100 511 909 Find your local number: https://zoom.us/u/acBVrM6HWR

View Event

161 - 180 of 4924