3 * Public API for core memory test functions for MemtestG80
4 * Includes functional and OO interfaces to GPU test functions.
6 * Author: Imran Haque, 2009
7 * Copyright 2009, Stanford University
9 * This file is licensed under the terms of the LGPL. Please see
10 * the COPYING file in the accompanying source distribution for
14 #ifndef _MEMTESTG80_CORE_H_
15 #define _MEMTESTG80_CORE_H_
17 #if defined (WINDOWS) || defined (WINNV)
19 inline unsigned int getTimeMilliseconds(void)
21 return GetTickCount();
24 #define SLEEPMS(x) Sleep(x)
25 #elif defined (LINUX) || defined (OSX)
27 inline unsigned int getTimeMilliseconds(void)
30 gettimeofday(&tv, NULL);
31 return tv.tv_sec*1000 + tv.tv_usec/1000;
34 #define SLEEPMS(x) usleep(x*1000)
36 #error Must #define LINUX, WINDOWS, WINNV, or OSX
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)
44 //while (cudaStreamQuery(0) != cudaSuccess) {SLEEPMS(length);}
45 unsigned startTime = getTimeMilliseconds();
46 while (cudaStreamQuery(0) == cudaErrorNotReady)
48 if ((getTimeMilliseconds() - startTime) > limit)
56 #define SOFTWAIT() if (_pollStatus() != 0) {return 0xFFFFFFFE; } // -2
57 #define SOFTWAIT_LIM(lim) if (_pollStatus(1, lim) != 0) {return 0xFFFFFFFE; } // -2
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);}
65 // Use this macro to check for kernel errors
66 #define CHECK_LAUNCH_ERROR() if (cudaGetLastError() != cudaSuccess) {return 0xFFFFFFFF; /* -1 */}
69 typedef unsigned int uint;
71 // OO interface to MemtestG80 functions
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(); }
89 uint allocate(uint mbToTest);
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; }
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;
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);
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);
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);
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);
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);
130 // Memtest86 Test 4: tseq=10
131 __host__ uint gpuMovingInversionsRandom(const uint nBlocks, const uint nThreads, uint* base, uint N, uint* blockErrorCounts, uint* errorCounts);
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);
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);
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);