o
    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);
    ,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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 r6   K/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/cupyx/_texture.py_create_texture_object5   sT   



r8   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   float32r'   r8   !_affine_transform_2d_array_kernel!_affine_transform_3d_array_kernelzeros
isinstancetypendarray)r.   transformation_matrixoutput_shapeoutputr9   r:   border_valuer+   r   texture_objectkernelr6   r6   r7   affine_transformationp   sT   %





rL   )r   )NNr   r   r   )r   r   	cupy.cudar   r   ElementwiseKernelr@   rA   strr8   rL   r6   r6   r6   r7   <module>   sF    

=