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>

1 files changed, 13 insertions, 8 deletions

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: