o
    0 i5                     @   s   d dl Z d dl mZ d dlmZ d dlmZ ejddddd	d
Zejdddddd
Z	 ddededefddZ						 ddedefddZ
dS )    N)_core)texture)runtimez%U texObj, raw float32 m, uint64 widthzT transformed_imagea#  
    float3 pixel = make_float3(
        (float)(i / width),
        (float)(i % width),
        1.0f
    );
    float x = dot(pixel, make_float3(m[0],  m[1],  m[2])) + .5f;
    float y = dot(pixel, make_float3(m[3],  m[4],  m[5])) + .5f;
    transformed_image = tex2D<T>(texObj, y, x);
    Z,cupyx_texture_affine_transformation_2d_arrayz
    inline __host__ __device__ float dot(float3 a, float3 b)
    {
        return a.x * b.x + a.y * b.y + a.z * b.z;
    }
    )preamblez4U texObj, raw float32 m, uint64 height, uint64 widthzT transformed_volumea  
    float4 voxel = make_float4(
        (float)(i / (width * height)),
        (float)((i % (width * height)) / width),
        (float)((i % (width * height)) % width),
        1.0f
    );
    float x = dot(voxel, make_float4(m[0],  m[1],  m[2],  m[3])) + .5f;
    float y = dot(voxel, make_float4(m[4],  m[5],  m[6],  m[7])) + .5f;
    float z = dot(voxel, make_float4(m[8],  m[9],  m[10], m[11])) + .5f;
    transformed_volume = tex3D<T>(texObj, z, y, x);
    Z,cupyx_texture_affine_transformation_3d_arrayz
    inline __host__ __device__ float dot(float4 a, float4 b)
    {
        return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w;
    }
    address_modefilter_mode	read_modec                 C   sd  t | jt jrtj}n t | jt jrtj}nt | jt jr$tj	}nt
d| j |dkr4tj}n|dkr<tj}nt
d| d|dkrLtj}n|dkrTtj}nt
d| d|d	krdtj}n|d
krltj}nt
d| dt| jd ddd|}tj|g| jd d d R  }tjtj|d}tj|f| j |||fd}	t||	}
||  |
S )NzUnsupported data type nearestconstantzUnsupported address mode z (supported: constant, nearest)linearzUnsupported filter mode z (supported: nearest, linear)element_typeZnormalized_floatzUnsupported read mode z, (supported: element_type, normalized_float)   r   )ZcuArr)ZborderColors)cupyZ
issubdtypedtypeZunsignedintegerr   ZcudaChannelFormatKindUnsignedintegerZcudaChannelFormatKindSignedZfloatingZcudaChannelFormatKindFloat
ValueErrorZcudaAddressModeClampZcudaAddressModeBorderZcudaFilterModePointZcudaFilterModeLinearZcudaReadModeElementTypeZcudaReadModeNormalizedFloatr   ZChannelFormatDescriptoritemsizeZ	CUDAarrayshapeZResourceDescriptorZcudaResourceTypeArrayZTextureDescriptorndimZTextureObjectZ	copy_from)datar   r   r   border_colorZfmt_kindZtexture_fmtarrayZres_descZtex_descZtex_obj r   Z/home/app/PaddleOCR-VL-test/.venv_paddleocr/lib/python3.10/site-packages/cupyx/_texture.py_create_texture_object5   sT   



r   r   r
   interpolationmodec                 C   sB  | j }|dk s|dkrtd| j}|tjkrtd| d|dvr+td| d|j|d	 |d	 fkr:td
t| ||d|d}	|dkrJt}
nt}
|du rS| j}|du r_tj	||d}n1t
|ttjfr||tjkrttd| dtj	||d}nt
|tjr|j|krtdntd|
|	|g|d	d |R   |S )a  
    Apply an affine transformation.

    The method uses texture memory and supports only 2D and 3D float32 arrays
    without channel dimension.

    Args:
        data (cupy.ndarray): The input array or texture object.
        transformation_matrix (cupy.ndarray): Affine transformation matrix.
            Must be a homogeneous and have shape ``(ndim + 1, ndim + 1)``.
        output_shape (tuple of ints): Shape of output. If not specified,
            the input array shape is used. Default is None.
        output (cupy.ndarray or ~cupy.dtype): The array in which to place the
            output, or the dtype of the returned array. If not specified,
            creates the output array with shape of ``output_shape``. Default is
            None.
        interpolation (str): Specifies interpolation mode: ``'linear'`` or
            ``'nearest'``. Default is ``'linear'``.
        mode (str): Specifies addressing mode for points outside of the array:
            (`'constant'``, ``'nearest'``). Default is ``'constant'``.
        border_value: Specifies value to be used for coordinates outside
            of the array for ``'constant'`` mode. Default is 0.

    Returns:
        cupy.ndarray:
            The transformed input.

    .. seealso:: :func:`cupyx.scipy.ndimage.affine_transform`
          zdTexture memory affine transformation is defined only for 2D and 3D arrays without channel dimension.zRTexture memory affine transformation is available only for float32 data type (not ))r   r	   zUnsupported interpolation z (supported: linear, nearest)   z.Matrix must be have shape (ndim + 1, ndim + 1)r   )r   r   r   r   N)r   zOutput shapes do not matchz/Output must be None, cupy.ndarray or cupy.dtype)r   r   r   r   Zfloat32r   r   !_affine_transform_2d_array_kernel!_affine_transform_3d_array_kernelZzeros
isinstancetypeZndarray)r   Ztransformation_matrixZoutput_shapeoutputr   r   Zborder_valuer   r   Ztexture_objectZkernelr   r   r   affine_transformationp   sT   %





r'   )r   )NNr   r
   r   )r   r   Z	cupy.cudar   r   ZElementwiseKernelr"   r#   strr   r'   r   r   r   r   <module>   sF    

=