From 55cd4210ecff75523699337704f7561f8c11a26e Mon Sep 17 00:00:00 2001 From: Yong He Date: Wed, 7 Oct 2020 00:30:10 -0700 Subject: Fix C++ emit for `bit_cast` inst. (#1570) Co-authored-by: Yong He --- prelude/slang-cpp-types.h | 6 ++++++ prelude/slang-cuda-prelude.h | 6 ++++++ source/slang/slang-emit-cpp.cpp | 4 ++-- 3 files changed, 14 insertions(+), 2 deletions(-) diff --git a/prelude/slang-cpp-types.h b/prelude/slang-cpp-types.h index b0a9c68f6..9c8fb3dec 100644 --- a/prelude/slang-cpp-types.h +++ b/prelude/slang-cpp-types.h @@ -1036,6 +1036,12 @@ struct ComputeVaryingInput typedef void(*ComputeThreadFunc)(ComputeThreadVaryingInput* varyingInput, void* uniformEntryPointParams, void* uniformState); typedef void(*ComputeFunc)(ComputeVaryingInput* varyingInput, void* uniformEntryPointParams, void* uniformState); +template +TResult slang_bit_cast(TInput val) +{ + return *(TResult*)(&val); +} + #ifdef SLANG_PRELUDE_NAMESPACE } #endif diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h index aebcffc10..a975ec99c 100644 --- a/prelude/slang-cuda-prelude.h +++ b/prelude/slang-cuda-prelude.h @@ -1237,6 +1237,12 @@ __device__ uint3 operator*(uint3 a, dim3 b) return r; } +template +__inline__ __device__ TResult slang_bit_cast(TInput val) +{ + return *(TResult*)(&val); +} + /* !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! */ diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp index 1e2482bb5..6a82815a9 100644 --- a/source/slang/slang-emit-cpp.cpp +++ b/source/slang/slang-emit-cpp.cpp @@ -2238,9 +2238,9 @@ bool CPPSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOut } case kIROp_BitCast: { - m_writer->emit("(("); + m_writer->emit("(slang_bit_cast<"); emitType(inst->getDataType()); - m_writer->emit(")("); + m_writer->emit(">("); emitOperand(inst->getOperand(0), getInfo(EmitOp::General)); m_writer->emit("))"); return true; -- cgit v1.2.3