diff options
| author | Tim Foley <tfoleyNV@users.noreply.github.com> | 2020-06-05 10:00:29 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-06-05 10:00:29 -0700 |
| commit | 00db8212b3266dfc6f3b1fba2d0f1f0c6fe5ec95 (patch) | |
| tree | 070f9a3a27796d6edf4f905a3bfbd43a73f06338 /source/slang/slang-emit-cuda.cpp | |
| parent | 899824eb3dbb327f3fd8479ca9bc01f8ea49ed98 (diff) | |
| parent | 3bb780724830ae830657a47e4eba008a4c0f4ff7 (diff) | |
Merge branch 'master' into loop_attrib
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); } } } |