Various fixes to CUDA backend. (#1877)

- Fix emitting `StructuredBuffer<ISomething>::Load`, which triggers emitting for `IROp_WrapExistential` that is previously unhandled.

- Fix cuda layout around vectors, they should be aligned to 1,2,4,8,16 bytes instead of just using element type's alignment. That means `float4` has alignment of 16 instead of 4.

- Fix `SLANG_CUDA_HANDLE_ERROR` macro definition.

- Fix navis sometimes fail to find `Slang::kIROp_*` enum values when debugging external projects.

Co-authored-by: Yong He <yhe@nvidia.com>
Co-authored-by: jsmall-nvidia <jsmall@nvidia.com>

5 files changed, 61 insertions, 21 deletions

diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index 1b0fe7c44..47c584251 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -1794,6 +1794,17 @@ void CLikeSourceEmitter::defaultEmitInstExpr(IRInst* inst, const EmitOpInfo& inO
         }
         break;
 
+    case kIROp_WrapExistential:
+        {
+            // Normally `WrapExistential` shouldn't exist in user code at this point.
+            // The only exception is when the user is calling a stdlib generic
+            // function that has an existential type argument, for example
+            // `StructuredBuffer<ISomething>.Load()`.
+            // We can safely ignore the `wrapExistential` operation in this case.
+            emitOperand(inst->getOperand(0), outerPrec);
+        }
+        break;
+
     case kIROp_Select:
         {
             
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index 1f68623bb..90d0181f7 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -6019,3 +6019,15 @@ namespace Slang
     }
 } // namespace Slang
 
+#if SLANG_VC
+#ifdef _DEBUG
+// Natvis sometimes cannot find enum values.
+// Export symbols for them to make sure natvis works correctly when debugging external projects.
+SLANG_API const int __SlangIROpNameHint = Slang::kIROp_NameHintDecoration;
+SLANG_API const int __SlangIROpExport = Slang::kIROp_ExportDecoration;
+SLANG_API const int __SlangIROpImport = Slang::kIROp_ImportDecoration;
+SLANG_API const int __SlangIROpStringLit = Slang::kIROp_StringLit;
+SLANG_API const int __SlangIROpIntLit = Slang::kIROp_IntLit;
+#endif
+#endif
+
diff --git a/source/slang/slang-type-layout.cpp b/source/slang/slang-type-layout.cpp
index 2568547a4..02b5b5bdc 100644
--- a/source/slang/slang-type-layout.cpp
+++ b/source/slang/slang-type-layout.cpp
@@ -465,6 +465,18 @@ struct CUDALayoutRulesImpl : DefaultLayoutRulesImpl
         return arrayInfo;
     }
 
+    // Given `size` between [0, 16] return the smallest power-of-2 that is greater than or equal to `size`.
+    uint32_t getVectorAlignment(uint32_t size)
+    {
+        SLANG_ASSERT(size <= 16);
+        --size;
+        // Set every bit after the highest bit.
+        size |= (size >> 1);
+        size |= (size >> 2);
+        ++size;
+        return size;
+    }
+
     SimpleLayoutInfo GetVectorLayout(BaseType elementType, SimpleLayoutInfo elementInfo, size_t elementCount) override
     {
         // Special case bool
@@ -479,8 +491,8 @@ struct CUDALayoutRulesImpl : DefaultLayoutRulesImpl
         SimpleLayoutInfo vectorInfo;
         vectorInfo.kind = elementInfo.kind;
         vectorInfo.size = elementInfo.size * elementCount;
-        vectorInfo.alignment = elementInfo.alignment;
-    
+        vectorInfo.alignment = getVectorAlignment(
+            (uint32_t)(elementInfo.size.getFiniteValue() * elementCount));
         return vectorInfo;
     }
 
diff --git a/source/slang/slang.natvis b/source/slang/slang.natvis
index ea010c187..5af533715 100644
--- a/source/slang/slang.natvis
+++ b/source/slang/slang.natvis
@@ -83,20 +83,20 @@
 			  <If Condition="child == 0">
 				  <Break/>
 			  </If>
