`iUddlmZddlZddlmZddlmZddlmZddl m Z dZ ej ej ejgZejejejejgZejejgZejejejejgZeezezezZdeeeDZd eDZd eDZ d Z!d Z"ej#e!d de DdeDzdeDzdeDzZ$ej#e"d de DdeDzdeDzdeDzdeDzdeDzZ%dZ&dZ'dZ(dZ)d#dZ*d Z+ d$d"Z,dS)%)productN)_normalize_axis_index) get_typename)runtime) axis_slicecRt|}|dkrtjrd}nd}|S)Nfloat16__halfhalf)rris_hip)dtypetypenames q/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/cupyx/scipy/signal/_iir_utils.py _get_typenamer s7E""H9 >  HHH Ocng|]2\}}tj||tj|u.||f3S)cupy promote_typesr .0xys r rsI<<<A#Aq))TZ]]::!f:::rc,g|]}t|Srr)rts rrr s . . .1mA . . .rcPg|]#\}}t|t|f$Srrrs rrr!s1OOODAqM!$$mA&6&67OOOra #include #include #include template __global__ void compute_correction_factors( const int m, const int k, const T* b, U* out) { int idx = blockDim.x * blockIdx.x + threadIdx.x; if(idx >= k) { return; } U* out_start = out + idx * (k + m); U* out_off = out_start + k; for(int i = 0; i < m; i++) { U acc = 0.0; for(int j = 0; j < k; j++) { acc += ((U) b[j]) * out_off[i - j - 1]; } out_off[i] = acc; } } template __global__ void first_pass_iir( const int m, const int k, const int n, const int n_blocks, const int carries_stride, const T* factors, T* out, T* carries) { int orig_idx = blockDim.x * (blockIdx.x % n_blocks) + threadIdx.x; int num_row = blockIdx.x / n_blocks; int idx = 2 * orig_idx + 1; if(idx >= n) { return; } int group_num = idx / m; int group_pos = idx % m; T* out_off = out + num_row * n; T* carries_off = carries + num_row * carries_stride; T* group_start = out_off + m * group_num; T* group_carries = carries_off + k * group_num; int pos = group_pos; int up_bound = pos; int low_bound = pos; int rel_pos; for(int level = 1, iter = 1; level < m; level *=2, iter++) { int sz = min(pow(2.0f, ((float) iter)), ((float) m)); if(level > 1) { int factor = ceil(pos / ((float) sz)); up_bound = sz * factor - 1; low_bound = up_bound - level + 1; } if(level == 1) { pos = low_bound; } if(pos < low_bound) { pos += level / 2; } if(pos + m * group_num >= n) { break; } rel_pos = pos % level; T carry = 0.0; for(int i = 1; i <= min(k, level); i++) { T k_value = group_start[low_bound - i]; const T* k_factors = factors + (m + k) * (i - 1) + k; T factor = k_factors[rel_pos]; carry += k_value * factor; } group_start[pos] += carry; __syncthreads(); } if(pos >= m - k) { if(carries != NULL) { group_carries[pos - (m - k)] = group_start[pos]; } } } template __global__ void correct_carries( const int m, const int k, const int n_blocks, const int carries_stride, const int offset, const T* factors, T* carries) { int idx = threadIdx.x; int pos = idx + (m - k); T* row_carries = carries + carries_stride * blockIdx.x; for(int i = offset; i < n_blocks; i++) { T* this_carries = row_carries + k * (i + (1 - offset)); T* prev_carries = row_carries + k * (i - offset); T carry = 0.0; for(int j = 1; j <= k; j++) { const T* k_factors = factors + (m + k) * (j - 1) + k; T factor = k_factors[pos]; T k_value = prev_carries[k - j]; carry += factor * k_value; } this_carries[idx] += carry; __syncthreads(); } } template __global__ void second_pass_iir( const int m, const int k, const int n, const int carries_stride, const int n_blocks, const int offset, const T* factors, T* carries, T* out) { int idx = blockDim.x * (blockIdx.x % n_blocks) + threadIdx.x; idx += offset * m; int row_num = blockIdx.x / n_blocks; int n_group = idx / m; int pos = idx % m; if(idx >= n) { return; } T* out_off = out + row_num * n; T* carries_off = carries + row_num * carries_stride; const T* prev_carries = carries_off + (n_group - offset) * k; T carry = 0.0; for(int i = 1; i <= k; i++) { const T* k_factors = factors + (m + k) * (i - 1) + k; T factor = k_factors[pos]; T k_value = prev_carries[k - i]; carry += factor * k_value; } out_off[idx] += carry; } a #include #include #include template __global__ void pick_carries( const int m, const int n, const int carries_stride, const int n_blocks, const int offset, T* x, T* carries) { int idx = m * (blockIdx.x % n_blocks) + threadIdx.x + m - 2; int pos = threadIdx.x; int row_num = blockIdx.x / n_blocks; int n_group = idx / m; T* x_off = x + row_num * n; T* carries_off = carries + row_num * carries_stride; T* group_carries = carries_off + (n_group + (1 - offset)) * 2; if(idx >= n) { return; } group_carries[pos] = x_off[idx]; } template __global__ void compute_correction_factors_sos( const int m, const T* f_const, U* all_out) { extern __shared__ __align__(sizeof(T)) thrust::complex bc_d[2]; T* b_c = reinterpret_cast(bc_d); extern __shared__ __align__(sizeof(T)) thrust::complex off_d[4]; U* off_cache = reinterpret_cast(off_d); int idx = threadIdx.x; int num_section = blockIdx.x; const int n_const = 6; const int a_off = 3; const int k = 2; const int off_idx = 1; U* out = all_out + num_section * k * m; U* out_start = out + idx * m; const T* b = f_const + num_section * n_const + a_off + 1; b_c[idx] = b[idx]; __syncthreads(); U* this_cache = off_cache + k * idx; this_cache[off_idx - idx] = 1; this_cache[idx] = 0; for(int i = 0; i < m; i++) { U acc = 0.0; for(int j = 0; j < k; j++) { acc += -((U) b_c[j]) * this_cache[off_idx - j]; } this_cache[0] = this_cache[1]; this_cache[1] = acc; out_start[i] = acc; } } template __global__ void first_pass_iir_sos( const int m, const int n, const int n_blocks, const T* factors, T* out, T* carries) { extern __shared__ unsigned int thread_status[2]; extern __shared__ __align__(sizeof(T)) thrust::complex fc_d[2 * 1024]; T* factor_cache = reinterpret_cast(fc_d); int orig_idx = blockDim.x * (blockIdx.x % n_blocks) + threadIdx.x; int num_row = blockIdx.x / n_blocks; int idx = 2 * orig_idx + 1; const int k = 2; if(idx >= n) { return; } int group_num = idx / m; int group_pos = idx % m; T* out_off = out + num_row * n; T* carries_off = carries + num_row * n_blocks * k; T* group_start = out_off + m * group_num; T* group_carries = carries_off + group_num * k; const T* section_factors = factors; T* section_carries = group_carries; factor_cache[group_pos] = section_factors[group_pos]; factor_cache[group_pos - 1] = section_factors[group_pos - 1]; factor_cache[m + group_pos] = section_factors[m + group_pos]; factor_cache[m + group_pos - 1] = section_factors[m + group_pos - 1]; __syncthreads(); int pos = group_pos; int up_bound = pos; int low_bound = pos; int rel_pos; for(int level = 1, iter = 1; level < m; level *= 2, iter++) { int sz = min(pow(2.0f, ((float) iter)), ((float) m)); if(level > 1) { int factor = ceil(pos / ((float) sz)); up_bound = sz * factor - 1; low_bound = up_bound - level + 1; } if(level == 1) { pos = low_bound; } if(pos < low_bound) { pos += level / 2; } if(pos + m * group_num >= n) { break; } rel_pos = pos % level; T carry = 0.0; for(int i = 1; i <= min(k, level); i++) { T k_value = group_start[low_bound - i]; const T* k_factors = factor_cache + m * (i - 1); T factor = k_factors[rel_pos]; carry += k_value * factor; } group_start[pos] += carry; __syncthreads(); } if(pos >= m - k) { if(carries != NULL) { section_carries[pos - (m - k)] = group_start[pos]; } } } template __global__ void correct_carries_sos( const int m, const int n_blocks, const int carries_stride, const int offset, const T* factors, T* carries) { extern __shared__ __align__(sizeof(T)) thrust::complex fcd3[4]; T* factor_cache = reinterpret_cast(fcd3); int idx = threadIdx.x; const int k = 2; int pos = idx + (m - k); T* row_carries = carries + carries_stride * blockIdx.x; factor_cache[2 * idx] = factors[pos]; factor_cache[2 * idx + 1] = factors[m + pos]; __syncthreads(); for(int i = offset; i < n_blocks; i++) { T* this_carries = row_carries + k * (i + (1 - offset)); T* prev_carries = row_carries + k * (i - offset); T carry = 0.0; for(int j = 1; j <= k; j++) { // const T* k_factors = factors + m * (j - 1); // T factor = k_factors[pos]; T factor = factor_cache[2 * idx + (j - 1)]; T k_value = prev_carries[k - j]; carry += factor * k_value; } this_carries[idx] += carry; __syncthreads(); } } template __global__ void second_pass_iir_sos( const int m, const int n, const int carries_stride, const int n_blocks, const int offset, const T* factors, T* carries, T* out) { extern __shared__ __align__(sizeof(T)) thrust::complex fcd2[2 * 1024]; T* factor_cache = reinterpret_cast(fcd2); extern __shared__ __align__(sizeof(T)) thrust::complex c_d[2]; T* carries_cache = reinterpret_cast(c_d); int idx = blockDim.x * (blockIdx.x % n_blocks) + threadIdx.x; idx += offset * m; int row_num = blockIdx.x / n_blocks; int n_group = idx / m; int pos = idx % m; const int k = 2; T* out_off = out + row_num * n; T* carries_off = carries + row_num * carries_stride; const T* prev_carries = carries_off + (n_group - offset) * k; if(pos < k) { carries_cache[pos] = prev_carries[pos]; } if(idx >= n) { return; } factor_cache[pos] = factors[pos]; factor_cache[pos + m] = factors[pos + m]; __syncthreads(); T carry = 0.0; for(int i = 1; i <= k; i++) { const T* k_factors = factor_cache + m * (i - 1); T factor = k_factors[pos]; T k_value = carries_cache[k - i]; carry += factor * k_value; } out_off[idx] += carry; } template __global__ void fir_sos( const int m, const int n, const int carries_stride, const int n_blocks, const int offset, const T* sos, T* carries, T* out) { extern __shared__ __align__(sizeof(T)) thrust::complex fir_cc[1024 + 2]; T* fir_cache = reinterpret_cast(fir_cc); extern __shared__ __align__(sizeof(T)) thrust::complex fir_b[3]; T* b = reinterpret_cast(fir_b); int idx = blockDim.x * (blockIdx.x % n_blocks) + threadIdx.x; int row_num = blockIdx.x / n_blocks; int n_group = idx / m; int pos = idx % m; const int k = 2; T* out_row = out + row_num * n; T* out_off = out_row + n_group * m; T* carries_off = carries + row_num * carries_stride; T* this_carries = carries_off + k * (n_group + (1 - offset)); T* group_carries = carries_off + (n_group - offset) * k; if(pos <= k) { b[pos] = sos[pos]; } if(pos < k) { if(offset && n_group == 0) { fir_cache[pos] = 0; } else { fir_cache[pos] = group_carries[pos]; } } if(idx >= n) { return; } fir_cache[pos + k] = out_off[pos]; __syncthreads(); T acc = 0.0; for(int i = k; i >= 0; i--) { acc += fir_cache[pos + i] * b[k - i]; } out_off[pos] = acc; } )z -std=c++11c&g|]\}}d|d|dS)zcompute_correction_factors<, >rrs rrrs?333!Q>A=====333rcg|]}d|d S)zcorrect_carries.s"EEE=++EEErr s^EE}EEEKyy%%H0=LY,,,,,,9K   - -F Mrctj||d}|j}|d|jd}|jjs|}||fS)Nrmoveaxisshapereshapeflags c_contiguouscopyraxisx_shapes r collapse_2drKsY ar""AgG "agbk""A 7  FFHH g:rctj||dzd}|j}||jdd|jd}|jjs|}||fS)Nr@rrArHs rcollapse_2d_restrNse a2&&AgG !'!*b!'"+..A 7  FFHH g:rc|j}tj||}tj|dddtj||f|f}t t d||}||fd||||f|S)Nr r@compute_correction_factorsrM)sizereyec_emptyr> IIR_MODULE)ablock_szr k correction corr_kernels rrQrQs A!5)))J44R4$*a]%@@@@BJ"0*aAAKKdXq!Z8999 rr@c >|tj|j|j}||}|||}|j}|j}t ||}|j}||} |dkr(t||\}}|t||\}} tj ||d} |jdkrdn |jd} | |zdz |z} | | z}tj ||}tj |dddtj ||f|f}tj | | |f|}ttd||}ttd| }ttd | }ttd | }||fd ||||f||f|d zf||| | | |z|| |f||jdkrtj|| d|jf}n-|jd kr"|| d|jd}|jdkr|}ntj||fd }|jjs|}| dks|Rt+|du}| |z }| d|z z|z}|| f|f||| ||||f|| |zf|f||| |||||| f |dkrK| |} tj| d|} | jjs| } | S)NrMTr rGrrPr@rQfirst_pass_iirsecond_pass_iircorrect_carriesrR)rI)r result_typer astyperCndimrrSrKarrayrTrUrVr>rW broadcast_torD concatenaterErFrGintrB)rrXrIzir rYrJx_ndimrZn_outnum_rowsn_blocks total_blocksr[carriesr\first_pass_kernelsecond_pass_kernelcarry_correction_kernelstarting_groupblocks_to_mergecarries_strides r apply_iirrzs } !'22 A ~ YYu  gG VF v . .D A A zz D)) 7 >D))EB *Qe$ / / /CFaKKqqQWQZHH q X-Hh&L!5)))J44R4$*a]%@@@@BJj 8Qu...G#0*aAAK(5EsKK)*6GMM.%s,,KdXq!Z8999|oA '7Ax(a!31222 ~ 7a<<"2!RW'=>>BB W\\Ha"66B <1  GG&G}1===G}) %llnnG!||r~R4Z"^3"a.&89Q> K! q(NN  " # # #   ' )H; q!^_ Z# 7 8 8 8 zzkk'""mCT**y% ((**C Jrc|jd}tj|d|f|}ttd||}||fd|||f|S)NrrcrPcompute_correction_factors_sos)rc)rCrrVr>IIR_SOS_MODULE)sosrYr n_sectionsr[r\s rr|r|esf1JZH5UCCCJ"8*cKKKK thZ%@AAA rTc |tj|j|j}||}|||}|j}|j} |jd} t || }d} ||} d} | dkrt||\}}|t||\}} |tj ||d}|jdkrdn |jd}| |zdz |z}||z}t|||}tj ||| f|}|}d}|0tj |}tj ||dz| f|}ttd|}ttd|}ttd |}ttd |}ttd |}t|du}||z }|d|z z| z}|||zf| f|| |||||ft!| D])}||}|9||ddddf}||dddddf<t#|| dz | ||ddddf<|r|||zf|f|| ||||||f||f|dzf|| |||||f|dks|f|)||ddddf}||dddddf<||ddddddf<||f| f|||||||f|||zf|f|| |||||||f|r|||zf| f|| |||||f|t#|| dz | ||ddddf<+| dkrK||}tj|d |}|jjs|}|^|| }t/| dkrtj|d |}|jjs|}|||fS|S) NrrcrMTr_rPfirst_pass_iir_sossecond_pass_iir_soscorrect_carries_sosfir_sos pick_carriesr@)rrdr rerCrfrrKrNrgr|rV empty_liker>r}rjrangerrDrBrErFrGlen) rr~rIrkr rY apply_firrorJrlrrZrmzi_shaperprqrrr[rs all_carrieszi_outrtrurv fir_kernelcarries_kernelrwrxrysb section_zis r apply_iir_sosrns@ } #)44 **U  C ~ YYu  gG VF1J v . .D A AH zz D)) 7 ~'D11 H {j%d333FaKKqqQWQZHH q X-Hh&L/XuEEJj 8Qu...GK F ~$$j x!|Q 'u666 ),c33)-s44.-s44!.)SAAJ%nncJJNt__N/O!n"45:NNHx')A4a>&''':  &9&9 F >Aqqq"1"HJ#-K1aaa )#q1ua88F1aaa!8   . J8+-{ !^X~;- . . .  Oh!m- q(JqM3 @ B B B a<<2>~111abb\ '1 AAAq!!!G$(/ AAAqrr111H% # # aT8^^qM;0 1 1 1  O+-{1no"JqM;E F F F  $ NH$&1nhk# $ $ $ >)#q1ua88F1aaa8  zzkk'""mCT**y% ((**C ~)) x==1  ]62t44F|( #[[]]F ~F{ Jr)r@NNr])r@NNr]TN)- itertoolsrrcupy._core.internalrcupy._core._scalarrcupy_backends.cuda.apircupyx.scipy.signal._arraytoolsrrr float32float64 FLOAT_TYPESint8int16int32int64 INT_TYPES complex64 complex128 COMPLEX_TYPESuint8uint16uint32uint64UNSIGNED_TYPESTYPES TYPE_PAIRS TYPE_NAMESTYPE_PAIR_NAMES IIR_KERNELIIR_SOS_KERNEL RawModulerWr}r>rKrNrQrzr|rrrrrs 555555++++++******555555   |T\4<8 Y DJ ; 1 *dk4; Di.0=@<<!6!6<<< / . . . . OOJOOOY vYvT^ _33"1333BBzBBBCBAjAAABCBzBBB CDDD   33"1333..:.../65*555654444 5 65*555 6 *)j))) *+++PPPPfBF&*qqqqqqr