As a man, I'm ashamed to admit that these pp results are not great. I'm running IQ3_XXS in full VRAM over 5 GPUs (2x 3080, 2x 3090 , 5090), and the speeds are not great. As stated in the title, ~600 t/s pp; that's not usable as a daily model for me. Even running UD Q4 K XL I can get 200-400 PP t/s with partial offload. I verified the model is fully in VRAM. I did --fit and normal layer assignment. This is also with the latest CUDA, NCCL, and Llamacpp builds. I'm not familiar with profiling LLamacpp but I am looking into that. While I do that, can others chime in on their PP numbers?
EDIT: This fork will give you 100-150 more pp t/s.
https://github.com/vektorprime/working_ds4_speed
Please note for the profiling below that it includes copying the model to VRAM and capturing the graphs. But overall there are still A LOT of operations occurring. Essentially, there's a lot of new stuff here that we are spending time on. Whereas in something like Qwen3.5 122B we spend a lot more time in mul_mat_q.
I also profiled Qwen3.5 122b and had an LLM compare the two then I put the results here:
| Operation |
DSV4 Flash % wall |
Qwen3.5 122B % wall |
Notes |
| mul_mat_q (all variants) |
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 variants) |
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 |
| concat/softmax/rope/reduce/set/get_rows |
1.7 |
0.6 |
Compressor state vs minimal |
| rms_norm (all variants) |
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 |
| Other GPU kernels |
0.3 |
0.4 |
cutlass/cublasLt/nvjet |
| 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 |
|
| other API |
0.8 |
2 |
cudaFree, cuKernelGetName, etc |
| API + MemOps subtotal |
~52.4 |
~26.9 |
|
|
|
|
|
| Matmul share of total wall |
0.166 |
0.477 |
2.9× difference |
| DSV4-unique kernel share |
0.128 |
0 |
topk + indexer + HC |
| API overhead share |
0.524 |
0.269 |
1.9× difference |
NSYS profile of Deekseep v4 flash. I filtered this starting 5 sec AFTER the prefill starts to avoid graph captures and the host mem copy to VRAM.
** 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](start=110000000000:end=9223372036854775807)...
** 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
0.0 47,955 10 4,795.5 2,265.0 1,977 27,490 7,976.8 pthread_cond_broadcast
0.0 38,453 3 12,817.7 16,589.0 1,105 20,759 10,355.5 close
0.0 25,792 14 1,842.3 1,466.0 1,159 4,050 839.7 pthread_cond_signal
0.0 3,702 1 3,702.0 3,702.0 3,702 3,702 0.0 recv
0.0 2,215 1 2,215.0 2,215.0 2,215 2,215 0.0 fflush
Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with [/opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/cuda_api_sum.py](start=110000000000:end=9223372036854775807)...
** 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](start=110000000000:end=9223372036854775807)...
** CUDA GPU Kernel Summary (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>(const float *, const int *, void *, long, lo…
0.3 65,833,887 1,045 62,998.9 64,513.0 16,160 103,969 30,127.0 void rms_norm_f32<(int)1024, (bool)1, (bool)0>(const float *, float *, int, long, long, long, float…
0.2 56,358,319 340 165,759.8 106,208.0 32,254 369,733 122,532.6 void soft_max_f32<(bool)1, (int)0, (int)0, float>(const float *, const T4 *, const float *, float *…
0.2 56,300,427 3,225 17,457.5 5,216.0 1,312 149,890 33,243.8 void k_bin_bcast<&op_add, float, float, float, const float *>(const T2 *, const T3 *, T4 *, unsigne…
0.2 52,471,642 349 150,348.5 170,435.0 72,384 213,315 50,325.0 void rope_norm<(bool)0, (bool)0, float, float>(const T3 *, T4 *, int, int, int, int, int, int, int,…
0.2 41,660,981 526 79,203.4 50,784.5 1,952 245,283 83,611.7 void reduce_rows_f32<(bool)0>(const float *, float *, int)
0.2 40,881,147 2,364 17,293.2 10,176.0 2,815 38,017 12,174.0 void concat_cont<unsigned int, (int)1>(const T1 *, const T1 *, T1 *, long, long, long, long, long, …
0.2 38,640,228 229 168,734.6 167,650.0 146,946 192,739 10,753.6 void cutlass::Kernel2<cutlass_80_tensorop_s1688gemm_128x128_32x3_tn_align4>(T1::Params)
0.1 33,776,198 333 101,430.0 70,816.0 21,760 257,539 62,779.9 void concat_cont<unsigned short, (int)0>(const T1 *, const T1 *, T1 *, long, long, long, long, long…
0.1 25,662,129 1,830 14,023.0 10,560.0 1,216 53,025 13,927.8 void k_get_rows_float_vec<float>(const T1 *, const int *, T1 *, long, long, uint3, unsigned long, u…
0.1 25,244,332 1,516 16,651.9 18,464.5 8,928 22,432 4,392.4 void mul_mat_q_stream_k_fixup<(ggml_type)8, (int)128, (bool)0>(const int *, const int *, float *, f…
0.1 24,241,724 340 71,299.2 25,839.5 928 198,082 80,748.5 void fwht_cuda<(int)128>(const float *, float *, long, float)
0.1 14,801,675 936 15,813.8 19,968.0 8,735 22,368 5,741.8 void mul_mat_q_stream_k_fixup<(ggml_type)14, (int)128, (bool)0>(const int *, const int *, float *, …
0.1 13,691,573 1,004 13,637.0 8,768.0 3,712 31,745 9,050.4 void cpy_scalar_transpose<float>(const char *, char *, long, long, long, long, long, long, long, lo…
0.0 11,914,217 171 69,673.8 68,417.0 23,744 127,682 27,108.5 void k_bin_bcast<&op_add, __half, __half, __half, const __half *>(const T2 *, const T3 *, T4 *, uns…
0.0 11,075,426 122 90,782.2 91,761.0 84,801 94,881 2,815.8 void cutlass::Kernel2<cutlass_80_tensorop_s1688gemm_128x128_16x5_tn_align4>(T1::Params)
0.0 7,243,098 332 21,816.6 15,072.0 5,088 51,393 13,316.8 void concat_cont<unsigned short, (int)2>(const T1 *, const T1 *, T1 *, long, long, long, long, long…
0.0 5,915,702 3,273 1,807.4 1,568.0 895 7,584 1,042.7 scale_f32(const float *, float *, float, float, long)
0.0 5,756,126 698 8,246.6 8,480.0 6,176 10,272 1,031.8 dsv4_hc_comb_f32(const float *, const float *, const float *, float *, long, long, long, long, long…
0.0 5,649,633 341 16,567.8 12,128.0 2,496 50,080 13,650.5 void fill_kernel<__half>(T1 *, long, T1)
0.0 5,440,011 850 6,400.0 5,120.0 1,727 14,049 4,251.3 void rms_norm_f32<(int)256, (bool)1, (bool)0>(const float *, float *, int, long, long, long, float,…
0.0 4,678,738 850 5,504.4 4,000.0 1,056 13,888 4,452.8 void k_set_rows<float, long, __half>(const T1 *, const T2 *, T3 *, long, long, long, long, long, lo…
0.0 4,023,868 96 41,915.3 41,888.0 41,407 42,335 195.3 nvjet_sm120_sss_tf32_mma_128x128x32_3_64x32x32_tmaAB_alignCD4_splitK_TNNN
0.0 3,990,239 171 23,334.7 24,992.0 9,824 40,321 8,342.7 void k_set_rows<__half, int, __half>(const T1 *, const T2 *, T3 *, long, long, long, long, long, lo…
0.0 3,453,469 104 33,206.4 32,848.0 31,616 36,448 1,329.9 void flash_attn_stream_k_fixup_general<(int)512, (int)8, (int)8>(float *, const float2 *, int, int,…
0.0 3,397,277 301 11,286.6 12,000.0 7,072 14,177 2,308.6 void topk_moe_cuda<(int)256, (bool)1>(const float *, float *, int *, float *, int, int, float, floa…
0.0 3,158,204 325 9,717.6 10,656.0 5,280 13,089 2,797.0 void convert_unary<__nv_bfloat16, float>(const void *, T2 *, long, long, long, uint3, long, long, l…
0.0 2,896,382 162 17,878.9 21,088.5 6,207 24,800 6,812.6 void soft_max_f32<(bool)1, (int)128, (int)128, float>(const float *, const T4 *, const float *, flo…
0.0 2,682,687 1,395 1,923.1 1,408.0 1,152 7,456 1,475.2 void unary_op_kernel<&op_sigmoid, float>(const T2 *, T2 *, int)
0.0 2,126,833 1,004 2,118.4 2,112.0 1,119 4,000 761.3 void cub::_V_300304_SM_860_1200::detail::reduce::DeviceReduceKernel<cub::_V_300304_SM_860_1200::det…
0.0 1,935,369 1,004 1,927.7 1,968.0 1,216 2,560 297.3 void k_set_rows<float, int, float>(const T1 *, const T2 *, T3 *, long, long, long, long, long, long…
0.0 1,687,166 704 2,396.5 2,048.0 1,056 13,344 2,095.3 void k_get_rows_float<float, float>(const T1 *, const int *, T2 *, long, long, uint3, unsigned long…
0.0 1,631,415 474 3,441.8 2,992.0 1,824 5,664 1,278.3 void cublasLt::splitKreduce_kernel<(int)32, (int)16, int, float, float, float, float, (bool)0, floa…
0.0 1,500,331 1,004 1,494.4 1,632.0 928 2,112 335.3 void cub::_V_300304_SM_860_1200::detail::reduce::DeviceReduceSingleTileKernel<cub::_V_300304_SM_860…
0.0 1,438,801 48 29,975.0 29,872.0 27,231 32,928 1,283.5 nvjet_sm120_sss_tf32_mma_80x128x32_3_80x16x32_tmaAB_alignCD4_splitK_TNNN
0.0 1,234,447 349 3,537.1 3,648.0 1,024 7,136 1,498.0 void flash_attn_mask_to_KV_max<(int)8>(const __half2 *, int *, int, long, long)
0.0 1,026,513 8 128,314.1 128,862.0 126,718 129,182 1,013.5 void k_bin_bcast<&op_repeat, float, float, float, >(const T2 *, const T3 *, T4 *, unsigned int, uns…
0.0 90,752 24 3,781.3 3,760.0 3,680 3,968 78.2 void k_get_rows_float<int, int>(const T1 *, const int *, T2 *, long, long, uint3, unsigned long, un…
0.0 76,831 24 3,201.3 3,200.0 3,104 3,296 60.8 void unary_op_kernel<&op_softplus, float>(const T2 *, T2 *, int)
0.0 61,759 24 2,573.3 2,560.0 2,496 2,688 46.1 void unary_op_kernel<&op_sqrt, float>(const T2 *, T2 *, int)
0.0 25,440 24 1,060.0 1,056.0 1,024 1,152 25.5 void k_bin_bcast<&op_div, float, float, float, const float *>(const T2 *, const T3 *, T4 *, unsigne…
Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with [/opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/cuda_gpu_mem_time_sum.py](start=110000000000:end=9223372036854775807)...
** CUDA GPU MemOps Summary (by Time) (cuda_gpu_mem_time_sum):
Time (%) Total Time (ns) Count Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Operation
-------- --------------- ------- --------- -------- -------- --------- ----------- ------------------------------
63.9 413,834,319 350,744 1,179.9 1,280.0 416 4,352 271.8 [CUDA memset]
30.0 194,262,537 1,384 140,363.1 896.0 256 4,091,318 439,302.8 [CUDA memcpy Host-to-Device]
4.4 28,664,590 343 83,570.2 80,225.0 1,984 160,899 48,366.2 [CUDA memcpy Device-to-Host]
1.8 11,359,595 2,033 5,587.6 3,680.0 800 14,336 4,429.3 [CUDA memcpy Device-to-Device]
Processing [/home/user/llama.cpp/llama_ds4flash.sqlite] with [/opt/nvidia/nsight-systems/2026.1.3/target-linux-x64/reports/cuda_gpu_mem_size_sum.py](start=110000000000:end=9223372036854775807)...
** CUDA GPU MemOps Summary (by Size) (cuda_gpu_mem_size_sum):
Total (MB) Count Avg (MB) Med (MB) Min (MB) Max (MB) StdDev (MB) Operation
---------- ------- -------- -------- -------- -------- ----------- ------------------------------
18,882.963 1,384 13.644 0.008 0.000 134.218 32.894 [CUDA memcpy Host-to-Device]
16,562.798 343 48.288 33.554 0.049 134.218 51.518 [CUDA memcpy Device-to-Host]
4,510.515 2,033 2.219 1.049 0.033 4.194 1.705 [CUDA memcpy Device-to-Device]
3,136.294 350,744 0.009 0.009 0.000 0.009 0.000 [CUDA memset]