-			  <If Condition="child->m_op == Slang::kIROp_NameHintDecoration">
+			  <If Condition="child->m_op == __SlangIROpNameHint">
 				  <Item Name="[name]">((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.chars,[((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.numChars]s8</Item>
 			  </If>
-			  <If Condition="child->m_op == Slang::kIROp_ExportDecoration">
+			  <If Condition="child->m_op == __SlangIROpExport">
 				  <Item Name="[exportName]">((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.chars,[((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.numChars]s8</Item>
 			  </If>
-			  <If Condition="child->m_op == Slang::kIROp_ImportDecoration">
+			  <If Condition="child->m_op == __SlangIROpImport">
 				  <Item Name="[importName]">((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.chars,[((Slang::IRStringLit*)(((Slang::IRUse*)(child + 1))->usedValue))->value.stringVal.numChars]s8</Item>
 			  </If>
 			  <Exec>child = child->next</Exec>
 		  </Loop>
       </CustomListItems>
-      <Item Name="[value]" Condition="m_op == Slang::kIROp_StringLit">((IRStringLit*)this)->value.stringVal.chars,[((IRStringLit*)this)->value.stringVal.numChars]s8</Item>
-      <Item Name="[value]" Condition="m_op == Slang::kIROp_IntLit">((IRIntLit*)this)->value.intVal</Item>
+      <Item Name="[value]" Condition="m_op == __SlangIROpStringLit">((IRStringLit*)this)->value.stringVal.chars,[((IRStringLit*)this)->value.stringVal.numChars]s8</Item>
+      <Item Name="[value]" Condition="m_op == __SlangIROpIntLit">((IRIntLit*)this)->value.intVal</Item>
 	  <!--
       <Synthetic Name="[operands]">
         <DisplayString>{{count = {operandCount}}}</DisplayString>
@@ -121,14 +121,14 @@
 			<Exec>child = pOperandInst->m_decorationsAndChildren.first</Exec>
 		    <Exec>nameDecoration = 0</Exec>
 		    <Loop Condition="child != 0">
-			    <If Condition="child->m_op == Slang::kIROp_NameHintDecoration">
+			    <If Condition="child->m_op == __SlangIROpNameHint">
 				    <Exec>nameDecoration = child</Exec>
 				    <Break/>
 			    </If>
-			    <If Condition="child->m_op == Slang::kIROp_ExportDecoration &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != Slang::kIROp_NameHintDecoration)">
+			    <If Condition="child->m_op == __SlangIROpExport &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != __SlangIROpNameHint)">
 				    <Exec>nameDecoration = child</Exec>
 			    </If>
-			    <If Condition="child->m_op == Slang::kIROp_ImportDecoration &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != Slang::kIROp_NameHintDecoration)">
+			    <If Condition="child->m_op == __SlangIROpImport &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != __SlangIROpNameHint)">
 				    <Exec>nameDecoration = child</Exec>
 			    </If>
 			    <Exec>child = child->next</Exec>
@@ -150,14 +150,14 @@
 				  <Exec>child = pItem->m_decorationsAndChildren.first </Exec>
 			      <Exec>nameDecoration = 0</Exec>
 				  <Loop Condition="child != 0">
-				      <If Condition="child->m_op == Slang::kIROp_NameHintDecoration">
+				      <If Condition="child->m_op == __SlangIROpNameHint">
 						  <Exec>nameDecoration = child</Exec>
 						  <Break/>
 				      </If>
-				      <If Condition="child->m_op == Slang::kIROp_ExportDecoration &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != Slang::kIROp_NameHintDecoration)">
+				      <If Condition="child->m_op == __SlangIROpExport &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != __SlangIROpNameHint)">
 						  <Exec>nameDecoration = child</Exec>
 				      </If>
-					  <If Condition="child->m_op == Slang::kIROp_ImportDecoration &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != Slang::kIROp_NameHintDecoration)">
+					  <If Condition="child->m_op == __SlangIROpImport &amp;&amp; (nameDecoration == 0 || nameDecoration->m_op != __SlangIROpNameHint)">
 						  <Exec>nameDecoration = child</Exec>
 					  </If>
 				      <Exec>child = child->next</Exec>
diff --git a/tools/gfx/cuda/render-cuda.cpp b/tools/gfx/cuda/render-cuda.cpp
index 64da1721d..ed7f44ed2 100644
--- a/tools/gfx/cuda/render-cuda.cpp
+++ b/tools/gfx/cuda/render-cuda.cpp
@@ -98,7 +98,7 @@ static SlangResult _handleCUDAError(cudaError_t error, const char* file, int lin
     return CUDAErrorInfo(file, line, cudaGetErrorName(error), cudaGetErrorString(error)).handle();
 }
 
-#    define SLANG_CUDA_HANDLE_ERROR(x) _handleCUDAError(_res, __FILE__, __LINE__)
+#    define SLANG_CUDA_HANDLE_ERROR(x) _handleCUDAError(x, __FILE__, __LINE__)
 
 #    define SLANG_CUDA_RETURN_ON_FAIL(x)              \
         {                                             \
@@ -431,7 +431,7 @@ public:
     Slang::RefPtr<MemoryCUDAResource> m_bufferResource;
     Slang::RefPtr<CUDAResourceView> m_bufferView;
     Slang::List<uint8_t> m_cpuBuffer;
-    void setCount(Index count)
+    Result setCount(Index count)
     {
         if (isHostOnly)
         {
@@ -444,7 +444,7 @@ public:
                 m_bufferView->proxyBuffer = m_cpuBuffer.getBuffer();
                 m_bufferView->desc = viewDesc;
             }
-            return;
+            return SLANG_OK;
         }
 
         if (!m_bufferResource)
@@ -454,7 +454,9 @@ public:
             desc.sizeInBytes = count;
             m_bufferResource = new MemoryCUDAResource(desc);
             if (count)
-                cudaMalloc(&m_bufferResource->m_cudaMemory, (size_t)count);
+            {
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMalloc(&m_bufferResource->m_cudaMemory, (size_t)count));
+            }
             IResourceView::Desc viewDesc = {};
             viewDesc.type = IResourceView::Type::UnorderedAccess;
             m_bufferView = new CUDAResourceView();
@@ -467,20 +469,21 @@ public:
             void* newMemory = nullptr;
             if (count)
             {
-                cudaMalloc(&newMemory, (size_t)count);
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMalloc(&newMemory, (size_t)count));
             }
             if (oldSize)
             {
-                cudaMemcpy(
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
                     newMemory,
                     m_bufferResource->m_cudaMemory,
                     Math::Min((size_t)count, oldSize),
-                    cudaMemcpyDefault);
+                    cudaMemcpyDefault));
             }
             cudaFree(m_bufferResource->m_cudaMemory);
             m_bufferResource->m_cudaMemory = newMemory;
             m_bufferResource->getDesc()->sizeInBytes = count;
         }
+        return SLANG_OK;
     }
 
     Slang::Index getCount()
@@ -1021,7 +1024,9 @@ public:
 
         virtual SLANG_NO_THROW void SLANG_MCALL wait() override
         {
-            cuStreamSynchronize(stream);
+            auto resultCode = cuStreamSynchronize(stream);
+            if (resultCode != cudaSuccess)
+                SLANG_CUDA_HANDLE_ERROR(resultCode);
         }
 
     public: