From c787c4b82ba76f87069911f203eb192060b5264f Mon Sep 17 00:00:00 2001 From: Yong He Date: Mon, 28 Aug 2023 21:24:49 -0700 Subject: Add `target_switch` and `intrinsic_asm` statement. (#3154) * Add `target_switch` and `__intrinsic_asm` statement. * Cleanup. * WaveGetActiveMask, WaveGetActiveMask, WaveCountBits. * WaveIsFirstLane. * More wave intrinsics. * wave intrinsics. * merge fix. * Fix. * Fix. * Update test. * update test. * Fix. --------- Co-authored-by: Yong He --- source/slang/slang-emit-cuda.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'source/slang/slang-emit-cuda.cpp') diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 3e974f60e..fa0e3c7aa 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -436,17 +436,17 @@ void CUDASourceEmitter::_emitInitializerList(IRType* elementType, IRUse* operand m_writer->emit("\n}"); } -void CUDASourceEmitter::emitIntrinsicCallExprImpl(IRCall* inst, IRTargetIntrinsicDecoration* targetIntrinsic, EmitOpInfo const& inOuterPrec) +void CUDASourceEmitter::emitIntrinsicCallExprImpl(IRCall* inst, UnownedStringSlice intrinsicDefinition, EmitOpInfo const& inOuterPrec) { // This works around the problem, where some intrinsics that require the "half" type enabled don't use the half/float16_t type. // For example `f16tof32` can operate on float16_t *and* uint. If the input is uint, although we are // using the half feature (as far as CUDA is concerned), the half/float16_t type is not visible/directly used. - if (targetIntrinsic->getDefinition().startsWith(toSlice("__half"))) + if (intrinsicDefinition.startsWith(toSlice("__half"))) { m_extensionTracker->requireBaseType(BaseType::Half); } - Super::emitIntrinsicCallExprImpl(inst, targetIntrinsic, inOuterPrec); + Super::emitIntrinsicCallExprImpl(inst, intrinsicDefinition, inOuterPrec); } bool CUDASourceEmitter::tryEmitInstStmtImpl(IRInst* inst) -- cgit v1.2.3