libhoomd/utils/TextureTools.h

   1 /*
   2 Highly Optimized Object-oriented Many-particle Dynamics -- Blue Edition
   3 (HOOMD-blue) Open Source Software License Copyright 2009-2014 The Regents of
   4 the University of Michigan All rights reserved.
   5
   6 HOOMD-blue may contain modifications ("Contributions") provided, and to which
   7 copyright is held, by various Contributors who have granted The Regents of the
   8 University of Michigan the right to modify and/or distribute such Contributions.
   9
  10 You may redistribute, use, and create derivate works of HOOMD-blue, in source
  11 and binary forms, provided you abide by the following conditions:
  12
  13 * Redistributions of source code must retain the above copyright notice, this
  14 list of conditions, and the following disclaimer both in the code and
  15 prominently in any materials provided with the distribution.
  16
  17 * Redistributions in binary form must reproduce the above copyright notice, this
  18 list of conditions, and the following disclaimer in the documentation and/or
  19 other materials provided with the distribution.
  20
  21 * All publications and presentations based on HOOMD-blue, including any reports
  22 or published results obtained, in whole or in part, with HOOMD-blue, will
  23 acknowledge its use according to the terms posted at the time of submission on:
  24 http://codeblue.umich.edu/hoomd-blue/citations.html
  25
  26 * Any electronic documents citing HOOMD-Blue will link to the HOOMD-Blue website:
  27 http://codeblue.umich.edu/hoomd-blue/
  28
  29 * Apart from the above required attributions, neither the name of the copyright
  30 holder nor the names of HOOMD-blue's contributors may be used to endorse or
  31 promote products derived from this software without specific prior written
  32 permission.
  33
  34 Disclaimer
  35
  36 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDER AND CONTRIBUTORS ``AS IS'' AND
  37 ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
  38 WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE, AND/OR ANY
  39 WARRANTIES THAT THIS SOFTWARE IS FREE OF INFRINGEMENT ARE DISCLAIMED.
  40
  41 IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
  42 INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
  43 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
  44 DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
  45 LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
  46 OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF
  47 ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  48 */
  49
  50 // Maintainer: joaander
  51
  52 #ifndef __HOOMD_TEXTURE_TOOLS_H__
  53 #define __HOOMD_TEXTURE_TOOLS_H__
  54
  55 /*! \file TextureTools.h
  56     \brief Utilities for working with textures
  57
  58     TextureTools.h exists to aid in defining Scalar textures which may be either float or double. It aims to simplify
  59     code that reads from these textures so that the amount of conditional code is simplified to be entirely within
  60     this header.
  61
  62     Planning for the future (__ldg), the fetch methods will also take in a pointer to the memory. That way, the initial
  63     work done to convert the texture loads over to the single/double will also make it easy to change over to __ldg
  64     in a single spot.
  65 */
  66
  67 #include "HOOMDMath.h"
  68
  69 #ifdef NVCC
  70
  71 //! Fetch an unsigned int from texture memory.
  72 /*! This function should be called whenever a CUDA kernel wants to retrieve a
  73     unsigned int value from texture memory.
  74
  75     \param ptr Pointer to bound memory
  76     \param tex_ref Texture in which the desired values are stored.
  77     \param ii Index at which to look.
  78 */
  79 __device__ inline unsigned int texFetchUint(const unsigned int *ptr, texture<unsigned int, 1> tex_ref, unsigned int ii)
  80     {
  81     #if __CUDA_ARCH__ >= 350
  82     return __ldg(ptr+ii);
  83     #else
  84     return tex1Dfetch(tex_ref, ii);
  85     #endif
  86     }
  87
  88 #ifdef SINGLE_PRECISION
  89
  90 typedef texture<Scalar, 1, cudaReadModeElementType> scalar_tex_t;
  91 typedef texture<Scalar2, 1, cudaReadModeElementType> scalar2_tex_t;
  92 typedef texture<Scalar4, 1, cudaReadModeElementType> scalar4_tex_t;
  93
  94 //! Fetch a Scalar value from texture memory.
  95 /*! This function should be called whenever a CUDA kernel wants to retrieve a
  96     Scalar value from texture memory.
  97
  98     \param ptr Pointer to bound memory
  99     \param tex_ref Texture in which the desired values are stored.
 100     \param ii Index at which to look.
 101 */
 102 __device__ inline Scalar texFetchScalar(const Scalar *ptr, texture<Scalar, 1> tex_ref, unsigned int ii)
 103     {
 104     #if __CUDA_ARCH__ >= 350
 105     return __ldg(ptr+ii);
 106     #else
 107     return tex1Dfetch(tex_ref, ii);
 108     #endif
 109     }
 110
 111 //! Fetch a Scalar2 value from texture memory.
 112 /*! This function should be called whenever a CUDA kernel wants to retrieve a
 113     Scalar2 value from texture memory.
 114
 115     \param ptr Pointer to bound memory
 116     \param tex_ref Texture in which the desired values are stored.
 117     \param ii Index at which to look.
 118 */
 119 __device__ inline Scalar2 texFetchScalar2(const Scalar2 *ptr, texture<Scalar2, 1> tex_ref, unsigned int ii)
 120     {
 121     #if __CUDA_ARCH__ >= 350
 122     return __ldg(ptr+ii);
 123     #else
 124     return tex1Dfetch(tex_ref, ii);
 125     #endif
 126     }
 127
 128 //! Fetch a Scalar4 value from texture memory.
 129 /*! This function should called whenever a CUDA kernel wants to retrieve a
 130     Scalar4 value from texture memory.
 131
 132     \param ptr Pointer to bound memory
 133     \param tex_ref Texture in which the desired values are stored.
 134     \param ii Index at which to look.
 135 */
 136 __device__ inline Scalar4 texFetchScalar4(const Scalar4 *ptr, texture<Scalar4, 1> tex_ref, unsigned int ii)
 137     {
 138     #if __CUDA_ARCH__ >= 350
 139     return __ldg(ptr+ii);
 140     #else
 141     return tex1Dfetch(tex_ref, ii);
 142     #endif
 143     }
 144
 145 #else
 146 typedef texture<int2, 1, cudaReadModeElementType> scalar_tex_t;
 147 typedef texture<int4, 1, cudaReadModeElementType> scalar2_tex_t;
 148 typedef texture<int4, 1, cudaReadModeElementType> scalar4_tex_t;
 149
 150 //! Fetch a Scalar value from texture memory.
 151 /*! This function should be called whenever a CUDA kernel wants to retrieve a
 152     Scalar value from texture memory.
 153
 154     \param ptr Pointer to bound memory
 155     \param tex_ref Texture in which the desired values are stored.
 156     \param ii Index at which to look.
 157 */
 158 __device__ inline Scalar texFetchScalar(const Scalar *ptr, texture<int2, 1> tex_ref, unsigned int ii)
 159     {
 160     #if __CUDA_ARCH__ >= 350
 161     return __ldg(ptr+ii);
 162     #else
 163     int2 val = tex1Dfetch(tex_ref, ii);
 164     return Scalar(__hiloint2double(val.y, val.x));
 165     #endif
 166     }
 167
 168 //! Fetch a Scalar2 value from texture memory.
 169 /*! This function should be called whenever a CUDA kernel wants to retrieve a
 170     Scalar2 value from texture memory.
 171
 172     \param ptr Pointer to bound memory
 173     \param tex_ref Texture in which the desired values are stored.
 174     \param ii Index at which to look.
 175 */
 176 __device__ inline Scalar2 texFetchScalar2(const Scalar2* ptr, texture<int4, 1> tex_ref, unsigned int ii)
 177     {
 178     #if __CUDA_ARCH__ >= 350
 179     return __ldg(ptr+ii);
 180     #else
 181     int4 val = tex1Dfetch(tex_ref, ii);
 182     return make_scalar2(__hiloint2double(val.y, val.x),
 183                         __hiloint2double(val.w, val.z));
 184     #endif
 185     }
 186
 187 //! Fetch a Scalar4 value from texture memory.
 188 /*! This function should be called whenever a CUDA kernel wants to retrieve a
 189     Scalar4 value from texture memory.
 190
 191     \param ptr Pointer to bound memory
 192     \param tex_ref Texture in which the desired values are stored.
 193     \param ii Index at which to look.
 194 */
 195 __device__ inline Scalar4 texFetchScalar4(const Scalar4 *ptr, texture<int4, 1> tex_ref, unsigned int ii)
 196     {
 197     unsigned int idx = 2*ii;
 198     #if __CUDA_ARCH__ >= 350
 199     int4 part1 = __ldg(((int4 *)ptr)+idx);;
 200     int4 part2 = __ldg(((int4 *)ptr)+idx+1);;
 201     #else
 202     int4 part1 = tex1Dfetch(tex_ref, idx);
 203     int4 part2 = tex1Dfetch(tex_ref, idx+1);
 204     #endif
 205     return make_scalar4(__hiloint2double(part1.y, part1.x),
 206                         __hiloint2double(part1.w, part1.z),
 207                         __hiloint2double(part2.y, part2.x),
 208                         __hiloint2double(part2.w, part2.z));
 209     }
 210 #endif
 211 #endif
 212
 213
 214
 215 #endif // __HOOMD_MATH_H__