00001 /*!@file CUDA/cuda-drawops.h CUDA/GPU optimized drawing operations code */ 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_drawops.h $ 00035 // $Id: cuda_drawops.h 12962 2010-03-06 02:13:53Z irock $ 00036 // 00037 00038 #ifndef CUDA_DRAWOPS_H_DEFINED 00039 #define CUDA_DRAWOPS_H_DEFINED 00040 00041 #include <cuda.h> 00042 #include "CUDA/cutil.h" 00043 #include "cudadefs.h" 00044 00045 // Fill in rectangle with intensity value 00046 __global__ void cuda_global_drawFilledRect(float *dst, int top, int left, int bottom, int right, const float intensity, const int w, const int h, const int tile_width, const int tile_height) 00047 { 00048 const int x_pos = left + IMUL(blockIdx.x,tile_width) + threadIdx.x; 00049 const int y_pos = top + IMUL(blockIdx.y,tile_height) + threadIdx.y; 00050 const int idx = IMUL(y_pos,w) + x_pos; 00051 int x_max = MIN(w,right+1); 00052 int y_max = MIN(h,bottom+1); 00053 if(x_pos < x_max && y_pos < y_max) 00054 { 00055 dst[idx] = intensity; 00056 } 00057 } 00058 00059 // Fill in rectangle with rgb color value 00060 __global__ void cuda_global_drawFilledRectRGB(float3_t *dst, int top, int left, int bottom, int right, const float c1, const float c2, const float c3, const int w, const int h, const int tile_width, const int tile_height) 00061 { 00062 const int x_pos = left + IMUL(blockIdx.x,tile_width) + threadIdx.x; 00063 const int y_pos = top + IMUL(blockIdx.y,tile_height) + threadIdx.y; 00064 const int idx = IMUL(y_pos,w) + x_pos; 00065 int x_max = MIN(w,right+1); 00066 int y_max = MIN(h,bottom+1); 00067 if(x_pos < x_max && y_pos < y_max) 00068 { 00069 dst[idx].p[0] = c1; 00070 dst[idx].p[1] = c2; 00071 dst[idx].p[2] = c3; 00072 } 00073 } 00074 00075 #endif
1.6.3