AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

Add constraints to texture indirect functions

Ver Código-Fonte 
Similar to the previous patch, this change adds type constraints to texture indirect functions. Since we don't have to deduce the return type for these, we simply just have to check if the user provided a valid channel type.

Change-Id: Ia094bd6126e01df2ea90902c9aa59cb6cfe85773
Esse commit está contido em:
Vladislav Sytchenko
2020-03-17 13:09:33 -04:00
pai ef2415edc7
commit 117f0ab102
2 arquivos alterados com 178 adições e 59 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
include/hip/hcc_detail/texture_fetch_functions.h
+13 -13
Ver Arquivo
@@ -35,7 +35,7 @@ THE SOFTWARE.
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template<typename T>
struct __hip_is_channel_type
struct __hip_is_tex_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
@@ -50,17 +50,17 @@ struct __hip_is_channel_type
template<
typename T,
unsigned int rank>
struct __hip_is_channel_type<HIP_vector_type<T, rank>>
struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_channel_type<T>::value &&
__hip_is_tex_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template<typename T>
struct __hip_is_normalized_channel_type
struct __hip_is_tex_normalized_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
@@ -72,10 +72,10 @@ struct __hip_is_normalized_channel_type
template<
typename T,
unsigned int rank>
struct __hip_is_normalized_channel_type<HIP_vector_type<T, rank>>
struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_normalized_channel_type<T>::value &&
__hip_is_tex_normalized_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
@@ -99,7 +99,7 @@ template <typename T>
struct __hip_tex_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
{
using type = T;
};
@@ -110,7 +110,7 @@ template<
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
};
@@ -119,7 +119,7 @@ template<typename T>
struct __hip_tex_ret<
T,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float;
};
@@ -130,7 +130,7 @@ template<
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
};
@@ -333,7 +333,7 @@ template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
@@ -344,7 +344,7 @@ template<
struct __hip_tex2dgather_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
@@ -353,7 +353,7 @@ template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_normalized_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float4;
};
include/hip/hcc_detail/texture_indirect_functions.h
+165 -46
Ver Arquivo
@@ -28,11 +28,40 @@ THE SOFTWARE.
#include <hip/hip_texture_types.h>
#include <hip/hcc_detail/ockl_image.h>
#include <type_traits>
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template <typename T>
template<typename T>
struct __hip_is_itex_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_itex_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -40,13 +69,17 @@ static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
{
*ptr = tex1Dfetch<T>(textureObject, x);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1D(hipTextureObject_t textureObject, float x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -54,13 +87,17 @@ static __device__ T tex1D(hipTextureObject_t textureObject, float x)
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
{
*ptr = tex1D<T>(textureObject, x);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -68,13 +105,17 @@ static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
{
*ptr = tex2D<T>(textureObject, x, y);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -82,13 +123,17 @@ static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, fl
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = tex3D<T>(textureObject, x, y, z);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -96,13 +141,17 @@ static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, layer);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -110,13 +159,17 @@ static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, floa
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, y, layer);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -124,13 +177,17 @@ static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = texCubemap<T>(textureObject, x, y, z);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -138,13 +195,17 @@ static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -173,13 +234,17 @@ static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float
return {};
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, comp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -187,13 +252,17 @@ static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float le
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
{
*ptr = tex1DLod<T>(textureObject, x, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -201,13 +270,17 @@ static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
{
*ptr = tex2DLod<T>(textureObject, x, y, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -215,13 +288,17 @@ static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = tex3DLod<T>(textureObject, x, y, z, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -229,13 +306,17 @@ static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, i
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
{
*ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -243,13 +324,17 @@ static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
*ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
}
template <class T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -257,13 +342,17 @@ static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, flo
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = texCubemapLod<T>(textureObject, x, y, z, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -273,13 +362,17 @@ static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, fl
return {};
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -287,13 +380,17 @@ static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
*ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -301,13 +398,17 @@ static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float d
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
*ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -315,13 +416,17 @@ static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
*ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -329,13 +434,17 @@ static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -343,13 +452,17 @@ static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
*ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -357,13 +470,17 @@ static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x,
return *reinterpret_cast<T*>(&tmp);
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
*ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
}
template <class T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -373,7 +490,9 @@ static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, flo
return {};
}
template <typename T>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);