1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
8 //
9 //
10 // License Agreement
11 // For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
17 //
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
20 //
21 // * Redistribution's of source code must retain the above copyright notice,
22 // this list of conditions and the following disclaimer.
23 //
24 // * Redistribution's in binary form must reproduce the above copyright notice,
25 // this list of conditions and the following disclaimer in the documentation
26 // and/or other materials provided with the distribution.
27 //
28 // * The name of the copyright holders may not be used to endorse or promote products
29 // derived from this software without specific prior written permission.
30 //
31 // This software is provided by the copyright holders and contributors "as is" and
32 // any express or implied warranties, including, but not limited to, the implied
33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
34 // In no event shall the Intel Corporation or contributors be liable for any direct,
35 // indirect, incidental, special, exemplary, or consequential damages
36 // (including, but not limited to, procurement of substitute goods or services;
37 // loss of use, data, or profits; or business interruption) however caused
38 // and on any theory of liability, whether in contract, strict liability,
39 // or tort (including negligence or otherwise) arising in any way out of
40 // the use of this software, even if advised of the possibility of such damage.
41 //
42 //M*/
43
44 #pragma once
45
46 #ifndef __OPENCV_CUDEV_UTIL_ATOMIC_HPP__
47 #define __OPENCV_CUDEV_UTIL_ATOMIC_HPP__
48
49 #include "../common.hpp"
50
51 namespace cv { namespace cudev {
52
53 //! @addtogroup cudev
54 //! @{
55
56 // atomicAdd
57
atomicAdd(int * address,int val)58 __device__ __forceinline__ int atomicAdd(int* address, int val)
59 {
60 return ::atomicAdd(address, val);
61 }
62
atomicAdd(uint * address,uint val)63 __device__ __forceinline__ uint atomicAdd(uint* address, uint val)
64 {
65 return ::atomicAdd(address, val);
66 }
67
atomicAdd(float * address,float val)68 __device__ __forceinline__ float atomicAdd(float* address, float val)
69 {
70 #if CV_CUDEV_ARCH >= 200
71 return ::atomicAdd(address, val);
72 #else
73 int* address_as_i = (int*) address;
74 int old = *address_as_i, assumed;
75 do {
76 assumed = old;
77 old = ::atomicCAS(address_as_i, assumed,
78 __float_as_int(val + __int_as_float(assumed)));
79 } while (assumed != old);
80 return __int_as_float(old);
81 #endif
82 }
83
atomicAdd(double * address,double val)84 __device__ static double atomicAdd(double* address, double val)
85 {
86 #if CV_CUDEV_ARCH >= 130
87 unsigned long long int* address_as_ull = (unsigned long long int*) address;
88 unsigned long long int old = *address_as_ull, assumed;
89 do {
90 assumed = old;
91 old = ::atomicCAS(address_as_ull, assumed,
92 __double_as_longlong(val + __longlong_as_double(assumed)));
93 } while (assumed != old);
94 return __longlong_as_double(old);
95 #else
96 (void) address;
97 (void) val;
98 return 0.0;
99 #endif
100 }
101
102 // atomicMin
103
atomicMin(int * address,int val)104 __device__ __forceinline__ int atomicMin(int* address, int val)
105 {
106 return ::atomicMin(address, val);
107 }
108
atomicMin(uint * address,uint val)109 __device__ __forceinline__ uint atomicMin(uint* address, uint val)
110 {
111 return ::atomicMin(address, val);
112 }
113
atomicMin(float * address,float val)114 __device__ static float atomicMin(float* address, float val)
115 {
116 #if CV_CUDEV_ARCH >= 120
117 int* address_as_i = (int*) address;
118 int old = *address_as_i, assumed;
119 do {
120 assumed = old;
121 old = ::atomicCAS(address_as_i, assumed,
122 __float_as_int(::fminf(val, __int_as_float(assumed))));
123 } while (assumed != old);
124 return __int_as_float(old);
125 #else
126 (void) address;
127 (void) val;
128 return 0.0f;
129 #endif
130 }
131
atomicMin(double * address,double val)132 __device__ static double atomicMin(double* address, double val)
133 {
134 #if CV_CUDEV_ARCH >= 130
135 unsigned long long int* address_as_ull = (unsigned long long int*) address;
136 unsigned long long int old = *address_as_ull, assumed;
137 do {
138 assumed = old;
139 old = ::atomicCAS(address_as_ull, assumed,
140 __double_as_longlong(::fmin(val, __longlong_as_double(assumed))));
141 } while (assumed != old);
142 return __longlong_as_double(old);
143 #else
144 (void) address;
145 (void) val;
146 return 0.0;
147 #endif
148 }
149
150 // atomicMax
151
atomicMax(int * address,int val)152 __device__ __forceinline__ int atomicMax(int* address, int val)
153 {
154 return ::atomicMax(address, val);
155 }
156
atomicMax(uint * address,uint val)157 __device__ __forceinline__ uint atomicMax(uint* address, uint val)
158 {
159 return ::atomicMax(address, val);
160 }
161
atomicMax(float * address,float val)162 __device__ static float atomicMax(float* address, float val)
163 {
164 #if CV_CUDEV_ARCH >= 120
165 int* address_as_i = (int*) address;
166 int old = *address_as_i, assumed;
167 do {
168 assumed = old;
169 old = ::atomicCAS(address_as_i, assumed,
170 __float_as_int(::fmaxf(val, __int_as_float(assumed))));
171 } while (assumed != old);
172 return __int_as_float(old);
173 #else
174 (void) address;
175 (void) val;
176 return 0.0f;
177 #endif
178 }
179
atomicMax(double * address,double val)180 __device__ static double atomicMax(double* address, double val)
181 {
182 #if CV_CUDEV_ARCH >= 130
183 unsigned long long int* address_as_ull = (unsigned long long int*) address;
184 unsigned long long int old = *address_as_ull, assumed;
185 do {
186 assumed = old;
187 old = ::atomicCAS(address_as_ull, assumed,
188 __double_as_longlong(::fmax(val, __longlong_as_double(assumed))));
189 } while (assumed != old);
190 return __longlong_as_double(old);
191 #else
192 (void) address;
193 (void) val;
194 return 0.0;
195 #endif
196 }
197
198 //! @}
199
200 }}
201
202 #endif
203