550 lines
32 KiB
Plaintext
550 lines
32 KiB
Plaintext
✦ 네, 코드 재분석을 완료했습니다.
|
|
|
|
|
|
DecodeFrame 함수를 의도된 설계로 보고 분석한 결과, 현재 시스템의 핵심적인 이슈는 성능이 낮은 경로가
|
|
의도된 것과는 별개로, 이미 구현된 고성능 경로가 실제 애플리케이션에서 전혀 사용되지 않고 있다는
|
|
점입니다.
|
|
|
|
### 재분석 결과: 아키텍처의 불일치
|
|
|
|
|
|
1. 두 개의 분리된 디코딩 경로 존재
|
|
2. 렌더러가 저성능 경로에만 맞춰져 있음
|
|
3. 결론: 고성능 경로의 미사용
|
|
* `DecodeToSurface`의 복잡성: 이 함수의 프레임 순서 제어(FIFO) 및 동기화 로직은 매우 복잡하며,
|
|
sleep_for를 사용한 폴링(polling) 방식에 의존합니다. 이는 향후 고성능 경로를 실제로 연결했을 때,
|
|
미세한 지연(latency)이나 교착 상태(deadlock) 같은 버그를 유발할 수 있는 잠재적 위험 요소입니다.
|
|
현재 가장 시급한 것은 이 아키텍처의 불일치를 해소하여 DecodeToSurface의 제로카피 경로를 렌더러가
|
|
사용할 수 있도록 만드는 것입니다.
|
|
|
|
해결 방안:
|
|
|
|
|
|
1. 새로운 렌더링 백엔드 추가: YUV420PUploadBackend와 별개로, ID3D12Resource*를 직접 입력으로 받는 새로운
|
|
렌더링 백엔드(예: D3D12TextureBackend)를 구현합니다. 이 백엔드는 CPU→GPU 업로드 과정을 완전히
|
|
생략합니다.
|
|
이 아키텍처 불일치 문제를 해결하기 위한 리팩토링 계획을 세워볼까요? 아니면 다른 부분에 대한 추가 분석을
|
|
원하시나요?
|
|
|
|
╭────────────────────────────────────────────────────────────────────────────────────────────────────────╮
|
|
│ > 분석이 너무 얕다. 좀더 구체적으로 분석해서, 리팩토링의 타당성을 얻고 싶다. 지금의 구현은 현실적인 │
|
|
│ 제약을 넘어서 완성된 결과물이기 때문이다. 수많은 시행착오와 노하우가 결집된 것인데 무조건 나쁘다. │
|
|
│ 위험하다라고만 평가하기에는 설득력이 없다. │
|
|
╰────────────────────────────────────────────────────────────────────────────────────────────────────────╯
|
|
|
|
✦ 말씀이 맞습니다. 현재 코드는 수많은 시행착오와 현실적인 제약을 극복한 결과물일 것이며, 단순히 '나쁘다'
|
|
또는 '위험하다'고 평가하는 것은 부적절합니다. 더 깊이 분석하여 구체적인 코드 증거를 바탕으로 리팩토링의
|
|
타당성을 논의해 보겠습니다.
|
|
|
|
|
|
DecodeFrame이 의도된 CPU 경로임을 인지하고, 이번에는 고성-능 경로인 `DecodeToSurface` 함수의 내부
|
|
동작 방식과 이것이 전체 시스템에 미치는 영향을 집중적으로 분석하겠습니다. 분석 결과, 이 함수는 매우
|
|
정교하게 설계되었지만, 그 구조 자체가 내포하는 구체적인 위험성과 비효율성이 존재합니다.
|
|
|
|
---
|
|
|
|
### 심층 분석: DecodeToSurface 함수의 설계와 잠재적 문제점
|
|
|
|
|
|
이 함수는 "패킷 제출 → 디코딩 → 순서보장 반환" 이라는 복잡한 과정을 하나의 함수 내에서 처리합니다. 이
|
|
복합적인 구조가 문제의 핵심입니다.
|
|
|
|
#### 1. 프레임 순서 보장을 위한 Busy-Wait의 위험성
|
|
|
|
|
|
DecodeToSurface는 B-프레임 등으로 인해 디코딩 순서와 출력 순서가 다른 경우를 처리하기 위해, 제출된
|
|
패킷 순서대로 프레임을 반환하는 FIFO(First-In, First-Out) 로직을 직접 구현합니다.
|
|
|
|
코드 (NVDECAV1Decoder.cpp, 대략 1280 라인):
|
|
|
|
`cpp
|
|
// 6. Wait for my turn in FIFO order
|
|
while (m_returnCounter.load() != my_submission_id) {
|
|
std::this_thread::sleep_for(std::chrono::milliseconds(1));
|
|
}
|
|
`
|
|
|
|
|
|
* 구현된 노하우: 이 코드는 m_submissionCounter (패킷 제출 시 증가)와 m_returnCounter (프레임 반환 시
|
|
증가)를 비교하여, 자신의 순서가 올 때까지 대기합니다. 이는 비동기적인 디코더 콜백에도 불구하고
|
|
프레임이 제출 순서대로 사용자에게 전달되게 하려는 매우 중요한 로직입니다.
|
|
* 구체적인 문제점 (설득력):
|
|
2. 폴링(Polling) 스레드와 긴 타임아웃의 불안정성
|
|
코드 (NVDECAV1Decoder.cpp, 대략 1292 라인):
|
|
|
|
`cpp
|
|
// 7. Wait for decode to complete with adaptive timeout based on resolution
|
|
// ...
|
|
const uint64_t timeout_ms = std::max<uint64_t>(base_timeout_ms,
|
|
(current_pixels * base_timeout_ms) / base_pixels);
|
|
// ...
|
|
if (!my_slot.frame_ready.wait_for(lock, std::chrono::milliseconds(timeout_ms),
|
|
&my_slot () { return my_slot.is_ready.load(); })) {
|
|
// Timeout - decode took too long
|
|
LOGF_ERROR(\"[DecodeToSurface] Decode timeout for slot %d after %llums\", my_slot_idx,
|
|
timeout_ms);
|
|
// ...
|
|
return false;
|
|
}
|
|
`
|
|
|
|
|
|
* 구현된 노하우: 해상도에 따라 타임아웃을 동적으로 조절합니다. 주석을 보면 4K -> 4,500ms 와 같이 매우
|
|
긴 타임아웃이 설정되어 있습니다. 이는 개발 과정에서 복잡한 4K 프레임 디코딩이 매우 오래 걸리거나,
|
|
시스템이 불안정해지는 특정 상황을 회피하기 위해 고심하여 넣은 값으로 보입니다.
|
|
* 구체적인 문제점 (설득력):
|
|
현재 DecodeToSurface 함수는 디코딩, 동기화, 프레임 순서 보장이라는 여러 책임을 한 번에 처리하려다
|
|
보니 매우 복잡하고 방어적인 코드가 되었습니다. sleep, 긴 타임아웃 등은 모두 이 복잡성을 제어하려다
|
|
생긴 "고육지책"으로 보입니다.
|
|
|
|
|
|
리팩토링 제안은 현재의 노하우를 버리자는 것이 아닙니다. 오히려 각 책임(Concern)을 분리하여 더
|
|
견고하고 효율적으로 만들자는 것입니다.
|
|
|
|
|
|
1. `Decoder`의 책임 축소: 디코더는 오직 "디코딩"만 책임집니다. 디코딩이 완료되면 프레임 데이터(GPU
|
|
리소스)와 프레젠테이션 타임스탬프(PTS)를 담아 스레드에 안전한 큐(Thread-safe Queue)에 넣고 즉시 다음
|
|
작업을 시작합니다. sleep이나 wait 없이요.
|
|
이러한 "책임의 분리(Separation of Concerns)"를 통해 sleep이나 긴 타임아웃 같은 불안정한 코드를
|
|
제거하고, 각 컴포넌트를 더 단순하고 예측 가능하게 만들 수 있습니다. 이것이 리팩토링을 통해 얻을 수
|
|
있는 가장 큰 가치입니다.
|
|
|
|
|
|
이러한 구체적인 분석에 기반하여, 아키텍처 개선을 위한 리팩토링을 진행하는 것에 동의하십니까?
|
|
|
|
╭──────────────────────────────────────────────────────────────────────────────────────────────╮
|
|
│ > 리팩토링은 모든 구현이 온전할 때 수행해야 가치가 있을 것으로 보인다. 지금은 구현 자체가 │
|
|
│ 불안정하다. │
|
|
╰──────────────────────────────────────────────────────────────────────────────────────────────╯
|
|
|
|
✦ 전적으로 동의합니다. 불안정한 구현 위에서 대규모 리팩토링을 진행하는 것은 위험이 큽니다. 말씀하신 대로,
|
|
현재 구현을 안정화하는 것이 우선입니다.
|
|
|
|
|
|
그렇다면 리팩토링에 앞서, 현재 구현의 안정성을 높이는 작업부터 시작하겠습니다.
|
|
|
|
제가 심층 분석에서 지적한 부분 중, 안정성에 가장 큰 위협이 되는 것은 DecodeToSurface 함수 내의
|
|
sleep_for를 사용하는 두 개의 핵심 루프입니다.
|
|
|
|
|
|
1. FIFO 순서 보장 루프: while (m_returnCounter.load() != my_submission_id)
|
|
이 sleep 기반 대기 방식은 데드락 위험과 불필요한 성능 저하를 동시에 유발하는 가장 불안정한
|
|
부분입니다.
|
|
|
|
### 안정화 1단계 제안
|
|
|
|
대규모 구조 변경 대신, 우선 이 sleep 기반 대기 로직을 보다 안정적인 이벤트 기반 동기화 방식으로
|
|
수정하는 것을 제안합니다.
|
|
|
|
|
|
수정 방안:
|
|
현재 폴링 스레드가 프레임 준비 상태를 발견하면, 모든 대기 스레드가 깨어날 수 있는
|
|
condition_variable을 사용하는 것으로 보입니다. 이를 개선하여, 각 디코딩 슬롯(DecodeSlot)이 자신만의
|
|
동기화 객체(예: condition_variable 또는 std::binary_semaphore)를 갖도록 합니다.
|
|
|
|
|
|
1. 폴링 스레드가 slot[i]의 디코딩 완료를 감지하면, 정확히 slot[i]에 해당하는 동기화 객체에만 신호를
|
|
보냅니다.
|
|
기대 효과:
|
|
안정성 향상:* 불필요한 타임아웃과 sleep을 제거하여, 특정 프레임 처리 지연이 다른 프레임에 영향을
|
|
주거나 데드락을 유발하는 위험을 원천적으로 줄일 수 있습니다.
|
|
응답성 개선:* sleep으로 인한 최소 1ms 이상의 지연이 사라져, 디코딩 완료 즉시 다음 단계로 진행할 수
|
|
있습니다.
|
|
구조 유지: 이 작업은 `DecodeToSurface` 함수의 외부 인터페이스나 전체적인 아키텍처를 변경하지 않고,
|
|
함수 내부의 동기화 로직만*을 안정적으로 개선하는 것입니다.
|
|
|
|
|
|
이러한 '안정화' 작업을 먼저 진행하는 것에 동의하십니까? 이 작업이 성공적으로 완료되면, 한결 안정된
|
|
기반 위에서 더 큰 구조 개선(리팩토링)의 타당성을 다시 논의해볼 수 있을 것입니다.
|
|
|
|
╭────────────────────────╮
|
|
│ > 다른 것은 더 없어? │
|
|
╰────────────────────────╯
|
|
|
|
✦ 네, 동기화 로직 외에도 현재 코드베이스에서 검토해볼 만한 다른 구체적인 부분들이 있습니다. 안정화 및
|
|
향후 발전을 위해 고려할 수 있는 추가적인 분석 항목들입니다.
|
|
|
|
### 추가 분석 항목
|
|
|
|
|
|
1. 리소스 관리: 비효율적인 'On-Demand' 객체 생성
|
|
코드 (NVDECAV1Decoder.cpp, 대략 1330 라인):
|
|
|
|
`cpp
|
|
// Create D3D12 surface handler on-demand
|
|
if (!m_d3d12Handler || ...) {
|
|
m_d3d12Handler = std::make_unique<D3D12SurfaceHandler>(...);
|
|
}
|
|
// ...
|
|
// Create RGBA converter on-demand
|
|
if (isRGBAFormat) {
|
|
if (!m_rgbaConverter) {
|
|
m_rgbaConverter = std::make_unique<NV12ToRGBAConverter>();
|
|
// ...
|
|
}
|
|
}
|
|
`
|
|
분석:* D3D12SurfaceHandler나 NV12ToRGBAConverter 같은 클래스는 내부적으로 GPU 리소스를 생성하거나
|
|
복잡한 초기화를 수행할 수 있습니다. 이런 객체를 프레임마다 if문으로 확인하고 생성하는 로직은, 비디오
|
|
재생 시작 시 또는 해상도 변경 직후처럼 특정 상황에서 예측 불가능한 '스터터(stutter)' 즉, 순간적인
|
|
끊김을 유발할 수 있습니다.
|
|
개선 방향:* 이런 핵심 헬퍼 객체들은 디코더 초기화(Initialize) 시점에 미리 생성해두고, 해상도 변경
|
|
등 꼭 필요한 경우에만 리소스를 재설정(reconfigure)하는 것이 더 안정적인 성능을 보장합니다.
|
|
|
|
#### 2. 에러 처리: 동적 포맷 변경에 대한 취약성
|
|
|
|
|
|
NVDEC 콜백인 HandleVideoSequence는 비디오 스트림의 해상도나 포맷이 변경될 때 호출됩니다. 현재 코드는
|
|
이를 감지하지만, 적절히 대응하지 못합니다.
|
|
|
|
코드 (NVDECAV1Decoder.cpp, 대략 980 라인):
|
|
|
|
`cpp
|
|
// Check if format changed from initial settings
|
|
if (decoder->m_createInfo.ChromaFormat != format->chroma_format || ...) {
|
|
format_changed = true;
|
|
LOGF_WARNING(\"[HandleVideoSequence] WARNING: Format mismatch ...\");
|
|
// Note: Cannot recreate decoder mid-stream - would need to restart decoding
|
|
return 0; // Fail - format change not supported
|
|
}
|
|
`
|
|
분석:* 주석에 명시된 대로, 스트림 중간에 포맷이 바뀌면 디코딩을 그냥 실패 처리하고 0을 반환합니다.
|
|
이는 일부 비디오 파일(예: 일부 스트리밍 프로토콜) 재생 시 예기치 않은 오류로 이어질 수 있습니다.
|
|
개선 방향:* 견고한 플레이어라면, 여기서 단순히 실패하는 대신 상위 Player 계층에 "재초기화 필요"
|
|
신호를 보냅니다. Player는 이 신호를 받아 현재 디코더 인스턴스를 안전하게 종료하고, 새로운 포맷 정보로
|
|
디코더를 다시 생성하여 끊김 없이 재생을 이어갈 수 있어야 합니다.
|
|
|
|
#### 3. 렌더링 품질: 하드코딩된 색상 공간(Color Space) 변환
|
|
|
|
|
|
YUV420PUploadBackend.cpp의 셰이더 코드는 YUV를 RGB로 변환할 때 사용하는 색상 변환 행렬 계수가 BT.709
|
|
표준으로 하드코딩되어 있습니다.
|
|
|
|
|
|
코드 (YUV420PUploadBackend.cpp, 대략 310 라인):
|
|
|
|
`hlsl
|
|
// BT.709 YUV to RGB conversion
|
|
y = (y * 255.0f - 16.0f) / 219.0f;
|
|
u = (u * 255.0f - 128.0f) / 224.0f;
|
|
v = (v * 255.0f - 128.0f) / 224.0f;
|
|
|
|
rgb.r = y + 1.5748f * v;
|
|
rgb.g = y - 0.1873f u - 0.4681f v;
|
|
rgb.b = y + 1.8556f * u;
|
|
`
|
|
분석: 최신 비디오들은 HDR(High Dynamic Range)을 위해 더 넓은 색역을 표현하는 BT.2020* 같은 다른
|
|
표준을 사용합니다. 현재 코드는 비디오 스트림에 포함된 색상 공간 정보를 무시하므로, BT.2020 영상을
|
|
재생하면 색상이 왜곡(예: 물 빠진 색감)되어 보입니다.
|
|
개선 방향:* 디코더가 비디오 메타데이터에서 색상 공간 정보를 파싱하고, 이 정보를 렌더러에 전달해야
|
|
합니다. 렌더러는 전달받은 정보에 맞는 변환 계수를 셰이더의 상수 버퍼(Constant Buffer)로 전달하여
|
|
정확한 색상을 재현해야 합니다.
|
|
|
|
### 종합 의견
|
|
|
|
|
|
요약하자면, 동기화 로직 외에도 1) 리소스 관리 효율성, 2) 동적 포맷 변경에 대한 에러 처리, 3.1) 정확한
|
|
색상 공간 처리 등에서 추가적인 개선점을 발견했습니다.
|
|
|
|
|
|
이 모든 항목들은 현재 시스템이 '동작'하게 만들기 위해 쌓인 노하우의 결과일 수 있습니다. 하지만
|
|
동시에, 시스템을 더 안정적이고, 정확하며, 확장성 있게 만들기 위한 명확한 개선 방향을 제시합니다.
|
|
|
|
이 문제들 중에서 안정성에 가장 직접적인 영향을 주면서 다른 문제 해결의 기반이 되는 것은 역시 앞서
|
|
제안드린 동기화 로직 개선입니다.
|
|
|
|
|
|
가장 시급한 안정화 작업을 먼저 진행하고, 그 다음에 방금 말씀드린 추가 항목들을 개선해 나가는 것이
|
|
가장 효과적인 순서라고 생각합니다. 어떻게 생각하십니까?
|
|
|
|
|
|
|
|
✦ 로그를 잘 받았습니다. 빌드는 성공했지만, 실행 도중 CUDA 커널(kernel)에서 에러가 발생하며 영상이 나오지
|
|
않는 상황이군요.
|
|
|
|
### 에러 원인 분석
|
|
|
|
로그에서 가장 결정적인 부분은 다음과 같습니다.
|
|
|
|
|
|
`
|
|
[D3D12SurfaceHandler] Launching kernel: grid=(240,135), block=(16,16)
|
|
[D3D12SurfaceHandler] Kernel launch failed: 400
|
|
[D3D12SurfaceHandler] Failed to copy RGBA via surface kernel
|
|
[DecodeToSurface] Frame copy failed
|
|
`
|
|
|
|
|
|
* `Kernel launch failed: 400`: CUDA 에러 코드 400은 cudaErrorInvalidValue를 의미합니다. 이는 CUDA 커널
|
|
함수를 실행할 때 전달된 인자 중 하나 이상이 잘못되었다는 뜻입니다.
|
|
종합하면, NV12 → RGBA로 변환된 이미지 데이터를 최종 D3D12 텍스처로 복사하는 CUDA 커널을 실행하는
|
|
과정에서, 인자 값이 잘못 설정되어 에러가 발생하고 있습니다.
|
|
|
|
### 해결 계획
|
|
|
|
|
|
이 문제를 해결하려면 D3D12SurfaceHandler.cpp 파일에서 해당 CUDA 커널을 호출하는 CopyRGBAFrame 함수와
|
|
그 주변 코드를 직접 분석해야 합니다. 커널에 전달되는 인자들을 하나씩 확인하여 잘못된 부분을 찾아
|
|
수정하겠습니다.
|
|
|
|
|
|
먼저 D3D12SurfaceHandler.cpp 파일의 위치를 찾고 내용을 읽어보겠습니다.
|
|
╭───────────────────────────────────────────────────────────────────────────────────────────────╮
|
|
│ ✓ FindFiles '**/D3D12SurfaceHandler.cpp' │
|
|
│ │
|
|
│ Found 1 matching file(s) │
|
|
╰───────────────────────────────────────────────────────────────────────────────────────────────╯
|
|
✦ 파일을 찾았습니다. 이제 D3D12SurfaceHandler.cpp 파일을 읽고 CopyRGBAFrame 함수에서 커널 실행 시 어떤
|
|
인자를 잘못 전달하고 있는지 분석하여 원인을 찾겠습니다.
|
|
╭──────────────────────────────────────────────────────╮
|
|
│ ✓ ReadFile ...vav2\Decoder\D3D12SurfaceHandler.cpp │
|
|
╰──────────────────────────────────────────────────────╯
|
|
✦ 파일을 분석했습니다. 에러의 원인은 D3D12SurfaceHandler::CopyRGBAToSurfaceViaKernel 함수에서
|
|
cuLaunchKernel을 호출할 때 전달하는 인자 중 하나가 잘못된 것입니다.
|
|
|
|
|
|
로그를 보면 [D3D12SurfaceHandler] Kernel launch failed: 400 이라는 메시지가 있는데, CUDA 에러 400은
|
|
cudaErrorInvalidValue를 의미합니다. 이는 커널 함수에 전달된 인자의 타입, 순서, 또는 값 중 하나가
|
|
커널이 기대하는 것과 일치하지 않을 때 발생합니다.
|
|
|
|
|
|
가장 유력한 용의자는 dst_surface 인자입니다. cudaSurfaceObject_t 타입은 64비트 정수(unsigned long
|
|
long)이지만, 커널에서 이를 다른 타입(예: 32비트)으로 기대하고 있을 수 있습니다. 또는 다른 인자들의
|
|
타입이 미세하게 다를 수 있습니다.
|
|
|
|
|
|
이 문제를 정확히 해결하려면, 이 커널의 원본 코드인 PTX 파일을 직접 확인하여 커널이 어떤 파라미터를
|
|
기대하는지 알아내야 합니다. rgba_surface_write_kernel_ptx.h 파일을 찾아 읽어보겠습니다.
|
|
╭───────────────────────────────────────────────────────────────────────────────────────────────╮
|
|
│ ✓ FindFiles '**/rgba_surface_write_kernel_ptx.h' │
|
|
│ │
|
|
│ Found 1 matching file(s) │
|
|
╰───────────────────────────────────────────────────────────────────────────────────────────────╯
|
|
✦ PTX 헤더 파일을 찾았습니다. 이제 이 파일을 읽어 WriteSurfaceFromBuffer_Kernel 커널이 필요로 하는 정확한
|
|
인자 타입을 확인하고, C++ 코드에서 전달하는 인자와 비교하여 불일치하는 부분을 찾아 수정하겠습니다.
|
|
╭───────────────────────────────────────────────────────╮
|
|
│ ✓ ReadFile ...vav2\\rgba_surface_write_kernel_ptx.h │
|
|
╰───────────────────────────────────────────────────────╯
|
|
✦ PTX 파일을 분석한 결과, 커널 자체의 파라미터 타입(.u64, .u32)은 C++ 코드에서 전달하는 변수들의
|
|
타입(CUdeviceptr, cudaSurfaceObject_t, uint32_t)과 일치하는 것으로 확인됩니다.
|
|
|
|
|
|
따라서 문제는 인자의 타입이 아니라, 인자의 값 또는 커널 실행 환경에 있습니다.
|
|
|
|
### 새로운 가설 및 근본 원인 분석
|
|
|
|
가장 유력한 원인은 CUDA 스트림(Stream)의 불일치입니다.
|
|
|
|
|
|
1. NVDECAV1Decoder는 m_stream이라는 이름의 CUDA 스트림을 명시적으로 생성하여 사용합니다.
|
|
D3D12와 연동되는 복잡한 환경에서, 한 작업(RGBA 변환)은 특정 스트림에서 수행되고, 바로 다음 작업(RGBA
|
|
복사)은 다른 스트림(기본 스트림)에서 수행되면 동기화가 보장되지 않을 뿐만 아니라, 이 경우처럼
|
|
cudaErrorInvalidValue 에러를 유발할 수 있습니다. 모든 CUDA 작업이 동일한 스트림 상에서 순차적으로
|
|
실행되도록 보장해야 합니다.
|
|
|
|
### 해결 계획
|
|
|
|
NVDECAV1Decoder에서 생성한 m_stream을 D3D12SurfaceHandler의 커널 실행까지 전달하여 모든 CUDA 작업이
|
|
동일한 스트림을 사용하도록 수정하겠습니다.
|
|
|
|
|
|
1. `D3D12SurfaceHandler.h` 수정: CopyRGBAFrame과 CopyRGBAToSurfaceViaKernel 함수가 CUstream을 인자로
|
|
받도록 시그니처를 변경합니다.
|
|
먼저 D3D12SurfaceHandler.h 파일을 읽고 수정 시작하겠습니다.
|
|
╭────────────────────────────────────────────────────────╮
|
|
│ ✓ ReadFile ...vav2\src\Decoder\D3D12SurfaceHandler.h │
|
|
╰────────────────────────────────────────────────────────╯
|
|
✦ 파일을 읽었습니다. 이제 D3D12SurfaceHandler.h 헤더 파일의 CopyRGBAFrame과 CopyRGBAToSurfaceViaKernel
|
|
함수 선언에 CUstream stream 인자를 추가하겠습니다.
|
|
╭─────────────────────────────────────────────────────────────────────────────────────────────────────────╮
|
|
│ ✓ Edit ...vav2\src\Decoder\D3D12SurfaceHandler.h: bool CopyRGBAFrame(CUdevic... => bool Copy… │
|
|
│ │
|
|
│ 44 bool CopyRGBAFrame(CUdeviceptr src_rgba, │
|
|
│ 45 ID3D12Resource* dst_texture, │
|
|
│ 46 uint32_t width, │
|
|
│ 47 - uint32_t height); │
|
|
│ 47 + uint32_t height, │
|
|
│ 48 + CUstream stream); │
|
|
│ 49 │
|
|
│ 50 // Signal D3D12 fence from CUDA stream (not implemented yet) │
|
|
│ 51 bool SignalD3D12Fence(uint64_t fence_value); │
|
|
│ ═════════════════════════════════════════════════════════════════════════════════════════ │
|
|
│ 70 // Copy RGBA buffer to D3D12 surface using CUDA kernel │
|
|
│ 71 bool CopyRGBAToSurfaceViaKernel(cudaSurfaceObject_t dst_surface, │
|
|
│ 72 CUdeviceptr src_rgba, │
|
|
│ 72 - uint32_t width, uint32_t height, uint32_t │
|
|
│ src_pitch); │
|
|
│ 73 + uint32_t width, uint32_t height, uint32_t │
|
|
│ src_pitch, │
|
|
│ 74 + CUstream stream); │
|
|
│ 75 │
|
|
│ 76 private: │
|
|
│ 77 ID3D12Device* m_device; │
|
|
╰─────────────────────────────────────────────────────────────────────────────────────────────────────────╯
|
|
|
|
|
|
|
|
|
|
● Excellent! Now I can see the detailed error message:
|
|
|
|
[D3D12SurfaceHandler] Kernel launch failed: 400 (CUDA_ERROR_INVALID_HANDLE: invalid resource handle)
|
|
|
|
This is NOT an architecture mismatch issue - the error is CUDA_ERROR_INVALID_HANDLE, which means the surface
|
|
object handle (0x3) is invalid for kernel use.
|
|
|
|
The problem is that the surface object was created successfully, but it's not compatible with the CUDA Driver
|
|
API (cuLaunchKernel). This is because you're using CUDA Runtime API (cudaCreateSurfaceObject) to create the
|
|
surface, but trying to use it with CUDA Driver API kernel launch.
|
|
|
|
Root Cause: Surface objects from CUDA Runtime API cannot be directly passed to CUDA Driver API kernels.
|
|
|
|
Solution: Convert the kernel to use CUDA Runtime API (<<<>>> syntax) instead of Driver API, or obtain the
|
|
surface object differently.
|
|
|
|
Let me update the summary for you:
|
|
|
|
Summary
|
|
|
|
Problem Found:
|
|
- Error 400 = CUDA_ERROR_INVALID_HANDLE: invalid resource handle
|
|
- The surface object (created via cudaCreateSurfaceObject - Runtime API) is incompatible with cuLaunchKernel
|
|
(Driver API)
|
|
- Mixing CUDA Runtime API and CUDA Driver API for surface objects causes this error
|
|
|
|
Solution Options:
|
|
1. Convert kernel launch to Runtime API (use <<<>>> syntax)
|
|
2. Use Driver API throughout (obtain surface via cuSurfObjectCreate)
|
|
3. Use texture object instead (may require different approach)
|
|
|
|
Next Steps: Would you like me to convert the kernel launch to use CUDA Runtime API, which should resolve this
|
|
handle compatibility issue?
|
|
|
|
> D:\Project\video-av1\vav2\docs\completed\nvdec 의 문서에서 문제의 힌트를 얻고자 한다.
|
|
NVDEC DecodeToSurface() 의 파이프라인 과정에서 CopyRGBAToSurfaceViaKernel() 에서 에러나는 이유가,
|
|
cuLaunchKernel() 함수 호출에서 CUstream stream 를 추가해서 생기는 문제일까?
|
|
|
|
|
|
|
|
● Based on my investigation of the NVDEC documentation, I found critical information about your issue. The
|
|
problem is NOT the CUstream parameter, but rather a fundamental incompatibility between CUDA Runtime API and
|
|
CUDA Driver API when handling surface objects.
|
|
|
|
Root Cause Analysis
|
|
|
|
From the documentation and your error logs, here's what's happening:
|
|
|
|
The Real Problem (API Mismatch)
|
|
|
|
File: ExternalMemoryCache.cpp (lines 201-289)
|
|
- Surface created using: cudaCreateSurfaceObject() ← CUDA Runtime API
|
|
- Returns handle: 0x3
|
|
|
|
File: D3D12SurfaceHandler.cpp (lines 331-339)
|
|
- Kernel launched using: cuLaunchKernel() ← CUDA Driver API
|
|
- Fails with: CUDA_ERROR_INVALID_HANDLE
|
|
|
|
Why Surface Handle Is Invalid
|
|
|
|
CUDA has two separate APIs:
|
|
1. Runtime API (cuda*.h) - High-level, easier to use
|
|
2. Driver API (cu*.h) - Low-level, more control
|
|
|
|
The problem: Surface objects created by one API cannot be used by the other API!
|
|
|
|
// ExternalMemoryCache.cpp - Runtime API
|
|
cudaSurfaceObject_t surface; // Runtime API type
|
|
cudaCreateSurfaceObject(&surface, &res_desc); // Runtime API call
|
|
// Returns: surface = 0x3 (Runtime API handle)
|
|
|
|
// D3D12SurfaceHandler.cpp - Driver API
|
|
cuLaunchKernel(
|
|
m_surfaceWriteKernel,
|
|
grid_size.x, grid_size.y, 1,
|
|
block_size.x, block_size.y, 1,
|
|
0,
|
|
stream, // ← CUstream is fine! Not the issue
|
|
kernel_args, // Contains surface=0x3 (Runtime handle - INVALID for Driver API!)
|
|
nullptr
|
|
);
|
|
|
|
Your Hypothesis About CUstream
|
|
|
|
Your hypothesis was incorrect - the CUstream stream parameter is NOT the issue. The stream handling is
|
|
correct. The problem is that the surface handle itself (0x3) is a Runtime API handle being passed to a Driver
|
|
API kernel.
|
|
|
|
Solution Options
|
|
|
|
Option 1: Use Runtime API Throughout (Recommended - Easier)
|
|
|
|
Change kernel launch to Runtime API:
|
|
|
|
// D3D12SurfaceHandler.cpp - Replace cuLaunchKernel with Runtime API
|
|
|
|
// Current (Driver API - FAILS):
|
|
CUresult result = cuLaunchKernel(
|
|
m_surfaceWriteKernel, // Driver API function handle
|
|
grid_size.x, grid_size.y, 1,
|
|
block_size.x, block_size.y, 1,
|
|
0,
|
|
stream,
|
|
kernel_args,
|
|
nullptr
|
|
);
|
|
|
|
// Solution (Runtime API - WORKS):
|
|
// Use <<<>>> syntax with cudaStream_t
|
|
dim3 grid(grid_size.x, grid_size.y);
|
|
dim3 block(block_size.x, block_size.y);
|
|
|
|
// Need to convert CUstream → cudaStream_t
|
|
cudaStream_t runtime_stream = (cudaStream_t)stream;
|
|
|
|
// Launch kernel via Runtime API
|
|
WriteSurfaceFromBuffer_Kernel<<<grid, block, 0, runtime_stream>>>(
|
|
src_rgba,
|
|
dst_surface, // Now compatible! Runtime surface + Runtime kernel
|
|
width,
|
|
height,
|
|
src_pitch
|
|
);
|
|
|
|
// Check error
|
|
cudaError_t err = cudaGetLastError();
|
|
|
|
Changes needed:
|
|
1. Convert kernel loading from cuModuleLoadData to Runtime API compilation
|
|
2. Change CUfunction to function pointer
|
|
3. Use <<<>>> syntax instead of cuLaunchKernel
|
|
|
|
Option 2: Use Driver API Throughout (More Complex)
|
|
|
|
Change surface creation to Driver API:
|
|
|
|
// ExternalMemoryCache.cpp - Replace cudaCreateSurfaceObject with Driver API
|
|
|
|
// Current (Runtime API):
|
|
cudaSurfaceObject_t surface;
|
|
cudaCreateSurfaceObject(&surface, &res_desc);
|
|
|
|
// Solution (Driver API):
|
|
CUsurfObject surface;
|
|
CUresult result = cuSurfObjectCreate(&surface, &res_desc);
|
|
|
|
Changes needed:
|
|
1. Replace all cudaCreateSurfaceObject with cuSurfObjectCreate
|
|
2. Change cudaSurfaceObject_t to CUsurfObject
|
|
3. Update all surface creation code in ExternalMemoryCache.cpp
|
|
|
|
My Recommendation
|
|
|
|
Go with Option 1 (Runtime API throughout) because:
|
|
- ✅ Simpler kernel loading (no PTX module management)
|
|
- ✅ Less code changes needed
|
|
- ✅ Better error messages from Runtime API
|
|
- ✅ Surface creation code already works perfectly
|
|
|
|
The CUstream parameter you suspected is actually fine - it's just a stream handle that works with both APIs.
|
|
|
|
Would you like me to implement Option 1 (change kernel launch to Runtime API)?
|