diff options
| author | Yong He <yonghe@outlook.com> | 2022-02-25 20:49:31 -0800 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2022-02-25 20:49:31 -0800 |
| commit | c31577953d5041c82375c22d847c2eba06106c58 (patch) | |
| tree | bc685a8b63fc13cb85d160ae13df950056ca6e91 /source/slang/slang-emit-cuda.cpp | |
| parent | 8990d270e3a0c01b1f7abbf4f79556c5ef82a096 (diff) | |
Improved SCCP, inlining and resource specialization passes, legalize `ImageSubscript` for GLSL (#2146)
Diffstat (limited to 'source/slang/slang-emit-cuda.cpp')
| -rw-r--r-- | source/slang/slang-emit-cuda.cpp | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 27eb75d34..f7850179d 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -895,6 +895,24 @@ void CUDASourceEmitter::emitSimpleFuncImpl(IRFunc* func) CLikeSourceEmitter::emitSimpleFuncImpl(func); } +void CUDASourceEmitter::emitSimpleValueImpl(IRInst* inst) +{ + // Make sure we convert float to half when emitting a half literal to avoid + // overload ambiguity errors from CUDA. + if (inst->getOp() == kIROp_FloatLit) + { + if (inst->getDataType()->getOp() == kIROp_HalfType) + { + m_writer->emit("__half("); + CLikeSourceEmitter::emitSimpleValueImpl(inst); + m_writer->emit(")"); + return; + } + } + CLikeSourceEmitter::emitSimpleValueImpl(inst); +} + + void CUDASourceEmitter::emitSemanticsImpl(IRInst* inst) { Super::emitSemanticsImpl(inst); |
