hostdepend/X86MAC64/util/X86MAC64/cuda/samples/0_Simple/simpleAtomicIntrinsics/simpleAtomicIntrinsics_kernel.cuh

   1 /*
   2  * Copyright 1993-2013 NVIDIA Corporation.  All rights reserved.
   3  *
   4  * Please refer to the NVIDIA end user license agreement (EULA) associated
   5  * with this source code for terms and conditions that govern your use of
   6  * this software. Any use, reproduction, disclosure, or distribution of
   7  * this software and related documentation outside the terms of the EULA
   8  * is strictly prohibited.
   9  *
  10  */
  11
  12 /* Simple kernel demonstrating atomic functions in device code. */
  13
  14 #ifndef _SIMPLEATOMICS_KERNEL_H_
  15 #define _SIMPLEATOMICS_KERNEL_H_
  16
  17 ////////////////////////////////////////////////////////////////////////////////
  18 //! Simple test kernel for atomic instructions
  19 //! @param g_idata  input data in global memory
  20 //! @param g_odata  output data in global memory
  21 ////////////////////////////////////////////////////////////////////////////////
  22 __global__ void
  23 testKernel(int *g_odata)
  24 {
  25     // access thread id
  26     const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
  27
  28     // Test various atomic instructions
  29
  30     // Arithmetic atomic instructions
  31
  32     // Atomic addition
  33     atomicAdd(&g_odata[0], 10);
  34
  35     // Atomic subtraction (final should be 0)
  36     atomicSub(&g_odata[1], 10);
  37
  38     // Atomic exchange
  39     atomicExch(&g_odata[2], tid);
  40
  41     // Atomic maximum
  42     atomicMax(&g_odata[3], tid);
  43
  44     // Atomic minimum
  45     atomicMin(&g_odata[4], tid);
  46
  47     // Atomic increment (modulo 17+1)
  48     atomicInc((unsigned int *)&g_odata[5], 17);
  49
  50     // Atomic decrement
  51     atomicDec((unsigned int *)&g_odata[6], 137);
  52
  53     // Atomic compare-and-swap
  54     atomicCAS(&g_odata[7], tid-1, tid);
  55
  56     // Bitwise atomic instructions
  57
  58     // Atomic AND
  59     atomicAnd(&g_odata[8], 2*tid+7);
  60
  61     // Atomic OR
  62     atomicOr(&g_odata[9], 1 << tid);
  63
  64     // Atomic XOR
  65     atomicXor(&g_odata[10], tid);
  66 }
  67
  68 #endif // #ifndef _SIMPLEATOMICS_KERNEL_H_