한 남자로서, 이 PP (Prompt Processing) 결과가 그리 좋지 않다는 점을 인정하기가 부끄럽습니다. 저는 5개의 GPU (2x 3080, 2x 3090, 5090)를 통해 IQ3_XXS를 전체 VRAM (Video RAM)에 올려 실행 중인데, 속도가 좋지 않습니다. 제목에 언급했듯이, 약 600 t/s PP입니다. 저에게 이것은 일상적인 모델로 사용하기에 적합하지 않은 수준입니다. 심지어 UD Q4 K XL을 부분 오프로드 (partial offload)로 실행할 때도 200-400 PP t/s를 얻을 수 있습니다. 저는 모델이 VRAM에 완전히 올라가 있음을 확인했습니다. --fit 및 일반적인 레이어 할당 (layer assignment)을 수행했습니다. 이는 최신 CUDA, NCCL, 그리고 Llamacpp 빌드를 사용한 결과입니다. Llamacpp의 프로파일링 (profiling)에는 익숙하지 않지만, 현재 조사 중입니다. 제가 조사하는 동안, 다른 분들은 본인의 PP 수치를 공유해 주실 수 있나요?

수정: 이 포크 (fork)를 사용하면 PP t/s가 100-150 정도 더 나옵니다.
https://github.com/vektorprime/working_ds4_speed

수정 2: 이제 이 포크는 메인 버전보다 약 200 pp t/s 더 빠르게 작동합니다.

수정 3: 이제 포크에서 1000 PP T/S에 도달했습니다!
초당 1003.28 토큰 (tokens per second) <-

아래의 프로파일링에는 모델을 VRAM으로 복사하고 그래프를 캡처하는 과정이 포함되어 있음을 참고해 주세요. 하지만 전반적으로 여전히 매우 많은 연산 (operations)이 발생하고 있습니다. 본질적으로, 여기에는 우리가 시간을 할애하고 있는 많은 새로운 요소들이 있습니다. 반면 Qwen3.5 122B와 같은 모델에서는 mul_mat_q에서 훨씬 더 많은 시간을 소비합니다.

또한 Qwen3.5 122B를 프로파일링하고 LLM을 사용하여 두 모델을 비교한 뒤, 그 결과를 여기에 정리했습니다:

연산 (Operation) DSV4 Flash % wall Qwen3.5 122B % wall 비고 (Notes)

mul_mat_q (모든 변형) 16.6 47.7 MoE + projections (투영)
flash_attn_ext_f16 13.4 11 DSV4 512-dim, Qwen 256-dim
gated_delta_net / ssm_conv — 7.9 Qwen GDN/SSM fused (융합)
cub::DeviceTopK (3개 변형) 8.9 — DSV4 indexer top-k
lightning_indexer_kernel_wmma 3 — DSV4 indexer scan
dsv4_hc_post 0.7 — Hyper-connection post
dsv4_hc_pre 0.2 — Hyper-connection pre
dsv4_hc_comb <0.1 — Sinkhorn comb (fused (융합))
fwht_cuda <0.1 0.1 Hadamard rotations (Hadamard 회전)
concat/softmax/rope/reduce/set/get_rows 1.7 0.6 Compressor state vs minimal
rms_norm (모든 변형) 0.9 1.1
quantize_mmq_q8_1 0.6 1
broadcast ops (add/mul/clamp/silu/sigmoid) 1.2 2.5
mul_mat_q_stream_k_fixup <0.1 0.6
k_argsort — 0.2
topk_moe_cuda <0.1 <0.1 MoE gate routing (MoE 게이트 라우팅)
기타 GPU 커널 (Other GPU kernels) 0.3 0.4 cutlass/cublasLt/nvjet
GPU 커널 소계 (GPU kernels subtotal) ~47.6 ~73.1

cudaLaunchKernel 30 1.5 Warmup/capture vs replay
cudaStreamSynchronize 11.8 7.9 Per-segment syncs (세그먼트별 동기화)
cudaEventSynchronize — 12.6 Pipeline sync (파이프라인 동기화)
cudaMemsetAsync 7.9 — KQ mask zero-fill
cudaMemcpy H2D 0.4 1.7
cudaMemcpy D2H 0.1 0.2
cudaMemcpyAsync / Peer / 2D 0.2 0.2
cudaEventRecord <0.1 0.2
cudaGraphLaunch <0.1 0.2
기타 API (other API) 0.8 2 cudaFree, cuKernelGetName 등
API + MemOps 소계 (API + MemOps subtotal) ~52.4 ~26.9

