Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[WGSL] __ref does not work as expected #5596

Closed
saipraveenb25 opened this issue Nov 19, 2024 · 0 comments · Fixed by #5597
Closed

[WGSL] __ref does not work as expected #5596

saipraveenb25 opened this issue Nov 19, 2024 · 0 comments · Fixed by #5597
Assignees
Labels
goal:quality & productivity Quality issues and issues that impact our productivity coding day to day inside slang

Comments

@saipraveenb25
Copy link
Collaborator

__ref parameters do not work correctly. Objects passed as reference to function should not be placed in a local variable.

The following test replicates this behavior

// groupshared-ref-param.slang

//TEST:SIMPLE(filecheck=CHECK): -target wgsl -entry computeMain -stage compute


//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
RWStructuredBuffer<int> outputBuffer;

groupshared uint sharedVal;

[ForceInline]
uint workgroupUniformLoad(__ref uint val)
{
    __intrinsic_asm "workgroupUniformLoad(&($0))";
}

// Expect sharedVal to be passed by reference.
//
// CHECK: fn computeMain({{.*}}({{.*}})
// CHECK: {{.*}}workgroupUniformLoad(&((sharedVal_{{[a-zA-Z0-9_]*}}))){{.*}};
// CHECK: }


[numthreads(1, 1, 1)]
void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
{
    int idx = dispatchThreadID.x;

    sharedVal = 1;
    outputBuffer[0] = workgroupUniformLoad(sharedVal);
}

WGSL output:

@binding(0) @group(0) var<storage, read_write> outputBuffer_0 : array<i32>;

var<workgroup> sharedVal_0 : u32;

@compute
@workgroup_size(1, 1, 1)
fn computeMain(@builtin(global_invocation_id) dispatchThreadID_0 : vec3<u32>)
{
    sharedVal_0 = u32(1);
    var _S1 : u32 = u32(1);
    var _S2 : u32 = (workgroupUniformLoad(&((_S1))));
    sharedVal_0 = _S1;
    outputBuffer_0[i32(0)] = i32(_S2);
    return;
}

HLSL output (this works as expected):

#pragma pack_matrix(column_major)
#ifdef SLANG_HLSL_ENABLE_NVAPI
#include "nvHLSLExtns.h"
#endif

#ifndef __DXC_VERSION_MAJOR
// warning X3557: loop doesn't seem to do anything, forcing loop to unroll
#pragma warning(disable : 3557)
#endif


#line 9 "tests/language-feature/pointer/groupshared-ref-param.slang"
RWStructuredBuffer<int > outputBuffer_0 : register(u0);

static groupshared uint sharedVal_0;


#line 34
[numthreads(1, 1, 1)]
void computeMain(int3 dispatchThreadID_0 : SV_DispatchThreadID)
{

    sharedVal_0 = 1U;
    uint _S1 = (workgroupUniformLoad(&((sharedVal_0))));

#line 39
    outputBuffer_0[int(0)] = int(_S1);
    return;
}
@csyonghe csyonghe self-assigned this Nov 19, 2024
@csyonghe csyonghe added the goal:quality & productivity Quality issues and issues that impact our productivity coding day to day inside slang label Nov 19, 2024
@csyonghe csyonghe added this to the Q4 2024 (Fall) milestone Nov 19, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
goal:quality & productivity Quality issues and issues that impact our productivity coding day to day inside slang
Projects
None yet
Development

Successfully merging a pull request may close this issue.

2 participants