feat(encode/nvenc): SPIR-V cursor blend over Vulkan-allocated input slots — retire the PTX kernels
A vendored PTX blob is JIT'd against the driver's ISA ceiling, so the cursor-blend module silently dies on drivers older than the generating toolkit (CUDA_ERROR_UNSUPPORTED_PTX_VERSION/INVALID_PTX, 222/218 — the KWin leg's invisible composite cursor on driver 595/CUDA 13.2 vs a CUDA 13.3 blob). SPIR-V has no such coupling, and the in-tree precedent already exists twice (vulkan_video's CSC blend, VkBridge's exportable OPAQUE_FD → cuImportExternalMemory bridge). New pf_zerocopy::vkslot::VkSlotBlend: the direct-SDK NVENC encoder now allocates its input ring as exportable Vulkan buffers CUDA-imports (same contiguous InputSurface layouts, pitch = row bytes rounded to 256), and the cursor composite is a compute dispatch over the cursor's rectangle (cursor_blend.comp, vendored .spv; spec-constant selects ARGB/NV12/YUV444; BT.709 limited, matching the retired .cu). The surface SSBO is uint[] with every invocation owning whole words — no 8-bit-storage device dependency. Cursor-bearing frames force the existing CPU-synced submit path so the CUDA copy → Vulkan dispatch (fence-waited) → NVENC encode ordering is CPU-established; cursorless frames keep the stream-ordered fast path untouched. Any bring-up/alloc/registration failure falls back wholesale to plain pitched CUDA surfaces (never a mixed or short ring): sessions always encode, composite mode just loses the cursor, warned once. cursor_blend.cu / cursor_blend.ptx and the CursorBlend PTX loader are deleted. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This commit is contained in:
Claude Fable 5
@@ -1,105 +0,0 @@
|
||||
// Cursor-overlay blend kernels for the CUDA/NVENC path (cursor-as-metadata). The cursor bitmap is
|
||||
// straight-alpha RGBA, row-packed (stride = curW*4). Blended into the encoder-OWNED NVENC input
|
||||
// surface — never the compositor's dmabuf. One thread per cursor pixel (ARGB / YUV444) or per 2x2
|
||||
// chroma block (NV12). Coefficients are BT.709 limited, matching rgb2yuv.comp so the cursor colour
|
||||
// matches the rest of the frame regardless of which zero-copy backend encodes it.
|
||||
//
|
||||
// Build (regenerate cursor_blend.ptx after editing):
|
||||
// nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx
|
||||
// PTX is JIT'd by the driver forward to the actual GPU, so a compute_75 (Turing) baseline runs on
|
||||
// every Turing-or-newer NVENC GPU. (CUDA 13's nvcc no longer targets pre-Turing archs.)
|
||||
|
||||
typedef unsigned char u8;
|
||||
|
||||
__device__ __forceinline__ u8 blend8(int dst, int src, int a) {
|
||||
return (u8)((src * a + dst * (255 - a)) / 255);
|
||||
}
|
||||
|
||||
// Packed 4-byte surface. NVENC's ARGB format stores bytes B,G,R,A in memory; the cursor is R,G,B,A.
|
||||
extern "C" __global__ void blend_argb(
|
||||
u8* surf, int pitch, int surfW, int surfH,
|
||||
const u8* cur, int curW, int curH, int ox, int oy)
|
||||
{
|
||||
int cx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int cy = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (cx >= curW || cy >= curH) return;
|
||||
int px = ox + cx, py = oy + cy;
|
||||
if (px < 0 || py < 0 || px >= surfW || py >= surfH) return;
|
||||
const u8* s = cur + (size_t)(cy * curW + cx) * 4;
|
||||
int a = s[3];
|
||||
if (a == 0) return;
|
||||
u8* d = surf + (size_t)py * pitch + (size_t)px * 4;
|
||||
d[0] = blend8(d[0], s[2], a); // B <- cursor B
|
||||
d[1] = blend8(d[1], s[1], a); // G <- cursor G
|
||||
d[2] = blend8(d[2], s[0], a); // R <- cursor R
|
||||
}
|
||||
|
||||
// Planar YUV444: three full-res planes stacked at base, base+plane, base+2*plane (plane=pitch*surfH).
|
||||
extern "C" __global__ void blend_yuv444(
|
||||
u8* base, int pitch, int surfW, int surfH,
|
||||
const u8* cur, int curW, int curH, int ox, int oy)
|
||||
{
|
||||
int cx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int cy = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (cx >= curW || cy >= curH) return;
|
||||
int px = ox + cx, py = oy + cy;
|
||||
if (px < 0 || py < 0 || px >= surfW || py >= surfH) return;
|
||||
const u8* s = cur + (size_t)(cy * curW + cx) * 4;
|
||||
int a = s[3];
|
||||
if (a == 0) return;
|
||||
float R = s[0], G = s[1], B = s[2];
|
||||
int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f);
|
||||
int U = (int)(128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B + 0.5f);
|
||||
int V = (int)(128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B + 0.5f);
|
||||
size_t plane = (size_t)pitch * surfH;
|
||||
u8* yp = base + (size_t)py * pitch + px;
|
||||
u8* up = base + plane + (size_t)py * pitch + px;
|
||||
u8* vp = base + 2 * plane + (size_t)py * pitch + px;
|
||||
*yp = blend8(*yp, Y, a);
|
||||
*up = blend8(*up, U, a);
|
||||
*vp = blend8(*vp, V, a);
|
||||
}
|
||||
|
||||
// NV12: full-res Y plane + interleaved half-res UV plane. One thread per 2x2 luma block; each blends
|
||||
// up to four Y samples and one (alpha-weighted) UV sample.
|
||||
extern "C" __global__ void blend_nv12(
|
||||
u8* yb, int yPitch, u8* uvb, int uvPitch, int surfW, int surfH,
|
||||
const u8* cur, int curW, int curH, int ox, int oy)
|
||||
{
|
||||
int bx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int by = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int base_cx = bx * 2, base_cy = by * 2;
|
||||
if (base_cx >= curW || base_cy >= curH) return;
|
||||
float ua = 0.0f, va = 0.0f, wa = 0.0f;
|
||||
int cnt = 0;
|
||||
for (int j = 0; j < 2; j++) {
|
||||
for (int i = 0; i < 2; i++) {
|
||||
int cx = base_cx + i, cy = base_cy + j;
|
||||
if (cx >= curW || cy >= curH) continue;
|
||||
int px = ox + cx, py = oy + cy;
|
||||
if (px < 0 || py < 0 || px >= surfW || py >= surfH) continue;
|
||||
const u8* s = cur + (size_t)(cy * curW + cx) * 4;
|
||||
int a = s[3];
|
||||
if (a == 0) continue;
|
||||
float R = s[0], G = s[1], B = s[2];
|
||||
int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f);
|
||||
u8* yp = yb + (size_t)py * yPitch + px;
|
||||
*yp = blend8(*yp, Y, a);
|
||||
ua += (128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B) * a;
|
||||
va += (128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B) * a;
|
||||
wa += a;
|
||||
cnt++;
|
||||
}
|
||||
}
|
||||
if (wa <= 0.0f || cnt == 0) return;
|
||||
// The chroma sample covering this block's top-left surface pixel.
|
||||
int uvx = (ox + base_cx) / 2;
|
||||
int uvy = (oy + base_cy) / 2;
|
||||
if (uvx < 0 || uvy < 0 || uvx * 2 >= surfW || uvy * 2 >= surfH) return;
|
||||
int U = (int)(ua / wa + 0.5f);
|
||||
int V = (int)(va / wa + 0.5f);
|
||||
int amean = (int)(wa / cnt + 0.5f);
|
||||
u8* uv = uvb + (size_t)uvy * uvPitch + (size_t)uvx * 2;
|
||||
uv[0] = blend8(uv[0], U, amean);
|
||||
uv[1] = blend8(uv[1], V, amean);
|
||||
}
|
||||
@@ -1,584 +0,0 @@
|
||||
//
|
||||
// Generated by NVIDIA NVVM Compiler
|
||||
//
|
||||
// Compiler Build ID: CL-38244171
|
||||
// Cuda compilation tools, release 13.3, V13.3.73
|
||||
// Based on NVVM 7.0.1
|
||||
//
|
||||
|
||||
// .version lowered from the generating toolkit's 9.3 (CUDA 13.3) by hand: a driver refuses PTX
|
||||
// whose ISA version is newer than its JIT (CUDA_ERROR_UNSUPPORTED_PTX_VERSION, 222 — driver
|
||||
// 595.58 = CUDA 13.2 rejects a 9.3 stamp), while a stamp OLDER than the body's actual syntax is
|
||||
// rejected as CUDA_ERROR_INVALID_PTX (218 — 8.0 was tried and failed against this body). 9.2 is
|
||||
// the highest ISA a CUDA-13.2 driver JITs and the lowest this 13.3-emitted body accepts.
|
||||
// TODO(portability): drivers older than the 13.2 era will refuse 9.2 too — the real fix is
|
||||
// regenerating cursor_blend.cu with an older toolkit (CUDA 12.x) so the body itself is old, then
|
||||
// stamping that toolkit's native version. Re-test on the oldest driver box after any regen.
|
||||
.version 9.2
|
||||
.target sm_75
|
||||
.address_size 64
|
||||
|
||||
// .globl blend_argb
|
||||
|
||||
.visible .entry blend_argb(
|
||||
.param .u64 blend_argb_param_0,
|
||||
.param .u32 blend_argb_param_1,
|
||||
.param .u32 blend_argb_param_2,
|
||||
.param .u32 blend_argb_param_3,
|
||||
.param .u64 blend_argb_param_4,
|
||||
.param .u32 blend_argb_param_5,
|
||||
.param .u32 blend_argb_param_6,
|
||||
.param .u32 blend_argb_param_7,
|
||||
.param .u32 blend_argb_param_8
|
||||
)
|
||||
{
|
||||
.reg .pred %p<10>;
|
||||
.reg .b16 %rs<2>;
|
||||
.reg .b32 %r<34>;
|
||||
.reg .b64 %rd<17>;
|
||||
|
||||
|
||||
ld.param.u64 %rd2, [blend_argb_param_0];
|
||||
ld.param.u32 %r5, [blend_argb_param_1];
|
||||
ld.param.u32 %r6, [blend_argb_param_2];
|
||||
ld.param.u32 %r7, [blend_argb_param_3];
|
||||
ld.param.u64 %rd3, [blend_argb_param_4];
|
||||
ld.param.u32 %r8, [blend_argb_param_5];
|
||||
ld.param.u32 %r11, [blend_argb_param_6];
|
||||
ld.param.u32 %r9, [blend_argb_param_7];
|
||||
ld.param.u32 %r10, [blend_argb_param_8];
|
||||
mov.u32 %r12, %ntid.x;
|
||||
mov.u32 %r13, %ctaid.x;
|
||||
mov.u32 %r14, %tid.x;
|
||||
mad.lo.s32 %r1, %r13, %r12, %r14;
|
||||
mov.u32 %r15, %ntid.y;
|
||||
mov.u32 %r16, %ctaid.y;
|
||||
mov.u32 %r17, %tid.y;
|
||||
mad.lo.s32 %r2, %r16, %r15, %r17;
|
||||
setp.ge.s32 %p1, %r1, %r8;
|
||||
setp.ge.s32 %p2, %r2, %r11;
|
||||
or.pred %p3, %p1, %p2;
|
||||
@%p3 bra $L__BB0_4;
|
||||
|
||||
add.s32 %r3, %r1, %r9;
|
||||
add.s32 %r4, %r2, %r10;
|
||||
or.b32 %r18, %r4, %r3;
|
||||
setp.lt.s32 %p4, %r18, 0;
|
||||
setp.ge.s32 %p5, %r3, %r6;
|
||||
or.pred %p6, %p5, %p4;
|
||||
setp.ge.s32 %p7, %r4, %r7;
|
||||
or.pred %p8, %p7, %p6;
|
||||
@%p8 bra $L__BB0_4;
|
||||
|
||||
mad.lo.s32 %r19, %r2, %r8, %r1;
|
||||
mul.wide.s32 %rd4, %r19, 4;
|
||||
cvta.to.global.u64 %rd5, %rd3;
|
||||
add.s64 %rd1, %rd5, %rd4;
|
||||
ld.global.u8 %rs1, [%rd1+3];
|
||||
setp.eq.s16 %p9, %rs1, 0;
|
||||
@%p9 bra $L__BB0_4;
|
||||
|
||||
cvt.u32.u16 %r20, %rs1;
|
||||
mul.wide.s32 %rd6, %r4, %r5;
|
||||
mul.wide.s32 %rd7, %r3, 4;
|
||||
add.s64 %rd8, %rd6, %rd7;
|
||||
cvta.to.global.u64 %rd9, %rd2;
|
||||
add.s64 %rd10, %rd9, %rd8;
|
||||
ld.global.u8 %r21, [%rd10];
|
||||
ld.global.u8 %r22, [%rd1+2];
|
||||
xor.b32 %r23, %r20, 255;
|
||||
mul.lo.s32 %r24, %r23, %r21;
|
||||
mad.lo.s32 %r25, %r22, %r20, %r24;
|
||||
mul.wide.u32 %rd11, %r25, -2139062143;
|
||||
shr.u64 %rd12, %rd11, 39;
|
||||
st.global.u8 [%rd10], %rd12;
|
||||
ld.global.u8 %r26, [%rd10+1];
|
||||
ld.global.u8 %r27, [%rd1+1];
|
||||
mul.lo.s32 %r28, %r23, %r26;
|
||||
mad.lo.s32 %r29, %r27, %r20, %r28;
|
||||
mul.wide.u32 %rd13, %r29, -2139062143;
|
||||
shr.u64 %rd14, %rd13, 39;
|
||||
st.global.u8 [%rd10+1], %rd14;
|
||||
ld.global.u8 %r30, [%rd10+2];
|
||||
ld.global.u8 %r31, [%rd1];
|
||||
mul.lo.s32 %r32, %r23, %r30;
|
||||
mad.lo.s32 %r33, %r31, %r20, %r32;
|
||||
mul.wide.u32 %rd15, %r33, -2139062143;
|
||||
shr.u64 %rd16, %rd15, 39;
|
||||
st.global.u8 [%rd10+2], %rd16;
|
||||
|
||||
$L__BB0_4:
|
||||
ret;
|
||||
|
||||
}
|
||||
// .globl blend_yuv444
|
||||
.visible .entry blend_yuv444(
|
||||
.param .u64 blend_yuv444_param_0,
|
||||
.param .u32 blend_yuv444_param_1,
|
||||
.param .u32 blend_yuv444_param_2,
|
||||
.param .u32 blend_yuv444_param_3,
|
||||
.param .u64 blend_yuv444_param_4,
|
||||
.param .u32 blend_yuv444_param_5,
|
||||
.param .u32 blend_yuv444_param_6,
|
||||
.param .u32 blend_yuv444_param_7,
|
||||
.param .u32 blend_yuv444_param_8
|
||||
)
|
||||
{
|
||||
.reg .pred %p<10>;
|
||||
.reg .b16 %rs<5>;
|
||||
.reg .f32 %f<16>;
|
||||
.reg .b32 %r<49>;
|
||||
.reg .b64 %rd<14>;
|
||||
|
||||
|
||||
ld.param.u64 %rd2, [blend_yuv444_param_0];
|
||||
ld.param.u32 %r5, [blend_yuv444_param_1];
|
||||
ld.param.u32 %r6, [blend_yuv444_param_2];
|
||||
ld.param.u32 %r7, [blend_yuv444_param_3];
|
||||
ld.param.u64 %rd3, [blend_yuv444_param_4];
|
||||
ld.param.u32 %r8, [blend_yuv444_param_5];
|
||||
ld.param.u32 %r11, [blend_yuv444_param_6];
|
||||
ld.param.u32 %r9, [blend_yuv444_param_7];
|
||||
ld.param.u32 %r10, [blend_yuv444_param_8];
|
||||
mov.u32 %r12, %ntid.x;
|
||||
mov.u32 %r13, %ctaid.x;
|
||||
mov.u32 %r14, %tid.x;
|
||||
mad.lo.s32 %r1, %r13, %r12, %r14;
|
||||
mov.u32 %r15, %ntid.y;
|
||||
mov.u32 %r16, %ctaid.y;
|
||||
mov.u32 %r17, %tid.y;
|
||||
mad.lo.s32 %r2, %r16, %r15, %r17;
|
||||
setp.ge.s32 %p1, %r1, %r8;
|
||||
setp.ge.s32 %p2, %r2, %r11;
|
||||
or.pred %p3, %p1, %p2;
|
||||
@%p3 bra $L__BB1_4;
|
||||
|
||||
add.s32 %r3, %r1, %r9;
|
||||
add.s32 %r4, %r2, %r10;
|
||||
or.b32 %r18, %r4, %r3;
|
||||
setp.lt.s32 %p4, %r18, 0;
|
||||
setp.ge.s32 %p5, %r3, %r6;
|
||||
or.pred %p6, %p5, %p4;
|
||||
setp.ge.s32 %p7, %r4, %r7;
|
||||
or.pred %p8, %p7, %p6;
|
||||
@%p8 bra $L__BB1_4;
|
||||
|
||||
mad.lo.s32 %r19, %r2, %r8, %r1;
|
||||
mul.wide.s32 %rd4, %r19, 4;
|
||||
cvta.to.global.u64 %rd5, %rd3;
|
||||
add.s64 %rd1, %rd5, %rd4;
|
||||
ld.global.u8 %rs1, [%rd1+3];
|
||||
setp.eq.s16 %p9, %rs1, 0;
|
||||
@%p9 bra $L__BB1_4;
|
||||
|
||||
cvt.u32.u16 %r20, %rs1;
|
||||
ld.global.u8 %rs2, [%rd1];
|
||||
cvt.rn.f32.u16 %f1, %rs2;
|
||||
ld.global.u8 %rs3, [%rd1+1];
|
||||
cvt.rn.f32.u16 %f2, %rs3;
|
||||
ld.global.u8 %rs4, [%rd1+2];
|
||||
cvt.rn.f32.u16 %f3, %rs4;
|
||||
fma.rn.f32 %f4, %f1, 0f3E3AFB7F, 0f41800000;
|
||||
fma.rn.f32 %f5, %f2, 0f3F1D3C36, %f4;
|
||||
fma.rn.f32 %f6, %f3, 0f3D7DF3B6, %f5;
|
||||
add.f32 %f7, %f6, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r21, %f7;
|
||||
fma.rn.f32 %f8, %f1, 0fBDCE075F, 0f43000000;
|
||||
fma.rn.f32 %f9, %f2, 0fBEAD5CFB, %f8;
|
||||
fma.rn.f32 %f10, %f3, 0f3EE0DED3, %f9;
|
||||
add.f32 %f11, %f10, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r22, %f11;
|
||||
fma.rn.f32 %f12, %f1, 0f3EE0DED3, 0f43000000;
|
||||
fma.rn.f32 %f13, %f2, 0fBECC3C9F, %f12;
|
||||
fma.rn.f32 %f14, %f3, 0fBD25119D, %f13;
|
||||
add.f32 %f15, %f14, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r23, %f15;
|
||||
mul.wide.s32 %rd6, %r4, %r5;
|
||||
cvt.s64.s32 %rd7, %r3;
|
||||
add.s64 %rd8, %rd6, %rd7;
|
||||
cvta.to.global.u64 %rd9, %rd2;
|
||||
add.s64 %rd10, %rd9, %rd8;
|
||||
ld.global.u8 %r24, [%rd10];
|
||||
mul.lo.s32 %r25, %r21, %r20;
|
||||
xor.b32 %r26, %r20, 255;
|
||||
mad.lo.s32 %r27, %r26, %r24, %r25;
|
||||
mul.hi.s32 %r28, %r27, -2139062143;
|
||||
add.s32 %r29, %r28, %r27;
|
||||
shr.u32 %r30, %r29, 31;
|
||||
shr.u32 %r31, %r29, 7;
|
||||
add.s32 %r32, %r31, %r30;
|
||||
st.global.u8 [%rd10], %r32;
|
||||
mul.wide.s32 %rd11, %r7, %r5;
|
||||
add.s64 %rd12, %rd10, %rd11;
|
||||
ld.global.u8 %r33, [%rd12];
|
||||
mul.lo.s32 %r34, %r22, %r20;
|
||||
mad.lo.s32 %r35, %r26, %r33, %r34;
|
||||
mul.hi.s32 %r36, %r35, -2139062143;
|
||||
add.s32 %r37, %r36, %r35;
|
||||
shr.u32 %r38, %r37, 31;
|
||||
shr.u32 %r39, %r37, 7;
|
||||
add.s32 %r40, %r39, %r38;
|
||||
st.global.u8 [%rd12], %r40;
|
||||
add.s64 %rd13, %rd12, %rd11;
|
||||
ld.global.u8 %r41, [%rd13];
|
||||
mul.lo.s32 %r42, %r23, %r20;
|
||||
mad.lo.s32 %r43, %r26, %r41, %r42;
|
||||
mul.hi.s32 %r44, %r43, -2139062143;
|
||||
add.s32 %r45, %r44, %r43;
|
||||
shr.u32 %r46, %r45, 31;
|
||||
shr.u32 %r47, %r45, 7;
|
||||
add.s32 %r48, %r47, %r46;
|
||||
st.global.u8 [%rd13], %r48;
|
||||
|
||||
$L__BB1_4:
|
||||
ret;
|
||||
|
||||
}
|
||||
// .globl blend_nv12
|
||||
.visible .entry blend_nv12(
|
||||
.param .u64 blend_nv12_param_0,
|
||||
.param .u32 blend_nv12_param_1,
|
||||
.param .u64 blend_nv12_param_2,
|
||||
.param .u32 blend_nv12_param_3,
|
||||
.param .u32 blend_nv12_param_4,
|
||||
.param .u32 blend_nv12_param_5,
|
||||
.param .u64 blend_nv12_param_6,
|
||||
.param .u32 blend_nv12_param_7,
|
||||
.param .u32 blend_nv12_param_8,
|
||||
.param .u32 blend_nv12_param_9,
|
||||
.param .u32 blend_nv12_param_10
|
||||
)
|
||||
{
|
||||
.reg .pred %p<43>;
|
||||
.reg .b16 %rs<17>;
|
||||
.reg .f32 %f<108>;
|
||||
.reg .b32 %r<123>;
|
||||
.reg .b64 %rd<35>;
|
||||
|
||||
|
||||
ld.param.u64 %rd11, [blend_nv12_param_0];
|
||||
ld.param.u32 %r21, [blend_nv12_param_1];
|
||||
ld.param.u64 %rd10, [blend_nv12_param_2];
|
||||
ld.param.u32 %r22, [blend_nv12_param_3];
|
||||
ld.param.u32 %r23, [blend_nv12_param_4];
|
||||
ld.param.u32 %r24, [blend_nv12_param_5];
|
||||
ld.param.u64 %rd12, [blend_nv12_param_6];
|
||||
ld.param.u32 %r25, [blend_nv12_param_7];
|
||||
ld.param.u32 %r26, [blend_nv12_param_8];
|
||||
ld.param.u32 %r27, [blend_nv12_param_9];
|
||||
ld.param.u32 %r28, [blend_nv12_param_10];
|
||||
cvta.to.global.u64 %rd1, %rd11;
|
||||
cvta.to.global.u64 %rd2, %rd12;
|
||||
mov.u32 %r29, %ntid.x;
|
||||
mov.u32 %r30, %ctaid.x;
|
||||
mov.u32 %r31, %tid.x;
|
||||
mad.lo.s32 %r32, %r30, %r29, %r31;
|
||||
mov.u32 %r33, %ntid.y;
|
||||
mov.u32 %r34, %ctaid.y;
|
||||
mov.u32 %r35, %tid.y;
|
||||
mad.lo.s32 %r36, %r34, %r33, %r35;
|
||||
shl.b32 %r1, %r32, 1;
|
||||
shl.b32 %r2, %r36, 1;
|
||||
setp.ge.s32 %p1, %r1, %r25;
|
||||
setp.ge.s32 %p2, %r2, %r26;
|
||||
or.pred %p3, %p1, %p2;
|
||||
mov.f32 %f102, 0f00000000;
|
||||
mov.f32 %f103, 0f00000000;
|
||||
mov.f32 %f104, 0f00000000;
|
||||
@%p3 bra $L__BB2_19;
|
||||
|
||||
cvt.s64.s32 %rd3, %r21;
|
||||
add.s32 %r3, %r2, %r28;
|
||||
setp.ge.s32 %p4, %r3, %r24;
|
||||
mul.lo.s32 %r4, %r2, %r25;
|
||||
mul.wide.s32 %rd4, %r3, %r21;
|
||||
add.s32 %r5, %r1, %r27;
|
||||
or.b32 %r38, %r5, %r3;
|
||||
setp.lt.s32 %p5, %r38, 0;
|
||||
mov.u32 %r121, 0;
|
||||
setp.ge.s32 %p6, %r5, %r23;
|
||||
or.pred %p7, %p6, %p5;
|
||||
or.pred %p8, %p4, %p7;
|
||||
@%p8 bra $L__BB2_4;
|
||||
|
||||
add.s32 %r40, %r1, %r4;
|
||||
mul.wide.s32 %rd13, %r40, 4;
|
||||
add.s64 %rd5, %rd2, %rd13;
|
||||
ld.global.u8 %rs1, [%rd5+3];
|
||||
setp.eq.s16 %p9, %rs1, 0;
|
||||
@%p9 bra $L__BB2_4;
|
||||
|
||||
cvt.u32.u16 %r42, %rs1;
|
||||
ld.global.u8 %rs5, [%rd5];
|
||||
cvt.rn.f32.u16 %f31, %rs5;
|
||||
ld.global.u8 %rs6, [%rd5+1];
|
||||
cvt.rn.f32.u16 %f32, %rs6;
|
||||
ld.global.u8 %rs7, [%rd5+2];
|
||||
cvt.rn.f32.u16 %f33, %rs7;
|
||||
fma.rn.f32 %f34, %f31, 0f3E3AFB7F, 0f41800000;
|
||||
fma.rn.f32 %f35, %f32, 0f3F1D3C36, %f34;
|
||||
fma.rn.f32 %f36, %f33, 0f3D7DF3B6, %f35;
|
||||
add.f32 %f37, %f36, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r43, %f37;
|
||||
cvt.s64.s32 %rd14, %r5;
|
||||
add.s64 %rd15, %rd4, %rd14;
|
||||
add.s64 %rd16, %rd1, %rd15;
|
||||
ld.global.u8 %r44, [%rd16];
|
||||
mul.lo.s32 %r45, %r43, %r42;
|
||||
xor.b32 %r46, %r42, 255;
|
||||
mad.lo.s32 %r47, %r46, %r44, %r45;
|
||||
mul.hi.s32 %r48, %r47, -2139062143;
|
||||
add.s32 %r49, %r48, %r47;
|
||||
shr.u32 %r50, %r49, 31;
|
||||
shr.u32 %r51, %r49, 7;
|
||||
add.s32 %r52, %r51, %r50;
|
||||
st.global.u8 [%rd16], %r52;
|
||||
fma.rn.f32 %f38, %f31, 0fBDCE075F, 0f43000000;
|
||||
fma.rn.f32 %f39, %f32, 0fBEAD5CFB, %f38;
|
||||
fma.rn.f32 %f40, %f33, 0f3EE0DED3, %f39;
|
||||
cvt.rn.f32.u16 %f104, %rs1;
|
||||
fma.rn.f32 %f102, %f40, %f104, 0f00000000;
|
||||
fma.rn.f32 %f41, %f31, 0f3EE0DED3, 0f43000000;
|
||||
fma.rn.f32 %f42, %f32, 0fBECC3C9F, %f41;
|
||||
fma.rn.f32 %f43, %f33, 0fBD25119D, %f42;
|
||||
fma.rn.f32 %f103, %f43, %f104, 0f00000000;
|
||||
mov.u32 %r121, 1;
|
||||
|
||||
$L__BB2_4:
|
||||
add.s32 %r7, %r1, 1;
|
||||
setp.ge.s32 %p10, %r7, %r25;
|
||||
@%p10 bra $L__BB2_8;
|
||||
|
||||
add.s32 %r8, %r7, %r27;
|
||||
or.b32 %r53, %r8, %r3;
|
||||
setp.lt.s32 %p12, %r53, 0;
|
||||
setp.ge.s32 %p13, %r8, %r23;
|
||||
or.pred %p14, %p13, %p12;
|
||||
or.pred %p15, %p4, %p14;
|
||||
@%p15 bra $L__BB2_8;
|
||||
|
||||
add.s32 %r54, %r7, %r4;
|
||||
mul.wide.s32 %rd17, %r54, 4;
|
||||
add.s64 %rd6, %rd2, %rd17;
|
||||
ld.global.u8 %rs2, [%rd6+3];
|
||||
setp.eq.s16 %p16, %rs2, 0;
|
||||
@%p16 bra $L__BB2_8;
|
||||
|
||||
cvt.u32.u16 %r55, %rs2;
|
||||
ld.global.u8 %rs8, [%rd6];
|
||||
cvt.rn.f32.u16 %f44, %rs8;
|
||||
ld.global.u8 %rs9, [%rd6+1];
|
||||
cvt.rn.f32.u16 %f45, %rs9;
|
||||
ld.global.u8 %rs10, [%rd6+2];
|
||||
cvt.rn.f32.u16 %f46, %rs10;
|
||||
fma.rn.f32 %f47, %f44, 0f3E3AFB7F, 0f41800000;
|
||||
fma.rn.f32 %f48, %f45, 0f3F1D3C36, %f47;
|
||||
fma.rn.f32 %f49, %f46, 0f3D7DF3B6, %f48;
|
||||
add.f32 %f50, %f49, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r56, %f50;
|
||||
cvt.s64.s32 %rd18, %r8;
|
||||
add.s64 %rd19, %rd4, %rd18;
|
||||
add.s64 %rd20, %rd1, %rd19;
|
||||
ld.global.u8 %r57, [%rd20];
|
||||
mul.lo.s32 %r58, %r56, %r55;
|
||||
xor.b32 %r59, %r55, 255;
|
||||
mad.lo.s32 %r60, %r59, %r57, %r58;
|
||||
mul.hi.s32 %r61, %r60, -2139062143;
|
||||
add.s32 %r62, %r61, %r60;
|
||||
shr.u32 %r63, %r62, 31;
|
||||
shr.u32 %r64, %r62, 7;
|
||||
add.s32 %r65, %r64, %r63;
|
||||
st.global.u8 [%rd20], %r65;
|
||||
fma.rn.f32 %f51, %f44, 0fBDCE075F, 0f43000000;
|
||||
fma.rn.f32 %f52, %f45, 0fBEAD5CFB, %f51;
|
||||
fma.rn.f32 %f53, %f46, 0f3EE0DED3, %f52;
|
||||
cvt.rn.f32.u16 %f54, %rs2;
|
||||
fma.rn.f32 %f102, %f53, %f54, %f102;
|
||||
fma.rn.f32 %f55, %f44, 0f3EE0DED3, 0f43000000;
|
||||
fma.rn.f32 %f56, %f45, 0fBECC3C9F, %f55;
|
||||
fma.rn.f32 %f57, %f46, 0fBD25119D, %f56;
|
||||
fma.rn.f32 %f103, %f57, %f54, %f103;
|
||||
add.f32 %f104, %f104, %f54;
|
||||
add.s32 %r121, %r121, 1;
|
||||
|
||||
$L__BB2_8:
|
||||
add.s32 %r11, %r2, 1;
|
||||
setp.ge.s32 %p17, %r11, %r26;
|
||||
add.s32 %r12, %r11, %r28;
|
||||
add.s32 %r13, %r4, %r25;
|
||||
cvt.s64.s32 %rd21, %r12;
|
||||
mul.lo.s64 %rd7, %rd21, %rd3;
|
||||
@%p17 bra $L__BB2_12;
|
||||
|
||||
setp.ge.s32 %p18, %r12, %r24;
|
||||
or.b32 %r66, %r5, %r12;
|
||||
setp.lt.s32 %p19, %r66, 0;
|
||||
or.pred %p21, %p6, %p19;
|
||||
or.pred %p22, %p18, %p21;
|
||||
@%p22 bra $L__BB2_12;
|
||||
|
||||
add.s32 %r67, %r1, %r13;
|
||||
mul.wide.s32 %rd22, %r67, 4;
|
||||
add.s64 %rd8, %rd2, %rd22;
|
||||
ld.global.u8 %rs3, [%rd8+3];
|
||||
setp.eq.s16 %p23, %rs3, 0;
|
||||
@%p23 bra $L__BB2_12;
|
||||
|
||||
cvt.u32.u16 %r68, %rs3;
|
||||
ld.global.u8 %rs11, [%rd8];
|
||||
cvt.rn.f32.u16 %f58, %rs11;
|
||||
ld.global.u8 %rs12, [%rd8+1];
|
||||
cvt.rn.f32.u16 %f59, %rs12;
|
||||
ld.global.u8 %rs13, [%rd8+2];
|
||||
cvt.rn.f32.u16 %f60, %rs13;
|
||||
fma.rn.f32 %f61, %f58, 0f3E3AFB7F, 0f41800000;
|
||||
fma.rn.f32 %f62, %f59, 0f3F1D3C36, %f61;
|
||||
fma.rn.f32 %f63, %f60, 0f3D7DF3B6, %f62;
|
||||
add.f32 %f64, %f63, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r69, %f64;
|
||||
cvt.s64.s32 %rd23, %r5;
|
||||
add.s64 %rd24, %rd7, %rd23;
|
||||
add.s64 %rd25, %rd1, %rd24;
|
||||
ld.global.u8 %r70, [%rd25];
|
||||
mul.lo.s32 %r71, %r69, %r68;
|
||||
xor.b32 %r72, %r68, 255;
|
||||
mad.lo.s32 %r73, %r72, %r70, %r71;
|
||||
mul.hi.s32 %r74, %r73, -2139062143;
|
||||
add.s32 %r75, %r74, %r73;
|
||||
shr.u32 %r76, %r75, 31;
|
||||
shr.u32 %r77, %r75, 7;
|
||||
add.s32 %r78, %r77, %r76;
|
||||
st.global.u8 [%rd25], %r78;
|
||||
fma.rn.f32 %f65, %f58, 0fBDCE075F, 0f43000000;
|
||||
fma.rn.f32 %f66, %f59, 0fBEAD5CFB, %f65;
|
||||
fma.rn.f32 %f67, %f60, 0f3EE0DED3, %f66;
|
||||
cvt.rn.f32.u16 %f68, %rs3;
|
||||
fma.rn.f32 %f102, %f67, %f68, %f102;
|
||||
fma.rn.f32 %f69, %f58, 0f3EE0DED3, 0f43000000;
|
||||
fma.rn.f32 %f70, %f59, 0fBECC3C9F, %f69;
|
||||
fma.rn.f32 %f71, %f60, 0fBD25119D, %f70;
|
||||
fma.rn.f32 %f103, %f71, %f68, %f103;
|
||||
add.f32 %f104, %f104, %f68;
|
||||
add.s32 %r121, %r121, 1;
|
||||
|
||||
$L__BB2_12:
|
||||
or.pred %p26, %p17, %p10;
|
||||
@%p26 bra $L__BB2_16;
|
||||
|
||||
setp.ge.s32 %p27, %r12, %r24;
|
||||
add.s32 %r16, %r7, %r27;
|
||||
or.b32 %r79, %r16, %r12;
|
||||
setp.lt.s32 %p28, %r79, 0;
|
||||
setp.ge.s32 %p29, %r16, %r23;
|
||||
or.pred %p30, %p29, %p28;
|
||||
or.pred %p31, %p27, %p30;
|
||||
@%p31 bra $L__BB2_16;
|
||||
|
||||
add.s32 %r80, %r7, %r13;
|
||||
mul.wide.s32 %rd26, %r80, 4;
|
||||
add.s64 %rd9, %rd2, %rd26;
|
||||
ld.global.u8 %rs4, [%rd9+3];
|
||||
setp.eq.s16 %p32, %rs4, 0;
|
||||
@%p32 bra $L__BB2_16;
|
||||
|
||||
cvt.u32.u16 %r81, %rs4;
|
||||
ld.global.u8 %rs14, [%rd9];
|
||||
cvt.rn.f32.u16 %f72, %rs14;
|
||||
ld.global.u8 %rs15, [%rd9+1];
|
||||
cvt.rn.f32.u16 %f73, %rs15;
|
||||
ld.global.u8 %rs16, [%rd9+2];
|
||||
cvt.rn.f32.u16 %f74, %rs16;
|
||||
fma.rn.f32 %f75, %f72, 0f3E3AFB7F, 0f41800000;
|
||||
fma.rn.f32 %f76, %f73, 0f3F1D3C36, %f75;
|
||||
fma.rn.f32 %f77, %f74, 0f3D7DF3B6, %f76;
|
||||
add.f32 %f78, %f77, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r82, %f78;
|
||||
cvt.s64.s32 %rd27, %r16;
|
||||
add.s64 %rd28, %rd7, %rd27;
|
||||
add.s64 %rd29, %rd1, %rd28;
|
||||
ld.global.u8 %r83, [%rd29];
|
||||
mul.lo.s32 %r84, %r82, %r81;
|
||||
xor.b32 %r85, %r81, 255;
|
||||
mad.lo.s32 %r86, %r85, %r83, %r84;
|
||||
mul.hi.s32 %r87, %r86, -2139062143;
|
||||
add.s32 %r88, %r87, %r86;
|
||||
shr.u32 %r89, %r88, 31;
|
||||
shr.u32 %r90, %r88, 7;
|
||||
add.s32 %r91, %r90, %r89;
|
||||
st.global.u8 [%rd29], %r91;
|
||||
fma.rn.f32 %f79, %f72, 0fBDCE075F, 0f43000000;
|
||||
fma.rn.f32 %f80, %f73, 0fBEAD5CFB, %f79;
|
||||
fma.rn.f32 %f81, %f74, 0f3EE0DED3, %f80;
|
||||
cvt.rn.f32.u16 %f82, %rs4;
|
||||
fma.rn.f32 %f102, %f81, %f82, %f102;
|
||||
fma.rn.f32 %f83, %f72, 0f3EE0DED3, 0f43000000;
|
||||
fma.rn.f32 %f84, %f73, 0fBECC3C9F, %f83;
|
||||
fma.rn.f32 %f85, %f74, 0fBD25119D, %f84;
|
||||
fma.rn.f32 %f103, %f85, %f82, %f103;
|
||||
add.f32 %f104, %f104, %f82;
|
||||
add.s32 %r121, %r121, 1;
|
||||
|
||||
$L__BB2_16:
|
||||
setp.eq.s32 %p33, %r121, 0;
|
||||
setp.le.f32 %p34, %f104, 0f00000000;
|
||||
or.pred %p35, %p34, %p33;
|
||||
@%p35 bra $L__BB2_19;
|
||||
|
||||
shr.u32 %r92, %r5, 31;
|
||||
add.s32 %r93, %r5, %r92;
|
||||
shr.s32 %r19, %r93, 1;
|
||||
setp.lt.s32 %p36, %r3, -1;
|
||||
setp.lt.s32 %p37, %r5, -1;
|
||||
or.pred %p38, %p37, %p36;
|
||||
and.b32 %r94, %r93, -2;
|
||||
setp.ge.s32 %p39, %r94, %r23;
|
||||
or.pred %p40, %p38, %p39;
|
||||
shr.u32 %r95, %r3, 31;
|
||||
add.s32 %r96, %r3, %r95;
|
||||
shr.s32 %r20, %r96, 1;
|
||||
and.b32 %r97, %r96, -2;
|
||||
setp.ge.s32 %p41, %r97, %r24;
|
||||
or.pred %p42, %p40, %p41;
|
||||
@%p42 bra $L__BB2_19;
|
||||
|
||||
div.rn.f32 %f86, %f102, %f104;
|
||||
add.f32 %f87, %f86, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r98, %f87;
|
||||
div.rn.f32 %f88, %f103, %f104;
|
||||
add.f32 %f89, %f88, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r99, %f89;
|
||||
cvt.rn.f32.s32 %f90, %r121;
|
||||
div.rn.f32 %f91, %f104, %f90;
|
||||
add.f32 %f92, %f91, 0f3F000000;
|
||||
cvt.rzi.s32.f32 %r100, %f92;
|
||||
mul.wide.s32 %rd30, %r20, %r22;
|
||||
mul.wide.s32 %rd31, %r19, 2;
|
||||
add.s64 %rd32, %rd30, %rd31;
|
||||
cvta.to.global.u64 %rd33, %rd10;
|
||||
add.s64 %rd34, %rd33, %rd32;
|
||||
ld.global.u8 %r101, [%rd34];
|
||||
mul.lo.s32 %r102, %r100, %r98;
|
||||
mov.u32 %r103, 255;
|
||||
sub.s32 %r104, %r103, %r100;
|
||||
mad.lo.s32 %r105, %r104, %r101, %r102;
|
||||
mul.hi.s32 %r106, %r105, -2139062143;
|
||||
add.s32 %r107, %r106, %r105;
|
||||
shr.u32 %r108, %r107, 31;
|
||||
shr.u32 %r109, %r107, 7;
|
||||
add.s32 %r110, %r109, %r108;
|
||||
st.global.u8 [%rd34], %r110;
|
||||
ld.global.u8 %r111, [%rd34+1];
|
||||
mul.lo.s32 %r112, %r100, %r99;
|
||||
mad.lo.s32 %r113, %r104, %r111, %r112;
|
||||
mul.hi.s32 %r114, %r113, -2139062143;
|
||||
add.s32 %r115, %r114, %r113;
|
||||
shr.u32 %r116, %r115, 31;
|
||||
shr.u32 %r117, %r115, 7;
|
||||
add.s32 %r118, %r117, %r116;
|
||||
st.global.u8 [%rd34+1], %r118;
|
||||
|
||||
$L__BB2_19:
|
||||
ret;
|
||||
|
||||
}
|
||||
|
||||
@@ -69,6 +69,7 @@ use super::{AuChunk, ChromaFormat, Codec, EncodedFrame, Encoder, EncoderCaps};
|
||||
use anyhow::{anyhow, bail, Context, Result};
|
||||
use pf_frame::{CapturedFrame, FramePayload};
|
||||
use pf_zerocopy::cuda::{self, InputSurface};
|
||||
use pf_zerocopy::vkslot::{SlotFormat, VkSlotBlend, VkSlotRef};
|
||||
use std::collections::VecDeque;
|
||||
use std::ffi::c_void;
|
||||
use std::ptr;
|
||||
@@ -76,16 +77,6 @@ use std::sync::mpsc;
|
||||
|
||||
use nvidia_video_codec_sdk::sys::nvEncodeAPI as nv;
|
||||
|
||||
/// Prebuilt PTX for the cursor-overlay blend kernels (cursor-as-metadata). Source is
|
||||
/// `cursor_blend.cu` beside this file; regenerate with
|
||||
/// `nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx` after editing. JIT'd by the
|
||||
/// driver, so it runs on any Turing-or-newer GPU. ⚠️ The `.version` stamp is load-bearing (see
|
||||
/// the comment in the .ptx): a driver refuses PTX with an ISA newer than its JIT (error 222),
|
||||
/// and a stamp older than the body's real syntax is INVALID_PTX (218) — so a new toolkit's
|
||||
/// output silently kills cursor compositing on older-driver boxes. Prefer regenerating with the
|
||||
/// OLDEST toolchain that compiles the .cu, and re-test on the oldest driver box.
|
||||
const CURSOR_PTX: &[u8] = include_bytes!("cursor_blend.ptx");
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Runtime-loaded NVENC entry table (Linux). Same shape as the Windows backend's `EncodeApi`, minus
|
||||
// the async-event entry points (Windows-only). Resolved once from `libnvidia-encode.so.1` — the two
|
||||
@@ -432,13 +423,54 @@ fn buffer_format(buf: &cuda::DeviceBuffer) -> nv::NV_ENC_BUFFER_FORMAT {
|
||||
}
|
||||
}
|
||||
|
||||
/// One encoder-owned CUDA input surface + its NVENC registration. The surface is copied into each
|
||||
/// One encoder-owned input surface + its NVENC registration. The surface is copied into each
|
||||
/// use (device→device) and the registration is created once at session init, unregistered at teardown.
|
||||
struct RingSlot {
|
||||
surface: InputSurface,
|
||||
surface: SlotSurface,
|
||||
reg: nv::NV_ENC_REGISTERED_PTR,
|
||||
}
|
||||
|
||||
/// The ring slot's backing allocation: Vulkan external memory CUDA-imported (the normal case —
|
||||
/// blendable by the SPIR-V cursor pass, see `vkslot.rs`) or a plain pitched CUDA allocation (the
|
||||
/// fallback when Vulkan bring-up fails: sessions still encode, composite mode just has no
|
||||
/// cursor). Both present the same `(ptr, pitch, height)` NVENC-registration vocabulary.
|
||||
enum SlotSurface {
|
||||
Cuda(InputSurface),
|
||||
/// Backing objects live in the encoder's [`VkSlotBlend`] (freed by its `free_slots`); the
|
||||
/// ref itself is Copy and carries the registered geometry.
|
||||
Vk(VkSlotRef),
|
||||
}
|
||||
|
||||
/// The [`SlotFormat`] for an NVENC buffer format (the ring-build + blend vocabulary).
|
||||
fn slot_fmt_of(fmt: nv::NV_ENC_BUFFER_FORMAT) -> SlotFormat {
|
||||
match fmt {
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => SlotFormat::Yuv444,
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => SlotFormat::Nv12,
|
||||
_ => SlotFormat::Argb,
|
||||
}
|
||||
}
|
||||
|
||||
impl SlotSurface {
|
||||
fn ptr(&self) -> pf_zerocopy::cuda::CUdeviceptr {
|
||||
match self {
|
||||
SlotSurface::Cuda(s) => s.ptr,
|
||||
SlotSurface::Vk(r) => r.ptr,
|
||||
}
|
||||
}
|
||||
fn pitch(&self) -> usize {
|
||||
match self {
|
||||
SlotSurface::Cuda(s) => s.pitch,
|
||||
SlotSurface::Vk(r) => r.pitch,
|
||||
}
|
||||
}
|
||||
fn height(&self) -> u32 {
|
||||
match self {
|
||||
SlotSurface::Cuda(s) => s.height,
|
||||
SlotSurface::Vk(r) => r.height,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// `doNotWait` sampling cadence inside [`Encoder::poll_chunk`] — the probe measured ~200 µs
|
||||
/// between slice completions on the 5070 Ti, so 50 µs keeps the added per-chunk delivery delay
|
||||
/// well under one slice time without hammering the driver.
|
||||
@@ -534,16 +566,18 @@ pub struct NvencCudaEncoder {
|
||||
split_mode: u32,
|
||||
/// The last reference-frame range we invalidated — dedupes repeated RFI requests for one loss.
|
||||
last_rfi_range: Option<(i64, i64)>,
|
||||
/// Cursor-as-metadata GPU blend (loaded lazily on the first frame that carries a cursor, once the
|
||||
/// CUDA context is current). `None` until then or if the module load fails; `cursor_tried` stops
|
||||
/// re-attempting a failed load every frame. `cursor_serial` tracks the uploaded bitmap.
|
||||
cursor: Option<cuda::CursorBlend>,
|
||||
/// Cursor-as-metadata GPU blend: the Vulkan device + SPIR-V compute pass the ring's
|
||||
/// external-memory slots are allocated through (`vkslot.rs`) — the driver-portable
|
||||
/// replacement for the retired PTX kernels. Brought up once at session init (`cursor_tried`
|
||||
/// stops re-attempts); `None` = bring-up failed, the ring fell back to plain CUDA
|
||||
/// allocations and composite mode degrades to no cursor. `cursor_serial` tracks the
|
||||
/// uploaded bitmap.
|
||||
vk_blend: Option<VkSlotBlend>,
|
||||
cursor_tried: bool,
|
||||
cursor_serial: u64,
|
||||
/// Suppress-until-success latches for the per-frame cursor upload/blend warns: a persistent
|
||||
/// failure sits in the submit() hot path, so warn once per failure streak (reset on success)
|
||||
/// rather than on every cursor-bearing frame, which would evict the log ring.
|
||||
cursor_upload_warned: bool,
|
||||
/// Suppress-until-success latch for the per-frame blend warn: a persistent failure sits in
|
||||
/// the submit() hot path, so warn once per failure streak (reset on success) rather than on
|
||||
/// every cursor-bearing frame, which would evict the log ring.
|
||||
cursor_blend_warned: bool,
|
||||
/// One-shot latch for [`diagnose_failed_open`](Self::diagnose_failed_open) so a rebuild-retry
|
||||
/// burst (the session loop's bounded encoder resets) logs the diagnosis once, not per attempt.
|
||||
@@ -649,10 +683,9 @@ impl NvencCudaEncoder {
|
||||
frame_idx: 0,
|
||||
force_kf: false,
|
||||
pending_anchor: false,
|
||||
cursor: None,
|
||||
vk_blend: None,
|
||||
cursor_tried: false,
|
||||
cursor_serial: u64::MAX,
|
||||
cursor_upload_warned: false,
|
||||
cursor_blend_warned: false,
|
||||
diagnosed: false,
|
||||
inited: false,
|
||||
@@ -743,7 +776,12 @@ impl NvencCudaEncoder {
|
||||
// (the forfeit contract), and the next session re-latches the arming at init.
|
||||
self.subframe_chunks = false;
|
||||
self.chunk = None;
|
||||
self.ring.clear(); // drops the InputSurfaces, freeing their CUDA allocations
|
||||
self.ring.clear(); // drops the CUDA InputSurfaces; Vk slots are freed just below
|
||||
if let Some(vk) = &mut self.vk_blend {
|
||||
// The Vulkan-backed slots' memory (and its CUDA mapping) — the device itself stays
|
||||
// up for the next session's ring (`cursor_tried` keeps bring-up one-shot).
|
||||
vk.free_slots();
|
||||
}
|
||||
self.bitstreams.clear();
|
||||
self.pending.clear();
|
||||
self.encoder = ptr::null_mut();
|
||||
@@ -1132,36 +1170,102 @@ impl NvencCudaEncoder {
|
||||
|
||||
// Encoder-owned input-surface ring: allocate + register POOL surfaces in the negotiated
|
||||
// format. Registered once here, mapped per submit, unregistered at teardown.
|
||||
for _ in 0..POOL {
|
||||
let surface = match self.buffer_fmt {
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => {
|
||||
InputSurface::alloc_yuv444(self.width, self.height)
|
||||
}
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => {
|
||||
InputSurface::alloc_nv12(self.width, self.height)
|
||||
}
|
||||
_ => InputSurface::alloc_rgb(self.width, self.height),
|
||||
// Preferred backing = Vulkan external memory CUDA-imported (`vkslot.rs`), so the
|
||||
// SPIR-V cursor blend can composite into the very bytes NVENC encodes; any bring-up
|
||||
// or per-slot failure falls back to plain pitched CUDA allocations (sessions always
|
||||
// encode — composite mode just loses the cursor, warned below).
|
||||
if !self.cursor_tried {
|
||||
self.cursor_tried = true;
|
||||
match VkSlotBlend::new() {
|
||||
Ok(v) => self.vk_blend = Some(v),
|
||||
Err(e) => tracing::warn!(
|
||||
error = %format!("{e:#}"),
|
||||
"NVENC (Linux): Vulkan slot-blend bring-up failed — plain CUDA input \
|
||||
surfaces, cursor compositing unavailable"
|
||||
),
|
||||
}
|
||||
.context("alloc NVENC input surface")?;
|
||||
let mut rr = nv::NV_ENC_REGISTER_RESOURCE {
|
||||
version: nv::NV_ENC_REGISTER_RESOURCE_VER,
|
||||
resourceType:
|
||||
nv::NV_ENC_INPUT_RESOURCE_TYPE::NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR,
|
||||
width: self.width,
|
||||
height: self.height,
|
||||
pitch: surface.pitch as u32,
|
||||
resourceToRegister: surface.ptr as *mut c_void,
|
||||
bufferFormat: self.buffer_fmt,
|
||||
bufferUsage: nv::NV_ENC_BUFFER_USAGE::NV_ENC_INPUT_IMAGE,
|
||||
..Default::default()
|
||||
};
|
||||
(api().register_resource)(self.encoder, &mut rr)
|
||||
.nv_ok()
|
||||
.map_err(|e| nvenc_status::call_err("register_resource (CUDADEVICEPTR)", e))?;
|
||||
self.ring.push(RingSlot {
|
||||
surface,
|
||||
reg: rr.registeredResource,
|
||||
});
|
||||
}
|
||||
let slot_fmt = slot_fmt_of(self.buffer_fmt);
|
||||
// Two attempts: the full ring on Vulkan slots, else (any failure) the full ring on
|
||||
// plain CUDA — never a mixed ring (it would blend on some slots only: a flickering
|
||||
// cursor) and never a short one.
|
||||
'ring: for use_vk in [self.vk_blend.is_some(), false] {
|
||||
if !use_vk && self.vk_blend.is_some() {
|
||||
// Second attempt: retire the Vulkan side wholesale first.
|
||||
for s in self.ring.drain(..) {
|
||||
let _ = (api().unregister_resource)(self.encoder, s.reg);
|
||||
}
|
||||
if let Some(vk) = &mut self.vk_blend {
|
||||
vk.free_slots();
|
||||
}
|
||||
self.vk_blend = None;
|
||||
}
|
||||
for _ in 0..POOL {
|
||||
let surface = if use_vk {
|
||||
let vk = self.vk_blend.as_mut().expect("use_vk implies Some");
|
||||
match vk.alloc_slot(slot_fmt, self.width, self.height) {
|
||||
Ok(r) => SlotSurface::Vk(r),
|
||||
Err(e) => {
|
||||
tracing::warn!(
|
||||
error = %format!("{e:#}"),
|
||||
"NVENC (Linux): Vulkan slot alloc failed — rebuilding the \
|
||||
ring on plain CUDA surfaces (cursor compositing \
|
||||
unavailable)"
|
||||
);
|
||||
continue 'ring;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
SlotSurface::Cuda(
|
||||
match self.buffer_fmt {
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => {
|
||||
InputSurface::alloc_yuv444(self.width, self.height)
|
||||
}
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => {
|
||||
InputSurface::alloc_nv12(self.width, self.height)
|
||||
}
|
||||
_ => InputSurface::alloc_rgb(self.width, self.height),
|
||||
}
|
||||
.context("alloc NVENC input surface")?,
|
||||
)
|
||||
};
|
||||
let mut rr = nv::NV_ENC_REGISTER_RESOURCE {
|
||||
version: nv::NV_ENC_REGISTER_RESOURCE_VER,
|
||||
resourceType:
|
||||
nv::NV_ENC_INPUT_RESOURCE_TYPE::NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR,
|
||||
width: self.width,
|
||||
height: self.height,
|
||||
pitch: surface.pitch() as u32,
|
||||
resourceToRegister: surface.ptr() as *mut c_void,
|
||||
bufferFormat: self.buffer_fmt,
|
||||
bufferUsage: nv::NV_ENC_BUFFER_USAGE::NV_ENC_INPUT_IMAGE,
|
||||
..Default::default()
|
||||
};
|
||||
match (api().register_resource)(self.encoder, &mut rr).nv_ok() {
|
||||
Ok(()) => {}
|
||||
Err(e) if use_vk => {
|
||||
// NVENC refusing the imported pointer is a Vulkan-side condition
|
||||
// too — same wholesale fallback.
|
||||
tracing::warn!(
|
||||
error = ?e,
|
||||
"NVENC (Linux): registering a Vulkan-imported slot failed — \
|
||||
rebuilding the ring on plain CUDA surfaces"
|
||||
);
|
||||
continue 'ring;
|
||||
}
|
||||
Err(e) => {
|
||||
return Err(nvenc_status::call_err(
|
||||
"register_resource (CUDADEVICEPTR)",
|
||||
e,
|
||||
))
|
||||
}
|
||||
}
|
||||
self.ring.push(RingSlot {
|
||||
surface,
|
||||
reg: rr.registeredResource,
|
||||
});
|
||||
}
|
||||
break 'ring; // full ring built
|
||||
}
|
||||
|
||||
self.inited = true;
|
||||
@@ -1246,9 +1350,9 @@ impl NvencCudaEncoder {
|
||||
/// IO-stream binding (stream-ordered submit — see the gate in [`Encoder::submit`]).
|
||||
fn copy_into_slot(&self, buf: &cuda::DeviceBuffer, slot: usize, sync: bool) -> Result<()> {
|
||||
let s = &self.ring[slot].surface;
|
||||
let base = s.ptr;
|
||||
let pitch = s.pitch;
|
||||
let hh = s.height as u64;
|
||||
let base = s.ptr();
|
||||
let pitch = s.pitch();
|
||||
let hh = s.height() as u64;
|
||||
match self.buffer_fmt {
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => {
|
||||
if !buf.yuv444 {
|
||||
@@ -1400,79 +1504,57 @@ impl Encoder for NvencCudaEncoder {
|
||||
// `async_rt` must be absent too: in two-thread mode the frame may be recycled right after
|
||||
// submit returns while the stream still holds its copy (belt-and-braces — an escalated
|
||||
// session was rebuilt without the binding, so `stream_ordered` is false there anyway).
|
||||
let ordered = self.stream_ordered && self.async_rt.is_none() && self.pending.is_empty();
|
||||
// Cursor-bearing frames additionally force the CPU-synced path: the Vulkan blend sits
|
||||
// between the CUDA copy and the encode, and its cross-API ordering is fence/CPU-
|
||||
// established, not stream-ordered. Frames without a cursor (games hide it; client-draws
|
||||
// sessions strip it) keep the stream-ordered fast path untouched.
|
||||
let ordered = self.stream_ordered
|
||||
&& self.async_rt.is_none()
|
||||
&& self.pending.is_empty()
|
||||
&& captured.cursor.is_none();
|
||||
let t0 = std::time::Instant::now();
|
||||
|
||||
// Copy the captured buffer into this slot's input surface before encoding it.
|
||||
self.copy_into_slot(buf, slot, !ordered)?;
|
||||
let t_copy = t0.elapsed();
|
||||
|
||||
// Cursor-as-metadata: blend the overlay into this slot's OWNED input surface (a tiny kernel
|
||||
// over the cursor's rect — never the compositor's dmabuf). The PTX module loads lazily on the
|
||||
// first cursor frame now that the CUDA context is current; any failure degrades to no cursor,
|
||||
// never a dropped frame.
|
||||
// Cursor-as-metadata: blend the overlay into this slot's OWNED input surface via the
|
||||
// SPIR-V compute pass (a dispatch over the cursor's rect — never the compositor's
|
||||
// dmabuf). Cursor-bearing frames forced `ordered = false` above, so the CUDA copy has
|
||||
// completed before the Vulkan dispatch and the fence-waited dispatch completes before
|
||||
// the encode below — the cross-API ordering is CPU-established. Any failure degrades to
|
||||
// no cursor, never a dropped frame.
|
||||
if let Some(ov) = &captured.cursor {
|
||||
if !self.cursor_tried {
|
||||
self.cursor_tried = true;
|
||||
match cuda::CursorBlend::new(CURSOR_PTX) {
|
||||
Ok(cb) => {
|
||||
// Success is as diagnosis-critical as failure: a silent no-cursor session
|
||||
// must be attributable to "overlay never arrived", not "module never said".
|
||||
tracing::info!("NVENC (Linux): cursor blend module loaded");
|
||||
self.cursor = Some(cb);
|
||||
}
|
||||
Err(e) => tracing::warn!(
|
||||
error = %format!("{e:#}"),
|
||||
"NVENC (Linux): cursor blend module load failed — cursor not composited"
|
||||
),
|
||||
}
|
||||
}
|
||||
if let Some(cb) = &self.cursor {
|
||||
if let (Some(vk), SlotSurface::Vk(vref)) =
|
||||
(self.vk_blend.as_mut(), &self.ring[slot].surface)
|
||||
{
|
||||
if self.cursor_serial != ov.serial {
|
||||
match cb.upload(ov.rgba.as_slice(), ov.w, ov.h) {
|
||||
Ok(()) => {
|
||||
self.cursor_serial = ov.serial;
|
||||
self.cursor_upload_warned = false;
|
||||
}
|
||||
Err(e) => {
|
||||
if !self.cursor_upload_warned {
|
||||
self.cursor_upload_warned = true;
|
||||
tracing::warn!(
|
||||
error = %format!("{e:#}"),
|
||||
serial = ov.serial,
|
||||
"NVENC (Linux): cursor upload failed — cursor not composited"
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
vk.upload_cursor(ov.rgba.as_slice(), ov.w, ov.h);
|
||||
self.cursor_serial = ov.serial;
|
||||
}
|
||||
let s = &self.ring[slot].surface;
|
||||
// surfW = content width; surfH = the surface's allocated height (matches
|
||||
// `copy_into_slot`'s plane math). Cursor pixels past the content are in cropped
|
||||
// padding rows — harmless.
|
||||
let (w, h) = (self.width, s.height);
|
||||
let r = match self.buffer_fmt {
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => {
|
||||
cb.blend_yuv444(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered)
|
||||
}
|
||||
nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => {
|
||||
cb.blend_nv12(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered)
|
||||
}
|
||||
_ => cb.blend_argb(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered),
|
||||
};
|
||||
// surfW = content width; the blend derives plane strides from the slot's luma
|
||||
// height. Cursor pixels past the content land in cropped padding rows — harmless.
|
||||
let r = vk.blend_ref(
|
||||
vref,
|
||||
slot_fmt_of(self.buffer_fmt),
|
||||
self.width,
|
||||
ov.w,
|
||||
ov.h,
|
||||
ov.x,
|
||||
ov.y,
|
||||
);
|
||||
if let Err(e) = r {
|
||||
if !self.cursor_blend_warned {
|
||||
self.cursor_blend_warned = true;
|
||||
tracing::warn!(
|
||||
error = %format!("{e:#}"),
|
||||
"NVENC (Linux): cursor blend launch failed — cursor not composited"
|
||||
"NVENC (Linux): cursor blend dispatch failed — cursor not composited"
|
||||
);
|
||||
}
|
||||
} else {
|
||||
self.cursor_blend_warned = false;
|
||||
// TEMP KWin composite probe (DROP BEFORE MERGE): prove the blend launches and
|
||||
// with what geometry — the on-glass symptom is a loaded module and a silent,
|
||||
// invisible cursor.
|
||||
// TEMP KWin composite probe (DROP BEFORE MERGE): prove the blend dispatches
|
||||
// and with what geometry.
|
||||
{
|
||||
use std::sync::atomic::{AtomicU64, Ordering as ProbeOrd};
|
||||
static PROBE_BLEND: AtomicU64 = AtomicU64::new(0);
|
||||
@@ -1485,14 +1567,18 @@ impl Encoder for NvencCudaEncoder {
|
||||
ov_y = ov.y,
|
||||
ov_w = ov.w,
|
||||
ov_h = ov.h,
|
||||
surf_w = w,
|
||||
surf_h = h,
|
||||
visible = ov.visible,
|
||||
"cursor blend probe: kernel launched"
|
||||
"cursor blend probe: vulkan dispatch submitted"
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
} else if !self.cursor_blend_warned {
|
||||
self.cursor_blend_warned = true;
|
||||
tracing::warn!(
|
||||
"NVENC (Linux): cursor overlay present but no Vulkan blend (bring-up failed \
|
||||
earlier) — cursor not composited"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1540,7 +1626,7 @@ impl Encoder for NvencCudaEncoder {
|
||||
version: nv::NV_ENC_PIC_PARAMS_VER,
|
||||
inputWidth: self.width,
|
||||
inputHeight: self.height,
|
||||
inputPitch: self.ring[slot].surface.pitch as u32,
|
||||
inputPitch: self.ring[slot].surface.pitch() as u32,
|
||||
inputBuffer: mp.mappedResource,
|
||||
bufferFmt: mp.mappedBufferFmt,
|
||||
outputBitstream: self.bitstreams[slot],
|
||||
|
||||
Reference in New Issue
Block a user