전체 wall 시간 대비 Matmul 비중 (Matmul share of total wall) 0.166 0.477 2.9배 차이
DSV4 고유 커널 비중 (DSV4-unique kernel share) 0.128 0 topk + indexer + HC
API 오버헤드 비중 (API overhead share) 0.524 0.269 1.9배 차이

Deepseek v4 flash의 NSYS 프로파일입니다. 그래프 캡처와 VRAM으로의 호스트 메모리 복사를 피하기 위해, prefill(프리필)이 시작된 후 5초가 지난 시점부터 필터링했습니다.

** NVTX Range Summary (nvtx_sum): Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Style Range -------- --------------- --------- -------- -------- -------- ---------- ----------- ------- -------------------------- 99.9 23,517,017,181 349,261 67,333.6 28,776.0 24,075 35,067,107 629,495.6 PushPop :cub::DeviceTopK::MaxPairs 0.1 15,104,814 1,004 15,044.6 14,122.5 11,576 32,460 2,906.0 PushPop :cub::DeviceReduce::Sum Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with /opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/osrt_sum.py... ** OS Runtime Summary (osrt_sum): Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name -------- --------------- --------- ---------------- ---------------- ------------- --------------- ---------------- ---------------------- 59.8 938,741,260,257 23 40,814,837,402.5 3,852,725,158.0 4,571,408 142,829,769,032 57,945,550,617.2 pthread_cond_wait 15.6 245,398,284,449 5,558 44,152,264.2 10,092,237.5 1,050 100,287,684 43,792,633.3 poll 13.6 213,028,000,123 363 586,853,994.8 500,070,864.0 500,048,924 4,000,086,753 544,982,083.7 pthread_cond_timedwait 7.5 117,355,440,079 37 3,171,768,650.8 1,000,068,679.0 509,768,866 60,000,274,987 10,358,166,691.2 pthread_cond_clockwait 3.5 54,626,648,415 2 27,313,324,207.5 27,313,324,207.5 5,852,098,841 48,774,549,574 30,350,755,978.4 accept4 0.0 157,947,698 269 587,166.2 38,553.0 8,784 21,271,489 2,496,230.0 ioctl 0.0 81,587,427 3 27,195,809.0 30,273,276.0 352,517 50,961,634 25,444,523.6 mmap 0.0 7,950,471 10 795,047.1 63,148.0 41,219 3,931,046 1,300,088.1 pthread_join 0.0 1,739,916 52 33,459.9 16,988.5 1,034 394,126 64,848.5 pthread_mutex_lock 0.0 778,438 20 38,921.9 39,685.0 10,518 83,093 14,240.4 send 0.0 226,340 11 20,576.4 8,889.0 5,664 81,985 27,757.8 munmap 0.0 66,859 3 22,286.3 13,314.0 11,459 42,086 17,172.1 shutdown 0.0 48,842 28 1,744.4 1,704.5 1,251 2,390 388.3 fputs

Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with /opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/cuda_api_sum.py... ** CUDA API Summary (cuda_api_sum): Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name -------- --------------- --------- ------------ ------------ -------- ----------- ------------ ------------------------- 59.2 15,868,026,900 1,438,692 11,029.5 3,515.0 2,857 33,592,937 240,705.8 cudaLaunchKernel 23.2 6,210,846,084 2,487 2,497,324.5 3,061.0 896 116,417,479 10,397,557.2 cudaStreamSynchronize 15.5 4,154,936,825 349,979 11,872.0 3,716.0 2,209 25,512,068 256,118.3 cudaMemsetAsync 0.8 207,425,433 1,438,692 144.2 140.0 105 24,531 55.4 cuKernelGetName 0.4 112,437,810 3 37,479,270.0 43,601,674.0 919,979 67,916,157 33,915,112.3 cudaFreeHost 0.3 85,680,813 343 249,798.3 157,930.0 7,627 1,159,978 265,836.5 cudaMemcpyPeerAsync 0.3 80,402,307 60 1,340,038.4 93,277.5 8,039 21,640,675 4,053,874.0 cudaFree 0.1 31,981,065 7,706 4,150.2 3,720.0 2,913 34,738 1,451.4 cudaLaunchKernelExC 0.0 11,758,896 1,714 6,860.5 5,005.0 2,425 41,309 5,386.1 cudaMemcpyAsync 0.0 8,199,214 5 1,639,842.8 1,897,337.0 737,767 2,482,765 687,145.1 cuMemUnmap 0.0 8,168,763 1,360 6,006.4 5,021.5 3,834 20,476 2,557.9 cudaMemcpy2DAsync 0.0 4,786,186 1,048 4,567.0 3,999.5 2,970 22,192 1,837.5 cuLaunchKernel 0.0 2,830,984 1,535 1,844.3 1,245.0 689 11,146 1,439.0 cudaEventRecord 0.0 2,623,086 131 20,023.6 17,178.0 7,606 55,372 10,223.8 cudaGraphLaunch 0.0 1,036,131 6,434 161.0 146.0 107 989 63.2 cuStreamGetCaptureInfo_v2 0.0 1,016,774 343 2,964.4 2,520.0 1,444 18,398 1,396.6 cudaStreamWaitEvent

