Kokkos Core Kernels Package: Kokkos

00001 /*
00002 //@HEADER
00003 // ************************************************************************
00004 //
00005 //                             Kokkos
00006 //         Manycore Performance-Portable Multidimensional Arrays
00007 //
00008 //              Copyright (2012) Sandia Corporation
00009 //
00010 // Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
00011 // the U.S. Government retains certain rights in this software.
00012 //
00013 // Redistribution and use in source and binary forms, with or without
00014 // modification, are permitted provided that the following conditions are
00015 // met:
00016 //
00017 // 1. Redistributions of source code must retain the above copyright
00018 // notice, this list of conditions and the following disclaimer.
00019 //
00020 // 2. Redistributions in binary form must reproduce the above copyright
00021 // notice, this list of conditions and the following disclaimer in the
00022 // documentation and/or other materials provided with the distribution.
00023 //
00024 // 3. Neither the name of the Corporation nor the names of the
00025 // contributors may be used to endorse or promote products derived from
00026 // this software without specific prior written permission.
00027 //
00028 // THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
00029 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
00030 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
00031 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
00032 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
00033 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
00034 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
00035 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
00036 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
00037 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
00038 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
00039 //
00040 // Questions?  Contact  H. Carter Edwards (hcedwar@sandia.gov)
00041 //
00042 // ************************************************************************
00043 //@HEADER
00044 */
00045 
00046 #ifndef KOKKOS_MACROS_HPP
00047 #define KOKKOS_MACROS_HPP
00048 
00049 //----------------------------------------------------------------------------
00065 #ifndef KOKKOS_DONT_INCLUDE_CORE_CONFIG_H
00066 #include <KokkosCore_config.h>
00067 #endif
00068 
00069 //----------------------------------------------------------------------------
00099 //----------------------------------------------------------------------------
00100 
00101 #if defined( KOKKOS_HAVE_CUDA ) && defined( __CUDACC__ )
00102 
00103 /*  Compiling with a CUDA compiler.
00104  *
00105  *  Include <cuda.h> to pick up the CUDA_VERSION macro defined as:
00106  *    CUDA_VERSION = ( MAJOR_VERSION * 1000 ) + ( MINOR_VERSION * 10 )
00107  *
00108  *  When generating device code the __CUDA_ARCH__ macro is defined as:
00109  *    __CUDA_ARCH__ = ( MAJOR_CAPABILITY * 100 ) + ( MINOR_CAPABILITY * 10 )
00110  */
00111 
00112 #include <cuda_runtime.h>
00113 #include <cuda.h>
00114 
00115 #if ! defined( CUDA_VERSION )
00116 #error "#include <cuda.h> did not define CUDA_VERSION"
00117 #endif
00118 
00119 #if ( CUDA_VERSION < 4010 )
00120 #error "Cuda version 4.1 or greater required"
00121 #endif
00122 
00123 #if defined( __CUDA_ARCH__ ) && ( __CUDA_ARCH__ < 200 )
00124 /*  Compiling with CUDA compiler for device code. */
00125 #error "Cuda device capability >= 2.0 is required"
00126 #endif
00127 
00128 #endif /* #if defined( KOKKOS_HAVE_CUDA ) && defined( __CUDACC__ ) */
00129 
00130 /*--------------------------------------------------------------------------*/
00131 /* Language info: C++, CUDA, OPENMP */
00132 
00133 #if defined( __CUDA_ARCH__ )
00134   // Compiling Cuda code to 'ptx'
00135 
00136   #define KOKKOS_FORCEINLINE_FUNCTION  __device__  __host__  __forceinline__
00137   #define KOKKOS_INLINE_FUNCTION       __device__  __host__  inline
00138   #define KOKKOS_FUNCTION              __device__  __host__
00139 
00140 #endif /* #if defined( __CUDA_ARCH__ ) */
00141 
00142 #if defined( _OPENMP )
00143 
00144   /*  Compiling with OpenMP.
00145    *  The value of _OPENMP is an integer value YYYYMM
00146    *  where YYYY and MM are the year and month designation
00147    *  of the supported OpenMP API version.
00148    */
00149 
00150 #endif /* #if defined( _OPENMP ) */
00151 
00152 /*--------------------------------------------------------------------------*/
00153 /* Mapping compiler built-ins to KOKKOS_COMPILER_*** macros */
00154 
00155 #if defined( __NVCC__ )
00156   // NVIDIA compiler is being used.
00157   // Code is parsed and separated into host and device code.
00158   // Host code is compiled again with another compiler.
00159   // Device code is compile to 'ptx'.
00160   #define KOKKOS_COMPILER_NVCC __NVCC__
00161 
00162   #if defined( KOKKOS_HAVE_CXX11 )
00163     // CUDA supports (inofficially) C++11 in device code starting with 
00164     // version 6.5. This includes auto type and device code internal
00165     // lambdas.
00166     #if ( CUDA_VERSION < 6050 )
00167       #error "NVCC does not support C++11"
00168     #endif
00169   #endif
00170 #else
00171   #if defined( KOKKOS_HAVE_CXX11 )
00172     // CUDA (including version 6.5) does not support giving lambdas as
00173     // arguments to global functions. Thus its not currently possible
00174     // to dispatch lambdas from the host.
00175     #define KOKKOS_HAVE_CXX11_DISPATCH_LAMBDA 1
00176   #endif
00177 #endif /* #if defined( __NVCC__ ) */
00178 
00179 
00180 #if ! defined( __CUDA_ARCH__ ) /* Not compiling Cuda code to 'ptx'. */
00181 
00182 #if defined( __INTEL_COMPILER )
00183   #define KOKKOS_COMPILER_INTEL __INTEL_COMPILER
00184 #elif defined( __ICC )
00185   // Old define
00186   #define KOKKOS_COMPILER_INTEL __ICC
00187 #elif defined( __ECC ) 
00188   // Very old define
00189   #define KOKKOS_COMPILER_INTEL __ECC
00190 #endif
00191 
00192 #if defined( _CRAYC )
00193   #define KOKKOS_COMPILER_CRAYC _CRAYC
00194 #endif
00195 
00196 #if defined( __IBMCPP__ )
00197   // IBM C++
00198   #define KOKKOS_COMPILER_IBM __IBMCPP__
00199 #elif defined( __IBMC__ )
00200   #define KOKKOS_COMPILER_IBM __IBMC__
00201 #endif
00202 
00203 #if defined( __APPLE_CC__ )
00204   #define KOKKOS_COMPILER_APPLECC __APPLE_CC__
00205 #endif
00206 
00207 #if defined( __clang__ )
00208   #define KOKKOS_COMPILER_CLANG __clang_major__*100+__clang_minor__*10+__clang_patchlevel__
00209 #endif
00210 
00211 #if ! defined( __clang__ ) && ! defined( KOKKOS_COMPILER_INTEL ) &&defined( __GNUC__ )
00212   #define KOKKOS_COMPILER_GNU __GNUC__*100+__GNUC_MINOR__*10+__GNUC_PATCHLEVEL__
00213 #endif
00214 
00215 #if defined( __PGIC__ ) && ! defined( __GNUC__ )
00216   #define KOKKOS_COMPILER_PGI __PGIC__*100+__PGIC_MINOR__*10+__PGIC_PATCHLEVEL__
00217 #endif
00218 
00219 #endif /* #if ! defined( __CUDA_ARCH__ ) */
00220 
00221 /*--------------------------------------------------------------------------*/
00222 /*--------------------------------------------------------------------------*/
00223 /* Intel compiler macros */
00224 
00225 #if defined( KOKKOS_COMPILER_INTEL )
00226 
00227   #define KOKKOS_HAVE_PRAGMA_UNROLL 1
00228   #define KOKKOS_HAVE_PRAGMA_IVDEP 1
00229   #define KOKKOS_HAVE_PRAGMA_LOOPCOUNT 1
00230   #define KOKKOS_HAVE_PRAGMA_VECTOR 1
00231   #define KOKKOS_HAVE_PRAGMA_SIMD 1
00232 
00233   #if ( 1200 <= KOKKOS_COMPILER_INTEL ) && ! defined( KOKKOS_ENABLE_ASM )
00234     #define KOKKOS_ENABLE_ASM 1
00235   #endif
00236 
00237   #define KOKKOS_FORCEINLINE_FUNCTION  __forceinline
00238 
00239   #if defined( __MIC__ )
00240     // Compiling for Xeon Phi
00241   #endif
00242 
00243 #endif
00244 
00245 /*--------------------------------------------------------------------------*/
00246 /* Cray compiler macros */
00247 
00248 #if defined( KOKKOS_COMPILER_CRAYC )
00249 
00250 
00251 #endif
00252 
00253 /*--------------------------------------------------------------------------*/
00254 /* IBM Compiler macros */
00255 
00256 #if defined( KOKKOS_COMPILER_IBM )
00257 
00258   #define KOKKOS_HAVE_PRAGMA_UNROLL 1
00259   //#define KOKKOS_HAVE_PRAGMA_IVDEP 1
00260   //#define KOKKOS_HAVE_PRAGMA_LOOPCOUNT 1
00261   //#define KOKKOS_HAVE_PRAGMA_VECTOR 1
00262   //#define KOKKOS_HAVE_PRAGMA_SIMD 1
00263 
00264 #endif
00265 
00266 /*--------------------------------------------------------------------------*/
00267 
00268 #if defined( KOKKOS_COMPILER_CLANG )
00269 
00270   //#define KOKKOS_HAVE_PRAGMA_UNROLL 1
00271   //#define KOKKOS_HAVE_PRAGMA_IVDEP 1
00272   //#define KOKKOS_HAVE_PRAGMA_LOOPCOUNT 1
00273   //#define KOKKOS_HAVE_PRAGMA_VECTOR 1
00274   //#define KOKKOS_HAVE_PRAGMA_SIMD 1
00275 
00276   #define KOKKOS_FORCEINLINE_FUNCTION  inline __attribute__((always_inline))
00277 
00278 #endif
00279 
00280 /*--------------------------------------------------------------------------*/
00281 
00282 #if defined( KOKKOS_COMPILER_GNU ) 
00283 
00284   //#define KOKKOS_HAVE_PRAGMA_UNROLL 1
00285   //#define KOKKOS_HAVE_PRAGMA_IVDEP 1
00286   //#define KOKKOS_HAVE_PRAGMA_LOOPCOUNT 1
00287   //#define KOKKOS_HAVE_PRAGMA_VECTOR 1
00288   //#define KOKKOS_HAVE_PRAGMA_SIMD 1
00289 
00290   #define KOKKOS_FORCEINLINE_FUNCTION inline __attribute__((always_inline))
00291 
00292   #if ! defined( KOKKOS_ENABLE_ASM ) && \
00293       ! ( defined( __powerpc) || \
00294           defined(__powerpc__) || \
00295           defined(__powerpc64__) || \
00296           defined(__POWERPC__) || \
00297           defined(__ppc__) || \
00298           defined(__ppc64__) )
00299     #define KOKKOS_ENABLE_ASM 1
00300   #endif
00301 
00302 #endif
00303 
00304 /*--------------------------------------------------------------------------*/
00305 
00306 #if defined( KOKKOS_COMPILER_PGI )
00307 
00308   #define KOKKOS_HAVE_PRAGMA_UNROLL 1
00309   #define KOKKOS_HAVE_PRAGMA_IVDEP 1
00310   //#define KOKKOS_HAVE_PRAGMA_LOOPCOUNT 1
00311   #define KOKKOS_HAVE_PRAGMA_VECTOR 1
00312   //#define KOKKOS_HAVE_PRAGMA_SIMD 1
00313 
00314 #endif
00315 
00316 /*--------------------------------------------------------------------------*/
00317 
00318 #if defined( KOKKOS_COMPILER_NVCC )
00319 
00320   #if defined(__CUDA_ARCH__ )
00321     #define KOKKOS_HAVE_PRAGMA_UNROLL 1
00322   #endif
00323 
00324 #endif
00325 
00326 /*--------------------------------------------------------------------------*/
00327 /* Select compiler dependent interface for atomics */
00328 
00329 #if ! defined( KOKKOS_ATOMICS_USE_CUDA ) || \
00330     ! defined( KOKKOS_ATOMICS_USE_GNU ) || \
00331     ! defined( KOKKOS_ATOMICS_USE_INTEL ) || \
00332     ! defined( KOKKOS_ATOMICS_USE_OPENMP31 )
00333 
00334 /* Atomic selection is not pre-defined, choose from language and compiler. */
00335 
00336 #if defined( __CUDA_ARCH__ )
00337 
00338   #define KOKKOS_ATOMICS_USE_CUDA
00339 
00340 #elif defined( KOKKOS_COMPILER_GNU ) || defined( KOKKOS_COMPILER_CLANG )
00341 
00342   #define KOKKOS_ATOMICS_USE_GNU
00343 
00344 #elif defined( KOKKOS_COMPILER_INTEL ) || defined( KOKKOS_COMPILER_CRAYC )
00345 
00346   #define KOKKOS_ATOMICS_USE_INTEL
00347 
00348 #elif defined( _OPENMP ) && ( 201107 <= _OPENMP )
00349 
00350   #define KOKKOS_ATOMICS_USE_OMP31
00351 
00352 #else
00353 
00354   #error "Compiler does not support atomic operations"
00355 
00356 #endif
00357 
00358 #endif
00359 
00360 //----------------------------------------------------------------------------
00363 #if ! defined( KOKKOS_FORCEINLINE_FUNCTION )
00364 #define KOKKOS_FORCEINLINE_FUNCTION  inline
00365 #endif
00366 
00367 #if ! defined( KOKKOS_INLINE_FUNCTION )
00368 #define KOKKOS_INLINE_FUNCTION  inline
00369 #endif
00370 
00371 #if ! defined( KOKKOS_FUNCTION )
00372 #define KOKKOS_FUNCTION 
00373 #endif
00374 
00375 //----------------------------------------------------------------------------
00380 #if 1 < ( ( defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_CUDA ) ? 1 : 0 ) + \
00381           ( defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_OPENMP ) ? 1 : 0 ) + \
00382           ( defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_THREADS ) ? 1 : 0 ) + \
00383           ( defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_SERIAL ) ? 1 : 0 ) )
00384 
00385 #error "More than one KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_* specified" ;
00386 
00387 #endif
00388 
00392 #if   defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_CUDA )
00393 #elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_OPENMP )
00394 #elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_THREADS )
00395 #elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_SERIAL )
00396 #elif defined ( KOKKOS_HAVE_CUDA )
00397 #define KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_CUDA
00398 #elif defined ( KOKKOS_HAVE_OPENMP )
00399 #define KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_OPENMP
00400 #elif defined ( KOKKOS_HAVE_PTHREAD )
00401 #define KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_THREADS
00402 #else
00403 #define KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_SERIAL
00404 #endif
00405 
00406 //----------------------------------------------------------------------------
00409 #if defined( __CUDACC__ ) && defined( __CUDA_ARCH__ )
00410 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA
00411 #else
00412 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST
00413 #endif
00414 
00415 //----------------------------------------------------------------------------
00416 //----------------------------------------------------------------------------
00417 
00418 #endif /* #ifndef KOKKOS_MACROS_HPP */
00419