summaryrefslogtreecommitdiffstats
path: root/source/slang/slang-emit-cuda.cpp
diff options
context:
space:
mode:
authorYong He <yonghe@outlook.com>2022-02-25 20:49:31 -0800
committerGitHub <noreply@github.com>2022-02-25 20:49:31 -0800
commitc31577953d5041c82375c22d847c2eba06106c58 (patch)
treebc685a8b63fc13cb85d160ae13df950056ca6e91 /source/slang/slang-emit-cuda.cpp
parent8990d270e3a0c01b1f7abbf4f79556c5ef82a096 (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.cpp18
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);