util/X86MAC64/cuda/samples/3_Imaging/imageDenoising/imageDenoising_knn_kernel.cuh

   1 /*
   2  * Copyright 1993-2013 NVIDIA Corporation.  All rights reserved.
   3  *
   4  * Please refer to the NVIDIA end user license agreement (EULA) associated
   5  * with this source code for terms and conditions that govern your use of
   6  * this software. Any use, reproduction, disclosure, or distribution of
   7  * this software and related documentation outside the terms of the EULA
   8  * is strictly prohibited.
   9  *
  10  */
  11
  12
  13
  14 ////////////////////////////////////////////////////////////////////////////////
  15 // KNN kernel
  16 ////////////////////////////////////////////////////////////////////////////////
  17 __global__ void KNN(
  18     TColor *dst,
  19     int imageW,
  20     int imageH,
  21     float Noise,
  22     float lerpC
  23 )
  24 {
  25     const int ix = blockDim.x * blockIdx.x + threadIdx.x;
  26     const int iy = blockDim.y * blockIdx.y + threadIdx.y;
  27     //Add half of a texel to always address exact texel centers
  28     const float x = (float)ix + 0.5f;
  29     const float y = (float)iy + 0.5f;
  30
  31     if (ix < imageW && iy < imageH)
  32     {
  33         //Normalized counter for the weight threshold
  34         float fCount = 0;
  35         //Total sum of pixel weights
  36         float sumWeights = 0;
  37         //Result accumulator
  38         float3 clr = {0, 0, 0};
  39         //Center of the KNN window
  40         float4 clr00 = tex2D(texImage, x, y);
  41
  42         //Cycle through KNN window, surrounding (x, y) texel
  43         for (float i = -KNN_WINDOW_RADIUS; i <= KNN_WINDOW_RADIUS; i++)
  44             for (float j = -KNN_WINDOW_RADIUS; j <= KNN_WINDOW_RADIUS; j++)
  45             {
  46                 float4     clrIJ = tex2D(texImage, x + j, y + i);
  47                 float distanceIJ = vecLen(clr00, clrIJ);
  48
  49                 //Derive final weight from color distance
  50                 float   weightIJ = __expf(- (distanceIJ * Noise + (i * i + j * j) * INV_KNN_WINDOW_AREA));
  51
  52                 //Accumulate (x + j, y + i) texel color with computed weight
  53                 clr.x += clrIJ.x * weightIJ;
  54                 clr.y += clrIJ.y * weightIJ;
  55                 clr.z += clrIJ.z * weightIJ;
  56
  57                 //Sum of weights for color normalization to [0..1] range
  58                 sumWeights     += weightIJ;
  59
  60                 //Update weight counter, if KNN weight for current window texel
  61                 //exceeds the weight threshold
  62                 fCount         += (weightIJ > KNN_WEIGHT_THRESHOLD) ? INV_KNN_WINDOW_AREA : 0;
  63             }
  64
  65         //Normalize result color by sum of weights
  66         sumWeights = 1.0f / sumWeights;
  67         clr.x *= sumWeights;
  68         clr.y *= sumWeights;
  69         clr.z *= sumWeights;
  70
  71         //Choose LERP quotent basing on how many texels
  72         //within the KNN window exceeded the weight threshold
  73         float lerpQ = (fCount > KNN_LERP_THRESHOLD) ? lerpC : 1.0f - lerpC;
  74
  75         //Write final result to global memory
  76         clr.x = lerpf(clr.x, clr00.x, lerpQ);
  77         clr.y = lerpf(clr.y, clr00.y, lerpQ);
  78         clr.z = lerpf(clr.z, clr00.z, lerpQ);
  79         dst[imageW * iy + ix] = make_color(clr.x, clr.y, clr.z, 0);
  80     };
  81 }
  82
  83 extern "C"
  84 void cuda_KNN(
  85     TColor *d_dst,
  86     int imageW,
  87     int imageH,
  88     float Noise,
  89     float lerpC
  90 )
  91 {
  92     dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  93     dim3 grid(iDivUp(imageW, BLOCKDIM_X), iDivUp(imageH, BLOCKDIM_Y));
  94
  95     KNN<<<grid, threads>>>(d_dst, imageW, imageH, Noise, lerpC);
  96 }
  97
  98
  99 ////////////////////////////////////////////////////////////////////////////////
 100 // Stripped KNN kernel, only highlighting areas with different LERP directions
 101 ////////////////////////////////////////////////////////////////////////////////
 102 __global__ void KNNdiag(
 103     TColor *dst,
 104     int imageW,
 105     int imageH,
 106     float Noise,
 107     float lerpC
 108 )
 109 {
 110     const int ix = blockDim.x * blockIdx.x + threadIdx.x;
 111     const int iy = blockDim.y * blockIdx.y + threadIdx.y;
 112     //Add half of a texel to always address exact texel centers
 113     const float x = (float)ix + 0.5f;
 114     const float y = (float)iy + 0.5f;
 115
 116     if (ix < imageW && iy < imageH)
 117     {
 118         //Normalized counter for the weight threshold
 119         float  fCount = 0;
 120         //Center of the KNN window
 121         float4  clr00 = tex2D(texImage, x, y);
 122
 123         //Cycle through KNN window, surrounding (x, y) texel
 124         for (float i = -KNN_WINDOW_RADIUS; i <= KNN_WINDOW_RADIUS; i++)
 125             for (float j = -KNN_WINDOW_RADIUS; j <= KNN_WINDOW_RADIUS; j++)
 126             {
 127                 float4     clrIJ = tex2D(texImage, x + j, y + i);
 128                 float distanceIJ = vecLen(clr00, clrIJ);
 129
 130                 //Derive final weight from color and geometric distance
 131                 float weightIJ  = __expf(- (distanceIJ * Noise + (i * i + j * j) * INV_KNN_WINDOW_AREA));
 132
 133                 //Update weight counter, if KNN weight for current window texel
 134                 //exceeds the weight threshold
 135                 fCount += (weightIJ > KNN_WEIGHT_THRESHOLD) ? INV_KNN_WINDOW_AREA : 0.0f;
 136             }
 137
 138         //Choose LERP quotent basing on how many texels
 139         //within the KNN window exceeded the weight threshold
 140         float lerpQ = (fCount > KNN_LERP_THRESHOLD) ? 1.0f : 0;
 141
 142         //Write final result to global memory
 143         dst[imageW * iy + ix] = make_color(lerpQ, 0, (1.0f - lerpQ), 0);
 144     };
 145 }
 146
 147 extern "C"
 148 void cuda_KNNdiag(
 149     TColor *d_dst,
 150     int imageW,
 151     int imageH,
 152     float Noise,
 153     float lerpC
 154 )
 155 {
 156     dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
 157     dim3 grid(iDivUp(imageW, BLOCKDIM_X), iDivUp(imageH, BLOCKDIM_Y));
 158
 159     KNNdiag<<<grid, threads>>>(d_dst, imageW, imageH, Noise, lerpC);
 160 }