Skip to content

Commit bd6dbaf

Browse files
csyonghecsyonghe
authored
Compile append and consume structured buffers to glsl. (shader-slang#3142)
* Compile append and consume structured buffers to glsl. * Fix. * Update CI config. --------- Co-authored-by: Yong He <yhe@nvidia.com>
1 parent f94b2f7 commit bd6dbaf

40 files changed

+809
-239
lines changed

‎.github/workflows/windows-selfhosted.yml

+10-7
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,9 @@ jobs:
1212
build:
1313
runs-on: [Windows, self-hosted]
1414
timeout-minutes: 100
15-
15+
continue-on-error: true
1616
strategy:
17+
fail-fast: false
1718
matrix:
1819
configuration: ['Release']
1920
platform: ['x64']
@@ -38,11 +39,13 @@ jobs:
3839
MSBuild.exe slang.sln -v:m -m -property:Configuration=${{matrix.configuration}} -property:Platform=${{matrix.platform}} -property:WindowsTargetPlatformVersion=10.0.19041.0 -maxcpucount:12
3940
- name: test-spirv-direct
4041
run: |
41-
set PATH=%PATH%;.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\
42-
".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe" tests/ -use-test-server -emit-spirv-directly -expected-failure-list tests/expected-failure.txt -api vk 2>&1
43-
shell: cmd
42+
$ErrorActionPreference = "SilentlyContinue"
43+
$env:Path += ';.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\'
44+
.\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe tests/ -use-test-server -emit-spirv-directly -expected-failure-list tests/expected-failure.txt -api vk
45+
4446
- name: test
4547
run: |
46-
set PATH=%PATH%;.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\
47-
".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe" -use-test-server -api vk 2>&1
48-
shell: cmd
48+
$ErrorActionPreference = "SilentlyContinue"
49+
$env:Path += ';.\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\'
50+
.\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\slang-test.exe -use-test-server -api all-cpu
51+

‎build/visual-studio/slang/slang.vcxproj

+2
Original file line numberDiff line numberDiff line change
@@ -408,6 +408,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla
408408
<ClInclude Include="..\..\..\source\slang\slang-ir-liveness.h" />
409409
<ClInclude Include="..\..\..\source\slang\slang-ir-loop-inversion.h" />
410410
<ClInclude Include="..\..\..\source\slang\slang-ir-loop-unroll.h" />
411+
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.h" />
411412
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-binding-query.h" />
412413
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-bit-cast.h" />
413414
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-buffer-element-type.h" />
@@ -618,6 +619,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla
618619
<ClCompile Include="..\..\..\source\slang\slang-ir-liveness.cpp" />
619620
<ClCompile Include="..\..\..\source\slang\slang-ir-loop-inversion.cpp" />
620621
<ClCompile Include="..\..\..\source\slang\slang-ir-loop-unroll.cpp" />
622+
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.cpp" />
621623
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-binding-query.cpp" />
622624
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-bit-cast.cpp" />
623625
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-buffer-element-type.cpp" />

‎build/visual-studio/slang/slang.vcxproj.filters

+6
Original file line numberDiff line numberDiff line change
@@ -312,6 +312,9 @@
312312
<ClInclude Include="..\..\..\source\slang\slang-ir-loop-unroll.h">
313313
<Filter>Header Files</Filter>
314314
</ClInclude>
315+
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.h">
316+
<Filter>Header Files</Filter>
317+
</ClInclude>
315318
<ClInclude Include="..\..\..\source\slang\slang-ir-lower-binding-query.h">
316319
<Filter>Header Files</Filter>
317320
</ClInclude>
@@ -938,6 +941,9 @@
938941
<ClCompile Include="..\..\..\source\slang\slang-ir-loop-unroll.cpp">
939942
<Filter>Source Files</Filter>
940943
</ClCompile>
944+
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-append-consume-structured-buffer.cpp">
945+
<Filter>Source Files</Filter>
946+
</ClCompile>
941947
<ClCompile Include="..\..\..\source\slang\slang-ir-lower-binding-query.cpp">
942948
<Filter>Source Files</Filter>
943949
</ClCompile>

‎source/slang/hlsl.meta.slang

+24-2
Original file line numberDiff line numberDiff line change
@@ -6,16 +6,31 @@ typedef uint UINT;
66
[ForceInline] float3 __asFloat3(float2 v) { return float3(v, 0); }
77
[ForceInline] float3 __asFloat3(float3 v) { return v; }
88

9+
__generic<T>
10+
__intrinsic_op($(kIROp_StructuredBufferGetDimensions))
11+
uint2 __structuredBufferGetDimensions(AppendStructuredBuffer<T> buffer);
12+
13+
__generic<T>
14+
__intrinsic_op($(kIROp_StructuredBufferGetDimensions))
15+
uint2 __structuredBufferGetDimensions(ConsumeStructuredBuffer<T> buffer);
16+
917
__generic<T>
1018
__magic_type(HLSLAppendStructuredBufferType)
1119
__intrinsic_type($(kIROp_HLSLAppendStructuredBufferType))
1220
struct AppendStructuredBuffer
1321
{
22+
__intrinsic_op($(kIROp_StructuredBufferAppend))
1423
void Append(T value);
1524

25+
[ForceInline]
1626
void GetDimensions(
1727
out uint numStructs,
18-
out uint stride);
28+
out uint stride)
29+
{
30+
let result = __structuredBufferGetDimensions(this);
31+
numStructs = result.x;
32+
stride = result.y;
33+
}
1934
};
2035

