
    Spf                        d dl Z d dl mZ d dlZd dlZd dlZd dlmc mZ ddl	m
Z
 ddlmZ d dlmZ dD ]&Z	  ej        e dd	
          Z n# e$ r dZY #w xY werJe                                                                D ]#\  ZZedk    rd ndZ ej        eede           $dD ]&Z	  ej        e dd	
          Z n# e$ r dZY #w xY wer@e                                                                D ]\  ZZ ej        eedd           d Zd Z eede          Z eede          Zd Z eede          Z dS )    N)partial   )custom_call)GpuLibNotLinkedError)
xla_client)z.cudajax_cuda12_pluginz._linalgjaxlib)packagecu_cholesky_updateCUDA)platformapi_version)z.rocmjax_rocm60_pluginROCMc                 B    t          j        t          j        | d          S )Nr   )	functoolsreduceoperatormul)xss    Q/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jaxlib/gpu_linalg.py<lambda>r   =   s    9#HL"a88     c                .   t          j        |j                  }|j        }t           j                            d          }|j        |k    s
J |            |st                      t          t          t          |          dz
  dd                    }|}t          |          }	||	d<   t           j                            |	|          }
t          |  dd|
g|gt          t           j                            ||                    |g|g          j        S )z?Kernel for the transformation of pivots to permutations on GPU.    r   _lu_pivots_to_permutation   )permutation_size)r   result_typesoperandsbackend_configoperand_layoutsresult_layouts)irRankedTensorTypetypeshapeIntegerTypeget_signlesselement_typer   tuplerangelenlistgetr   dictIntegerAttrresults)r   
gpu_linalgpivotsr   typdimsi32_typepivots_layoutpermutations_layoutpermutations_dimspermutations_types              r   _lu_pivots_to_permutation_hlor=   @   s%   
FK((#	$^((,,(		X	%	%	%s	%	%	%	 !

 
  c$ii!mR4455-%4jj*B)--.?JJ		,,,%&x>--h8HII   %o)*

 

 

 
r   cuhipc           	         ~ t          j        |j                  }|j        }|d         |d         k    sJ |d         }|st	                      t          j        |          }|                    ||          }	t          d||gt           j        	                    ||f|j
                  t           j        	                    |f|j
                  gddd|	          j        dd         S )zCholesky update.r   r   r   )r   r   )r!   r    operand_output_aliasesr"   N)r%   r&   r'   r(   r   npdtype build_cholesky_update_descriptorr   r0   r+   r3   )
r   r4   r_matrixw_vectorrC   r_typer7   nnp_typeopaques
             r   _cholesky_update_hlorK   c   s    x}--&	$	aDG				
1g!	 !

 
  HUOO'66wBB&	H%


!
!1a&&*=
>
>


!
!1$(;
<
< "#q\\	
 	
 	
 BQB	
 	r   )!r   r   	importlibnumpyrB   r   jaxlib.mlir.irmlirr%   hlo_helpersr   gpu_common_utilsr   r	   r   cuda_module_nameimport_module_cuda_linalgImportErrorregistrationsitems_name_valuer   register_custom_call_targetrocm_module_name_hip_linalg_prodr=   cuda_lu_pivots_to_permutationhip_lu_pivots_to_permutationrK   cuda_cholesky_update r   r   <module>rb      s                               $ $ $ $ $ $ 2 2 2 2 2 2      6 
 

*9*%%%x  L 
E 
   LLL
  #113399;;  meV 444!!!K*J*vK     7 
 

)))%%%x  K 
E 
   KKK
  "002288::  meV*J*vA     	98  8 !((Et(4!6 !6 &w!5+ 7  7 
  4 w3T<HH   s#   AAA,CCC