atomic.hpp 5.9 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202
  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. #pragma once
  44. #ifndef OPENCV_CUDEV_UTIL_ATOMIC_HPP
  45. #define OPENCV_CUDEV_UTIL_ATOMIC_HPP
  46. #include "../common.hpp"
  47. namespace cv { namespace cudev {
  48. //! @addtogroup cudev
  49. //! @{
  50. // atomicAdd
  51. __device__ __forceinline__ int atomicAdd(int* address, int val)
  52. {
  53. return ::atomicAdd(address, val);
  54. }
  55. __device__ __forceinline__ uint atomicAdd(uint* address, uint val)
  56. {
  57. return ::atomicAdd(address, val);
  58. }
  59. __device__ __forceinline__ float atomicAdd(float* address, float val)
  60. {
  61. #if CV_CUDEV_ARCH >= 200
  62. return ::atomicAdd(address, val);
  63. #else
  64. int* address_as_i = (int*) address;
  65. int old = *address_as_i, assumed;
  66. do {
  67. assumed = old;
  68. old = ::atomicCAS(address_as_i, assumed,
  69. __float_as_int(val + __int_as_float(assumed)));
  70. } while (assumed != old);
  71. return __int_as_float(old);
  72. #endif
  73. }
  74. __device__ static double atomicAdd(double* address, double val)
  75. {
  76. #if CV_CUDEV_ARCH >= 130
  77. unsigned long long int* address_as_ull = (unsigned long long int*) address;
  78. unsigned long long int old = *address_as_ull, assumed;
  79. do {
  80. assumed = old;
  81. old = ::atomicCAS(address_as_ull, assumed,
  82. __double_as_longlong(val + __longlong_as_double(assumed)));
  83. } while (assumed != old);
  84. return __longlong_as_double(old);
  85. #else
  86. CV_UNUSED(address);
  87. CV_UNUSED(val);
  88. return 0.0;
  89. #endif
  90. }
  91. // atomicMin
  92. __device__ __forceinline__ int atomicMin(int* address, int val)
  93. {
  94. return ::atomicMin(address, val);
  95. }
  96. __device__ __forceinline__ uint atomicMin(uint* address, uint val)
  97. {
  98. return ::atomicMin(address, val);
  99. }
  100. __device__ static float atomicMin(float* address, float val)
  101. {
  102. #if CV_CUDEV_ARCH >= 120
  103. int* address_as_i = (int*) address;
  104. int old = *address_as_i, assumed;
  105. do {
  106. assumed = old;
  107. old = ::atomicCAS(address_as_i, assumed,
  108. __float_as_int(::fminf(val, __int_as_float(assumed))));
  109. } while (assumed != old);
  110. return __int_as_float(old);
  111. #else
  112. CV_UNUSED(address);
  113. CV_UNUSED(val);
  114. return 0.0f;
  115. #endif
  116. }
  117. __device__ static double atomicMin(double* address, double val)
  118. {
  119. #if CV_CUDEV_ARCH >= 130
  120. unsigned long long int* address_as_ull = (unsigned long long int*) address;
  121. unsigned long long int old = *address_as_ull, assumed;
  122. do {
  123. assumed = old;
  124. old = ::atomicCAS(address_as_ull, assumed,
  125. __double_as_longlong(::fmin(val, __longlong_as_double(assumed))));
  126. } while (assumed != old);
  127. return __longlong_as_double(old);
  128. #else
  129. CV_UNUSED(address);
  130. CV_UNUSED(val);
  131. return 0.0;
  132. #endif
  133. }
  134. // atomicMax
  135. __device__ __forceinline__ int atomicMax(int* address, int val)
  136. {
  137. return ::atomicMax(address, val);
  138. }
  139. __device__ __forceinline__ uint atomicMax(uint* address, uint val)
  140. {
  141. return ::atomicMax(address, val);
  142. }
  143. __device__ static float atomicMax(float* address, float val)
  144. {
  145. #if CV_CUDEV_ARCH >= 120
  146. int* address_as_i = (int*) address;
  147. int old = *address_as_i, assumed;
  148. do {
  149. assumed = old;
  150. old = ::atomicCAS(address_as_i, assumed,
  151. __float_as_int(::fmaxf(val, __int_as_float(assumed))));
  152. } while (assumed != old);
  153. return __int_as_float(old);
  154. #else
  155. CV_UNUSED(address);
  156. CV_UNUSED(val);
  157. return 0.0f;
  158. #endif
  159. }
  160. __device__ static double atomicMax(double* address, double val)
  161. {
  162. #if CV_CUDEV_ARCH >= 130
  163. unsigned long long int* address_as_ull = (unsigned long long int*) address;
  164. unsigned long long int old = *address_as_ull, assumed;
  165. do {
  166. assumed = old;
  167. old = ::atomicCAS(address_as_ull, assumed,
  168. __double_as_longlong(::fmax(val, __longlong_as_double(assumed))));
  169. } while (assumed != old);
  170. return __longlong_as_double(old);
  171. #else
  172. CV_UNUSED(address);
  173. CV_UNUSED(val);
  174. return 0.0;
  175. #endif
  176. }
  177. //! @}
  178. }}
  179. #endif