diff options
| author | Yong He <yonghe@outlook.com> | 2020-10-07 00:30:10 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-10-07 00:30:10 -0700 |
| commit | 55cd4210ecff75523699337704f7561f8c11a26e (patch) | |
| tree | 5f205b93f7cc7da8f7e2ee315bd4f8ea11f90100 | |
| parent | 4ad2e52662a00f7d8b25be6d451bba33ba62947f (diff) | |
Fix C++ emit for `bit_cast` inst. (#1570)
Co-authored-by: Yong He <yhe@nvidia.com>
| -rw-r--r-- | prelude/slang-cpp-types.h | 6 | ||||
| -rw-r--r-- | prelude/slang-cuda-prelude.h | 6 | ||||
| -rw-r--r-- | 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<typename TResult, typename TInput> +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<typename TResult, typename TInput> +__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; |