Spaces:
Runtime error
Runtime error
| // Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. | |
| // | |
| // NVIDIA CORPORATION and its licensors retain all intellectual property | |
| // and proprietary rights in and to this software, related documentation | |
| // and any modifications thereto. Any use, reproduction, disclosure or | |
| // distribution of this software and related documentation without an express | |
| // license agreement from NVIDIA CORPORATION is strictly prohibited. | |
| #include "common.h" | |
| #include "texture.h" | |
| //------------------------------------------------------------------------ | |
| // Memory access and math helpers. | |
| static __device__ __forceinline__ void accum_from_mem(float* a, int s, float b, float c) { a[0] += b * c; } | |
| static __device__ __forceinline__ void accum_from_mem(float* a, int s, float2 b, float c) { a[0] += b.x * c; a[s] += b.y * c; } | |
| static __device__ __forceinline__ void accum_from_mem(float* a, int s, float4 b, float c) { a[0] += b.x * c; a[s] += b.y * c; a[2*s] += b.z * c; a[3*s] += b.w * c; } | |
| static __device__ __forceinline__ void accum_to_mem(float& a, float* b, int s) { a += b[0]; } | |
| static __device__ __forceinline__ void accum_to_mem(float2& a, float* b, int s) { float2 v = a; v.x += b[0]; v.y += b[s]; a = v; } | |
| static __device__ __forceinline__ void accum_to_mem(float4& a, float* b, int s) { float4 v = a; v.x += b[0]; v.y += b[s]; v.z += b[2*s]; v.w += b[3*s]; a = v; } | |
| static __device__ __forceinline__ bool isfinite_vec3(const float3& a) { return isfinite(a.x) && isfinite(a.y) && isfinite(a.z); } | |
| static __device__ __forceinline__ bool isfinite_vec4(const float4& a) { return isfinite(a.x) && isfinite(a.y) && isfinite(a.z) && isfinite(a.w); } | |
| template<class T> static __device__ __forceinline__ T lerp (const T& a, const T& b, float c) { return a + c * (b - a); } | |
| template<class T> static __device__ __forceinline__ T bilerp(const T& a, const T& b, const T& c, const T& d, const float2& e) { return lerp(lerp(a, b, e.x), lerp(c, d, e.x), e.y); } | |
| //------------------------------------------------------------------------ | |
| // Cube map wrapping for smooth filtering across edges and corners. At corners, | |
| // one of the texture coordinates will be negative. For correct interpolation, | |
| // the missing texel must take the average color of the other three. | |
| static __constant__ uint32_t c_cubeWrapMask1[48] = | |
| { | |
| 0x1530a440, 0x1133a550, 0x6103a110, 0x1515aa44, 0x6161aa11, 0x40154a04, 0x44115a05, 0x04611a01, | |
| 0x2630a440, 0x2233a550, 0x5203a110, 0x2626aa44, 0x5252aa11, 0x40264a04, 0x44225a05, 0x04521a01, | |
| 0x32608064, 0x3366a055, 0x13062091, 0x32328866, 0x13132299, 0x50320846, 0x55330a55, 0x05130219, | |
| 0x42508064, 0x4455a055, 0x14052091, 0x42428866, 0x14142299, 0x60420846, 0x66440a55, 0x06140219, | |
| 0x5230a044, 0x5533a055, 0x1503a011, 0x5252aa44, 0x1515aa11, 0x40520a44, 0x44550a55, 0x04150a11, | |
| 0x6130a044, 0x6633a055, 0x2603a011, 0x6161aa44, 0x2626aa11, 0x40610a44, 0x44660a55, 0x04260a11, | |
| }; | |
| static __constant__ uint8_t c_cubeWrapMask2[48] = | |
| { | |
| 0x26, 0x33, 0x11, 0x05, 0x00, 0x09, 0x0c, 0x04, 0x04, 0x00, 0x00, 0x05, 0x00, 0x81, 0xc0, 0x40, | |
| 0x02, 0x03, 0x09, 0x00, 0x0a, 0x00, 0x00, 0x02, 0x64, 0x30, 0x90, 0x55, 0xa0, 0x99, 0xcc, 0x64, | |
| 0x24, 0x30, 0x10, 0x05, 0x00, 0x01, 0x00, 0x00, 0x06, 0x03, 0x01, 0x05, 0x00, 0x89, 0xcc, 0x44, | |
| }; | |
| static __device__ __forceinline__ int4 wrapCubeMap(int face, int ix0, int ix1, int iy0, int iy1, int w) | |
| { | |
| // Calculate case number. | |
| int cx = (ix0 < 0) ? 0 : (ix1 >= w) ? 2 : 1; | |
| int cy = (iy0 < 0) ? 0 : (iy1 >= w) ? 6 : 3; | |
| int c = cx + cy; | |
| if (c >= 5) | |
| c--; | |
| c = (face << 3) + c; | |
| // Compute coordinates and faces. | |
| unsigned int m = c_cubeWrapMask1[c]; | |
| int x0 = (m >> 0) & 3; x0 = (x0 == 0) ? 0 : (x0 == 1) ? ix0 : iy0; | |
| int x1 = (m >> 2) & 3; x1 = (x1 == 0) ? 0 : (x1 == 1) ? ix1 : iy0; | |
| int x2 = (m >> 4) & 3; x2 = (x2 == 0) ? 0 : (x2 == 1) ? ix0 : iy1; | |
| int x3 = (m >> 6) & 3; x3 = (x3 == 0) ? 0 : (x3 == 1) ? ix1 : iy1; | |
| int y0 = (m >> 8) & 3; y0 = (y0 == 0) ? 0 : (y0 == 1) ? ix0 : iy0; | |
| int y1 = (m >> 10) & 3; y1 = (y1 == 0) ? 0 : (y1 == 1) ? ix1 : iy0; | |
| int y2 = (m >> 12) & 3; y2 = (y2 == 0) ? 0 : (y2 == 1) ? ix0 : iy1; | |
| int y3 = (m >> 14) & 3; y3 = (y3 == 0) ? 0 : (y3 == 1) ? ix1 : iy1; | |
| int f0 = ((m >> 16) & 15) - 1; | |
| int f1 = ((m >> 20) & 15) - 1; | |
| int f2 = ((m >> 24) & 15) - 1; | |
| int f3 = ((m >> 28) ) - 1; | |
| // Flips. | |
| unsigned int f = c_cubeWrapMask2[c]; | |
| int w1 = w - 1; | |
| if (f & 0x01) x0 = w1 - x0; | |
| if (f & 0x02) x1 = w1 - x1; | |
| if (f & 0x04) x2 = w1 - x2; | |
| if (f & 0x08) x3 = w1 - x3; | |
| if (f & 0x10) y0 = w1 - y0; | |
| if (f & 0x20) y1 = w1 - y1; | |
| if (f & 0x40) y2 = w1 - y2; | |
| if (f & 0x80) y3 = w1 - y3; | |
| // Done. | |
| int4 tcOut; | |
| tcOut.x = x0 + (y0 + f0 * w) * w; | |
| tcOut.y = x1 + (y1 + f1 * w) * w; | |
| tcOut.z = x2 + (y2 + f2 * w) * w; | |
| tcOut.w = x3 + (y3 + f3 * w) * w; | |
| return tcOut; | |
| } | |
| //------------------------------------------------------------------------ | |
| // Cube map indexing and gradient functions. | |
| // Map a 3D lookup vector into an (s,t) face coordinates (returned in first . | |
| // two parameters) and face index. | |
| static __device__ __forceinline__ int indexCubeMap(float& x, float& y, float z) | |
| { | |
| float ax = fabsf(x); | |
| float ay = fabsf(y); | |
| float az = fabsf(z); | |
| int idx; | |
| float c; | |
| if (az > fmaxf(ax, ay)) { idx = 4; c = z; } | |
| else if (ay > ax) { idx = 2; c = y; y = z; } | |
| else { idx = 0; c = x; x = z; } | |
| if (c < 0.f) idx += 1; | |
| float m = __frcp_rz(fabsf(c)) * .5; | |
| float m0 = __uint_as_float(__float_as_uint(m) ^ ((0x21u >> idx) << 31)); | |
| float m1 = (idx != 2) ? -m : m; | |
| x = x * m0 + .5; | |
| y = y * m1 + .5; | |
| if (!isfinite(x) || !isfinite(y)) | |
| return -1; // Invalid uv. | |
| x = fminf(fmaxf(x, 0.f), 1.f); | |
| y = fminf(fmaxf(y, 0.f), 1.f); | |
| return idx; | |
| } | |
| // Based on dA/d{s,t}, compute dA/d{x,y,z} at a given 3D lookup vector. | |
| static __device__ __forceinline__ float3 indexCubeMapGrad(float3 uv, float gu, float gv) | |
| { | |
| float ax = fabsf(uv.x); | |
| float ay = fabsf(uv.y); | |
| float az = fabsf(uv.z); | |
| int idx; | |
| float c; | |
| float c0 = gu; | |
| float c1 = gv; | |
| if (az > fmaxf(ax, ay)) { idx = 0x10; c = uv.z; c0 *= uv.x; c1 *= uv.y; } | |
| else if (ay > ax) { idx = 0x04; c = uv.y; c0 *= uv.x; c1 *= uv.z; } | |
| else { idx = 0x01; c = uv.x; c0 *= uv.z; c1 *= uv.y; } | |
| if (c < 0.f) idx += idx; | |
| float m = __frcp_rz(fabsf(c)); | |
| c0 = (idx & 0x34) ? -c0 : c0; | |
| c1 = (idx & 0x2e) ? -c1 : c1; | |
| float gl = (c0 + c1) * m; | |
| float gx = (idx & 0x03) ? gl : (idx & 0x20) ? -gu : gu; | |
| float gy = (idx & 0x0c) ? gl : -gv; | |
| float gz = (idx & 0x30) ? gl : (idx & 0x03) ? gu : gv; | |
| gz = (idx & 0x09) ? -gz : gz; | |
| float3 res = make_float3(gx, gy, gz) * (m * .5f); | |
| if (!isfinite_vec3(res)) | |
| return make_float3(0.f, 0.f, 0.f); // Invalid uv. | |
| return res; | |
| } | |
| // Based on dL/d(d{s,t}/s{X,Y}), compute dL/d(d{x,y,z}/d{X,Y}). This is just two | |
| // indexCubeMapGrad() functions rolled together. | |
| static __device__ __forceinline__ void indexCubeMapGrad4(float3 uv, float4 dw, float3& g0, float3& g1) | |
| { | |
| float ax = fabsf(uv.x); | |
| float ay = fabsf(uv.y); | |
| float az = fabsf(uv.z); | |
| int idx; | |
| float c, c0, c1; | |
| if (az > fmaxf(ax, ay)) { idx = 0x10; c = uv.z; c0 = uv.x; c1 = uv.y; } | |
| else if (ay > ax) { idx = 0x04; c = uv.y; c0 = uv.x; c1 = uv.z; } | |
| else { idx = 0x01; c = uv.x; c0 = uv.z; c1 = uv.y; } | |
| if (c < 0.f) idx += idx; | |
| float m = __frcp_rz(fabsf(c)); | |
| c0 = (idx & 0x34) ? -c0 : c0; | |
| c1 = (idx & 0x2e) ? -c1 : c1; | |
| float gl0 = (dw.x * c0 + dw.z * c1) * m; | |
| float gl1 = (dw.y * c0 + dw.w * c1) * m; | |
| float gx0 = (idx & 0x03) ? gl0 : (idx & 0x20) ? -dw.x : dw.x; | |
| float gx1 = (idx & 0x03) ? gl1 : (idx & 0x20) ? -dw.y : dw.y; | |
| float gy0 = (idx & 0x0c) ? gl0 : -dw.z; | |
| float gy1 = (idx & 0x0c) ? gl1 : -dw.w; | |
| float gz0 = (idx & 0x30) ? gl0 : (idx & 0x03) ? dw.x : dw.z; | |
| float gz1 = (idx & 0x30) ? gl1 : (idx & 0x03) ? dw.y : dw.w; | |
| if (idx & 0x09) | |
| { | |
| gz0 = -gz0; | |
| gz1 = -gz1; | |
| } | |
| g0 = make_float3(gx0, gy0, gz0) * (m * .5f); | |
| g1 = make_float3(gx1, gy1, gz1) * (m * .5f); | |
| if (!isfinite_vec3(g0) || !isfinite_vec3(g1)) | |
| { | |
| g0 = make_float3(0.f, 0.f, 0.f); // Invalid uv. | |
| g1 = make_float3(0.f, 0.f, 0.f); | |
| } | |
| } | |
| // Compute d{s,t}/d{X,Y} based on d{x,y,z}/d{X,Y} at a given 3D lookup vector. | |
| // Result is (ds/dX, ds/dY, dt/dX, dt/dY). | |
| static __device__ __forceinline__ float4 indexCubeMapGradST(float3 uv, float3 dvdX, float3 dvdY) | |
| { | |
| float ax = fabsf(uv.x); | |
| float ay = fabsf(uv.y); | |
| float az = fabsf(uv.z); | |
| int idx; | |
| float c, gu, gv; | |
| if (az > fmaxf(ax, ay)) { idx = 0x10; c = uv.z; gu = uv.x; gv = uv.y; } | |
| else if (ay > ax) { idx = 0x04; c = uv.y; gu = uv.x; gv = uv.z; } | |
| else { idx = 0x01; c = uv.x; gu = uv.z; gv = uv.y; } | |
| if (c < 0.f) idx += idx; | |
| if (idx & 0x09) | |
| { | |
| dvdX.z = -dvdX.z; | |
| dvdY.z = -dvdY.z; | |
| } | |
| float m = __frcp_rz(fabsf(c)); | |
| float dm = m * .5f; | |
| float mm = m * dm; | |
| gu *= (idx & 0x34) ? -mm : mm; | |
| gv *= (idx & 0x2e) ? -mm : mm; | |
| float4 res; | |
| if (idx & 0x03) | |
| { | |
| res = make_float4(gu * dvdX.x + dm * dvdX.z, | |
| gu * dvdY.x + dm * dvdY.z, | |
| gv * dvdX.x - dm * dvdX.y, | |
| gv * dvdY.x - dm * dvdY.y); | |
| } | |
| else if (idx & 0x0c) | |
| { | |
| res = make_float4(gu * dvdX.y + dm * dvdX.x, | |
| gu * dvdY.y + dm * dvdY.x, | |
| gv * dvdX.y + dm * dvdX.z, | |
| gv * dvdY.y + dm * dvdY.z); | |
| } | |
| else // (idx & 0x30) | |
| { | |
| res = make_float4(gu * dvdX.z + copysignf(dm, c) * dvdX.x, | |
| gu * dvdY.z + copysignf(dm, c) * dvdY.x, | |
| gv * dvdX.z - dm * dvdX.y, | |
| gv * dvdY.z - dm * dvdY.y); | |
| } | |
| if (!isfinite_vec4(res)) | |
| return make_float4(0.f, 0.f, 0.f, 0.f); | |
| return res; | |
| } | |
| // Compute d(d{s,t}/d{X,Y})/d{x,y,z}, i.e., how the pixel derivatives of 2D face | |
| // coordinates change w.r.t. 3D texture coordinate vector, returned as follows: | |
| // | d(ds/dX)/dx d(ds/dY)/dx d(dt/dX)/dx d(dt/dY)/dx | | |
| // | d(ds/dX)/dy d(ds/dY)/dy d(dt/dX)/dy d(dt/dY)/dy | | |
| // | d(ds/dX)/dz d(ds/dY)/dz d(dt/dX)/dz d(dt/dY)/dz | | |
| static __device__ __forceinline__ void indexCubeMapGrad2(float3 uv, float3 dvdX, float3 dvdY, float4& dx, float4& dy, float4& dz) | |
| { | |
| float ax = fabsf(uv.x); | |
| float ay = fabsf(uv.y); | |
| float az = fabsf(uv.z); | |
| int idx; | |
| float c, gu, gv; | |
| if (az > fmaxf(ax, ay)) { idx = 0x10; c = uv.z; gu = uv.x; gv = uv.y; } | |
| else if (ay > ax) { idx = 0x04; c = uv.y; gu = uv.x; gv = uv.z; } | |
| else { idx = 0x01; c = uv.x; gu = uv.z; gv = uv.y; } | |
| if (c < 0.f) idx += idx; | |
| if (idx & 0x09) | |
| { | |
| dvdX.z = -dvdX.z; | |
| dvdY.z = -dvdY.z; | |
| } | |
| float m = __frcp_rz(c); | |
| float dm = -m * fabsf(m) * .5; | |
| float mm = m * m * .5; | |
| float mu = (idx & 0x34) ? -mm : mm; | |
| float mv = (idx & 0x2e) ? -mm : mm; | |
| gu *= -2.0 * m * mu; | |
| gv *= -2.0 * m * mv; | |
| if (idx & 0x03) | |
| { | |
| dx.x = gu * dvdX.x + dm * dvdX.z; | |
| dx.y = gu * dvdY.x + dm * dvdY.z; | |
| dx.z = gv * dvdX.x - dm * dvdX.y; | |
| dx.w = gv * dvdY.x - dm * dvdY.y; | |
| dy.x = 0.f; | |
| dy.y = 0.f; | |
| dy.z = mv * dvdX.x; | |
| dy.w = mv * dvdY.x; | |
| dz.x = mu * dvdX.x; | |
| dz.y = mu * dvdY.x; | |
| dz.z = 0.f; | |
| dz.w = 0.f; | |
| } | |
| else if (idx & 0x0c) | |
| { | |
| dx.x = mu * dvdX.y; | |
| dx.y = mu * dvdY.y; | |
| dx.z = 0.f; | |
| dx.w = 0.f; | |
| dy.x = gu * dvdX.y + dm * dvdX.x; | |
| dy.y = gu * dvdY.y + dm * dvdY.x; | |
| dy.z = gv * dvdX.y + dm * dvdX.z; | |
| dy.w = gv * dvdY.y + dm * dvdY.z; | |
| dz.x = 0.f; | |
| dz.y = 0.f; | |
| dz.z = mv * dvdX.y; | |
| dz.w = mv * dvdY.y; | |
| } | |
| else // (idx & 0x30) | |
| { | |
| dx.x = mu * dvdX.z; | |
| dx.y = mu * dvdY.z; | |
| dx.z = 0.f; | |
| dx.w = 0.f; | |
| dy.x = 0.f; | |
| dy.y = 0.f; | |
| dy.z = mv * dvdX.z; | |
| dy.w = mv * dvdY.z; | |
| dz.x = gu * dvdX.z - fabsf(dm) * dvdX.x; | |
| dz.y = gu * dvdY.z - fabsf(dm) * dvdY.x; | |
| dz.z = gv * dvdX.z - dm * dvdX.y; | |
| dz.w = gv * dvdY.z - dm * dvdY.y; | |
| } | |
| } | |
| //------------------------------------------------------------------------ | |
| // General texture indexing. | |
| template <bool CUBE_MODE> | |
| static __device__ __forceinline__ int indexTextureNearest(const TextureKernelParams& p, float3 uv, int tz) | |
| { | |
| int w = p.texWidth; | |
| int h = p.texHeight; | |
| float u = uv.x; | |
| float v = uv.y; | |
| // Cube map indexing. | |
| if (CUBE_MODE) | |
| { | |
| // No wrap. Fold face index into tz right away. | |
| int idx = indexCubeMap(u, v, uv.z); // Rewrites u, v. | |
| if (idx < 0) | |
| return -1; // Invalid uv. | |
| tz = 6 * tz + idx; | |
| } | |
| else | |
| { | |
| // Handle boundary. | |
| if (p.boundaryMode == TEX_BOUNDARY_MODE_WRAP) | |
| { | |
| u = u - (float)__float2int_rd(u); | |
| v = v - (float)__float2int_rd(v); | |
| } | |
| } | |
| u = u * (float)w; | |
| v = v * (float)h; | |
| int iu = __float2int_rd(u); | |
| int iv = __float2int_rd(v); | |
| // In zero boundary mode, return texture address -1. | |
| if (!CUBE_MODE && p.boundaryMode == TEX_BOUNDARY_MODE_ZERO) | |
| { | |
| if (iu < 0 || iu >= w || iv < 0 || iv >= h) | |
| return -1; | |
| } | |
| // Otherwise clamp and calculate the coordinate properly. | |
| iu = min(max(iu, 0), w-1); | |
| iv = min(max(iv, 0), h-1); | |
| return iu + w * (iv + tz * h); | |
| } | |
| template <bool CUBE_MODE> | |
| static __device__ __forceinline__ float2 indexTextureLinear(const TextureKernelParams& p, float3 uv, int tz, int4& tcOut, int level) | |
| { | |
| // Mip level size. | |
| int2 sz = mipLevelSize(p, level); | |
| int w = sz.x; | |
| int h = sz.y; | |
| // Compute texture-space u, v. | |
| float u = uv.x; | |
| float v = uv.y; | |
| bool clampU = false; | |
| bool clampV = false; | |
| // Cube map indexing. | |
| int face = 0; | |
| if (CUBE_MODE) | |
| { | |
| // Neither clamp or wrap. | |
| face = indexCubeMap(u, v, uv.z); // Rewrites u, v. | |
| if (face < 0) | |
| { | |
| tcOut.x = tcOut.y = tcOut.z = tcOut.w = -1; // Invalid uv. | |
| return make_float2(0.f, 0.f); | |
| } | |
| u = u * (float)w - 0.5f; | |
| v = v * (float)h - 0.5f; | |
| } | |
| else | |
| { | |
| if (p.boundaryMode == TEX_BOUNDARY_MODE_WRAP) | |
| { | |
| // Wrap. | |
| u = u - (float)__float2int_rd(u); | |
| v = v - (float)__float2int_rd(v); | |
| } | |
| // Move to texel space. | |
| u = u * (float)w - 0.5f; | |
| v = v * (float)h - 0.5f; | |
| if (p.boundaryMode == TEX_BOUNDARY_MODE_CLAMP) | |
| { | |
| // Clamp to center of edge texels. | |
| u = fminf(fmaxf(u, 0.f), w - 1.f); | |
| v = fminf(fmaxf(v, 0.f), h - 1.f); | |
| clampU = (u == 0.f || u == w - 1.f); | |
| clampV = (v == 0.f || v == h - 1.f); | |
| } | |
| } | |
| // Compute texel coordinates and weights. | |
| int iu0 = __float2int_rd(u); | |
| int iv0 = __float2int_rd(v); | |
| int iu1 = iu0 + (clampU ? 0 : 1); // Ensure zero u/v gradients with clamped. | |
| int iv1 = iv0 + (clampV ? 0 : 1); | |
| u -= (float)iu0; | |
| v -= (float)iv0; | |
| // Cube map wrapping. | |
| bool cubeWrap = CUBE_MODE && (iu0 < 0 || iv0 < 0 || iu1 >= w || iv1 >= h); | |
| if (cubeWrap) | |
| { | |
| tcOut = wrapCubeMap(face, iu0, iu1, iv0, iv1, w); | |
| tcOut += 6 * tz * w * h; // Bring in tz. | |
| return make_float2(u, v); // Done. | |
| } | |
| // Fold cube map face into tz. | |
| if (CUBE_MODE) | |
| tz = 6 * tz + face; | |
| // Wrap overflowing texel indices. | |
| if (!CUBE_MODE && p.boundaryMode == TEX_BOUNDARY_MODE_WRAP) | |
| { | |
| if (iu0 < 0) iu0 += w; | |
| if (iv0 < 0) iv0 += h; | |
| if (iu1 >= w) iu1 -= w; | |
| if (iv1 >= h) iv1 -= h; | |
| } | |
| // Coordinates with tz folded in. | |
| int iu0z = iu0 + tz * w * h; | |
| int iu1z = iu1 + tz * w * h; | |
| tcOut.x = iu0z + w * iv0; | |
| tcOut.y = iu1z + w * iv0; | |
| tcOut.z = iu0z + w * iv1; | |
| tcOut.w = iu1z + w * iv1; | |
| // Invalidate texture addresses outside unit square if we are in zero mode. | |
| if (!CUBE_MODE && p.boundaryMode == TEX_BOUNDARY_MODE_ZERO) | |
| { | |
| bool iu0_out = (iu0 < 0 || iu0 >= w); | |
| bool iu1_out = (iu1 < 0 || iu1 >= w); | |
| bool iv0_out = (iv0 < 0 || iv0 >= h); | |
| bool iv1_out = (iv1 < 0 || iv1 >= h); | |
| if (iu0_out || iv0_out) tcOut.x = -1; | |
| if (iu1_out || iv0_out) tcOut.y = -1; | |
| if (iu0_out || iv1_out) tcOut.z = -1; | |
| if (iu1_out || iv1_out) tcOut.w = -1; | |
| } | |
| // All done. | |
| return make_float2(u, v); | |
| } | |
| //------------------------------------------------------------------------ | |
| // Mip level calculation. | |
| template <bool CUBE_MODE, bool BIAS_ONLY, int FILTER_MODE> | |
| static __device__ __forceinline__ void calculateMipLevel(int& level0, int& level1, float& flevel, const TextureKernelParams& p, int pidx, float3 uv, float4* pdw, float3* pdfdv) | |
| { | |
| // Do nothing if mips not in use. | |
| if (FILTER_MODE == TEX_MODE_NEAREST || FILTER_MODE == TEX_MODE_LINEAR) | |
| return; | |
| // Determine mip level based on UV pixel derivatives. If no derivatives are given (mip level bias only), leave as zero. | |
| if (!BIAS_ONLY) | |
| { | |
| // Get pixel derivatives of texture coordinates. | |
| float4 uvDA; | |
| float3 dvdX, dvdY; // Gradients use these later. | |
| if (CUBE_MODE) | |
| { | |
| // Fetch. | |
| float2 d0 = ((const float2*)p.uvDA)[3 * pidx + 0]; | |
| float2 d1 = ((const float2*)p.uvDA)[3 * pidx + 1]; | |
| float2 d2 = ((const float2*)p.uvDA)[3 * pidx + 2]; | |
| // Map d{x,y,z}/d{X,Y} into d{s,t}/d{X,Y}. | |
| dvdX = make_float3(d0.x, d1.x, d2.x); // d{x,y,z}/dX | |
| dvdY = make_float3(d0.y, d1.y, d2.y); // d{x,y,z}/dY | |
| uvDA = indexCubeMapGradST(uv, dvdX, dvdY); // d{s,t}/d{X,Y} | |
| } | |
| else | |
| { | |
| // Fetch. | |
| uvDA = ((const float4*)p.uvDA)[pidx]; | |
| } | |
| // Scaling factors. | |
| float uscl = p.texWidth; | |
| float vscl = p.texHeight; | |
| // d[s,t]/d[X,Y]. | |
| float dsdx = uvDA.x * uscl; | |
| float dsdy = uvDA.y * uscl; | |
| float dtdx = uvDA.z * vscl; | |
| float dtdy = uvDA.w * vscl; | |
| // Calculate footprint axis lengths. | |
| float A = dsdx*dsdx + dtdx*dtdx; | |
| float B = dsdy*dsdy + dtdy*dtdy; | |
| float C = dsdx*dsdy + dtdx*dtdy; | |
| float l2b = 0.5 * (A + B); | |
| float l2n = 0.25 * (A-B)*(A-B) + C*C; | |
| float l2a = sqrt(l2n); | |
| float lenMinorSqr = fmaxf(0.0, l2b - l2a); | |
| float lenMajorSqr = l2b + l2a; | |
| // Footprint vs. mip level gradient. | |
| if (pdw && FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_LINEAR) | |
| { | |
| float dw = 0.72134752f / (l2n + l2a * l2b); // Constant is 0.5/ln(2). | |
| float AB = dw * .5f * (A - B); | |
| float Cw = dw * C; | |
| float l2aw = dw * l2a; | |
| float d_f_ddsdX = uscl * (dsdx * (l2aw + AB) + dsdy * Cw); | |
| float d_f_ddsdY = uscl * (dsdy * (l2aw - AB) + dsdx * Cw); | |
| float d_f_ddtdX = vscl * (dtdx * (l2aw + AB) + dtdy * Cw); | |
| float d_f_ddtdY = vscl * (dtdy * (l2aw - AB) + dtdx * Cw); | |
| float4 d_f_dw = make_float4(d_f_ddsdX, d_f_ddsdY, d_f_ddtdX, d_f_ddtdY); | |
| if (!CUBE_MODE) | |
| *pdw = isfinite_vec4(d_f_dw) ? d_f_dw : make_float4(0.f, 0.f, 0.f, 0.f); | |
| // In cube maps, there is also a texture coordinate vs. mip level gradient. | |
| // Only output nonzero vectors if both are free of inf/Nan garbage. | |
| if (CUBE_MODE) | |
| { | |
| float4 dx, dy, dz; | |
| indexCubeMapGrad2(uv, dvdX, dvdY, dx, dy, dz); | |
| float3 d_dsdX_dv = make_float3(dx.x, dy.x, dz.x); | |
| float3 d_dsdY_dv = make_float3(dx.y, dy.y, dz.y); | |
| float3 d_dtdX_dv = make_float3(dx.z, dy.z, dz.z); | |
| float3 d_dtdY_dv = make_float3(dx.w, dy.w, dz.w); | |
| float3 d_f_dv = make_float3(0.f, 0.f, 0.f); | |
| d_f_dv += d_dsdX_dv * d_f_ddsdX; | |
| d_f_dv += d_dsdY_dv * d_f_ddsdY; | |
| d_f_dv += d_dtdX_dv * d_f_ddtdX; | |
| d_f_dv += d_dtdY_dv * d_f_ddtdY; | |
| bool finite = isfinite_vec4(d_f_dw) && isfinite_vec3(d_f_dv); | |
| *pdw = finite ? d_f_dw : make_float4(0.f, 0.f, 0.f, 0.f); | |
| *pdfdv = finite ? d_f_dv : make_float3(0.f, 0.f, 0.f); | |
| } | |
| } | |
| // Finally, calculate mip level. | |
| flevel = .5f * __log2f(lenMajorSqr); // May be inf/NaN, but clamp fixes it. | |
| } | |
| // Bias the mip level and clamp. | |
| if (p.mipLevelBias) | |
| flevel += p.mipLevelBias[pidx]; | |
| flevel = fminf(fmaxf(flevel, 0.f), (float)p.mipLevelMax); | |
| // Calculate levels depending on filter mode. | |
| level0 = __float2int_rd(flevel); | |
| // Leave everything else at zero if flevel == 0 (magnification) or when in linear-mipmap-nearest mode. | |
| if (FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_LINEAR && flevel > 0.f) | |
| { | |
| level1 = min(level0 + 1, p.mipLevelMax); | |
| flevel -= level0; // Fractional part. Zero if clamped on last level. | |
| } | |
| } | |
| //------------------------------------------------------------------------ | |
| // Texel fetch and accumulator helpers that understand cube map corners. | |
| template<class T> | |
| static __device__ __forceinline__ void fetchQuad(T& a00, T& a10, T& a01, T& a11, const float* pIn, int4 tc, bool corner) | |
| { | |
| // For invalid cube map uv, tc will be all negative, and all texel values will be zero. | |
| if (corner) | |
| { | |
| T avg = zero_value<T>(); | |
| if (tc.x >= 0) avg += (a00 = *((const T*)&pIn[tc.x])); | |
| if (tc.y >= 0) avg += (a10 = *((const T*)&pIn[tc.y])); | |
| if (tc.z >= 0) avg += (a01 = *((const T*)&pIn[tc.z])); | |
| if (tc.w >= 0) avg += (a11 = *((const T*)&pIn[tc.w])); | |
| avg *= 0.33333333f; | |
| if (tc.x < 0) a00 = avg; | |
| if (tc.y < 0) a10 = avg; | |
| if (tc.z < 0) a01 = avg; | |
| if (tc.w < 0) a11 = avg; | |
| } | |
| else | |
| { | |
| a00 = (tc.x >= 0) ? *((const T*)&pIn[tc.x]) : zero_value<T>(); | |
| a10 = (tc.y >= 0) ? *((const T*)&pIn[tc.y]) : zero_value<T>(); | |
| a01 = (tc.z >= 0) ? *((const T*)&pIn[tc.z]) : zero_value<T>(); | |
| a11 = (tc.w >= 0) ? *((const T*)&pIn[tc.w]) : zero_value<T>(); | |
| } | |
| } | |
| static __device__ __forceinline__ void accumQuad(float4 c, float* pOut, int level, int4 tc, bool corner, CA_TEMP_PARAM) | |
| { | |
| // For invalid cube map uv, tc will be all negative, and no accumulation will take place. | |
| if (corner) | |
| { | |
| float cb; | |
| if (tc.x < 0) cb = c.x; | |
| if (tc.y < 0) cb = c.y; | |
| if (tc.z < 0) cb = c.z; | |
| if (tc.w < 0) cb = c.w; | |
| cb *= 0.33333333f; | |
| if (tc.x >= 0) caAtomicAddTexture(pOut, level, tc.x, c.x + cb); | |
| if (tc.y >= 0) caAtomicAddTexture(pOut, level, tc.y, c.y + cb); | |
| if (tc.z >= 0) caAtomicAddTexture(pOut, level, tc.z, c.z + cb); | |
| if (tc.w >= 0) caAtomicAddTexture(pOut, level, tc.w, c.w + cb); | |
| } | |
| else | |
| { | |
| if (tc.x >= 0) caAtomicAddTexture(pOut, level, tc.x, c.x); | |
| if (tc.y >= 0) caAtomicAddTexture(pOut, level, tc.y, c.y); | |
| if (tc.z >= 0) caAtomicAddTexture(pOut, level, tc.z, c.z); | |
| if (tc.w >= 0) caAtomicAddTexture(pOut, level, tc.w, c.w); | |
| } | |
| } | |
| //------------------------------------------------------------------------ | |
| // Mip builder kernel. | |
| template<class T, int C> | |
| static __forceinline__ __device__ void MipBuildKernelTemplate(const TextureKernelParams p) | |
| { | |
| // Sizes. | |
| int2 sz_in = mipLevelSize(p, p.mipLevelOut - 1); | |
| int2 sz_out = mipLevelSize(p, p.mipLevelOut); | |
| // Calculate pixel position. | |
| int px = blockIdx.x * blockDim.x + threadIdx.x; | |
| int py = blockIdx.y * blockDim.y + threadIdx.y; | |
| int pz = blockIdx.z; | |
| if (px >= sz_out.x || py >= sz_out.y) | |
| return; | |
| // Pixel indices. | |
| int pidx_in0 = p.channels * (((px + sz_in.x * py) << 1) + (pz * sz_in.x * sz_in.y)); | |
| int pidx_in1 = pidx_in0 + p.channels * sz_in.x; // Next pixel down. | |
| int pidx_out = p.channels * (px + sz_out.x * (py + sz_out.y * pz)); | |
| // Input and output pointers. | |
| const float* pin = p.tex[p.mipLevelOut - 1]; | |
| float* pout = (float*)p.tex[p.mipLevelOut]; | |
| // Special case: Input texture height or width is 1. | |
| if (sz_in.x == 1 || sz_in.y == 1) | |
| { | |
| if (sz_in.y == 1) | |
| pidx_in1 = pidx_in0 + p.channels; // Next pixel on the right. | |
| for (int i=0; i < p.channels; i += C) | |
| { | |
| T v0 = *((const T*)&pin[pidx_in0 + i]); | |
| T v1 = *((const T*)&pin[pidx_in1 + i]); | |
| T avg = .5f * (v0 + v1); | |
| #if TEX_DEBUG_MIP_RETAIN_VARIANCE | |
| avg = (avg - .5f) * 1.41421356f + .5f; | |
| #endif | |
| *((T*)&pout[pidx_out + i]) = avg; | |
| } | |
| return; | |
| } | |
| for (int i=0; i < p.channels; i += C) | |
| { | |
| T v0 = *((const T*)&pin[pidx_in0 + i]); | |
| T v1 = *((const T*)&pin[pidx_in0 + i + p.channels]); | |
| T v2 = *((const T*)&pin[pidx_in1 + i]); | |
| T v3 = *((const T*)&pin[pidx_in1 + i + p.channels]); | |
| T avg = .25f * (v0 + v1 + v2 + v3); | |
| #if TEX_DEBUG_MIP_RETAIN_VARIANCE | |
| avg = (avg - .5f) * 2.f + .5f; | |
| #endif | |
| *((T*)&pout[pidx_out + i]) = avg; | |
| } | |
| } | |
| // Template specializations. | |
| __global__ void MipBuildKernel1(const TextureKernelParams p) { MipBuildKernelTemplate<float, 1>(p); } | |
| __global__ void MipBuildKernel2(const TextureKernelParams p) { MipBuildKernelTemplate<float2, 2>(p); } | |
| __global__ void MipBuildKernel4(const TextureKernelParams p) { MipBuildKernelTemplate<float4, 4>(p); } | |
| //------------------------------------------------------------------------ | |
| // Forward kernel. | |
| template <class T, int C, bool CUBE_MODE, bool BIAS_ONLY, int FILTER_MODE> | |
| static __forceinline__ __device__ void TextureFwdKernelTemplate(const TextureKernelParams p) | |
| { | |
| // Calculate pixel position. | |
| int px = blockIdx.x * blockDim.x + threadIdx.x; | |
| int py = blockIdx.y * blockDim.y + threadIdx.y; | |
| int pz = blockIdx.z; | |
| int tz = (p.texDepth == 1) ? 0 : pz; | |
| if (px >= p.imgWidth || py >= p.imgHeight || pz >= p.n) | |
| return; | |
| // Pixel index. | |
| int pidx = px + p.imgWidth * (py + p.imgHeight * pz); | |
| // Output ptr. | |
| float* pOut = p.out + pidx * p.channels; | |
| // Get UV. | |
| float3 uv; | |
| if (CUBE_MODE) | |
| uv = ((const float3*)p.uv)[pidx]; | |
| else | |
| uv = make_float3(((const float2*)p.uv)[pidx], 0.f); | |
| // Nearest mode. | |
| if (FILTER_MODE == TEX_MODE_NEAREST) | |
| { | |
| int tc = indexTextureNearest<CUBE_MODE>(p, uv, tz); | |
| tc *= p.channels; | |
| const float* pIn = p.tex[0]; | |
| // Copy if valid tc, otherwise output zero. | |
| for (int i=0; i < p.channels; i += C) | |
| *((T*)&pOut[i]) = (tc >= 0) ? *((const T*)&pIn[tc + i]) : zero_value<T>(); | |
| return; // Exit. | |
| } | |
| // Calculate mip level. In 'linear' mode these will all stay zero. | |
| float flevel = 0.f; // Fractional level. | |
| int level0 = 0; // Discrete level 0. | |
| int level1 = 0; // Discrete level 1. | |
| calculateMipLevel<CUBE_MODE, BIAS_ONLY, FILTER_MODE>(level0, level1, flevel, p, pidx, uv, 0, 0); | |
| // Get texel indices and pointer for level 0. | |
| int4 tc0 = make_int4(0, 0, 0, 0); | |
| float2 uv0 = indexTextureLinear<CUBE_MODE>(p, uv, tz, tc0, level0); | |
| const float* pIn0 = p.tex[level0]; | |
| bool corner0 = CUBE_MODE && ((tc0.x | tc0.y | tc0.z | tc0.w) < 0); | |
| tc0 *= p.channels; | |
| // Bilinear fetch. | |
| if (FILTER_MODE == TEX_MODE_LINEAR || FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_NEAREST) | |
| { | |
| // Interpolate. | |
| for (int i=0; i < p.channels; i += C, tc0 += C) | |
| { | |
| T a00, a10, a01, a11; | |
| fetchQuad<T>(a00, a10, a01, a11, pIn0, tc0, corner0); | |
| *((T*)&pOut[i]) = bilerp(a00, a10, a01, a11, uv0); | |
| } | |
| return; // Exit. | |
| } | |
| // Get texel indices and pointer for level 1. | |
| int4 tc1 = make_int4(0, 0, 0, 0); | |
| float2 uv1 = indexTextureLinear<CUBE_MODE>(p, uv, tz, tc1, level1); | |
| const float* pIn1 = p.tex[level1]; | |
| bool corner1 = CUBE_MODE && ((tc1.x | tc1.y | tc1.z | tc1.w) < 0); | |
| tc1 *= p.channels; | |
| // Trilinear fetch. | |
| for (int i=0; i < p.channels; i += C, tc0 += C, tc1 += C) | |
| { | |
| // First level. | |
| T a00, a10, a01, a11; | |
| fetchQuad<T>(a00, a10, a01, a11, pIn0, tc0, corner0); | |
| T a = bilerp(a00, a10, a01, a11, uv0); | |
| // Second level unless in magnification mode. | |
| if (flevel > 0.f) | |
| { | |
| T b00, b10, b01, b11; | |
| fetchQuad<T>(b00, b10, b01, b11, pIn1, tc1, corner1); | |
| T b = bilerp(b00, b10, b01, b11, uv1); | |
| a = lerp(a, b, flevel); // Interpolate between levels. | |
| } | |
| // Write. | |
| *((T*)&pOut[i]) = a; | |
| } | |
| } | |
| // Template specializations. | |
| __global__ void TextureFwdKernelNearest1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelNearest2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelNearest4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinear1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinear2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinear4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearest1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearest2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearest4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinear1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinear2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinear4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeNearest1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeNearest2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeNearest4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinear1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinear2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinear4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearest1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearest2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearest4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinear1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinear2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinear4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearestBO1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearestBO2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapNearestBO4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinearBO1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, false, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinearBO2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, false, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelLinearMipmapLinearBO4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, false, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearestBO1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearestBO2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapNearestBO4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinearBO1 (const TextureKernelParams p) { TextureFwdKernelTemplate<float, 1, true, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinearBO2 (const TextureKernelParams p) { TextureFwdKernelTemplate<float2, 2, true, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureFwdKernelCubeLinearMipmapLinearBO4 (const TextureKernelParams p) { TextureFwdKernelTemplate<float4, 4, true, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| //------------------------------------------------------------------------ | |
| // Gradient mip puller kernel. | |
| template<class T, int C> | |
| static __forceinline__ __device__ void MipGradKernelTemplate(const TextureKernelParams p) | |
| { | |
| // Calculate pixel position. | |
| int px = blockIdx.x * blockDim.x + threadIdx.x; | |
| int py = blockIdx.y * blockDim.y + threadIdx.y; | |
| int pz = blockIdx.z; | |
| if (px >= p.texWidth || py >= p.texHeight) | |
| return; | |
| // Number of wide elements. | |
| int c = p.channels; | |
| if (C == 2) c >>= 1; | |
| if (C == 4) c >>= 2; | |
| // Dynamically allocated shared memory for holding a texel. | |
| extern __shared__ float s_texelAccum[]; | |
| int sharedOfs = threadIdx.x + threadIdx.y * blockDim.x; | |
| int sharedStride = blockDim.x * blockDim.y; | |
| # define TEXEL_ACCUM(_i) (s_texelAccum + (sharedOfs + (_i) * sharedStride)) | |
| // Clear the texel. | |
| for (int i=0; i < p.channels; i++) | |
| *TEXEL_ACCUM(i) = 0.f; | |
| // Track texel position and accumulation weight over the mip stack. | |
| int x = px; | |
| int y = py; | |
| float w = 1.f; | |
| // Pull gradients from all levels. | |
| int2 sz = mipLevelSize(p, 0); // Previous level size. | |
| for (int level=1; level <= p.mipLevelMax; level++) | |
| { | |
| // Weight decay depends on previous level size. | |
| if (sz.x > 1) w *= .5f; | |
| if (sz.y > 1) w *= .5f; | |
| // Current level size and coordinates. | |
| sz = mipLevelSize(p, level); | |
| x >>= 1; | |
| y >>= 1; | |
| T* pIn = (T*)(p.gradTex[level] + (x + sz.x * (y + sz.y * pz)) * p.channels); | |
| for (int i=0; i < c; i++) | |
| accum_from_mem(TEXEL_ACCUM(i * C), sharedStride, pIn[i], w); | |
| } | |
| // Add to main texture gradients. | |
| T* pOut = (T*)(p.gradTex[0] + (px + p.texWidth * (py + p.texHeight * pz)) * p.channels); | |
| for (int i=0; i < c; i++) | |
| accum_to_mem(pOut[i], TEXEL_ACCUM(i * C), sharedStride); | |
| } | |
| // Template specializations. | |
| __global__ void MipGradKernel1(const TextureKernelParams p) { MipGradKernelTemplate<float, 1>(p); } | |
| __global__ void MipGradKernel2(const TextureKernelParams p) { MipGradKernelTemplate<float2, 2>(p); } | |
| __global__ void MipGradKernel4(const TextureKernelParams p) { MipGradKernelTemplate<float4, 4>(p); } | |
| //------------------------------------------------------------------------ | |
| // Gradient kernel. | |
| template <bool CUBE_MODE, bool BIAS_ONLY, int FILTER_MODE> | |
| static __forceinline__ __device__ void TextureGradKernelTemplate(const TextureKernelParams p) | |
| { | |
| // Temporary space for coalesced atomics. | |
| CA_DECLARE_TEMP(TEX_GRAD_MAX_KERNEL_BLOCK_WIDTH * TEX_GRAD_MAX_KERNEL_BLOCK_HEIGHT); | |
| // Calculate pixel position. | |
| int px = blockIdx.x * blockDim.x + threadIdx.x; | |
| int py = blockIdx.y * blockDim.y + threadIdx.y; | |
| int pz = blockIdx.z; | |
| int tz = (p.texDepth == 1) ? 0 : pz; | |
| if (px >= p.imgWidth || py >= p.imgHeight || pz >= p.n) | |
| return; | |
| // Pixel index. | |
| int pidx = px + p.imgWidth * (py + p.imgHeight * pz); | |
| // Early exit if output gradients are zero. | |
| const float* pDy = p.dy + pidx * p.channels; | |
| unsigned int dmax = 0u; | |
| if ((p.channels & 3) == 0) | |
| { | |
| for (int i=0; i < p.channels; i += 4) | |
| { | |
| uint4 dy = *((const uint4*)&pDy[i]); | |
| dmax |= (dy.x | dy.y | dy.z | dy.w); | |
| } | |
| } | |
| else | |
| { | |
| for (int i=0; i < p.channels; i++) | |
| dmax |= __float_as_uint(pDy[i]); | |
| } | |
| // Store zeros and exit. | |
| if (__uint_as_float(dmax) == 0.f) | |
| { | |
| if (CUBE_MODE) | |
| { | |
| if (FILTER_MODE != TEX_MODE_NEAREST) | |
| ((float3*)p.gradUV)[pidx] = make_float3(0.f, 0.f, 0.f); | |
| if (FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_LINEAR) | |
| { | |
| if (p.gradUVDA) | |
| { | |
| ((float2*)p.gradUVDA)[3 * pidx + 0] = make_float2(0.f, 0.f); | |
| ((float2*)p.gradUVDA)[3 * pidx + 1] = make_float2(0.f, 0.f); | |
| ((float2*)p.gradUVDA)[3 * pidx + 2] = make_float2(0.f, 0.f); | |
| } | |
| if (p.gradMipLevelBias) | |
| p.gradMipLevelBias[pidx] = 0.f; | |
| } | |
| } | |
| else | |
| { | |
| if (FILTER_MODE != TEX_MODE_NEAREST) | |
| ((float2*)p.gradUV)[pidx] = make_float2(0.f, 0.f); | |
| if (FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_LINEAR) | |
| { | |
| if (p.gradUVDA) | |
| ((float4*)p.gradUVDA)[pidx] = make_float4(0.f, 0.f, 0.f, 0.f); | |
| if (p.gradMipLevelBias) | |
| p.gradMipLevelBias[pidx] = 0.f; | |
| } | |
| } | |
| return; | |
| } | |
| // Get UV. | |
| float3 uv; | |
| if (CUBE_MODE) | |
| uv = ((const float3*)p.uv)[pidx]; | |
| else | |
| uv = make_float3(((const float2*)p.uv)[pidx], 0.f); | |
| // Nearest mode - texture gradients only. | |
| if (FILTER_MODE == TEX_MODE_NEAREST) | |
| { | |
| int tc = indexTextureNearest<CUBE_MODE>(p, uv, tz); | |
| if (tc < 0) | |
| return; // Outside texture. | |
| tc *= p.channels; | |
| float* pOut = p.gradTex[0]; | |
| // Accumulate texture gradients. | |
| for (int i=0; i < p.channels; i++) | |
| caAtomicAddTexture(pOut, 0, tc + i, pDy[i]); | |
| return; // Exit. | |
| } | |
| // Calculate mip level. In 'linear' mode these will all stay zero. | |
| float4 dw = make_float4(0.f, 0.f, 0.f, 0.f); | |
| float3 dfdv = make_float3(0.f, 0.f, 0.f); | |
| float flevel = 0.f; // Fractional level. | |
| int level0 = 0; // Discrete level 0. | |
| int level1 = 0; // Discrete level 1. | |
| calculateMipLevel<CUBE_MODE, BIAS_ONLY, FILTER_MODE>(level0, level1, flevel, p, pidx, uv, &dw, &dfdv); | |
| // UV gradient accumulators. | |
| float gu = 0.f; | |
| float gv = 0.f; | |
| // Get texel indices and pointers for level 0. | |
| int4 tc0 = make_int4(0, 0, 0, 0); | |
| float2 uv0 = indexTextureLinear<CUBE_MODE>(p, uv, tz, tc0, level0); | |
| const float* pIn0 = p.tex[level0]; | |
| float* pOut0 = p.gradTex[level0]; | |
| bool corner0 = CUBE_MODE && ((tc0.x | tc0.y | tc0.z | tc0.w) < 0); | |
| tc0 *= p.channels; | |
| // Texel weights. | |
| float uv011 = uv0.x * uv0.y; | |
| float uv010 = uv0.x - uv011; | |
| float uv001 = uv0.y - uv011; | |
| float uv000 = 1.f - uv0.x - uv001; | |
| float4 tw0 = make_float4(uv000, uv010, uv001, uv011); | |
| // Attribute weights. | |
| int2 sz0 = mipLevelSize(p, level0); | |
| float sclu0 = (float)sz0.x; | |
| float sclv0 = (float)sz0.y; | |
| // Bilinear mode - texture and uv gradients. | |
| if (FILTER_MODE == TEX_MODE_LINEAR || FILTER_MODE == TEX_MODE_LINEAR_MIPMAP_NEAREST) | |
| { | |
| for (int i=0; i < p.channels; i++, tc0 += 1) | |
| { | |
| float dy = pDy[i]; | |
| accumQuad(tw0 * dy, pOut0, level0, tc0, corner0, CA_TEMP); | |
| float a00, a10, a01, a11; | |
| fetchQuad<float>(a00, a10, a01, a11, pIn0, tc0, corner0); | |
| float ad = (a11 + a00 - a10 - a01); | |
| gu += dy * ((a10 - a00) + uv0.y * ad) * sclu0; | |
| gv += dy * ((a01 - a00) + uv0.x * ad) * sclv0; | |
| } | |
| // Store UV gradients and exit. | |
| if (CUBE_MODE) | |
| ((float3*)p.gradUV)[pidx] = indexCubeMapGrad(uv, gu, gv); | |
| else | |
| ((float2*)p.gradUV)[pidx] = make_float2(gu, gv); | |
| return; | |
| } | |
| // Accumulate fractional mip level gradient. | |
| float df = 0; // dL/df. | |
| // Get texel indices and pointers for level 1. | |
| int4 tc1 = make_int4(0, 0, 0, 0); | |
| float2 uv1 = indexTextureLinear<CUBE_MODE>(p, uv, tz, tc1, level1); | |
| const float* pIn1 = p.tex[level1]; | |
| float* pOut1 = p.gradTex[level1]; | |
| bool corner1 = CUBE_MODE && ((tc1.x | tc1.y | tc1.z | tc1.w) < 0); | |
| tc1 *= p.channels; | |
| // Texel weights. | |
| float uv111 = uv1.x * uv1.y; | |
| float uv110 = uv1.x - uv111; | |
| float uv101 = uv1.y - uv111; | |
| float uv100 = 1.f - uv1.x - uv101; | |
| float4 tw1 = make_float4(uv100, uv110, uv101, uv111); | |
| // Attribute weights. | |
| int2 sz1 = mipLevelSize(p, level1); | |
| float sclu1 = (float)sz1.x; | |
| float sclv1 = (float)sz1.y; | |
| // Trilinear mode. | |
| for (int i=0; i < p.channels; i++, tc0 += 1, tc1 += 1) | |
| { | |
| float dy = pDy[i]; | |
| float dy0 = (1.f - flevel) * dy; | |
| accumQuad(tw0 * dy0, pOut0, level0, tc0, corner0, CA_TEMP); | |
| // UV gradients for first level. | |
| float a00, a10, a01, a11; | |
| fetchQuad<float>(a00, a10, a01, a11, pIn0, tc0, corner0); | |
| float ad = (a11 + a00 - a10 - a01); | |
| gu += dy0 * ((a10 - a00) + uv0.y * ad) * sclu0; | |
| gv += dy0 * ((a01 - a00) + uv0.x * ad) * sclv0; | |
| // Second level unless in magnification mode. | |
| if (flevel > 0.f) | |
| { | |
| // Texture gradients for second level. | |
| float dy1 = flevel * dy; | |
| accumQuad(tw1 * dy1, pOut1, level1, tc1, corner1, CA_TEMP); | |
| // UV gradients for second level. | |
| float b00, b10, b01, b11; | |
| fetchQuad<float>(b00, b10, b01, b11, pIn1, tc1, corner1); | |
| float bd = (b11 + b00 - b10 - b01); | |
| gu += dy1 * ((b10 - b00) + uv1.y * bd) * sclu1; | |
| gv += dy1 * ((b01 - b00) + uv1.x * bd) * sclv1; | |
| // Mip level gradient. | |
| float a = bilerp(a00, a10, a01, a11, uv0); | |
| float b = bilerp(b00, b10, b01, b11, uv1); | |
| df += (b-a) * dy; | |
| } | |
| } | |
| // Store UV gradients. | |
| if (CUBE_MODE) | |
| ((float3*)p.gradUV)[pidx] = indexCubeMapGrad(uv, gu, gv) + (dfdv * df); | |
| else | |
| ((float2*)p.gradUV)[pidx] = make_float2(gu, gv); | |
| // Store mip level bias gradient. | |
| if (p.gradMipLevelBias) | |
| p.gradMipLevelBias[pidx] = df; | |
| // Store UV pixel differential gradients. | |
| if (!BIAS_ONLY) | |
| { | |
| // Final gradients. | |
| dw *= df; // dL/(d{s,y}/d{X,Y}) = df/(d{s,y}/d{X,Y}) * dL/df. | |
| // Store them. | |
| if (CUBE_MODE) | |
| { | |
| // Remap from dL/(d{s,t}/s{X,Y}) to dL/(d{x,y,z}/d{X,Y}). | |
| float3 g0, g1; | |
| indexCubeMapGrad4(uv, dw, g0, g1); | |
| ((float2*)p.gradUVDA)[3 * pidx + 0] = make_float2(g0.x, g1.x); | |
| ((float2*)p.gradUVDA)[3 * pidx + 1] = make_float2(g0.y, g1.y); | |
| ((float2*)p.gradUVDA)[3 * pidx + 2] = make_float2(g0.z, g1.z); | |
| } | |
| else | |
| ((float4*)p.gradUVDA)[pidx] = dw; | |
| } | |
| } | |
| // Template specializations. | |
| __global__ void TextureGradKernelNearest (const TextureKernelParams p) { TextureGradKernelTemplate<false, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureGradKernelLinear (const TextureKernelParams p) { TextureGradKernelTemplate<false, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureGradKernelLinearMipmapNearest (const TextureKernelParams p) { TextureGradKernelTemplate<false, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureGradKernelLinearMipmapLinear (const TextureKernelParams p) { TextureGradKernelTemplate<false, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureGradKernelCubeNearest (const TextureKernelParams p) { TextureGradKernelTemplate<true, false, TEX_MODE_NEAREST>(p); } | |
| __global__ void TextureGradKernelCubeLinear (const TextureKernelParams p) { TextureGradKernelTemplate<true, false, TEX_MODE_LINEAR>(p); } | |
| __global__ void TextureGradKernelCubeLinearMipmapNearest (const TextureKernelParams p) { TextureGradKernelTemplate<true, false, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureGradKernelCubeLinearMipmapLinear (const TextureKernelParams p) { TextureGradKernelTemplate<true, false, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureGradKernelLinearMipmapNearestBO (const TextureKernelParams p) { TextureGradKernelTemplate<false, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureGradKernelLinearMipmapLinearBO (const TextureKernelParams p) { TextureGradKernelTemplate<false, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| __global__ void TextureGradKernelCubeLinearMipmapNearestBO (const TextureKernelParams p) { TextureGradKernelTemplate<true, true, TEX_MODE_LINEAR_MIPMAP_NEAREST>(p); } | |
| __global__ void TextureGradKernelCubeLinearMipmapLinearBO (const TextureKernelParams p) { TextureGradKernelTemplate<true, true, TEX_MODE_LINEAR_MIPMAP_LINEAR>(p); } | |
| //------------------------------------------------------------------------ | |