summaryrefslogtreecommitdiffstats
path: root/source/slang
diff options
context:
space:
mode:
Diffstat (limited to 'source/slang')
-rw-r--r--source/slang/slang-emit-c-like.cpp16
-rw-r--r--source/slang/slang-emit-c-like.h1
-rw-r--r--source/slang/slang-emit-cpp.cpp9
-rw-r--r--source/slang/slang-emit-cpp.h1
-rw-r--r--source/slang/slang-emit-cuda.cpp8
-rw-r--r--source/slang/slang-emit-cuda.h2
-rw-r--r--source/slang/slang-emit-glsl.cpp10
-rw-r--r--source/slang/slang-emit-glsl.h2
-rw-r--r--source/slang/slang-emit-hlsl.cpp8
-rw-r--r--source/slang/slang-emit-hlsl.h1
10 files changed, 43 insertions, 15 deletions
diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index 3631040b8..b217a2a1b 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -2592,20 +2592,8 @@ void CLikeSourceEmitter::emitRegion(Region* inRegion)
//
if (auto loopControlDecoration = loopInst->findDecoration<IRLoopControlDecoration>())
{
- switch (loopControlDecoration->getMode())
- {
- case kIRLoopControl_Unroll:
- // Note: loop unrolling control is only available in HLSL, not GLSL
- if(getSourceStyle() == SourceStyle::HLSL)
- {
- m_writer->emit("[unroll]\n");
- }
- break;
-
- default:
- break;
- }
- }
+ emitLoopControlDecorationImpl(loopControlDecoration);
+ }
m_writer->emit("for(;;)\n{\n");
m_writer->indent();
diff --git a/source/slang/slang-emit-c-like.h b/source/slang/slang-emit-c-like.h
index ddca69462..e5a8de9ec 100644
--- a/source/slang/slang-emit-c-like.h
+++ b/source/slang/slang-emit-c-like.h
@@ -328,6 +328,7 @@ public:
virtual void emitParamTypeImpl(IRType* type, String const& name);
virtual void emitIntrinsicCallExprImpl(IRCall* inst, IRTargetIntrinsicDecoration* targetIntrinsic, EmitOpInfo const& inOuterPrec);
virtual void emitFunctionPreambleImpl(IRInst* inst) { SLANG_UNUSED(inst); }
+ virtual void emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) { SLANG_UNUSED(decl); }
// Only needed for glsl output with $ prefix intrinsics - so perhaps removable in the future
virtual void emitTextureOrTextureSamplerTypeImpl(IRTextureTypeBase* type, char const* baseName) { SLANG_UNUSED(type); SLANG_UNUSED(baseName); }
diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp
index bece6c2d0..8f59da214 100644
--- a/source/slang/slang-emit-cpp.cpp
+++ b/source/slang/slang-emit-cpp.cpp
@@ -1868,6 +1868,15 @@ void CPPSourceEmitter::emitIntrinsicCallExprImpl(
return Super::emitIntrinsicCallExprImpl(inst, targetIntrinsic, inOuterPrec);
}
+void CPPSourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl)
+{
+ if (decl->getMode() == kIRLoopControl_Unroll)
+ {
+ // This relies on a suitable definition in slang-cpp-prelude.h or defined in C++ compiler invocation.
+ m_writer->emit("SLANG_UNROLL\n");
+ }
+}
+
bool CPPSourceEmitter::_tryEmitInstExprAsIntrinsic(IRInst* inst, const EmitOpInfo& inOuterPrec)
{
HLSLIntrinsic* specOp = m_intrinsicSet.add(inst);
diff --git a/source/slang/slang-emit-cpp.h b/source/slang/slang-emit-cpp.h
index 99f180850..aeef30804 100644
--- a/source/slang/slang-emit-cpp.h
+++ b/source/slang/slang-emit-cpp.h
@@ -79,6 +79,7 @@ protected:
virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
virtual void emitIntrinsicCallExprImpl(IRCall* inst, IRTargetIntrinsicDecoration* targetIntrinsic, EmitOpInfo const& inOuterPrec) SLANG_OVERRIDE;
+ virtual void emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) SLANG_OVERRIDE;
// Replaceable for classes derived from CPPSourceEmitter
virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out);
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 91439d5d3..64cb240fc 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -341,6 +341,14 @@ void CUDASourceEmitter::emitCall(const HLSLIntrinsic* specOp, IRInst* inst, cons
return Super::emitCall(specOp, inst, operands, numOperands, inOuterPrec);
}
+void CUDASourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl)
+{
+ if (decl->getMode() == kIRLoopControl_Unroll)
+ {
+ m_writer->emit("#pragma unroll\n");
+ }
+}
+
bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec)
{
switch(inst->op)
diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h
index 392e4f98d..3d23fd80f 100644
--- a/source/slang/slang-emit-cuda.h
+++ b/source/slang/slang-emit-cuda.h
@@ -49,6 +49,8 @@ protected:
virtual void emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
virtual void emitFunctionPreambleImpl(IRInst* inst) SLANG_OVERRIDE { SLANG_UNUSED(inst); m_writer->emit("__device__ "); }
+ virtual void emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) SLANG_OVERRIDE;
+
//virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
diff --git a/source/slang/slang-emit-glsl.cpp b/source/slang/slang-emit-glsl.cpp
index b433b4d94..af1cfd897 100644
--- a/source/slang/slang-emit-glsl.cpp
+++ b/source/slang/slang-emit-glsl.cpp
@@ -668,6 +668,16 @@ void GLSLSourceEmitter::_maybeEmitGLSLFlatModifier(IRType* valueType)
}
}
+void GLSLSourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl)
+{
+ if (decl->getMode() == kIRLoopControl_Unroll)
+ {
+ // https://github.com/KhronosGroup/GLSL/blob/master/extensions/ext/GL_EXT_control_flow_attributes.txt
+ m_glslExtensionTracker->requireExtension(UnownedStringSlice::fromLiteral("GL_EXT_control_flow_attributes"));
+ m_writer->emit("[[unroll]]\n");
+ }
+}
+
void GLSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
{
switch (inst->op)
diff --git a/source/slang/slang-emit-glsl.h b/source/slang/slang-emit-glsl.h
index a88a3b39b..bc430c9a8 100644
--- a/source/slang/slang-emit-glsl.h
+++ b/source/slang/slang-emit-glsl.h
@@ -46,7 +46,7 @@ protected:
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
virtual void emitSimpleValueImpl(IRInst* inst) SLANG_OVERRIDE;
-
+ virtual void emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) SLANG_OVERRIDE;
void _emitGLSLTextureOrTextureSamplerType(IRTextureTypeBase* type, char const* baseName);
void _emitGLSLStructuredBuffer(IRGlobalParam* varDecl, IRHLSLStructuredBufferTypeBase* structuredBufferType);
diff --git a/source/slang/slang-emit-hlsl.cpp b/source/slang/slang-emit-hlsl.cpp
index a0e6e872d..489e17c8e 100644
--- a/source/slang/slang-emit-hlsl.cpp
+++ b/source/slang/slang-emit-hlsl.cpp
@@ -504,6 +504,14 @@ void HLSLSourceEmitter::emitVectorTypeNameImpl(IRType* elementType, IRIntegerVal
m_writer->emit(">");
}
+void HLSLSourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl)
+{
+ if (decl->getMode() == kIRLoopControl_Unroll)
+ {
+ m_writer->emit("[unroll]\n");
+ }
+}
+
void HLSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
{
switch (inst->op)
diff --git a/source/slang/slang-emit-hlsl.h b/source/slang/slang-emit-hlsl.h
index 05fa3bb11..f9b0ad0fb 100644
--- a/source/slang/slang-emit-hlsl.h
+++ b/source/slang/slang-emit-hlsl.h
@@ -33,6 +33,7 @@ protected:
virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
virtual void emitSimpleValueImpl(IRInst* inst) SLANG_OVERRIDE;
+ virtual void emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) SLANG_OVERRIDE;
// Emit a single `register` semantic, as appropriate for a given resource-type-specific layout info
// Keyword to use in the uniform case (`register` for globals, `packoffset` inside a `cbuffer`)
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 20:32:15 +0000