gpuenergymicrobench/memoryTests.h at master · johnastratton/gpuenergymicrobench · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
#ifndef MEMORYTESTS_H
#define MEMORYTESTS_H

//#include "arithmeticTests.h"

template <typename T>
__global__
void globalMemKernel(int n, int iterateNum, volatile T *x);


template <typename T>
__global__
void createData(int n, T *x);


template <typename T>
class MemoryTestBase {
public:

  T *d_x;
  int n;
  int iterNum;
  int numBlocks;
  int blockSize;
  int numBlockScale;
  int opsPerIteration; //number of operations in one iteration. Not including loop calculations

  MemoryTestBase(int blockSize, int iterNum);
  MemoryTestBase(int blockSize, int iterNum, int numBlockScale);

  ~MemoryTestBase();

  void kernelSetup(cudaDeviceProp deviceProp);

  //get the number of threads launched in the kernel. Must be
  //called after kernelSetup() or the neccisary fields may not be initialized
  int getNumThreads();

  //return the number of operations that are executed in the kernel's loop
  //for the specified number of operations.
  //Ex: 6 operations per iteration * 1000000 iterations = 6000000 operations
  int getOpsPerThread();

  void runKernel();

  void CUDA_ERROR(cudaError_t e);

};


template <typename T>
class GlobalMemTest : public MemoryTestBase<T> {
public:
  GlobalMemTest(int blockSize, int iterNum);
  GlobalMemTest(int blockSize, int iterNum, int numBlockScale);
 // void kernelSetup(cudaDeviceProp deviceProp);
  void runKernel();
};

#endif