00001 /*!@file CUDA/cuda_saliencyops.h CUDA/GPU optimized saliency calculations */ 00002 00003 // //////////////////////////////////////////////////////////////////// // 00004 // The iLab Neuromorphic Vision C++ Toolkit - Copyright (C) 2000-2005 // 00005 // by the University of Southern California (USC) and the iLab at USC. // 00006 // See http://iLab.usc.edu for information about this project. // 00007 // //////////////////////////////////////////////////////////////////// // 00008 // Major portions of the iLab Neuromorphic Vision Toolkit are protected // 00009 // under the U.S. patent ``Computation of Intrinsic Perceptual Saliency // 00010 // in Visual Environments, and Applications'' by Christof Koch and // 00011 // Laurent Itti, California Institute of Technology, 2001 (patent // 00012 // pending; application number 09/912,225 filed July 23, 2001; see // 00013 // http://pair.uspto.gov/cgi-bin/final/home.pl for current status). // 00014 // //////////////////////////////////////////////////////////////////// // 00015 // This file is part of the iLab Neuromorphic Vision C++ Toolkit. // 00016 // // 00017 // The iLab Neuromorphic Vision C++ Toolkit is free software; you can // 00018 // redistribute it and/or modify it under the terms of the GNU General // 00019 // Public License as published by the Free Software Foundation; either // 00020 // version 2 of the License, or (at your option) any later version. // 00021 // // 00022 // The iLab Neuromorphic Vision C++ Toolkit is distributed in the hope // 00023 // that it will be useful, but WITHOUT ANY WARRANTY; without even the // 00024 // implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR // 00025 // PURPOSE. See the GNU General Public License for more details. // 00026 // // 00027 // You should have received a copy of the GNU General Public License // 00028 // along with the iLab Neuromorphic Vision C++ Toolkit; if not, write // 00029 // to the Free Software Foundation, Inc., 59 Temple Place, Suite 330, // 00030 // Boston, MA 02111-1307 USA. // 00031 // //////////////////////////////////////////////////////////////////// // 00032 // 00033 // Primary maintainer for this file: 00034 // $HeadURL: svn://isvn.usc.edu/software/invt/trunk/saliency/src/CUDA/cuda_saliencyops.h $ 00035 // $Id: cuda_saliencyops.h 13227 2010-04-15 01:38:09Z dparks $ 00036 // 00037 00038 #ifndef CUDA_SALIENCYOPS_H_DEFINED 00039 #define CUDA_SALIENCYOPS_H_DEFINED 00040 00041 #include <cuda.h> 00042 #include "CUDA/cutil.h" 00043 #include "cudadefs.h" 00044 00045 __global__ void cuda_global_inertiaMap(float *dst, float s, float r_inv, int px, int py, int tile_width, int tile_height, int w, int h) 00046 { 00047 // Destination column 00048 const int dest_col = blockIdx.x*tile_width + threadIdx.x; 00049 // Destination row index 00050 const int dest_row = blockIdx.y*tile_height + threadIdx.y; 00051 // Destination index 00052 const int dest_idx = dest_row*w + dest_col; 00053 00054 if(dest_col < w && dest_row < h) 00055 { 00056 const int dsq = (px - dest_col)*(px - dest_col) + (py - dest_row)*(py - dest_row); 00057 dst[dest_idx] = s * exp(-dsq * r_inv); 00058 } 00059 } 00060 00061 00062 __global__ void cuda_global_inhibitionMap(float *dst, float factorOld, float factorNew, float radius, int px, int py, int tile_width, int tile_height, int w, int h) 00063 { 00064 // Destination column 00065 const int dest_col = blockIdx.x*tile_width + threadIdx.x; 00066 // Destination row index 00067 const int dest_row = blockIdx.y*tile_height + threadIdx.y; 00068 // Destination index 00069 const int dest_idx = dest_row*w + dest_col; 00070 00071 if(dest_col < w && dest_row < h) 00072 { 00073 const int dsq = (px - dest_col)*(px - dest_col) + (py - dest_row)*(py - dest_row); 00074 const float newval = 00075 dst[dest_idx] * factorOld 00076 + (factorNew * exp(- dsq / radius)); 00077 dst[dest_idx] = newval < 0.0F 00078 ? 0.0F 00079 : newval > 255.0F 00080 ? 255.0F 00081 : newval; 00082 } 00083 } 00084 00085 00086 00087 #endif
1.6.3