From eda50b4eb51b205d6574e30ea4c7dadda3ed9057 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 10 Apr 2026 11:31:35 +0200 Subject: [PATCH] fix(critical): pad backtest evaluator states buffer to state_dim_padded MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Root cause of cublasLtMatmul CUBLAS_STATUS_NOT_SUPPORTED: - states_buf allocated [batch * state_dim] (unpadded) - cuBLAS forward GEMM reads with stride state_dim_padded (pad128) - Buffer overflow: GEMM reads past buffer end cublasLtMatmul validates buffer sizes against layout descriptors and returns NOT_SUPPORTED for undersized buffers. cublasSgemm silently read garbage — this was the source of the "parallel test congestion errors" seen previously. Fix: - Allocate states_buf with state_dim_padded stride (3 allocation sites) - gather_states kernel: add padded_sd parameter, write with padded stride - DtoD copy: use padded row bytes - Both gather_states call sites updated with padded_sd arg Co-Authored-By: Claude Opus 4.6 (1M context) --- .../cuda_pipeline/backtest_gather_kernel.cu | 14 +- .../cuda_pipeline/gpu_backtest_evaluator.rs | 20 +- testing/cublaslt_debug | Bin 0 -> 34248 bytes testing/cublaslt_debug.cu | 246 ++++++++++++++++++ 4 files changed, 268 insertions(+), 12 deletions(-) create mode 100755 testing/cublaslt_debug create mode 100644 testing/cublaslt_debug.cu diff --git a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu index fdc8989e7..c200bee88 100644 --- a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu @@ -34,7 +34,7 @@ extern "C" __global__ void gather_states( const __nv_bfloat16* __restrict__ features, // [n_windows * max_len * feat_dim] (bf16 from data pipeline) const float* __restrict__ portfolio, // [n_windows * 8] (f32 from env kernel) - float* states_out, // [n_windows * state_dim] (f32 for cublasSgemm) + float* states_out, // [n_windows * padded_sd] (f32 for cublasLtMatmul, padded stride) int n_windows, int max_len, int feat_dim, @@ -42,7 +42,8 @@ extern "C" __global__ void gather_states( int current_step, float initial_capital, // host scalar — convert on first use float spread_cost, // host scalar — convert on first use - int ofi_dim // 0 = no OFI, 8 = standard OFI features + int ofi_dim, // 0 = no OFI, 8 = standard OFI features + int padded_sd // pad128(state_dim) — output row stride for cuBLAS K-tile alignment ) { // Shared memory tile for coalesced portfolio reads — 8 F32 values per thread __shared__ float shmem_portfolio[256 * 8]; @@ -63,7 +64,7 @@ extern "C" __global__ void gather_states( if (w >= n_windows) return; int feat_base = (w * max_len + current_step) * feat_dim; - int out_base = w * state_dim; + int out_base = w * padded_sd; // Read portfolio state from shared memory tile — all f32 (output is f32 for cublasSgemm) float f_value = shmem_portfolio[local_tid * 8 + 0]; @@ -222,15 +223,16 @@ extern "C" __global__ void gather_states( } } - // ── 5. Zero-pad remaining [filled .. state_dim) ── // + // ── 5. Zero-pad remaining [filled .. padded_sd) ── // + // Pads both the state_dim gap AND the cuBLAS K-tile alignment columns. int filled = market_dim + 8 + 16 + ofi_dim; - for (i = filled; i + 3 < state_dim; i += 4) { + for (i = filled; i + 3 < padded_sd; i += 4) { states_out[out_base + i] = 0.0f; states_out[out_base + i + 1] = 0.0f; states_out[out_base + i + 2] = 0.0f; states_out[out_base + i + 3] = 0.0f; } - for (; i < state_dim; i++) { + for (; i < padded_sd; i++) { states_out[out_base + i] = 0.0f; } diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 677651b46..6ad61a9a3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -551,8 +551,9 @@ impl GpuBacktestEvaluator { // 8-aligned for H100 tensor core HMMA dispatch. const PORTFOLIO_AND_MTF_DIM: usize = 24; // 8 portfolio + 16 multi-timeframe let state_dim = (feature_dim + PORTFOLIO_AND_MTF_DIM + 7) & !7; + let state_dim_padded = (state_dim + 127) & !127; // pad128 for cuBLAS K-tile alignment let states_buf = stream - .alloc_zeros::(n_windows * state_dim) + .alloc_zeros::(n_windows * state_dim_padded) .map_err(|e| MLError::ModelError(format!("states_buf alloc: {e}")))?; // Pre-allocate forward kernel output buffer (used by evaluate_dqn path) @@ -713,7 +714,8 @@ impl GpuBacktestEvaluator { // Safety: argument order matches `gather_states` signature exactly: // features (bf16), portfolio (f32), states_out (f32), n_windows, max_len, feat_dim, - // state_dim, current_step, initial_capital, spread_cost, ofi_dim + // state_dim, current_step, initial_capital, spread_cost, ofi_dim, padded_sd + let padded_sd_i32 = ((state_dim + 127) & !127) as i32; unsafe { self.stream .launch_builder(&self.gather_kernel) @@ -728,12 +730,14 @@ impl GpuBacktestEvaluator { .arg(&initial_capital) .arg(&spread_cost) .arg(&ofi_dim_i32) + .arg(&padded_sd_i32) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!("gather_states launch step {step}: {e}")))?; } - // Return a copy of the states buffer as CudaSlice. - let n_elems = self.n_windows * state_dim; + // Return a copy of the states buffer as CudaSlice (padded stride). + let state_dim_padded = (state_dim + 127) & !127; + let n_elems = self.n_windows * state_dim_padded; let dst = self.stream.alloc_zeros::(n_elems) .map_err(|e| MLError::ModelError(format!("alloc states step {step}: {e}")))?; let src_view = self.states_buf.slice(..n_elems); @@ -1042,7 +1046,8 @@ impl GpuBacktestEvaluator { let b3 = dqn_cfg.branch_3_size as i32; let v_min_f = dqn_cfg.v_min; let v_max_f = dqn_cfg.v_max; - let state_row_bytes = n * state_dim * std::mem::size_of::(); + let state_dim_padded = (state_dim + 127) & !127; + let state_row_bytes = n * state_dim_padded * std::mem::size_of::(); let total_chunks = (self.max_len + DQN_BACKTEST_CHUNK_SIZE - 1) / DQN_BACKTEST_CHUNK_SIZE; for chunk_start in (0..self.max_len).step_by(DQN_BACKTEST_CHUNK_SIZE) { @@ -1511,7 +1516,8 @@ impl GpuBacktestEvaluator { let ch_q_values = self.stream.alloc_zeros::(cn * total_actions) .map_err(|e| MLError::ModelError(format!("alloc chunked q_values: {e}")))?; - let ch_states_buf = self.stream.alloc_zeros::(cn * self.state_dim) + let state_dim_padded = (self.state_dim + 127) & !127; + let ch_states_buf = self.stream.alloc_zeros::(cn * state_dim_padded) .map_err(|e| MLError::ModelError(format!("alloc chunked states_buf: {e}")))?; let ch_actions_buf = self.stream.alloc_zeros::(cn) .map_err(|e| MLError::ModelError(format!("alloc chunked actions_buf: {e}")))?; @@ -1711,6 +1717,7 @@ impl GpuBacktestEvaluator { let initial_capital = self.config.initial_capital; let spread_cost = self.config.spread_cost; let ofi_dim_i32 = self.config.ofi_dim as i32; + let padded_sd_i32 = ((state_dim + 127) & !127) as i32; unsafe { self.stream @@ -1726,6 +1733,7 @@ impl GpuBacktestEvaluator { .arg(&initial_capital) .arg(&spread_cost) .arg(&ofi_dim_i32) + .arg(&padded_sd_i32) .launch(launch_cfg) .map_err(|e| { MLError::ModelError(format!("gather_states launch step {step}: {e}")) diff --git a/testing/cublaslt_debug b/testing/cublaslt_debug new file mode 100755 index 0000000000000000000000000000000000000000..7128bac0abb718d148072edf94351462818e32f9 GIT binary patch literal 34248 zcmeHwe{>wxb^okv*^2pL+1QE0FAv578Sq+?Wmz^DR@M(jv1Bin1%_Z&EA7hKAnhu< zvyRLOM4)WKcD;?8eofqOOKVb6oZk+C*5r#p28!f1aRedY^qchf(1d7-A&LW3P=flo z@5jvOX~)jtq<{3BeL9-^?&sZi?|paPyqS4BGw*YaEloupk0#SB?Rt&4;pH6hi2~zp zkpb~(wc2^uF4Hd4=768gFkbHC1f@KkD4I@71U?It>Mfy69(pl{HA$t1M5&%zDqp}$ zB-KhePxXo^tJ4NaXIqQbUBr1wo%NVKpO9B*>G>R&)LD;W({np7SLq4S?sC}};U|TQ z>J5r|D%C|iC6)0cHS}LA^s2N1Wt1dU>-C6wDsAA^G)c(@rM6$0p$}uvFX|0!;(AoN zhr^nr(%wy|M{)UMlL)7_RoGMGFZunV-I7v2p|q+u(pgizsyDQ-cFQnfLY-DL$a$%c$`*-olE3Jwn;*EP?-JFj}*ZC|%GY9^s+W8vd{D=yh{@w(F)dH9GD zn+!t1LYuX6*Sg^E!HInG&$`H80*$(*)y_QWqGt&jnos}DsG1L7+>%1yIkb=y7)QjqUSLe`6d^6 zpNrk2F8J#%`qM7*KXZ|P-v$4=i=T5{gsXfYeY2{Rtstx3K<9uAsei){^>eaT*bJlqX?(XQ}Wird0wooU7+ zoyi$GV5sX1#dWgCPBb65Jh2nt8n=!+8r76C`&g<`PR9!husP zWIE}<$+k@64qWaVDw}ZN=ulRgbl~$W6d4g8(tVBtFLB`KI`C2ljt*?4B@SHNvavGV zf#U$NQn>?{_YzcC;lMAjq;cQn!09?)CZ7XeEJ0Aifq&G2`yKel9QbwzzQlodIPi}< z@E!+#p#$%C;Nq5-l@2&?=S{<)1E*_xneKPsmr4+HuLED|!1p_F-GLu);Fme@gAUwz zopZ>6U+$1U?7){h@L>loh749Z>cFpb$d5VjatB_>79voHKp_H!2oxevh(I9%g$NWP zP>4Vw0)+?^BJjUC0>7Vs>3hb|sS+boeC!5HOQeR)qNy=s=(&<-S+l0r+;_dEO_hBf z+xbg<8p-!iN$%CDsj2S0jHd~w+>16(6HU2iY@Q~Va*x_PO)TZUYV$OimHV>I(*#rQ zKAWe>pIqGLX+kO2Y4bFZl-p|aG=Y@cX!A62l&i9NnlQ>Ovw4~*$}O~cnjp#***r}Q z}TsvL=I|~1%!oQ~Q|E};aEBtYVe_r9AQuu#V`0ptEKPdd)D*OWq|K|$-1%ZP;e3QaoukbYr?^XCK6#i0$|L8Qn@{Q*73wIdlSB;@xPWZPq zRt{GlGqM{-aTrf6z5t?cb$jP8{X9A$2Dt2oj{`BX#S4k7n=}_==zAHzGSJZXk*SmO zmkyHLvw}iCM}^rnuM%1Dfsvjtj=p(=addK);d$P8;RACK6g(&tluVuE{aV(i@o{j& zuTZQduWmPnHrz#7BmFD$TqCpL3t+}`?@vvQhcMbaUwj0-X9wz9_NjjEGbp4?JL@t8u3p57&Ej!7zK>bT>)PwKwHf!Gx*ub`kv_5?gM*QNd@m#W?q?)B z2-1>H4S+PHvnGL=#!@4l>Ot0N`9$TgkzsNfEuA{d1Uc)QolZUJC8UT(VTv_4m&6PJa%sOLKdRKXsI@S_`k@L;#=48uswl>Ww_R^u?#oA z-Hk6jd;?sd7ZzNWV-`{xL@Hbm}=67&<~*&1}dp7G-x}*BR+&p!}^`O>0h% zJ%Mv5>MyP&``IWEB1iig*-${O;NILncN0!SW%>e0YiZKwIVmW#GeB%ZKqA9*27F zLmBs>!s)|e%ZE+OhcoSoOp5k|<-_Z+lTMvrKG2r=KpFUO9LjhcLc#}}VyQnqgi8qY zhHnv(arg@;;&B+6r17P+cnfyb@8>T_Jq&H#^M5-``Gd%3_`#e_`QXS8pixU7r$b{E zbLHod7&<~*&0L?^M;ROt|AAk)>u^A9;ST8BfkvkTrIrIta9||$E!3prVUK2FAOE@ivmpF8+I5P@(msL}cs#64iO@ zkNqdrk%q`LH$>LR&yOKE8DrD_1sXvY3M4;PzMaj!+t^@eWMZ? z!?0r@zRl?gYhOQwxTI6#%rtG8Y0AL#eZur@-1O79h(sHINkq2syQt3F_^WrOnP$g( z*EG|9#dJM0eTtc;Ei+9Sn4T+4f0CPS;HD!^)76&gPchTHgQL}!X^v2NknVRzHF2yw$&bn_7J)(;bTGR%ZGwW}3FlG-Y6V zfiS&;o8HV#C!D6Qw@jNbji&WLKlP!12AfVzGQG5AdMN|FPgLE4AGVuq6`fAg`prfl;wsVB_d;Yw`J)2 zZ_y$77^xs}oOIEYmQ0o+){@yr`^?B>sn-~pebjG_%;PjXz`6aX)RJX&>-@;{1FrLb z%5zI=dW?KDs0~BU>FJb)lA$BC)spivtVU`$d1YoF+0C$Gb~>+d`k4Fu)Ksn=4K^~3 zhxqCKdT04Dl%utus^k&*7Z$MxDO4?)2Zj5!Fp0)cD_Z#YiWX91S~B0RVB~AOP1$mA zvT50UG%|`7PO+<>wyDSOCoKCQwX&_Qt!``GO{~=i0oVD_=3Hy~BsKdo@@wche>ydr z=}hgViYO>YK{{0eaOena*@#XVl)X=t(sYIK@(g? z-abvr$c~m=fa{@ZoWRM{kAVN?;tcbEMfMtAaF}-T_?J*BI~RogERZP*ndN_m9j(MJ zpe?gN8CWocg|Bl9DQ@8#L}VnBmW6Nrj_Mqt4uft` z!HDP->?mx=WQVaI8#4Q7WaC$&xIjG2_r8%mLjB;D%n=g0dh<<2HcMUUrpEMG z`UQ4IwT|HR)TE*3eCgC8X3Xj)9(a;YEde-mgtn}kPzKTeDS4TuZgQuWcaz6>;|^l( z0^$2U5!tx?u)xE2e1epb9WD70bdwE;!Efdwb1elUvu^-C8ku-0cYpfD_R+fZBqT3) zhHW$SIXckOoGpib>LR7A0T+WCIzn640Lq|jk=1~2tmF;&e`5dRE)ESiO++@}6$p#{ z&n;943*P~oPU*}7ZJ7niz`~;>%0gra3r~vu&n@hAT1Z+JzWH0$B)0#drG;5o&MeTD zS)dHc{Fa5+@EAAU4IJSXF6S1!P74=U7MfvUFUGPHxclNaXSn+ck*;ib{F8OJ*X?lJ z@5nST#m&4W0?{rh*!+QGjZ5$WG?2y_2tlYje|fv#ta$nZ~EJ4=nDI8myYSN6{AMim=Sq#XLonEarAZl=zB#+Czq~xCd1w98Z{!vcD%b(zBJ-QzLA#Y77 zZ_OHKYc#7hHKcbWRZBYX)Fh5X$|C0FMEHGa6*E=>0A~#nr2A6Jz>(*vJ`3XQFiSo0 z1#b4QFcpd*K1@UgF^$miAU=JH8L41KM#jgOg>#t&+A<52frXG|;XORgNEUv|Ev)4h zHaaa_X<66-3nM8XbPPg#B-MbNkL3)JL)W>C1tsMNGLqVE$=;3}o*+f74qJ9Mm+fJ) z{g$jrvKX00QUjEui-m0Qe=wg?gG}~*fsds2Ajf^!Ys)^vW%o1L1D3?YR_>rJdw|Ox zVzP&^UCmnbq$N9y+#oRca1WOqWwK+~((TMh>Lui87JMXi0y%c)kuCltmp#d3OR=SS zaOm(ed}Lhv4F%{eI>3>B(}|Oc4^kNhQ76acmS@=Tj2TBy&T7fd_Mjh?Lx9d>jHGl% zvb99$M}X{x3!(d#x+hx65_a<+Xg3c%S31OWdbsLqQ{>EN_rjI8G1X3T$=Au1$(gPY z^dVh&9J=8OyMu*^xaA7zgeycWS8$_Fjey*GLB$UBG@@LRX1IcNwcLbdRo}VHjJ^v?E=bpbY zjyl=mD1ITK*c=DTF7?#OzPDlDp6oO1nZ#|d<~aRii#GuvRbx2k!JF#@`y4b+pK&a; zwc!3|x_;_wwhfJRu!&1dDIg?!9i9xGDy5^QO!!Y>N*9k1n=Sq}xqx|0Y7AWY9v7N= z5&5r)3LB_Gw)i?GIWq1=%%GjkqmpLwA)TCn0{qj=DZ6$&@oNeOpU0)iRU@9s87&A zLI0-}d@zMOSS0DEPH$IIt8Z_p)A6=iEUpJly)3c5EVLd37e-yljV*O;Gs^HToS95)C<{S^C8M!@ zYjfMS=6b!UuDPYLK`*yNu7Lx|s0rZ}^RyoMs-(VUlV1KAsJJV+Lf2aBwzamm5ccZd z{3R(Azzk(8An=uWnh+GlfVdZ$mbd+^z3e?k765A4jsA9OpNc7=NL@m54*o$`nlhHf1p2~`fKxfeG!b^ZX!FV_n=&7=DmD;XA zZ)|4-Zw}H+i1Ds~iFZ~tZQh0r8}#~i3TmK!ORK+qTVr5#Ra2m;u5DXjThr<)9Rixg zxydsr)YIlQ#G+xk+|UMZpsoLn_8)j~M{<$W(U5sLSgVFav6LyUSWm<8K&3?EX?GDierYZ0VZpSp#9`7WK z+q1m&AN(WhW1`#|KX=aSmTyT{{)q)$k%6ZsO1Xw%NGx)2bgK~Q5ltX(&6)v8(A#Iku_ zd@g}Ll%UsoYv+0KWsIF+Q&02-)>O^&21B8EIFSe>@FfAgrg|QdR1*knzB$mkrJ=p0 zG0+Sno{MH(y9b4(_&@m4)YLHO(jGs+S9Ry$bZ&OoCz;6Pr1r394 z2R#T%=l4~Zl{f&p57Vr}px*&K3Hk~s4fbz?>Yz(8zw86O8?*!TAm|_{rA3@*p1Zbc zo`F)&Md!{fp>M(velh;{pO~7$G>xV;%q5q#<@l#Kp8Uzwl%I5*U)ps3l26P(=QAaP z+6@=3yK43FWk6Z|+d)gvHdc{9B+_eA1~~a76WOA;EI`{YB(U-P(k~R%pEu_wI0Tl; zzlHyE&QShClozAC9@MIT7q3r#pTd7R%D;%Ie5<@AJ8RGEq2gKb^OzvnS%mg~@ijaO zLnj^S$E+Q5NPx<#QT`)LXyd6FyWDU4xeevdqZ||Dc6o10ei?>j&qkJKi7g8Ytg$NWPP>4Vw z0)+?^B2b7xAp(U66e3WFKp_Huf(Xd(`^oS5$?x~!7L+CXyMdg?WD85O-2QGM$G3>` zb2*C}S)QR=dP>#Ta?0+_k(J*YyokeW-i5MkjsaADUy$ybDbZX5rMEwriV-~|>{-7c z2mzWJpmb7{%L$*m1W)tzl;r$5-jSG=gqU`K-Nf5-^Btmq=1C~Y-+A!<3`?V;p%Nbu zfslUtMd0w@4@=TMn`;0f4y;e zlM(s-b5&u5&JmxWenC3~?H6=V(7l2l5cH6s!-9?pdQ#8{LGg)umP!TH1+5U&C#YZ0 z4ng|`9Tar0pa%p!B64_n}%*KXoa9YLH&Yu2+FAi9SctktbYdlVNtjYJtn~ zNR~7sgD;$F{tk|if0Wws@4?Y5#|s8x^I80pemTxb{C;FS+Q+Q%N8*nllh6O}G5Jfh z?HycF9;g2bmh^vJ_%Fw;*Mz*>Pg1^!IedY&AGtP7I}f-AvzMhE9FgM!ecw8tKQzxw zap)29Of9OK0`C{NwEKU7lb_EjevS(K7Xq*6=CxOuo`qV8j{{QAEObH-zIa}$;L8QR zM8O*ct}A%Az{?f=};2zk9s!z{?f96c;YEe=8LH3W3)u_;mvJDfm`_ z8w%bdaKD0oPT<=WoW4Fy{&Xn#w}I1m^7eWjw=zz&<4pb%ZJ>h#GM;Y<`SXOly#8B= z3ltB2b?Vu~<>w1Ma=aZ>^jryfvU`Ql!+yg-9Yf$-1TL?qiaez7f*I#WIo_5F{c^nh zH1p>Y?S6qv{~r?i`-J|hNG;Mafy?V>ak&VK`N!>XrvIbbx%9w=lJbr7??4{E7HM)l z7Fp#ZaLj|NWO;omaeADc;-ib}?1iG>3MT(?YyBE2f9pI=(=}XucSwZORgjl;oh4gb z@L%A%H=q0t=*QW@TDRy5p{(CU&jY}9r$Lebnv47+xDd%__a9y4Um?6`fps4uP5r_} ze$oZM+{JFI3rnrIKy?eBv z=pwD8Ln53`xagUW`?-8^sC2=%0x!jNnorSxhl~7f7yPeX@aJ9dUjrw*dlkEaEG>6AuG?Mk`&{sE124t(t<*2nYmd3e|J((iaKXrOmRn(Wj9x3yGN@I^?lm`8oGGZHm|(GXiz zhMv5&mj&T#g0!1hyMroK@k&@FhLnXvsExV_t((z?pcy2w$~EvLC{`!g79KD)v#)=k z+YD4xR8_6+3|1h+)5;pGgrYoH9SU}bE4^L){nBD|7i$Zi748d0gF8XvSUcplmeueE zW!!cT^FPoXi>p#=nN%XUD|{wVzke5VeO&;?u+|Qi90_y<6J{V;g-4JBw{5P9?yBKW z6xD>nCKfYkS%(L113j^XnP_woz@ta2*91Bv=9vW6o<(2{9+*C(hBZ}Z6Igu~fts_K zsHr@ghO_yw_G~8Bo>jxyB3XMj6SXyGX+`bX_C@WwvuId%CLby*&>`3Z@UHt)h@2_h zTm;9jd#|Xhq^=oas{^T>vt?~0!FsLICtKD5=T+2_%8FogH-EEsxXIVw3)$GOi5)0jq!YqsOIj*b;PmNp8Wo8fj$gne4uB;COfeN`hz&j1h6!i ze^(>Mc4TO*q_aVOAu6SlKp$J(Di8@_ji)%PQo~Om`A#rw{1sz5J9pWmec&?@vqxy- z12snfN@klmA2lQ9Zn->FJ^FqpR-J;N!`VMZoOxJHi&jObWu#gvt?TWDe}Sg0b*+to z#?1}1;#vUfKn5BO0U8ZkwZM%nTQ=6U1hzCawKd|^l)8<0qEr(LRs|Yv-CWn&TrXZF zvX&L|Su0tIb*kjDR(wS^ySPOzWJQi)WiBj1=2%pW%Q=?2QWl7#1-X23;VW9>N}A!0 z}Y$v49$_@h4YdQ&#v3;ev`Uk7n&H zZryJ6nfwg}X|Xu&-n5S4Tvg32BGVQFbS~$Yr@5jdIF^i?VU$PDxptndN9@?ZZB-TY z!k2wX*vHo-tXv~wUmwG@*?{S_B~1UWzO9>Tx1nZzFcGfAct%&0_520C zZFmJLVEQ+y%M7Zk5Zc~2`VDr6b+CsLu>ilw6#GK1@n_dmI_Q~q4L^XL2j8?^)Jc!3 z)7p+0XQnTrXYH_Rc5HGz(`-#sp57Ky@3NSJKTX6YmLPAx3X#cN8W##VMMVh?gj)F#(j}rY}ep zM6SQr)bREpmSXfG>)ja>nMAlt^U_hOdFgHy(lHisysRj0zz}@ot&0QA8}1Qzu00{x zksQ~~Z(+GEnTHd>KCB$-m5O=CZVT`1!(FlF#lfR_aWl=h>G$v-hQrJ#=2>Y@l_y#c znUbXPxde`KS(zt#IsR!1SJtl;^(Cz!AxDz)780Tt8=dv#d#RExmx6-I`8^3YEA{2? zwUX*mP*Al#y-$sA55}bb@_7qMrT>`TXG!%x2ozIkR{Q1i7?SG3AhnU=EbGhp_y{lx znY3>S2toNg2SivR9nSP=F!Zd1tS_IBkaW8cmi|jQN$){9J$oT}`TT{X2Sk1HpPC~5 z7qFHUxU4Uq2a!}h4qBiN>C=|-)OC;?O{fX*dh<~;IL%=8|Qop8yp`^1d z&Gb%7>i+*5lsfC{Iv1Cey<_CiAm1gD^bw_g|1vHtsg#o{BroY9rT(C(FR4!kRw$75 zBz;1uKfIg^N*c!}VVz0Vli1V9INL9uUy>Blk<*eIzke3>EjN8aTu@ra)tTh}lh}`t zA^WnveEvgHYln*pYW)*x{Y_j@{k=_8SJkog&7Jn;?@*HNkrf41+y6FjXZ?~EE@)BN zFi~IipB_BH)GHeUNJKs_CaJo=WPM4`fpPjxQHEJQA0nxgBi)oFFX_inpN<__Uq1IF zsV@Ad+FlxW7FMqF-w1WyAN!FKCN4c}U)**xim3E{GSx@p;DD_KaJwYp^ zprC5~wWv>+^xr4yN$OYX>k5CJu&>&Y^;atL)K*GPs&o-n-rvE9Rx1ub9V|TqB>k7; pCw?xod`H$w2m5&TR_a`iB<)IQ0Wz}f5+mxN2aAvzl?n>0{Wq{(M3evk literal 0 HcmV?d00001 diff --git a/testing/cublaslt_debug.cu b/testing/cublaslt_debug.cu new file mode 100644 index 000000000..2890b46cd --- /dev/null +++ b/testing/cublaslt_debug.cu @@ -0,0 +1,246 @@ +/** + * Minimal cublasLtMatmul debug test. + * + * Tests the EXACT same configuration that fails in the Rust code: + * C[N, B] = op(W)[N, K] @ A[K, B] where op(W) = W^T + * + * Failing dimensions on H100: m=128, n=4096, k=256 + * Failing dimensions on 3050: m=128, n=512, k=64 + * + * Compile: + * nvcc -o cublaslt_debug cublaslt_debug.cu -lcublasLt -lcublas -lcudart + * + * Run: + * ./cublaslt_debug + */ + +#include +#include +#include +#include +#include +#include +#include + +#define CHECK_CUDA(call) \ + do { \ + cudaError_t err = (call); \ + if (err != cudaSuccess) { \ + fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \ + cudaGetErrorString(err)); \ + exit(1); \ + } \ + } while (0) + +#define CHECK_CUBLAS(call) \ + do { \ + cublasStatus_t st = (call); \ + if (st != CUBLAS_STATUS_SUCCESS) { \ + fprintf(stderr, "cuBLAS error at %s:%d: status=%d\n", \ + __FILE__, __LINE__, (int)st); \ + exit(1); \ + } \ + } while (0) + +// ── Test configuration ───────────────────────────────────────────── +struct TestCase { + int m, n, k; + const char* label; +}; + +void test_cublaslt_matmul(cublasLtHandle_t ltHandle, + cudaStream_t stream, + void* workspace, size_t wsSize, + const TestCase& tc, + cublasComputeType_t computeType, + const char* computeLabel) { + int m = tc.m, n = tc.n, k = tc.k; + printf(" [%s] (%d, %d, %d) compute=%s ... ", tc.label, m, n, k, computeLabel); + fflush(stdout); + + // Allocate device memory + float *d_W, *d_A, *d_C; + CHECK_CUDA(cudaMalloc(&d_W, (size_t)k * m * sizeof(float))); // W[m, k] row-major = [k, m] col-major + CHECK_CUDA(cudaMalloc(&d_A, (size_t)k * n * sizeof(float))); // A[n, k] row-major = [k, n] col-major + CHECK_CUDA(cudaMalloc(&d_C, (size_t)m * n * sizeof(float))); // C[n, m] row-major = [m, n] col-major + CHECK_CUDA(cudaMemset(d_W, 0, (size_t)k * m * sizeof(float))); + CHECK_CUDA(cudaMemset(d_A, 0, (size_t)k * n * sizeof(float))); + CHECK_CUDA(cudaMemset(d_C, 0, (size_t)m * n * sizeof(float))); + + float alpha = 1.0f, beta = 0.0f; + + // ── Approach 1: TRANSA=T (matching our Rust code) ────────────── + // GEMM: C[m, n] = W^T[m, k] @ A[k, n] + // Physical: W is [k, m] col-major (row-major W[m, k]), TRANSA=T → [m, k] + // A is [k, n] col-major (row-major A[n, k]), TRANSB=N → [k, n] + // C is [m, n] col-major + { + cublasLtMatmulDesc_t matmulDesc; + CHECK_CUBLAS(cublasLtMatmulDescCreate(&matmulDesc, computeType, CUDA_R_32F)); + + cublasOperation_t opT = CUBLAS_OP_T; + cublasOperation_t opN = CUBLAS_OP_N; + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(matmulDesc, + CUBLASLT_MATMUL_DESC_TRANSA, &opT, sizeof(opT))); + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(matmulDesc, + CUBLASLT_MATMUL_DESC_TRANSB, &opN, sizeof(opN))); + + // A = W: physical [k, m], ld=k. TRANSA=T → logical [m, k]. + cublasLtMatrixLayout_t Adesc, Bdesc, Cdesc, Ddesc; + CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Adesc, CUDA_R_32F, k, m, k)); + CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Bdesc, CUDA_R_32F, k, n, k)); + CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, CUDA_R_32F, m, n, m)); + CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Ddesc, CUDA_R_32F, m, n, m)); + + cublasLtMatmulPreference_t pref; + CHECK_CUBLAS(cublasLtMatmulPreferenceCreate(&pref)); + CHECK_CUBLAS(cublasLtMatmulPreferenceSetAttribute(pref, + CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &wsSize, sizeof(wsSize))); + + cublasLtMatmulHeuristicResult_t heurResult; + int returnedAlgoCount = 0; + cublasStatus_t heurStatus = cublasLtMatmulAlgoGetHeuristic( + ltHandle, matmulDesc, Adesc, Bdesc, Cdesc, Ddesc, + pref, 1, &heurResult, &returnedAlgoCount); + + if (heurStatus != CUBLAS_STATUS_SUCCESS || returnedAlgoCount == 0) { + printf("HEURISTIC FAILED (status=%d, count=%d)\n", (int)heurStatus, returnedAlgoCount); + } else { + printf("heuristic OK (ws=%zu) ", heurResult.workspaceSize); + fflush(stdout); + + cublasStatus_t matmulStatus = cublasLtMatmul( + ltHandle, matmulDesc, + &alpha, + d_W, Adesc, + d_A, Bdesc, + &beta, + d_C, Cdesc, + d_C, Ddesc, + &heurResult.algo, + workspace, wsSize, + stream); + + CHECK_CUDA(cudaStreamSynchronize(stream)); + + if (matmulStatus == CUBLAS_STATUS_SUCCESS) { + printf("MATMUL OK ✓\n"); + } else { + printf("MATMUL FAILED (status=%d) ✗\n", (int)matmulStatus); + } + } + + cublasLtMatmulPreferenceDestroy(pref); + cublasLtMatrixLayoutDestroy(Ddesc); + cublasLtMatrixLayoutDestroy(Cdesc); + cublasLtMatrixLayoutDestroy(Bdesc); + cublasLtMatrixLayoutDestroy(Adesc); + cublasLtMatmulDescDestroy(matmulDesc); + } + + // ── Approach 2: NO TRANSPOSE (swap A/B, row-major trick) ─────── + // GEMM: C'[m, n] = W'[m, k] @ A'[k, n] TRANSA=N, TRANSB=N + // But W stored row-major [m, k] → col-major [k, m]. So W' has rows=m, cols=k, ld=k. + // Wait — col-major [rows, cols] with ld: element (i,j) = data[j*ld + i]. + // Row-major W[m, k]: element (i,j) = data[i*k + j]. + // Col-major interpretation: rows=k, cols=m, ld=k. (j*k + i) maps to data[col*k + row]. + // This is [k, m] col-major. To get [m, k] col-major without transpose, + // we'd need ld=m, but ld must be >= rows=m. That works for an [m, k] layout with ld=m. + // But that's NOT how the data is stored. Row-major [m, k] IS col-major [k, m]. + // + // Correct no-transpose approach: swap the matrices. + // Row-major: C[n, m] = A[n, k] @ W[m, k]^T + // → col-major: C'[m, n] = W'[k, m]^T_col @ A'[k, n] + // But W'[k, m] transposed is [m, k]. That needs TRANSA=T again. + // + // Alternative: D = alpha * A * B + beta * C → m,n,k from D perspective. + // Row-major: out[batch, out_dim] = in[batch, in_dim] @ W[out_dim, in_dim]^T + // cuBLAS col-major: out'[out_dim, batch] = W'[in_dim, out_dim]^T[out_dim, in_dim] @ in'[in_dim, batch] + // So: M=out_dim, N=batch, K=in_dim, A=W with TRANSA=T, B=in with TRANSB=N. + // This is approach 1. There's no clean no-transpose version. + { + // Skip approach 2 — approach 1 is the correct formulation. + } + + // ── Approach 3: cublasSgemm for comparison ───────────────────── + { + cublasHandle_t handle; + CHECK_CUBLAS(cublasCreate(&handle)); + CHECK_CUBLAS(cublasSetStream(handle, stream)); + + cublasStatus_t sgemmStatus = cublasSgemm(handle, + CUBLAS_OP_T, CUBLAS_OP_N, + m, n, k, + &alpha, + d_W, k, // W[k, m] col-major, transpose → [m, k] + d_A, k, // A[k, n] col-major + &beta, + d_C, m); // C[m, n] col-major + + CHECK_CUDA(cudaStreamSynchronize(stream)); + printf(" [%s] (%d, %d, %d) cublasSgemm ... %s\n", tc.label, m, n, k, + sgemmStatus == CUBLAS_STATUS_SUCCESS ? "OK ✓" : "FAILED ✗"); + + cublasDestroy(handle); + } + + CHECK_CUDA(cudaFree(d_W)); + CHECK_CUDA(cudaFree(d_A)); + CHECK_CUDA(cudaFree(d_C)); +} + +int main() { + // Print GPU info + cudaDeviceProp prop; + CHECK_CUDA(cudaGetDeviceProperties(&prop, 0)); + printf("GPU: %s (SM %d.%d)\n", prop.name, prop.major, prop.minor); + printf("CUDA driver: "); + int driverVersion; + CHECK_CUDA(cudaDriverGetVersion(&driverVersion)); + printf("%d.%d\n", driverVersion / 1000, (driverVersion % 1000) / 10); + + // Create cublasLt handle + cublasLtHandle_t ltHandle; + CHECK_CUBLAS(cublasLtCreate(<Handle)); + + // Allocate workspace (32 MB) + size_t wsSize = 32 * 1024 * 1024; + void* workspace; + CHECK_CUDA(cudaMalloc(&workspace, wsSize)); + + // Create stream + cudaStream_t stream; + CHECK_CUDA(cudaStreamCreate(&stream)); + + // Test cases: the dimensions that fail in Rust + TestCase cases[] = { + {128, 64, 64, "training_small"}, + {128, 512, 64, "eval_chunk"}, + {128, 4096, 256, "h100_batch"}, + {256, 64, 256, "shared_h2"}, + {64, 64, 256, "shared_h1"}, + {51, 64, 128, "v_logits"}, + {3, 5, 4, "cudarc_test"}, + }; + int numCases = sizeof(cases) / sizeof(cases[0]); + + printf("\n=== CUBLAS_COMPUTE_32F_FAST_TF32 ===\n"); + for (int i = 0; i < numCases; i++) { + test_cublaslt_matmul(ltHandle, stream, workspace, wsSize, cases[i], + CUBLAS_COMPUTE_32F_FAST_TF32, "FAST_TF32"); + } + + printf("\n=== CUBLAS_COMPUTE_32F ===\n"); + for (int i = 0; i < numCases; i++) { + test_cublaslt_matmul(ltHandle, stream, workspace, wsSize, cases[i], + CUBLAS_COMPUTE_32F, "32F"); + } + + // Cleanup + CHECK_CUDA(cudaStreamDestroy(stream)); + CHECK_CUDA(cudaFree(workspace)); + CHECK_CUBLAS(cublasLtDestroy(ltHandle)); + + printf("\nDone.\n"); + return 0; +}