2136
__magic_type(HLSLByteAddressBufferType)
@@ -257,11 +272,18 @@ __magic_type(HLSLConsumeStructuredBufferType)
257272
__intrinsic_type($(kIROp_HLSLConsumeStructuredBufferType))
258273
struct ConsumeStructuredBuffer
259274
{
275+
__intrinsic_op($(kIROp_StructuredBufferConsume))
260276
T Consume();
261277

278+
[ForceInline]
262279
void GetDimensions(
263280
out uint numStructs,
264-
out uint stride);
281+
out uint stride)
282+
{
283+
let result = __structuredBufferGetDimensions(this);
284+
numStructs = result.x;
285+
stride = result.y;
286+
}
265287
};
266288

267289
__generic<T, let N : int>

‎source/slang/slang-emit-c-like.cpp

+91-1
Original file line numberDiff line numberDiff line change
@@ -455,6 +455,66 @@ void CLikeSourceEmitter::emitRTTIObject(IRRTTIObject* rttiObject)
455455
// This is only used in targets that support dynamic dispatching.
456456
}
457457

458+
void CLikeSourceEmitter::defaultEmitInstStmt(IRInst* inst)
459+
{
460+
switch (inst->getOp())
461+
{
462+
case kIROp_AtomicCounterIncrement:
463+
{
464+
auto oldValName = getName(inst);
465+
m_writer->emit("int ");
466+
m_writer->emit(oldValName);
467+
m_writer->emit(";\n");
468+
m_writer->emit("InterlockedAdd(");
469+
emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
470+
m_writer->emit(", 1, ");
471+
m_writer->emit(oldValName);
472+
m_writer->emit(");\n");
473+
}
474+
break;
475+
case kIROp_AtomicCounterDecrement:
476+
{
477+
auto oldValName = getName(inst);
478+
m_writer->emit("int ");
479+
m_writer->emit(oldValName);
480+
m_writer->emit(";\n");
481+
m_writer->emit("InterlockedAdd(");
482+
emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
483+
m_writer->emit(", -1, ");
484+
m_writer->emit(oldValName);
485+
m_writer->emit(");\n");
486+
}
487+
break;
488+
case kIROp_StructuredBufferGetDimensions:
489+
{
490+
auto count = _generateUniqueName(UnownedStringSlice("_elementCount"));
491+
auto stride = _generateUniqueName(UnownedStringSlice("_stride"));
492+
493+
m_writer->emit("uint ");
494+
m_writer->emit(count);
495+
m_writer->emit(";\n");
496+
m_writer->emit("uint ");
497+
m_writer->emit(stride);
498+
m_writer->emit(";\n");
499+
emitOperand(inst->getOperand(0), leftSide(getInfo(EmitOp::General), getInfo(EmitOp::Postfix)));
500+
m_writer->emit(".GetDimensions(");
501+
m_writer->emit(count);
502+
m_writer->emit(", ");
503+
m_writer->emit(stride);
504+
m_writer->emit(");\n");
505+
emitInstResultDecl(inst);
506+
m_writer->emit("uint2(");
507+
m_writer->emit(count);
508+
m_writer->emit(", ");
509+
m_writer->emit(stride);
510+
m_writer->emit(");\n");
511+
}
512+
break;
513+
default:
514+
diagnoseUnhandledInst(inst);
515+
}
516+
}
517+
458518

