Fixes for active mask synthesis + tests (#1370)

* Fixes for active mask synthesis + tests

There are two fixes here:

* The code generation that follows active mask synthesis was requiring CUDA SM architecture version 7.0 for one of the introduced instructions, but not all of them. This change centralizes the handling of upgrading the required CUDA SM architecture version, and makes sure that the instructions introduced by active mask synthesis request version 7.0.

* The tests for active mask synthesis were not flagged as requiring the `cuda_sm_7_0` feature when invoking `render-test-tool`, which meant they run but produce unexpected results when invoked on a GPU without the required semantics for functions like `__ballot_sync()`. This change adds the missing `-render-feature cuda_sm_7_0` to those tests.

* fixup: mark more tests that rely on implicit active mask

2 files changed, 15 insertions, 11 deletions

diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 639e7f737..25b06027d 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -520,6 +520,8 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
         }
         case kIROp_WaveMaskBallot:
         {
+            _requireCUDASMVersion(SemanticVersion(7, 0));
+
             m_writer->emit("__ballot_sync(");
             emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
             m_writer->emit(", ");
@@ -529,12 +531,7 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
         }
         case kIROp_WaveMaskMatch:
         {
-            SemanticVersion version;
-            version.set(7, 0);
-            if (version > m_extensionTracker->m_smVersion)
-            {
-                m_extensionTracker->m_smVersion = version;
-            }
+            _requireCUDASMVersion(SemanticVersion(7, 0));
 
             m_writer->emit("__match_any_sync(");
             emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
@@ -549,6 +546,14 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
     return Super::tryEmitInstExprImpl(inst, inOuterPrec);
 }
 
+void CUDASourceEmitter::_requireCUDASMVersion(SemanticVersion const& version)
+{
+    if (version > m_extensionTracker->m_smVersion)
+    {
+        m_extensionTracker->m_smVersion = version;
+    }
+}
+
 void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue)
 {
     // Does this function declare any requirements on GLSL version or
@@ -566,11 +571,7 @@ void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue)
         {
             SemanticVersion version;
             version.setFromInteger(SemanticVersion::IntegerType(smDecoration->getCUDASMVersion()));
-
-            if (version > m_extensionTracker->m_smVersion)
-            {
-                m_extensionTracker->m_smVersion = version;
-            }
+            _requireCUDASMVersion(version);
         }
     }
 }
diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h
index 669bf2d20..c0a5ac5dc 100644
--- a/source/slang/slang-emit-cuda.h
+++ b/source/slang/slang-emit-cuda.h
@@ -79,6 +79,9 @@ protected:
     void _emitInitializerList(IRType* elementType, IRUse* operands, Index operandCount);
     void _emitInitializerListValue(IRType* elementType, IRInst* value);
 
+        /// Ensure that the generated code is compiled for at least CUDA SM `version`
+    void _requireCUDASMVersion(SemanticVersion const& version);
+
     RefPtr<CUDAExtensionTracker> m_extensionTracker;
 };