o
    X۷iY                     @  s   d dl mZ 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ddZ
					 dd ddZdS )!    )annotationsN)_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);
    ,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);
    ,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str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_typenormalized_floatzUnsupported read mode z, (supported: element_type, normalized_float)   r   )cuArr)borderColors)cupy
issubdtypedtypeunsignedintegerr   cudaChannelFormatKindUnsignedintegercudaChannelFormatKindSignedfloatingcudaChannelFormatKindFloat
ValueErrorcudaAddressModeClampcudaAddressModeBordercudaFilterModePointcudaFilterModeLinearcudaReadModeElementTypecudaReadModeNormalizedFloatr   ChannelFormatDescriptoritemsize	CUDAarrayshapeResourceDescriptorcudaResourceTypeArrayTextureDescriptorndimTextureObject	copy_from)datar	   r   r   border_colorfmt_kindtexture_fmtarrayres_desctex_desctex_obj r8   D/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupyx/_texture.py_create_texture_object7   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   r1   N)r   zOutput shapes do not matchz/Output must be None, cupy.ndarray or cupy.dtype)r-   r   r   r   float32r)   r:   !_affine_transform_2d_array_kernel!_affine_transform_3d_array_kernelzeros
isinstancetypendarray)r0   transformation_matrixoutput_shapeoutputr;   r<   border_valuer-   r   texture_objectkernelr8   r8   r9   affine_transformationr   sT   %





rN   )r   )r	   r
   r   r
   r   r
   )NNr   r   r   )r;   r
   r<   r
   )
__future__r   r   r   	cupy.cudar   r   ElementwiseKernelrB   rC   r:   rN   r8   r8   r8   r9   <module>   s4    
=