src/gromacs/gmxlib/gpu_utils/memtestG80_core.h

   1 /*
   2  * memtestG80_core.h
   3  * Public API for core memory test functions for MemtestG80
   4  * Includes functional and OO interfaces to GPU test functions.
   5  *
   6  * Author: Imran Haque, 2009
   7  * Copyright 2009, Stanford University
   8  *
   9  * This file is licensed under the terms of the LGPL. Please see
  10  * the COPYING file in the accompanying source distribution for
  11  * full license terms.
  12  *
  13  */
  14 #ifndef _MEMTESTG80_CORE_H_
  15 #define _MEMTESTG80_CORE_H_
  16
  17 #if defined (WINDOWS) || defined (WINNV)
  18     #include <windows.h>
  19 inline unsigned int getTimeMilliseconds(void)
  20 {
  21     return GetTickCount();
  22 }
  23     #include <windows.h>
  24     #define SLEEPMS(x) Sleep(x)
  25 #elif defined (LINUX) || defined (OSX)
  26     #include <sys/time.h>
  27 inline unsigned int getTimeMilliseconds(void)
  28 {
  29     struct timeval tv;
  30     gettimeofday(&tv, NULL);
  31     return tv.tv_sec*1000 + tv.tv_usec/1000;
  32 }
  33     #include <unistd.h>
  34     #define SLEEPMS(x) usleep(x*1000)
  35 #else
  36     #error Must #define LINUX, WINDOWS, WINNV, or OSX
  37 #endif
  38
  39 // By default the driver will spinwait when blocked on a kernel call
  40 // Use the SOFTWAIT macro to replace this with a thread sleep and occasional poll
  41 // limit expresses the max time we're willing to stay in the sleep loop - default = 15sec
  42 inline int _pollStatus(unsigned length = 1, unsigned limit = 15000)
  43 {
  44     //while (cudaStreamQuery(0) != cudaSuccess) {SLEEPMS(length);}
  45     unsigned startTime = getTimeMilliseconds();
  46     while (cudaStreamQuery(0) == cudaErrorNotReady)
  47     {
  48         if ((getTimeMilliseconds() - startTime) > limit)
  49         {
  50             return -1;
  51         }
  52         SLEEPMS(length);
  53     }
  54     return 0;
  55 }
  56 #define SOFTWAIT() if (_pollStatus() != 0) {return 0xFFFFFFFE; }              // -2
  57 #define SOFTWAIT_LIM(lim) if (_pollStatus(1, lim) != 0) {return 0xFFFFFFFE; } // -2
  58 //#define SOFTWAIT()
  59 //#define SOFTWAIT(delay) if (_pollStatus(delay) != 0) return -2;
  60 //#define SOFTWAIT(delay,limit) if (_pollStatus(delay,limit) != 0) return -2;
  61 //#define SOFTWAIT() while (cudaStreamQuery(0) != cudaSuccess) {SLEEPMS(1);}
  62 //#define SOFTWAIT(x) while (cudaStreamQuery(0) != cudaSuccess) {SLEEPMS(x);}
  63 //#define SOFTWAIT()
  64
  65 // Use this macro to check for kernel errors
  66 #define CHECK_LAUNCH_ERROR() if (cudaGetLastError() != cudaSuccess) {return 0xFFFFFFFF; /* -1 */}
  67
  68
  69 typedef unsigned int uint;
  70
  71 // OO interface to MemtestG80 functions
  72 class memtestState
  73 {
  74     protected:
  75         const uint nBlocks;
  76         const uint nThreads;
  77         uint       loopIters;
  78         uint       megsToTest;
  79         int        lcgPeriod;
  80         uint     * devTestMem;
  81         uint     * devTempMem;
  82         uint     * hostTempMem;
  83         bool       allocated;
  84     public:
  85         uint       initTime;
  86         memtestState() : nBlocks(1024), nThreads(512), loopIters(0), megsToTest(0), allocated(false), devTestMem(NULL), devTempMem(NULL), hostTempMem(NULL), initTime(0), lcgPeriod(1024) {};
  87         ~memtestState() {deallocate(); }
  88
  89         uint allocate(uint mbToTest);
  90         void deallocate();
  91         bool isAllocated() const {return allocated; }
  92         uint size() const {return megsToTest; }
  93         void setLCGPeriod(int period) {lcgPeriod = period; }
  94         int getLCGPeriod() const {return lcgPeriod; }
  95
  96         bool gpuMemoryBandwidth(double &bandwidth, uint mbToTest, uint iters = 5);
  97         bool gpuWriteConstant(const uint constant) const;
  98         bool gpuVerifyConstant(uint &errorCount, const uint constant) const;
  99         bool gpuShortLCG0(uint &errorCount, const uint repeats) const;
 100         bool gpuShortLCG0Shmem(uint &errorCount, const uint repeats) const;
 101         bool gpuMovingInversionsOnesZeros(uint &errorCount) const;
 102         bool gpuWalking8BitM86(uint &errorCount, const uint shift) const;
 103         bool gpuWalking8Bit(uint &errorCount, const bool ones, const uint shift) const;
 104         bool gpuMovingInversionsRandom(uint &errorCount) const;
 105         bool gpuWalking32Bit(uint &errorCount, const bool ones, const uint shift) const;
 106         bool gpuRandomBlocks(uint &errorCount, const uint seed) const;
 107         bool gpuModuloX(uint &errorCount, const uint shift, const uint pattern, const uint modulus, const uint overwriteIters) const;
 108 };
 109
 110 // Utility functions
 111 __host__ double gpuMemoryBandwidth(uint* src, uint* dst, uint mbToTest, uint iters);
 112 __host__ void gpuWriteConstant(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint constant);
 113 __host__ uint gpuVerifyConstant(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint constant, uint* blockErrorCount, uint* errorCounts);
 114
 115 __host__ void cpuWriteConstant(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint constant);
 116 __host__ uint cpuVerifyConstant(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint constant);
 117
 118 // Logic tests
 119 __host__ uint gpuShortLCG0(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint repeats, const int period, uint* blockErrorCounts, uint* errorCounts);
 120 __host__ uint gpuShortLCG0Shmem(const uint nBlocks, const uint nThreads, uint* base, uint N, const uint repeats, const int period, uint* blockErrorCounts, uint* errorCounts);
 121
 122 // Memtest86 Test 2: tseq=0,4
 123 __host__ uint gpuMovingInversionsOnesZeros(const uint nBlocks, const uint nThreads, uint* base, uint N, uint* blockErrorCounts, uint* errorCounts);
 124
 125 // Memtest86 Test 3: tseq=1
 126 __host__ uint gpuWalking8BitM86(const uint nBlocks, const uint nThreads, uint* base, uint N, uint shift, uint* blockErrorCounts, uint* errorCounts);
 127 __host__ uint cpuWalking8BitM86(const uint nBlocks, const uint nThreads, uint* base, uint N, uint shift);
 128 __host__ uint gpuWalking8Bit(const uint nBlocks, const uint nThreads, uint* base, uint N, bool ones, uint shift, uint* blockErrorCount, uint* errorCounts);
 129
 130 // Memtest86 Test 4: tseq=10
 131 __host__ uint gpuMovingInversionsRandom(const uint nBlocks, const uint nThreads, uint* base, uint N, uint* blockErrorCounts, uint* errorCounts);
 132
 133 // Memtest86 Test 6: tseq=2
 134 __host__ uint gpuWalking32Bit(const uint nBlocks, const uint nThreads, uint* base, uint N, bool ones, uint shift, uint* blockErrorCount, uint* errorCounts);
 135 //
 136 // Memtest86 Test 7: tseq=9
 137 __host__ uint gpuRandomBlocks(const uint nBlocks, const uint nThreads, uint* base, uint N, uint seed, uint* blockErrorCount, uint* errorCounts);
 138
 139 // Memtest86 Test 8: tseq=3 (M86 uses modulus = 20)
 140 __host__ uint gpuModuloX(const uint nBlocks, const uint nThreads, uint* base, const uint N, uint shift, uint pattern1, const uint modulus, const uint iters, uint* blockErrorCount, uint* errorCounts);
 141
 142 #endif