0.0 790,325 1,048 754.1 505.5 206 3,619 577.4 cuKernelGetFunction 0.0 755,796 3,576 211.4 184.0 133 1,676 99.3 cuStreamGetGreenCtx 0.0 635,879 144 4,415.8 4,438.5 3,066 6,252 631.0 cuLaunchKernelEx 0.0 291,936 16 18,246.0 17,624.5 12,381 28,538 3,836.8 cudaGraphExecDestroy 0.0 182,120 144 1,264.7 949.0 545 17,819 1,496.4 cuKernelSetAttribute 0.0 134,865 5 26,973.0 14,345.0 13,946 77,028 27,986.9 cudaMemGetInfo 0.0 94,174 20 4,708.7 3,346.0 2,133 22,402 4,872.7 cudaDeviceSynchronize 0.0 81,982 95 863.0 699.0 568 6,618 670.6 cudaEventDestroy 0.0 72,234 16 4,514.6 4,293.5 2,350 7,471 1,505.6 cudaGraphDestroy 0.0 61,017 5 12,203.4 11,918.0 9,749 16,656 2,736.4 cuMemAddressFree 0.0 47,404 5 9,480.8 7,797.0 6,357 17,767 4,677.2 cudaStreamDestroy 0.0 2,662 5 532.4 397.0 328 886 250.2 cudaGetDeviceProperties Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with /opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/cuda_gpu_kern_sum.py...

