summaryrefslogtreecommitdiffstats
path: root/source
diff options
context:
space:
mode:
Diffstat (limited to 'source')
-rw-r--r--source/slang/hlsl.meta.slang38
-rw-r--r--source/slang/hlsl.meta.slang.h40
-rw-r--r--source/slang/slang-emit-cuda.cpp16
3 files changed, 89 insertions, 5 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index a646c1b77..669fbb440 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -347,10 +347,13 @@ __generic<T : __BuiltinType, let N : int, let M : int> bool all(matrix<T,N,M> x)
// Barrier for writes to all memory spaces (HLSL SM 5.0)
__target_intrinsic(glsl, "memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()")
+// TODO(JS): Doesn't seem to be weaker form of sync, so use this?
+__target_intrinsic(cuda, "__syncthreads()")
void AllMemoryBarrier();
// Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0)
__target_intrinsic(glsl, "memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()")
+__target_intrinsic(cuda, "__syncthreads()")
void AllMemoryBarrierWithGroupSync();
// Test if any components is non-zero (HLSL SM 1.0)
@@ -648,6 +651,7 @@ __target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBu
void DeviceMemoryBarrier();
__target_intrinsic(glsl, "memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()")
+__target_intrinsic(glsl, "__syncthreads()")
void DeviceMemoryBarrierWithGroupSync();
// Vector distance
@@ -812,99 +816,131 @@ float2 GetRenderTargetSamplePosition(int Index);
__target_intrinsic(glsl, "groupMemoryBarrier")
void GroupMemoryBarrier();
+
__target_intrinsic(glsl, "groupMemoryBarrier(), barrier()")
+__target_intrinsic(cuda, "__syncthreads()")
void GroupMemoryBarrierWithGroupSync();
// Atomics
__target_intrinsic(glsl, "$atomicAdd($A, $1)")
+__target_intrinsic(cuda, "atomicAdd(&$0, $1)")
void InterlockedAdd(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicAdd($A, $1)")
+__target_intrinsic(cuda, "atomicAdd((uint*)&$0, $1)")
void InterlockedAdd(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicAdd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAdd(&$0, $1))")
void InterlockedAdd(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicAdd($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicAdd((uint*)&$0, $1))")
void InterlockedAdd(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicAnd($A, $1)")
+__target_intrinsic(cuda, "atomicAnd(&$0, $1)")
void InterlockedAnd(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicAnd($A, $1)")
+__target_intrinsic(cuda, "atomicAnd((int*)&$0, $1)")
void InterlockedAnd(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicAnd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAnd(&$0, $1))")
void InterlockedAnd(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicAnd($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicAnd((int*)&$0, $1))")
void InterlockedAnd(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "($3 = $atomicCompSwap($A, $1, $2))")
+__target_intrinsic(cuda, "($3 = atomicCAS(&$0, $1, $2))")
void InterlockedCompareExchange(__ref int dest, int compare_value, int value, out int original_value);
__target_intrinsic(glsl, "($3 = $atomicCompSwap($A, $1, $2))")
+__target_intrinsic(cuda, "($3 = (uint)atomicCAS((int*)&$0, $1, $2))")
void InterlockedCompareExchange(__ref uint dest, uint compare_value, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicCompSwap($A, $1, $2)")
-void InterlockedCompareStore(__ref int dest, int compare_value, int value);
+__target_intrinsic(cuda, "atomicCAS(&$0, $1, $2)")
+void InterlockedCompareStore(__ref int dest, int compare_value, int value);
__target_intrinsic(glsl, "$atomicCompSwap($A, $1, $2)")
+__target_intrinsic(cuda, "atomicCAS((int*)&$0, $1, $2)")
void InterlockedCompareStore(__ref uint dest, uint compare_value, uint value);
__target_intrinsic(glsl, "($2 = $atomicExchange($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicExch(&$0, $1))")
void InterlockedExchange(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicExchange($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicExch((int*)&$0, $1))")
void InterlockedExchange(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicMax($A, $1)")
+__target_intrinsic(cuda, "atomicMax(&$0, $1)")
void InterlockedMax(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicMax($A, $1)")
+__target_intrinsic(cuda, "atomicMax((int*)&$0, $1)")
void InterlockedMax(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicMax($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicMax(&$0, $1))")
void InterlockedMax(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicMax($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicMax((int*)&$0, $1))")
void InterlockedMax(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicMin($A, $1)")
+__target_intrinsic(cuda, "atomicMin(&$0, $1)")
void InterlockedMin(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicMin($A, $1)")
+__target_intrinsic(cuda, "atomicMin((int*)&$0, $1)")
void InterlockedMin(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicMin($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicMin(&$0, $1))")
void InterlockedMin(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicMin($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicMin((int*)&$0, $1))")
void InterlockedMin(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicOr($A, $1)")
+__target_intrinsic(cuda, "atomicOr(&$0, $1)")
void InterlockedOr(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicOr($A, $1)")
+__target_intrinsic(cuda, "atomicOr((int*)&$0, $1)")
void InterlockedOr(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicOr($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicOr(&$0, $1))")
void InterlockedOr(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicOr($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicOr((int*)&$0, $1))")
void InterlockedOr(__ref uint dest, uint value, out uint original_value);
__target_intrinsic(glsl, "$atomicXor($A, $1)")
+__target_intrinsic(cuda, "atomicXor(&$0, $1)")
void InterlockedXor(__ref int dest, int value);
__target_intrinsic(glsl, "$atomicXor($A, $1)")
+__target_intrinsic(cuda, "atomicXor((int*)&$0, $1)")
void InterlockedXor(__ref uint dest, uint value);
__target_intrinsic(glsl, "($2 = $atomicXor($A, $1))")
+__target_intrinsic(cuda, "($2 = atomicXor(&$0, $1))")
void InterlockedXor(__ref int dest, int value, out int original_value);
__target_intrinsic(glsl, "($2 = $atomicXor($A, $1))")
+__target_intrinsic(cuda, "($2 = (uint)atomicXor((int*)&$0, $1))")
void InterlockedXor(__ref uint dest, uint value, out uint original_value);
// Is floating-point value finite?
diff --git a/source/slang/hlsl.meta.slang.h b/source/slang/hlsl.meta.slang.h
index b66f43103..f69d0cb3d 100644
--- a/source/slang/hlsl.meta.slang.h
+++ b/source/slang/hlsl.meta.slang.h
@@ -396,10 +396,13 @@ SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> bool all(matri
SLANG_RAW("\n")
SLANG_RAW("// Barrier for writes to all memory spaces (HLSL SM 5.0)\n")
SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()\")\n")
+SLANG_RAW("// TODO(JS): Doesn't seem to be weaker form of sync, so use this?\n")
+SLANG_RAW("__target_intrinsic(cuda, \"__syncthreads()\")\n")
SLANG_RAW("void AllMemoryBarrier();\n")
SLANG_RAW("\n")
SLANG_RAW("// Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0)\n")
SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"__syncthreads()\")\n")
SLANG_RAW("void AllMemoryBarrierWithGroupSync();\n")
SLANG_RAW("\n")
SLANG_RAW("// Test if any components is non-zero (HLSL SM 1.0)\n")
@@ -724,6 +727,7 @@ SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), memoryBarrierImage(), mem
SLANG_RAW("void DeviceMemoryBarrier();\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()\")\n")
+SLANG_RAW("__target_intrinsic(glsl, \"__syncthreads()\")\n")
SLANG_RAW("void DeviceMemoryBarrierWithGroupSync();\n")
SLANG_RAW("\n")
SLANG_RAW("// Vector distance\n")
@@ -888,99 +892,131 @@ SLANG_RAW("// Group memory barrier\n")
SLANG_RAW("__target_intrinsic(glsl, \"groupMemoryBarrier\")\n")
SLANG_RAW("void GroupMemoryBarrier();\n")
SLANG_RAW("\n")
+SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"groupMemoryBarrier(), barrier()\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"__syncthreads()\")\n")
SLANG_RAW("void GroupMemoryBarrierWithGroupSync();\n")
SLANG_RAW("\n")
SLANG_RAW("// Atomics\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicAdd($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicAdd(&$0, $1)\")\n")
SLANG_RAW("void InterlockedAdd(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicAdd($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicAdd((uint*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedAdd(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAdd($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAdd(&$0, $1))\")\n")
SLANG_RAW("void InterlockedAdd(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAdd($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicAdd((uint*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedAdd(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicAnd($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicAnd(&$0, $1)\")\n")
SLANG_RAW("void InterlockedAnd(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicAnd($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicAnd((int*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedAnd(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAnd($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAnd(&$0, $1))\")\n")
SLANG_RAW("void InterlockedAnd(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAnd($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAnd((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedAnd(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($3 = $atomicCompSwap($A, $1, $2))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($3 = atomicCAS(&$0, $1, $2))\")\n")
SLANG_RAW("void InterlockedCompareExchange(__ref int dest, int compare_value, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($3 = $atomicCompSwap($A, $1, $2))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($3 = (uint)atomicCAS((int*)&$0, $1, $2))\")\n")
SLANG_RAW("void InterlockedCompareExchange(__ref uint dest, uint compare_value, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicCompSwap($A, $1, $2)\")\n")
-SLANG_RAW("void InterlockedCompareStore(__ref int dest, int compare_value, int value);\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicCAS(&$0, $1, $2)\")\n")
+SLANG_RAW("void InterlockedCompareStore(__ref int dest, int compare_value, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicCompSwap($A, $1, $2)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicCAS((int*)&$0, $1, $2)\")\n")
SLANG_RAW("void InterlockedCompareStore(__ref uint dest, uint compare_value, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicExchange($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicExch(&$0, $1))\")\n")
SLANG_RAW("void InterlockedExchange(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicExchange($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicExch((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedExchange(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicMax($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicMax(&$0, $1)\")\n")
SLANG_RAW("void InterlockedMax(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicMax($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicMax((int*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedMax(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMax($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicMax(&$0, $1))\")\n")
SLANG_RAW("void InterlockedMax(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMax($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicMax((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedMax(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicMin($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicMin(&$0, $1)\")\n")
SLANG_RAW("void InterlockedMin(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicMin($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicMin((int*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedMin(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMin($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicMin(&$0, $1))\")\n")
SLANG_RAW("void InterlockedMin(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMin($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicMin((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedMin(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicOr($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicOr(&$0, $1)\")\n")
SLANG_RAW("void InterlockedOr(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicOr($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicOr((int*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedOr(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicOr($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicOr(&$0, $1))\")\n")
SLANG_RAW("void InterlockedOr(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicOr($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicOr((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedOr(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicXor($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicXor(&$0, $1)\")\n")
SLANG_RAW("void InterlockedXor(__ref int dest, int value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"$atomicXor($A, $1)\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"atomicXor((int*)&$0, $1)\")\n")
SLANG_RAW("void InterlockedXor(__ref uint dest, uint value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicXor($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicXor(&$0, $1))\")\n")
SLANG_RAW("void InterlockedXor(__ref int dest, int value, out int original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicXor($A, $1))\")\n")
+SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicXor((int*)&$0, $1))\")\n")
SLANG_RAW("void InterlockedXor(__ref uint dest, uint value, out uint original_value);\n")
SLANG_RAW("\n")
SLANG_RAW("// Is floating-point value finite?\n")
@@ -1589,7 +1625,7 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa)
sb << "};\n";
}
-SLANG_RAW("#line 1516 \"hlsl.meta.slang\"")
+SLANG_RAW("#line 1552 \"hlsl.meta.slang\"")
SLANG_RAW("\n")
SLANG_RAW("\n")
SLANG_RAW("\n")
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index c628e5caf..0bbaafa5b 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -470,7 +470,7 @@ void CUDASourceEmitter::emitRateQualifiersImpl(IRRate* rate)
{
if (as<IRGroupSharedRate>(rate))
{
- m_writer->emit("groupshared ");
+ m_writer->emit("__shared__ ");
}
}
@@ -599,6 +599,18 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
m_writer->emit("\n};\n\n");
}
+ // Output group shared variables
+
+ {
+ for (auto action : actions)
+ {
+ if (action.level == EmitAction::Level::Definition && action.inst->op == kIROp_GlobalVar && as<IRGroupSharedRate>(action.inst->getRate()))
+ {
+ emitGlobalInst(action.inst);
+ }
+ }
+ }
+
// Output the 'Context' which will be used for execution
{
m_writer->emit("struct Context\n{\n");
@@ -614,7 +626,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
// Output all the thread locals
for (auto action : actions)
{
- if (action.level == EmitAction::Level::Definition && action.inst->op == kIROp_GlobalVar)
+ if (action.level == EmitAction::Level::Definition && action.inst->op == kIROp_GlobalVar && !as<IRGroupSharedRate>(action.inst->getRate()))
{
emitGlobalInst(action.inst);
}
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 21:18:51 +0000