[Bug]: HEVC 10-bit 4:4:4 CUDA Array 首帧卡死与色度错位：注册 pitch 和原生 UV 布局的复现及修补 #1039

diff --git a/src/nvenc/win/impl/nvenc_d3d11_on_cuda.cpp b/src/nvenc/win/impl/nvenc_d3d11_on_cuda.cpp
index 959a1a3..ea89dce 100644
--- a/src/nvenc/win/impl/nvenc_d3d11_on_cuda.cpp
+++ b/src/nvenc/win/impl/nvenc_d3d11_on_cuda.cpp
@@ -15,6 +15,39 @@
 namespace NVENC_NAMESPACE {
 #else
 namespace nvenc {
+#endif
+
+#if NVENCAPI_MAJOR_VERSION * 100 + NVENCAPI_MINOR_VERSION >= 1301
+  // Interleave two stacked 16-bit chroma planes into the native NVENC UV array.
+  // Embedded PTX avoids adding an nvcc dependency. Samples remain bit-exact;
+  // this changes their storage layout without chroma subsampling or conversion.
+  static constexpr char cuda_interleave_chroma_ptx[] = R"ptx(
+.version 6.0
+.target sm_52
+.address_size 64
+.visible .entry pack_uv(
+ .param .u64 src, .param .u64 pitch, .param .u64 dstsurf,
+ .param .u32 width, .param .u32 height)
+{
+ .reg .pred %p<3>;
+ .reg .b32 %r<20>;
+ .reg .b64 %rd<20>;
+ ld.param.u64 %rd1,[src]; ld.param.u64 %rd2,[pitch]; ld.param.u64 %rd3,[dstsurf];
+ ld.param.u32 %r1,[width]; ld.param.u32 %r2,[height];
+ mov.u32 %r3,%ctaid.x; mov.u32 %r4,%ntid.x; mov.u32 %r5,%tid.x;
+ mad.lo.u32 %r6,%r3,%r4,%r5;
+ mov.u32 %r7,%ctaid.y; mov.u32 %r8,%ntid.y; mov.u32 %r9,%tid.y;
+ mad.lo.u32 %r10,%r7,%r8,%r9;
+ setp.ge.u32 %p1,%r6,%r1; setp.ge.u32 %p2,%r10,%r2; or.pred %p1,%p1,%p2; @%p1 bra DONE;
+ cvt.u64.u32 %rd4,%r10; mul.lo.u64 %rd5,%rd4,%rd2;
+ mul.wide.u32 %rd6,%r6,2; add.u64 %rd7,%rd1,%rd5; add.u64 %rd8,%rd7,%rd6;
+ cvt.u64.u32 %rd9,%r2; mul.lo.u64 %rd10,%rd9,%rd2; add.u64 %rd11,%rd8,%rd10;
+ ld.global.u16 %r11,[%rd8]; ld.global.u16 %r12,[%rd11]; shl.b32 %r13,%r12,16; or.b32 %r14,%r11,%r13;
+ shl.b32 %r15,%r6,2;
+ sust.b.2d.b32.trap [%rd3,{%r15,%r10}],%r14;
+ DONE: ret;
+}
+)ptx";
 #endif
 
   nvenc_d3d11_on_cuda::nvenc_d3d11_on_cuda(ID3D11Device *d3d_device, shared_dll dll):
@@ -135,7 +168,14 @@ namespace nvenc {
 #if NVENCAPI_MAJOR_VERSION * 100 + NVENCAPI_MINOR_VERSION >= 1301
     else if (!load_function(functions.cuArray3DCreate, "cuArray3DCreate_v2") ||
              !load_function(functions.cuArrayDestroy, "cuArrayDestroy") ||
-             !load_function(functions.cuArrayGetPlane, "cuArrayGetPlane")) {
+             !load_function(functions.cuArrayGetPlane, "cuArrayGetPlane") ||
+             !load_function(functions.cuSurfObjectCreate, "cuSurfObjectCreate") ||
+             !load_function(functions.cuSurfObjectDestroy, "cuSurfObjectDestroy") ||
+             !load_function(functions.cuModuleLoadData, "cuModuleLoadData") ||
+             !load_function(functions.cuModuleGetFunction, "cuModuleGetFunction") ||
+             !load_function(functions.cuModuleUnload, "cuModuleUnload") ||
+             !load_function(functions.cuLaunchKernel, "cuLaunchKernel") ||
+             !load_function(functions.cuStreamSynchronize, "cuStreamSynchronize")) {
       BOOST_LOG(info) << "NvEnc: CUDA array input functions unavailable, using CUDA device pointer input";
       functions.cuArray3DCreate = nullptr;
       functions.cuArrayDestroy = nullptr;
@@ -318,7 +358,9 @@ namespace nvenc {
       CUDA_ARRAY3D_DESCRIPTOR array_descriptor = {};
       array_descriptor.Width = encoder_params.width;
       array_descriptor.Height = encoder_params.height;
-      array_descriptor.Format = CU_AD_FORMAT_UINT16_PLANAR_444;
+      // NVDEC's native 16-bit 4:4:4 layout has a Y plane and a two-channel UV
+      // plane. Registering generic three-plane CUDA storage can corrupt chroma.
+      array_descriptor.Format = CU_AD_FORMAT_YUV444_16BIT_SEMIPLANAR;
       array_descriptor.NumChannels = planar_yuv_plane_count;
       array_descriptor.Flags = CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_VIDEO_ENCODE_DECODE;
 
@@ -328,7 +370,7 @@ namespace nvenc {
         return false;
       }
 
-      for (std::uint32_t plane = 0; plane < planar_yuv_plane_count; ++plane) {
+      for (std::uint32_t plane = 0; plane < 2; ++plane) {
         if (cuda_failed(cuda_functions.cuArrayGetPlane(&cuda_array_planes[plane], cuda_array_surface, plane))) {
           BOOST_LOG(warning) << "NvEnc: cuArrayGetPlane() failed for NVENC input plane " << plane << ": error " << last_cuda_error;
           destroy_cuda_array_input();
@@ -337,24 +379,86 @@ namespace nvenc {
       }
     }
 
-    // NVENC wants the byte width of the whole allocation. The header spells this
-    // as `CUDA_ARRAY3D_DESCRIPTOR::Width * NumChannels`, which only works out to
-    // bytes for 8-bit formats; UINT16_PLANAR_444 is 2 bytes per sample across 3
-    // planes, so the row stride is width * 3 * 2.
+    if (!create_cuda_chroma_converter()) {
+      destroy_cuda_array_input();
+      return false;
+    }
+
+    // Multi-planar array registration uses the byte width of one luma row.
+    // NumChannels is the logical channel count, not an additional pitch factor.
     if (!register_cuda_input(
           NV_ENC_INPUT_RESOURCE_TYPE_CUDAARRAY,
           cuda_array_surface,
-          encoder_params.width * planar_yuv_plane_count * planar_yuv_bytes_per_sample)) {
+          encoder_params.width * planar_yuv_bytes_per_sample)) {
       BOOST_LOG(warning) << "NvEnc: CUDA array registration failed: " << last_nvenc_error_string;
       destroy_cuda_array_input();
       return false;
     }
 
+    BOOST_LOG(info) << "NvEnc: native 16-bit 4:4:4 CUDA Array layout Y+UV, luma pitch="
+                    << encoder_params.width * planar_yuv_bytes_per_sample;
+    return true;
+  }
+
+  bool
+  nvenc_d3d11_on_cuda::create_cuda_chroma_converter() {
+    // The caller holds the CUDA context and cleans up any partial allocation.
+    if (!cuda_chroma_scratch &&
+        cuda_failed(cuda_functions.cuMemAllocPitch(
+          &cuda_chroma_scratch, &cuda_chroma_scratch_pitch,
+          encoder_params.width * planar_yuv_bytes_per_sample,
+          encoder_params.height * 2, 16))) {
+      BOOST_LOG(warning) << "NvEnc: chroma scratch allocation failed: error " << last_cuda_error;
+      return false;
+    }
+
+    if (!cuda_chroma_surface) {
+      CUDA_RESOURCE_DESC resource_desc = {};
+      resource_desc.resType = CU_RESOURCE_TYPE_ARRAY;
+      resource_desc.res.array.hArray = cuda_array_planes[1];
+      if (cuda_failed(cuda_functions.cuSurfObjectCreate(&cuda_chroma_surface, &resource_desc))) {
+        BOOST_LOG(warning) << "NvEnc: chroma surface creation failed: error " << last_cuda_error;
+        return false;
+      }
+    }
+
+    if (!cuda_chroma_module &&
+        cuda_failed(cuda_functions.cuModuleLoadData(&cuda_chroma_module, cuda_interleave_chroma_ptx))) {
+      BOOST_LOG(warning) << "NvEnc: chroma PTX loading failed: error " << last_cuda_error;
+      return false;
+    }
+    if (!cuda_chroma_kernel &&
+        cuda_failed(cuda_functions.cuModuleGetFunction(&cuda_chroma_kernel, cuda_chroma_module, "pack_uv"))) {
+      BOOST_LOG(warning) << "NvEnc: chroma kernel lookup failed: error " << last_cuda_error;
+      return false;
+    }
     return true;
   }
 
   void
   nvenc_d3d11_on_cuda::destroy_cuda_array_input() {
+    // Per-frame stream synchronization completes every launch before cleanup.
+    // Surface objects must be destroyed before their underlying CUDA arrays.
+    if (cuda_chroma_surface) {
+      if (cuda_failed(cuda_functions.cuSurfObjectDestroy(cuda_chroma_surface))) {
+        BOOST_LOG(error) << "NvEnc: chroma surface destruction failed: error " << last_cuda_error;
+      }
+      cuda_chroma_surface = 0;
+    }
+    if (cuda_chroma_scratch) {
+      if (cuda_failed(cuda_functions.cuMemFree(cuda_chroma_scratch))) {
+        BOOST_LOG(error) << "NvEnc: chroma scratch release failed: error " << last_cuda_error;
+      }
+      cuda_chroma_scratch = 0;
+      cuda_chroma_scratch_pitch = 0;
+    }
+    if (cuda_chroma_module) {
+      if (cuda_failed(cuda_functions.cuModuleUnload(cuda_chroma_module))) {
+        BOOST_LOG(error) << "NvEnc: chroma module unload failed: error " << last_cuda_error;
+      }
+      cuda_chroma_module = nullptr;
+      cuda_chroma_kernel = nullptr;
+    }
     if (cuda_array_surface) {
       if (cuda_failed(cuda_functions.cuArrayDestroy(cuda_array_surface))) {
         BOOST_LOG(error) << "NvEnc: cuArrayDestroy() failed: error " << last_cuda_error;
@@ -450,16 +554,47 @@ namespace nvenc {
 #if NVENCAPI_MAJOR_VERSION * 100 + NVENCAPI_MINOR_VERSION >= 1301
     if (cuda_array_surface) {
       copy_params.dstMemoryType = CU_MEMORYTYPE_ARRAY;
+      copy_params.dstArray = cuda_array_planes[0];
       copy_params.Height = encoder_params.height;
+      if (cuda_failed(cuda_functions.cuMemcpy2D(&copy_params))) {
+        BOOST_LOG(error) << "NvEnc: luma copy to CUDA array failed: error " << last_cuda_error;
+        return false;
+      }
 
-      for (std::uint32_t plane = 0; plane < planar_yuv_plane_count; ++plane) {
-        copy_params.srcY = encoder_params.height * plane;
-        copy_params.dstArray = cuda_array_planes[plane];
-        if (cuda_failed(cuda_functions.cuMemcpy2D(&copy_params))) {
-          BOOST_LOG(error) << "NvEnc: cuMemcpy2D() to CUDA array plane " << plane
-                           << " failed: error " << last_cuda_error;
-          return false;
-        }
+      // Keep the existing D3D11 R16 texture and its Y/U/V stacked layout. Copy
+      // U and V together to pitched GPU scratch, then interleave into native UV.
+      // This avoids surface-load restrictions on the mapped D3D11 resource.
+      CUDA_MEMCPY2D chroma_copy = {};
+      chroma_copy.srcMemoryType = CU_MEMORYTYPE_ARRAY;
+      chroma_copy.srcArray = input_texture_array;
+      chroma_copy.srcY = encoder_params.height;
+      chroma_copy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+      chroma_copy.dstDevice = cuda_chroma_scratch;
+      chroma_copy.dstPitch = cuda_chroma_scratch_pitch;
+      chroma_copy.WidthInBytes = encoder_params.width * planar_yuv_bytes_per_sample;
+      chroma_copy.Height = encoder_params.height * 2;
+      if (cuda_failed(cuda_functions.cuMemcpy2D(&chroma_copy))) {
+        BOOST_LOG(error) << "NvEnc: chroma copy to CUDA scratch failed: error " << last_cuda_error;
+        return false;
+      }
+
+      std::uint32_t width = encoder_params.width;
+      std::uint32_t height = encoder_params.height;
+      void *kernel_params[] = {
+        &cuda_chroma_scratch, &cuda_chroma_scratch_pitch,
+        &cuda_chroma_surface, &width, &height
+      };
+      if (cuda_failed(cuda_functions.cuLaunchKernel(
+            cuda_chroma_kernel, (width + 31) / 32, (height + 7) / 8, 1,
+            32, 8, 1, 0, nullptr, kernel_params, nullptr))) {
+        BOOST_LOG(error) << "NvEnc: chroma interleave launch failed: error " << last_cuda_error;
+        return false;
+      }
+      // Complete writes before unmapping the D3D resource and giving NVENC the
+      // array. Use the same default stream as the existing CUDA interop copies.
+      if (cuda_failed(cuda_functions.cuStreamSynchronize(nullptr))) {
+        BOOST_LOG(error) << "NvEnc: chroma interleave synchronization failed: error " << last_cuda_error;
+        return false;
       }
     }
     else
diff --git a/src/nvenc/win/impl/nvenc_d3d11_on_cuda.h b/src/nvenc/win/impl/nvenc_d3d11_on_cuda.h
index 17a8591..d0cb3ca 100644
--- a/src/nvenc/win/impl/nvenc_d3d11_on_cuda.h
+++ b/src/nvenc/win/impl/nvenc_d3d11_on_cuda.h
@@ -37,6 +37,14 @@ namespace nvenc {
     tcuArray3DCreate *cuArray3DCreate;
     tcuArrayDestroy *cuArrayDestroy;
     tcuArrayGetPlane *cuArrayGetPlane;
+    // CUsurfObject is a 64-bit handle; the bundled CUDA shim omits these APIs.
+    CUresult(CUDAAPI *cuSurfObjectCreate)(unsigned long long *, const CUDA_RESOURCE_DESC *);
+    CUresult(CUDAAPI *cuSurfObjectDestroy)(unsigned long long);
+    tcuModuleLoadData *cuModuleLoadData;
+    tcuModuleGetFunction *cuModuleGetFunction;
+    tcuModuleUnload *cuModuleUnload;
+    tcuLaunchKernel *cuLaunchKernel;
+    tcuStreamSynchronize *cuStreamSynchronize;
 #endif
     shared_dll dll;
   };
@@ -88,6 +96,9 @@ namespace nvenc {
     bool
     create_and_register_cuda_array_input();
 
+    bool
+    create_cuda_chroma_converter();
+
     void
     destroy_cuda_array_input();
 #endif
@@ -141,7 +152,13 @@ namespace nvenc {
     size_t cuda_surface_pitch = 0;
 #if NVENCAPI_MAJOR_VERSION * 100 + NVENCAPI_MINOR_VERSION >= 1301
     CUarray cuda_array_surface = nullptr;
-    CUarray cuda_array_planes[planar_yuv_plane_count] = {};
+    // The native video array contains Y and interleaved UV, without subsampling.
+    CUarray cuda_array_planes[2] = {};
+    CUdeviceptr cuda_chroma_scratch = 0;
+    size_t cuda_chroma_scratch_pitch = 0;
+    unsigned long long cuda_chroma_surface = 0;
+    CUmodule cuda_chroma_module = nullptr;
+    CUfunction cuda_chroma_kernel = nullptr;
 #endif
   };
 }