** CUDA GPU 커널 요약 (cuda_gpu_kern_sum): 시간(Time) (%) 총 시간 (Total Time) (ns) 인스턴스 (Instances) 평균 (Avg) (ns) 중앙값 (Med) (ns) 최소 (Min) (ns) 최대 (Max) (ns) 표준편차 (StdDev) (ns) 이름 (Name)** -------- --------------- --------- ------------ ------------ --------- ---------- ------------ ---------------------------------------------------------------------------------------------------- 28.0 7,090,472,754 349 20,316,540.8 15,223,112.0 3,575,898 54,619,022 12,959,412.1 void flash_attn_ext_f16<(int)512, (int)512, (int)8, (int)8, (bool)0, (bool)0>(const char *, const c… 14.6 3,707,932,015 608 6,098,572.4 6,951,152.5 2,671,811 9,289,001 2,164,262.5 void mul_mat_q<(ggml_type)18, (int)128, (bool)0>(const char *, const int *, const int *, const int … 11.9 3,017,599,294 700,054 4,310.5 4,544.0 640 7,264 970.5 void cub::_V_300304_SM_860_1200::detail::topk::DeviceTopKKernel<cub::_V_300304_SM_860_1200::detail:… 11.0 2,792,243,865 406 6,877,447.9 7,901,809.0 3,294,396 9,689,904 2,261,982.8 void mul_mat_q<(ggml_type)17, (int)128, (bool)0>(const char *, const int *, const int *, const int … 6.3 1,610,604,112 170 9,474,141.8 9,690,097.0 2,107,177 17,705,517 4,630,249.6 void lightning_indexer_kernel_wmma<(int)8, (int)32, (long)128, (long)64, (ggml_type)1>(const float … 5.6 1,430,234,513 2,600 550,090.2 195,330.0 18,815 1,742,132 591,166.2 void mul_mat_q<(ggml_type)8, (int)128, (bool)0>(const char *, const int *, const int *, const int *… 5.1 1,287,357,624 350,027 3,677.9 3,968.0 2,336 5,152 744.4 void cub::_V_300304_SM_860_1200::detail::topk::DeviceTopKKernel<cub::_V_300304_SM_860_1200::detail:… 1.7 440,834,512 1,717 256,746.9 88,607.0 1,151 782,665 295,774.0 void concat_non_cont<unsigned int, (int)0>(const char *, const char *, char *, long, long, long, lo… 1.7 435,604,190 1,304 334,052.3 414,437.0 86,559 556,391 147,693.9 void mul_mat_q<(ggml_type)14, (int)128, (bool)0>(const char *, const int *, const int *, const int … 1.6 416,290,262 350,027 1,189.3 1,280.0 735 3,456 297.0 void

cub::_V_300304_SM_860_1200::detail::topk::DeviceTopKLastFilterKernel<cub::_V_300304_SM_860_120… 1.4 362,571,747 665 545,220.7 620,552.0 301,692 755,721 162,632.6 dsv4_hc_post_f32(const float *, const float *, const float *, const float *, float *, long, long, l… 1.3 326,013,784 4,253 76,655.0 52,897.0 3,135 488,741 99,589.4 void quantize_mmq_q8_1<(mmq_q8_1_ds_layout)0, (bool)0>(const float *, const int *, void *, long, lo… 1.0 254,610,290 664 383,449.2 443,892.5 161,087 544,391 152,737.7 void rms_norm_f32<(int)1024, (bool)0, (bool)0>(const float *, float *, int, long, long, long, float… 0.8 208,508,768 348 599,163.1 664,982.5 327,581 807,337 184,455.3 void rms_norm_f32<(int)256, (bool)0, (bool)0>(const float *, float *, int, long, long, long, float,… 0.7 170,623,787 1,386 123,105.2 92,175.0 767 296,515 106,498.3 void op_clamp_kernel<float>(const T1 *, T1 *, T1, T1, int) 0.7 167,385,333 2,246 74,526.0 2,368.0 1,344 588,839 167,913.8 void k_bin_bcast<&op_mul, float, float, float, const float *>(const T2 *, const T3 *, T4 *, unsigne… 0.5 132,972,685 673 197,582.0 190,622.0 17,984 440,933 156,392.7 void unary_gated_op_kernel<&op_silu, float>(const T2 *, const T2 *, T2 *, long, long, long, long) 0.5 132,446,060 17 7,790,944.7 7,847,099.0 6,718,456 8,890,373 673,322.6 void mul_mat_q<(ggml_type)39, (int)128, (bool)0>(const char *, const int *, const int *, const int … 0.5 128,014,036 697 183,664.3 200,642.0 101,375 245,699 52,616.8 dsv4_hc_pre_f32(const float *, const float *, float *, long, long, long, long, long, long, long, lo… 0.5 123,128,168 349 352,802.8 365,092.0 251,774 471,014 71,825.0 void cpy_scalar<&cpy_1_scalar<float, float>>(const char *, char *, long, long, long, long, long, lo… 0.5 123,107,791 697 176,625.2 209,250.0 56,735 265,603 78,864.8 void cutlass::Kernel2<cutlass_80_tensorop_s1688gemm_64x128_32x3_tn_align4>(T1::Params) 0.4 98,654,288 16 6,165,893.0 6,156,784.5 5,953,070 6,544,534 159,662.6 void mul_mat_q<(ggml_type)21, (int)128, (bool)0>(const char

*, const int *, const int *, const int … 0.3 86,973,709 349 249,208.3 273,507.0 117,694 335,812 83,764.6 void k_bin_bcast<&op_add, float, float, float, const float *, const float *, const float *, const f… 0.3 86,654,923 1,047 82,765.0 72,449.0 53,600 120,127 19,901.7 void mm_ids_helper<(int)6>(const int *, int *, int *, int *, int, int, int, int, int, bool) 0.3 84,496,163 1,368 61,766.2 5,264.0 1,088 213,410 81,951.4 void rope_norm<(bool)1, (bool)0, float, float>(const T3 *, T4 *, int, int, int, int, int, int, int,… 0.3 67,801,522 698 97,136.9 111,553.0 23,872 136,866 39,861.3 void quantize_mmq_q8_1<(mmq_q8_1_ds_layout)0, (bool)1>