diff options
| author | Yong He <yonghe@outlook.com> | 2020-06-05 13:01:06 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-06-05 13:01:06 -0700 |
| commit | 92fc3aaa835315ff08750c7b5a7498b7228e2c33 (patch) | |
| tree | 162b2b4b9411ede11d6a53015713ab431283f6c6 /source/slang/slang-emit-cuda.cpp | |
| parent | 389be08822ba554327b5948266f54acf8a312fe7 (diff) | |
| parent | e3e1cf2045f14837cfecb14e252c0e1083787b93 (diff) | |
Merge branch 'master' into findtypebynamefix
Diffstat (limited to 'source/slang/slang-emit-cuda.cpp')
| -rw-r--r-- | source/slang/slang-emit-cuda.cpp | 23 |
1 files changed, 12 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); } } } |