jS0dZddlZddlZddlmcmZddlmZm Z  ddl Z ddl m Z ddlmZmZn #e$rdZ YnwxYwdZdZejjejejjejejjejiZejdZd efd Zd Zd e fd Z!d Z"d efdZ#dZ$dZ%dZ&d"dZd#dZ'dZ(dZ)e e j*dZ+e j*dZ,e j*gdde j-de j-fdZ.e j*gdde j-de j-fd Z/dd!l0m1Z1e1dSdS)$aBF16 convolutions on FP16 tensor cores, for GPUs with FP16 but no BF16 tensor cores (validated on Turing, compute capability 7.5), where PyTorch has no fast BF16 convolution (cuDNN has no BF16 kernels there: `slow_conv_dilated3d`, i.e. vol2col and a CUDA-core GEMM). A BF16 value converts exactly to FP16 (8 significant bits into 11) while its magnitude stays in FP16's normal range. The weight is scaled by a power of two (largest |w| in [0.5, 1)); the input by the power of two that keeps every input below 2^15 and every possible output below 2^15 (largest |x| times the largest L1 norm of an output channel's weights), so nothing overflows; cuDNN accumulates in FP32; the FP16 result is scaled back, the bias added in FP32 and rounded once to BF16, in place. Precision: values that end up below FP16's normal range lose bits, and the result rounds to FP16 (11 bits) before BF16 (8 bits): LTX-2.5's VAE decoder, encoder and spatial upsampler measured within +-3% of the PyTorch BF16 route's error against FP32, 4.8-6.8x faster. Used with every kernel precision (user decision, 2026-10-21: a few % more error than PyTorch's route is worth it where PyTorch has no fast BF16 convolution), an exception to the strict rule. install_globally(device), called once by wgp.py: on an eligible GPU the stock _conv_forward of Conv1d / Conv2d / Conv3d (which subclasses with causal paddings and caches call through super()) computes every BF16 CUDA call of the process this way; other dtypes, devices, autograd, compilation and graph capture keep PyTorch's route. install(module, device) does the same for the layers of one module (tests, standalone use); conv2d() is the functional entry for a layer computing a convolution itself.N)choosecuda_inference)bf16_bits_to_fp32fp32_to_bf16_bitsixcHtj|dkS)N)torchcudaget_device_capability)indexs BC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\scaled_fp16.py _supportedr$s : + +E 2 2f <<returnctj|}tduoO|jdkoDtjjduo1tjot|j S)z|Validated on compute capability 7.5 only (Volta has the same FP16 tensor cores, untested); needs Triton for the conversions.Nr) r devicetritontypeversionhipr is_availablerr)rs r supportedr)s \& ! !F   Q&+"7 QEMDs  = 5<+= > >t u}$$)G)GEM)G^c^wx~^^)G(00ENBB(( 56 : : :flen>\>\F||EL%(2D#EEEZ_ZfIfIf TZ5>%A%AejFlFlFnFn%At Lrc~t|sdStD]#}d|jvr|j|_t |_$dS)zRoute the BF16 calls of every Conv1d / Conv2d / Conv3d of the process (their stock _conv_forward, which subclasses call through super()) on an eligible GPU (supported(device)); once per process, from wgp.py. False (nothing patched) elsewhere.F_scaled_fp16_stockT)rrr(r!r@_global_conv_forward)rr#s rinstall_globallyrBSsN V  u55 s| 3 3%(%6C " 4C  4rc$t|||}|dStddtfdtD}t ||||jjjj j j d S)zAlayer's convolution on the route for an eligible call, else None.NBF16 ConvolutionT owned_tritonc3BK|]\}}t||VdSN)r ).0r#r$r"s r z_route..es6 P PWS"E39O9O Pb P P P P P Pr_scaled_fp16_stats) r>rnextrr_scaledstridepaddingdilationgroups padding_mode _reversed_padding_repeated_twice)r"r;r<r=routedr$s` r_routerU_s ufd + +F ~t D1111 P P P P 3 3 5 5 P P P P PB 2vvtU\5=%.Z_Zfhmhz95BV X XXrcxt||||}|%t|||||n|SrH)rUrr@)r"r;r<r=outs rrArAjs? vt , ,CIL4;; ) )% E E EZ]]rcHt||||}| ||||n|SrH)rU)r"originalr;r<r=rWs rr!r!os2 vt , ,C,/K88E64 ( ( (S@rct|||} | tj|||||||Stddt tj| ||||||dd|| S)zF.conv2d, on the scaled-FP16 route where _conv_forward would take it (a layer computing a 2-D convolution itself, for example with a slice of a 3-D weight); owner.__dict__[key] caches the bound of that weight (one measurement, as for a layer).NrDTrEzeros)r>Fconv2drrM) r;r<r=rNrOrPrQownerkeyrTs rr]r]tstufd + +F ~xvtVWhOOO D1111 18VVT67HfV]_cejlo p pprrKc t|j|}||}tj|\}}|dd}tt ttt|  }|t | d|zzdzf}||j|<|S)z(power of two applied to the weight, largest L1 norm of an output channel's scaled weights): measured once per layer (or per owner and key), at its first call on the GPU (one synchronization).Nrg@g?)r(getabsr frexpamaxfloatflattensummax _SHIFT_LIMITminintitem) r^r<r_stats magnitude_exponentl1shifts r _weight_shiftrs~s N  s # #E }JJLL k).."2"2"8"8":":;; 8 __   & &q ) ) - -a 0 0 5 5 7 7\M3|c(--//6J6J5J#K#KLLbggii((3%<7>JK#s Lrctt|||||j|j|j|j|j|j|d S)zlayer's convolution (its stride, padding, padding mode, dilation, groups) of BF16 CUDA input, weight and bias on FP16 tensor cores.rK)rMr%rNrOrPrQrRrS)r"r;r<r=s r convolutionrusP ;u%%ufdEL%-Y^Yginiuw|xJ95BV X XXrc t| || \} } |}tj|\}}tj|jtj|j}tjdtj|j}ttj | tf|tj|| d|tj|tj| |dt d}|}tj|jtj|j}ttj | tf|tj|| | dddddt |dkrt!j|| | d}}tjd d 5|||d||||}dddn #1swxYwYdx}}|jd|d  }}t'tj | tf||tj|dn|tj|| | |||dut |tjS) z}fn (F.conv1d/2d/3d) of BF16 CUDA input, weight and bias on FP16 tensor cores; owner.__dict__[key] caches the bound of weight.)r1rrrTNgFr[)moder) device_typeenabled)rr)rs contiguousr aminmaxemptyshapefloat16rint32_to_fp16_kernelrcdivnumel_BLOCKviewint16r\padautocast_to_bf16_kernelr3)r$r;r<r=rNrOrPrQrRreversed_paddingr^r_ weight_shiftrqxlowhighx16rrw16ychannelsspatials rrMrMs8$UFC88L" A a  IC +agU]18 D D DC KQX > > >EV[F3356qvvek7J7JCQRQXQXQZQZ\]_b_g_ghmhs_t_tvzvvAFALwMwMOQ8=tVMMM A    F +fl%- N N NCV[88:;FKK > > > >E bga'k4;; < <}U?S?S S Wadd2:..55555r)rrrrHAS_BIASc @tjdtj| ztjd| z} | |k} tj||z } | dz} tj|| z| dtjt| zt| | z z}|r/|ttj|| |z|zz| dz }tj || zt|| dS)Nr) rrr6rrrr2rrrr)Y16Y_BITSBIASrrrrrrrrrtotalhalfrs rrrs -""%%bh//%7")Au:M:MM{'%..</0z GC'M4 + + . .rz : :]4=P=P PS`afimamSnSn n  ] "2747g3E2Q+QSWYZ#[#[\\ \A '!#4Q#7#7>>>>>r)!install_triton_compilation_logger)Nrrrr)rK)2__doc__r)r torch.nn.functionalr8 functionalr\dispatchrrrtriton.languagelanguager bf16_bitsrr ImportErrorrriConv1dconv1dConv2dr]Conv3dconv3drcacherboolrr%rkr-r>rBrUrAr!rsrurMjitrr constexprrr%shared.kernels.triton_compilation_logrrrrspp ,,,,,,,,MMM      ????????? FFF  !(EHOQXuxXYX`a  ===QQQQQ  s            XXX^^^ AAA qqqq    XXX """2  ZBBZB ZQQZQVZ"<"<"<===6R\6bdbn666>=6 VZ"N"N"NOOO ?`b`l ?uwvB ? ? ?PO ?XWWWWW%%'''''Qs 0::