 
 #include "aten4idaten.h"
 
-// NOTE
-// ddx, ddy
-// http://mosapui.blog116.fc2.com/blog-entry-35.html
-// https://www.gamedev.net/forums/topic/478820-derivative-instruction-details-ddx-ddy-or-dfdx-dfdy-etc/
-// http://d.hatena.ne.jp/umonist/20110616/p1
-// http://monsho.blog63.fc2.com/blog-entry-105.html
-
-inline __device__ float ddx(
-    int x, int y,
-    int w, int h,
-    const float4* __restrict__ aovNormalDepth)
-{
-    // NOTE
-    // 2x2 pixelƂɌvZ.
-
-    int leftX = x; 
-    int rightX = x + 1;
-
-#if 0
-    if ((x & 0x01) == 1) {
-        leftX = x - 1;
-        rightX = x;
-    }
-#else
-    int offset = (x & 0x01);
-    leftX -= offset;
-    rightX -= offset;
-#endif
-
-    rightX = min(rightX, w - 1);
-
-    const int idxL = getIdx(leftX, y, w);
-    const int idxR = getIdx(rightX, y, w);
-
-
-    float left = aovNormalDepth[idxL].w;
-    float right = aovNormalDepth[idxR].w;
-
-    return right - left;
-}
-
-inline __device__ float ddy(
-    int x, int y,
-    int w, int h,
-    const float4* __restrict__ aovNormalDepth)
-{
-    // NOTE
-    // 2x2 pixelƂɌvZ.
-
-    int topY = y;
-    int bottomY = y + 1;
-
-#if 0
-    if ((y & 0x01) == 1) {
-        topY = y - 1;
-        bottomY = y;
-    }
-#else
-    int offset = (y & 0x01);
-    topY -= offset;
-    bottomY -= offset;
-#endif
-
-    bottomY = min(bottomY, h - 1);
-
-    int idxT = getIdx(x, topY, w);
-    int idxB = getIdx(x, bottomY, w);
-
-    float top = aovNormalDepth[idxT].w;
-    float bottom = aovNormalDepth[idxB].w;
-
-    return bottom - top;
-}
-
 inline __device__ float gaussFilter3x3(
     int ix, int iy,
     int w, int h,

    const float4* __restrict__ clrVarBuffer,
    float4* nextClrVarBuffer,
    int stepScale,
    int width, int height)
{
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;

        return;
    }

    float centerLum = AT_NAME::color::luminance(centerColor.x, centerColor.y, centerColor.z);

    // KEXtB^3x3