459519
void CLikeSourceEmitter::emitTypeImpl(IRType* type, const StringSliceLoc* nameAndLoc)
460520
{
@@ -1874,6 +1934,16 @@ void CLikeSourceEmitter::emitInstExpr(IRInst* inst, const EmitOpInfo& inOuterPre
18741934
defaultEmitInstExpr(inst, inOuterPrec);
18751935
}
18761936

1937+
void CLikeSourceEmitter::emitInstStmt(IRInst* inst)
1938+
{
1939+
// Try target specific impl first
1940+
if (tryEmitInstStmtImpl(inst))
1941+
{
1942+
return;
1943+
}
1944+
defaultEmitInstStmt(inst);
1945+
}
1946+
18771947
void CLikeSourceEmitter::diagnoseUnhandledInst(IRInst* inst)
18781948
{
18791949
getSink()->diagnose(inst, Diagnostics::unimplemented, "unexpected IR opcode during code emit");
@@ -2193,6 +2263,23 @@ void CLikeSourceEmitter::defaultEmitInstExpr(IRInst* inst, const EmitOpInfo& inO
21932263
}
21942264
break;
21952265

2266+
case kIROp_StructuredBufferAppend:
2267+
{
2268+
auto outer = getInfo(EmitOp::General);
2269+
emitOperand(inst->getOperand(0), leftSide(outer, getInfo(EmitOp::Postfix)));
2270+
m_writer->emit(".Append(");
2271+
emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
2272+
m_writer->emit(")");
2273+
}
2274+
break;
2275+
case kIROp_StructuredBufferConsume:
2276+
{
2277+
auto outer = getInfo(EmitOp::General);
2278+
emitOperand(inst->getOperand(0), leftSide(outer, getInfo(EmitOp::Postfix)));
2279+
m_writer->emit(".Consume()");
2280+
}
2281+
break;
2282+
21962283
case kIROp_Call:
21972284
{
21982285
emitCallExpr((IRCall*)inst, outerPrec);
@@ -2562,7 +2649,10 @@ void CLikeSourceEmitter::_emitInst(IRInst* inst)
25622649

25632650
// Insts that needs to be emitted as code blocks.
25642651
case kIROp_CudaKernelLaunch:
2565-
emitInstStmtImpl(inst);
2652+
case kIROp_AtomicCounterIncrement:
2653+
case kIROp_AtomicCounterDecrement:
2654+
case kIROp_StructuredBufferGetDimensions:
2655+
emitInstStmt(inst);
25662656
break;
25672657

25682658
case kIROp_LiveRangeStart:

‎source/slang/slang-emit-c-like.h

+4-1
Original file line numberDiff line numberDiff line change
@@ -549,7 +549,10 @@ class CLikeSourceEmitter: public SourceEmitterBase
549549

550550
virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) { SLANG_UNUSED(varDecl); SLANG_UNUSED(varType); return false; }
551551
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) { SLANG_UNUSED(inst); SLANG_UNUSED(inOuterPrec); return false; }
552-
virtual void emitInstStmtImpl(IRInst* inst) { SLANG_UNUSED(inst); }
552+
virtual bool tryEmitInstStmtImpl(IRInst* inst) { SLANG_UNUSED(inst); return false; }
553+
554+
void defaultEmitInstStmt(IRInst* inst);
555+
void emitInstStmt(IRInst* inst);
553556

554557
virtual void emitPostKeywordTypeAttributesImpl(IRInst* inst) { SLANG_UNUSED(inst); }
555558

‎source/slang/slang-emit-glsl.cpp

+36-2
Original file line numberDiff line numberDiff line change
@@ -201,8 +201,11 @@ void GLSLSourceEmitter::_emitGLSLStructuredBuffer(IRGlobalParam* varDecl, IRHLSL
201201
m_writer->emit("buffer ");
202202

203203
// Generate a dummy name for the block
204-
m_writer->emit("_S");
205-
m_writer->emit(m_uniqueIDCounter++);
204+
StringBuilder blockTypeName;
205+
blockTypeName << "StructuredBuffer_";
206+
getTypeNameHint(blockTypeName, structuredBufferType->getElementType());
207+
blockTypeName << "_t";
208+
m_writer->emit(_generateUniqueName(blockTypeName.produceString().getUnownedSlice()));
206209

207210
m_writer->emit(" {\n");
208211
m_writer->indent();
@@ -2007,6 +2010,37 @@ bool GLSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
20072010
return false;
20082011
}
20092012

2013+
bool GLSLSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
2014+
{
2015+
switch (inst->getOp())
2016+
{
2017+
case kIROp_AtomicCounterIncrement:
2018+
{
2019+
auto oldValName = getName(inst);
2020+
m_writer->emit("int ");
2021+
m_writer->emit(oldValName);
2022+
m_writer->emit(" = ");
2023+
m_writer->emit("atomicAdd(");
2024+
emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
2025+
m_writer->emit(", 1);\n");
2026+
return true;
2027+
}
2028+
case kIROp_AtomicCounterDecrement:
2029+
{
2030+
auto oldValName = getName(inst);
2031+
m_writer->emit("int ");
2032+
m_writer->emit(oldValName);
2033+
m_writer->emit(" = ");
2034+
m_writer->emit("atomicAdd(");
2035+
emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
2036+
m_writer->emit(", -1);\n");
2037+
return true;
2038+
}
2039+
default:
2040+
return false;
2041+
}
2042+
}
2043+
20102044
void GLSLSourceEmitter::handleRequiredCapabilitiesImpl(IRInst* inst)
20112045
{
20122046
// Does this function declare any requirements on GLSL version or

‎source/slang/slang-emit-glsl.h

+2
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,8 @@ class GLSLSourceEmitter : public CLikeSourceEmitter
4848

4949
virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
5050
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
51+
virtual bool tryEmitInstStmtImpl(IRInst* inst) SLANG_OVERRIDE;
52+
5153
virtual void emitGlobalInstImpl(IRInst* inst) override;
5254
void emitBufferPointerTypeDefinition(IRInst* ptrType);
5355

‎source/slang/slang-emit-torch.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -65,12 +65,12 @@ void emitTorchScalarTypeName(SourceWriter* m_writer, IRInst* type)
6565
}
6666
}
6767

