finding local mean in an image using mex-cuda

Question

I have an image named HSIImage, of size is 565x585, in which I have find the local mean and standard deviation at every pixel. For this I am using a window W of size 9x9, if we a re finding the mean of x(i,j) we need values in the W where x(i,j) is at its center.

For working on the corner and edge pixels, I am padding the HSIImage and naming it as HSIImage2.

MATLAB code

[m,n,~]  =  size(HSIImage);
HSIImage2=padarray(HSIImage,[4,4],'symmetric');

mean1    = zeros(m,n);
sd       = zeros(m,n);
phi_x=zeros(m,n);

for i=5:m+4
    for j=5:n+4
        mean1(i-4,j-4) = mean( mean(HSIImage2(i-4:i+4, j-4:j+4, 3) )); %sum / (4*4);
        sd(i-4,j-4) = std( std(HSIImage2(i-4:i+4, j-4:j+4, 3), 1)); 
    end
end
[phi_x2,mean2,sd2] = getPhi(HSIImage(:,:,3)',HSIImage2(:,:,3)',m,n);

Serial mean displayed as image. My cuda code for finding mean and sd is

__global__ void phi(double *d_HSIImage,double *d_HSIImage2, int row, int col, double *d_phi_x, double *d_mean, double *d_std)
{
int X = blockDim.x * blockIdx.x + threadIdx.x;
int Y = blockDim.y * blockIdx.y + threadIdx.y;
int i,j;
double sum = 0;

if(Y>3 && X>3 && Y<row+4 && X<col+4)
{
    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            sum= sum + d_HSIImage2[i*col+j];
        }
    }

    d_mean[(Y-4)*col+X-4] = sum/81;
    double mean = sum/81;
    sum = 0;

    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            int index = i*col+j;
            sum= sum + (d_HSIImage2[index]-mean) * (d_HSIImage2[index]-mean);
        }
    }
    d_std[(Y-4)*col+X-4] = sqrt(sum/81);
}
void mexFunction( int nlhs, mxArray *plhs[],int nrhs, const mxArray *prhs[])
{
double*     HSIImage;
double*     d_HSIImage;
double*     HSIImage2;
double*     d_HSIImage2;
double      row;
double      col;
double*     phi_x;
double*     d_phi_x;
double*     mean2;
double*     d_mean;
double*     d_std;
double*     sd2;

HSIImage    = (double*)mxGetPr(prhs[0]);
HSIImage2   = (double*)mxGetPr(prhs[1]);
row         = mxGetScalar(prhs[2]);
col         = mxGetScalar(prhs[3]);

plhs[0]     = mxCreateDoubleMatrix(row,col,mxREAL);
phi_x       = mxGetPr(plhs[0]);
plhs[1]     = mxCreateDoubleMatrix(row,col,mxREAL);
mean2       = mxGetPr(plhs[1]);
plhs[2]     = mxCreateDoubleMatrix(row,col,mxREAL);
sd2         = mxGetPr(plhs[2]);

dim3 grid(((col+8)/TILE_WIDTH)+1,((row+8)/TILE_WIDTH)+1,1);
dim3 block(TILE_WIDTH,TILE_WIDTH,1);

if ( cudaMalloc(&d_HSIImage,sizeof(double)*row*col)!= cudaSuccess )  
    mexErrMsgTxt("Memory allocating failure on the GPU");
if ( cudaMalloc(&d_HSIImage2,sizeof(double)*(row+8)*(col+8))!= cudaSuccess )  
    mexErrMsgTxt("Memory allocating failure on the GPU");
if ( cudaMalloc(&d_phi_x,sizeof(double)*row*col)!= cudaSuccess )  
    mexErrMsgTxt("Memory allocating failure on the GPU");
if ( cudaMalloc(&d_mean,sizeof(double)*row*col)!= cudaSuccess )  
    mexErrMsgTxt("Memory allocating failure on the GPU");
if ( cudaMalloc(&d_std,sizeof(double)*row*col)!= cudaSuccess )  
    mexErrMsgTxt("Memory allocating failure on the GPU");

cudaMemcpy(d_HSIImage,HSIImage,sizeof(double)*row*col,cudaMemcpyHostToDevice);
cudaMemcpy(d_HSIImage2,HSIImage2,sizeof(double)*(row+8)*(col+8),cudaMemcpyHostToDevice);

phi <<< grid,block >>> (d_HSIImage,d_HSIImage2,row,col,d_phi_x,d_mean,d_std);

cudaMemcpy(phi_x,d_phi_x,sizeof(double)*row*col,cudaMemcpyDeviceToHost);
cudaMemcpy(mean2,d_mean,sizeof(double)*row*col,cudaMemcpyDeviceToHost);
cudaMemcpy(sd2,d_std,sizeof(double)*row*col,cudaMemcpyDeviceToHost);

cudaFree(d_HSIImage);
cudaFree(d_HSIImage2);
cudaFree(d_phi_x);

}

