From 1c643167a9417e75082b3898425ab9d2d999f583 Mon Sep 17 00:00:00 2001
From: jsmall-nvidia <jsmall@nvidia.com>
Date: Tue, 4 May 2021 14:44:20 -0400
Subject: More CUDA Half support (#1833)

* #include an absolute path didn't work - because paths were taken to always be relative.

* Split out StringEscapeUtil.

* Added StringEscapeUtil.

* Fix typo in unix quoting type.

* Small comment improvements.

* Try to fix linux linking issue.

* Fix typo.

* Attempt to fix linux link issue.

* Update VS proj even though nothing really changed.

* Fix another typo issue.

* Fix for windows issue.
Fixed bug.

* Make separate Utils for escaping.

* Fix typo.

* Split out into StringEscapeHandler.

* Windows shell does handle removing quotes (so remove code to remove them).

* Handle unescaping if not initiating using the shell.

* Slight improvement around shell like decoding.

* Simplify command extraction.

* Add shared-library category type.

* Fix bug in command extraction.

* Typo in transcendental category.

* Enable unit-test on in smoke test category.

* Make parsing failing output as a failing test.

* Fixes for transcendental tests. Disable tests that do not work.

* Changed category parsing.

* Removed the TestResult parameter from _gatherTestsForFile.
Made testsList only output.

* Remove testing if all tests were disabled.

* Make args of CommandLine always unescaped.

* Add category.

* Don't need escaping on unix/linux.

* Remove some no longer used functions.

* Add requireSMVersion to CUDAExtensionTracker.

* half-calc.slang now works for CUDA.

* bit-cast-16-bit works on CUDA.

* WIP handling of CUDA vector<half> types.

* Half swizzle CUDA.

* Half vector test.

* Fix swizzle half bug.

* Fix compilation issue with narrowing to Index.

* Add unary ops.

* Add some vector scalar maths ops.

* Add half vector conversions for CUDA.

* Fix erroneous comment.
---
 source/slang/slang-emit-cuda.cpp | 28 ++++++++++++++++++++++++++++
 1 file changed, 28 insertions(+)

(limited to 'source')

diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 5f7eada68..b0c2cc02b 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -173,6 +173,33 @@ void CUDASourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic*
 
     if (auto vecType = as <IRVectorType>(specOp->returnType))
     {
+        // Converting to or from half vector types is implemented prelude as convert___half functions
+        // Get the from type -> if it's half we ignore
+
+        if (specOp->op == Op::ConstructConvert)
+        {
+            auto signatureType = specOp->signatureType;
+
+            // Need to have impl of convert_float, double, int, uint, in prelude
+
+            const auto paramCount = signatureType->getParamCount();
+            SLANG_UNUSED(paramCount);
+
+            // We have 2 'params' and param 1 is the source type
+            SLANG_ASSERT(paramCount == 2);
+            IRType* paramType = signatureType->getParamType(1);
+
+            auto vecParamType = as<IRVectorType>(paramType);
+
+            if (auto baseType = as<IRBasicType>(vecParamType->getElementType()))
+            {
+                if (baseType->getBaseType() == BaseType::Half)
+                {
+                    return;
+                }
+            }
+        }
+
         if (auto baseType = as<IRBasicType>(vecType->getElementType()))
         {
             if (baseType->getBaseType() == BaseType::Half)
@@ -187,6 +214,7 @@ void CUDASourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic*
                     case Op::Neg:
 
                     case Op::ConstructFromScalar:
+                    case Op::ConstructConvert:
 
                     case Op::Leq:
                     case Op::Less:
-- 
cgit v1.2.3