68-
void TorchCppSourceEmitter::emitInstStmtImpl(IRInst* inst)
68+
bool TorchCppSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
6969
{
7070
switch (inst->getOp())
7171
{
7272
default:
73-
return;
73+
return false;
7474
case kIROp_CudaKernelLaunch:
7575
{
7676
m_writer->emit("AT_CUDA_CHECK(cudaLaunchKernel(");
@@ -101,7 +101,7 @@ void TorchCppSourceEmitter::emitInstStmtImpl(IRInst* inst)
101101
emitOperand(inst->getOperand(4), getInfo(EmitOp::General));
102102
m_writer->emit(")));\n");
103103

104-
break;
104+
return true;
105105
}
106106
}
107107
}

‎source/slang/slang-emit-torch.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ class TorchCppSourceEmitter : public CPPSourceEmitter
1919

2020
protected:
2121
// CPPSourceEmitter overrides
22-
virtual void emitInstStmtImpl(IRInst* inst) override;
22+
virtual bool tryEmitInstStmtImpl(IRInst* inst) override;
2323

2424
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) override;
2525
virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out) override;

‎source/slang/slang-emit.cpp

+9
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@
3131
#include "slang-ir-legalize-varying-params.h"
3232
#include "slang-ir-link.h"
3333
#include "slang-ir-com-interface.h"
34+
#include "slang-ir-lower-append-consume-structured-buffer.h"
3435
#include "slang-ir-lower-binding-query.h"
3536
#include "slang-ir-lower-generics.h"
3637
#include "slang-ir-lower-tuple-types.h"
@@ -494,6 +495,14 @@ Result linkAndOptimizeIR(
494495

495496
validateIRModuleIfEnabled(codeGenContext, irModule);
496497

498+
// On non-HLSL targets, there isn't an implementation of `AppendStructuredBuffer`
499+
// and `ConsumeStructuredBuffer` types, so we lower them into normal struct types
500+
// of `RWStructuredBuffer` typed fields now.
501+
if (target != CodeGenTarget::HLSL)
502+
{
503+
lowerAppendConsumeStructuredBuffers(targetRequest, irModule, sink);
504+
}
505+
497506
// We don't need the legalize pass for C/C++ based types
498507
if(options.shouldLegalizeExistentialAndResourceTypes )
499508
{

‎source/slang/slang-ir-byte-address-legalize.cpp

+2
Original file line numberDiff line numberDiff line change
@@ -741,6 +741,8 @@ struct ByteAddressBufferLegalizationContext
741741
paramBuilder.setInsertBefore(byteAddressBufferParam);
742742

743743
auto structuredBufferParam = paramBuilder.createGlobalParam(structuredBufferParamType);
744+
if (auto nameHint = byteAddressBufferParam->findDecoration<IRNameHintDecoration>())
745+
paramBuilder.addNameHintDecoration(structuredBufferParam, nameHint->getName());
744746

745747
// The new parameter needs to be given a layout to match the existing
746748
// parameter, so that it is given the same `binding` in the generated code.

‎source/slang/slang-ir-inst-defs.h

+8
Original file line numberDiff line numberDiff line change
@@ -447,6 +447,14 @@ INST(RWStructuredBufferStore, rwstructuredBufferStore, 3, 0)
447447

448448
INST(RWStructuredBufferGetElementPtr, rwstructuredBufferGetElementPtr, 2, 0)
449449

450+
// Append/Consume-StructuredBuffer operations
451+
INST(StructuredBufferAppend, StructuredBufferAppend, 1, 0)
452+
INST(StructuredBufferConsume, StructuredBufferConsume, 1, 0)
453+
INST(StructuredBufferGetDimensions, StructuredBufferGetDimensions, 1, 0)
454+
455+
INST(AtomicCounterIncrement, AtomicCounterIncrement, 1, 0)
456+
INST(AtomicCounterDecrement, AtomicCounterDecrement, 1, 0)
457+
450458
INST(MeshOutputRef, meshOutputRef, 2, 0)
451459

452460
// Construct a vector from a scalar

0 commit comments

Comments
 (0)