its working fine when image is full of ones. but when I give regular image, there is lot of difference in serial(MATLAB) and parallel(CUDA) outputs(When mean1 and mean2 are compared). Please tell me the error.

I am launching with

dim3 grid(((col+8)/TILE_WIDTH)+1,((row+8)/TILE_WIDTH)+1,1);
dim3 block(TILE_WIDTH,TILE_WIDTH,1);

TILEWIDTH is 32. row=565, col=584. Parallel mean displayed as image

Answer 1

It is important to note Matlab's c api is column-major ordered, however as mentioned in the comments OP has made sure of the consistency. The problem is that the stride used to access the data did not include the pads of the image. Going from one row to another requires a stride of col+8 (8 being padding of 4 on each side.

changing

__global__ void phi(double *d_HSIImage,double *d_HSIImage2, int row, int col, double *d_phi_x, double *d_mean, double *d_std)
{
int X = blockDim.x * blockIdx.x + threadIdx.x;
int Y = blockDim.y * blockIdx.y + threadIdx.y;
int i,j;
double sum = 0;

if(Y>3 && X>3 && Y<row+4 && X<col+4)
{
    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            sum= sum + d_HSIImage2[i*col+j];
        }
    }

    d_mean[(Y-4)*col+X-4] = sum/81;
    double mean = sum/81;
    sum = 0;

    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            int index = i*col+j;
            sum= sum + (d_HSIImage2[index]-mean) * (d_HSIImage2[index]-mean);
        }
    }
    d_std[(Y-4)*col+X-4] = sqrt(sum/81);
}

to

__global__ void phi(double *d_HSIImage,double *d_HSIImage2, int row, int col, double *d_phi_x, double *d_mean, double *d_std)
{
int X = blockDim.x * blockIdx.x + threadIdx.x;
int Y = blockDim.y * blockIdx.y + threadIdx.y;
int i,j;
double sum = 0;

if(Y>3 && X>3 && Y<row+4 && X<col+4)
{
    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            sum= sum + d_HSIImage2[i*(col+8)+j];
        }
    }

    d_mean[(Y-4)*col+X-4] = sum/81;
    double mean = sum/81;
    sum = 0;

    for(i=Y-4;i<=Y+4;i++){
        for(j=X-4;j<=X+4;j++){
            int index = i*(col+8)+j;
            sum= sum + (d_HSIImage2[index]-mean) * (d_HSIImage2[index]-mean);
        }
    }
    d_std[(Y-4)*col+X-4] = sqrt(sum/81);
}

Should work, however, I have included a compilable example that I validated on a small sample, that should be easy to expand.

It is not optimized, but that wasn't part of your question. Optimization using shared memory would give a large performance boost.

#include <stdio.h>
#include <stdlib.h>
#include <iostream>

#include <cuda.h>

using namespace std;

__global__ void phi(double *img, int row, int col, double *d_mean){

  int X=blockDim.x*blockIdx.x+threadIdx.x+4;
  int Y=blockDim.y*blockIdx.y+threadIdx.y+4;
  double sum = 0;

  if(Y<row+4 && X<col+4){


    for(int i=-4; i<=4; ++i){
      for(int j=-4; j<=4; ++j){
        sum+=img[ (Y+j)*(col+8)+X+i];
      }
    }

    sum/=81;

    d_mean[(Y-4)*col+X-4]=sum;

  }

}

int main(int argc, char * argv[]) {

  int width=10, height=10;
  double *h_img=new double[(width+8)*(height+8)];

  for(int i=0; i<height+8; i++){
    for(int j=0; j<width+8; j++){
      h_img[i*(width+8)+j]=0.0;
    }
  }

  for(int i=0; i<height; i++){
    for(int j=0; j<width; j++){
      int index = (i+4)*(width+8)+j+4;
      h_img[index]=i*width+j;
    }
  }

  for(int i=0; i<height+8; i++){
    for(int j=0; j<width+8; j++){
      cout<<h_img[i*(width+8)+j]<<" ";
    }cout<<endl;
  }

  double *d_img;
  size_t size=sizeof(double)*(height+8)*(width*8);
  cudaMalloc(&d_img, size);
  cudaMemcpy(d_img, h_img, size, cudaMemcpyHostToDevice);

  size = sizeof(double)*height*width;
  double *d_avg;
  cudaMalloc(&d_avg, size);

  dim3 block(32, 32, 1);
  dim3 grid(width/32+1, height/32+1, 1);
  phi<<<grid, block>>>(d_img, height, width, d_avg);
  cudaDeviceSynchronize();
  double *h_avg=new double[width*height];
  cudaMemcpy(h_avg, d_avg, size, cudaMemcpyDeviceToHost);

  for(int i=0; i<height; i++){
    for(int j=0; j<width; j++){
      cout<<h_avg[i*width+j]<<" ";
    }cout<<endl;
  }

   return 0;
}

Answer 2

Here's my 2 cents regarding local mean and local std. You should check whether using matlab's optimized built-in functions (conv2 and stdfilt , with their gpu support) gives you better performance than a "simple" mex version. For example, to take the local mean, the fastest will be to use conv2 as follows:

local_mean_image=conv2(image,normalized_window,'same');

where in your case normalized_window=ones(9)./9^2;

For local std use stdfilt :

local_std_image = stdfilt(image, ones(9));

Both options are available for faster GPU performance, I use conv2 with Jacket routinely, and I saw the stdfilt supports gpuarray variables.

Answer 3

By observing the answers of @Christian Sarofeen and of @bla, I made some changes to my code and now I am able to find the mean exactly same as MATLAB. I posting this thinking that some one may use it in future(I am sending the image as is from MATLAB). Still finding standard deviation is little problem.

__global__ void phi(double *d_HSIImage,double *d_HSIImage2, int row, int col, double *d_phi_x, double *d_mean, double *d_std)
{
int X = blockDim.x * blockIdx.x + threadIdx.x;
int Y = blockDim.y * blockIdx.y + threadIdx.y;
int i,j;
double sum = 0;

if(Y>3 && X>3 && Y<row+4 && X<col+4)
{
    int index = (X-4)*row+Y-4;

    for(i=-4;i<=4;i++){
        for(j=-4;j<=4;j++){
            sum= sum + d_HSIImage2[(X+j)*(row+8)+(Y+i)];
        }
    }

    d_mean[index] = sum/81;
    double mean = 0;
    double temp_std[9] = {0} ;

    for(j=-4;j<=4;j++){
       sum = 0;
       for(i=-4;i<=4;i++){
           sum = sum + d_HSIImage2[(X+j)*(row+8)+(Y+i)];//vector mean
       }
       mean = sum/9;
       sum =0 ;
       for(i=-4;i<=4;i++){
          int index = (X+j)*(row+8)+(Y+i);
          sum= sum + (d_HSIImage2[index]-mean) * (d_HSIImage2[index]-mean);
        }
        temp_std[j+4] = (sqrt(sum/9));//vector std
    }
    sum =0 ;
    for(j=-4;j<=4;j++){
        sum = sum + temp_std[j+4];//mean of vectors
    }

    mean = sum/9;
    sum = 0 ;

    for(j=-4;j<=4;j++){
        sum = sum + (temp_std[j+4]-mean) * (temp_std[j+4]-mean);
    }

    d_std[index] = sqrt(sum);//std of vectors

    d_phi_x[index] = 1.0/(1.0+exp((d_mean[index]-d_HSIImage[index])/d_std[index]));
}
}