cuda_common.h
Go to the documentation of this file.
1 
2 #ifndef EIGEN_TEST_CUDA_COMMON_H
3 #define EIGEN_TEST_CUDA_COMMON_H
4 
5 #include <cuda.h>
6 #include <cuda_runtime.h>
7 #include <cuda_runtime_api.h>
8 #include <iostream>
9 
10 #ifndef __CUDACC__
12 #endif
13 
14 template<typename Kernel, typename Input, typename Output>
15 void run_on_cpu(const Kernel& ker, int n, const Input& in, Output& out)
16 {
17  for(int i=0; i<n; i++)
18  ker(i, in.data(), out.data());
19 }
20 
21 
22 template<typename Kernel, typename Input, typename Output>
23 __global__
24 void run_on_cuda_meta_kernel(const Kernel ker, int n, const Input* in, Output* out)
25 {
26  int i = threadIdx.x + blockIdx.x*blockDim.x;
27  if(i<n) {
28  ker(i, in, out);
29  }
30 }
31 
32 
33 template<typename Kernel, typename Input, typename Output>
34 void run_on_cuda(const Kernel& ker, int n, const Input& in, Output& out)
35 {
36  typename Input::Scalar* d_in;
37  typename Output::Scalar* d_out;
38  std::ptrdiff_t in_bytes = in.size() * sizeof(typename Input::Scalar);
39  std::ptrdiff_t out_bytes = out.size() * sizeof(typename Output::Scalar);
40 
41  cudaMalloc((void**)(&d_in), in_bytes);
42  cudaMalloc((void**)(&d_out), out_bytes);
43 
44  cudaMemcpy(d_in, in.data(), in_bytes, cudaMemcpyHostToDevice);
45  cudaMemcpy(d_out, out.data(), out_bytes, cudaMemcpyHostToDevice);
46 
47  // Simple and non-optimal 1D mapping assuming n is not too large
48  // That's only for unit testing!
49  dim3 Blocks(128);
50  dim3 Grids( (n+int(Blocks.x)-1)/int(Blocks.x) );
51 
52  cudaThreadSynchronize();
53  run_on_cuda_meta_kernel<<<Grids,Blocks>>>(ker, n, d_in, d_out);
54  cudaThreadSynchronize();
55 
56  // check inputs have not been modified
57  cudaMemcpy(const_cast<typename Input::Scalar*>(in.data()), d_in, in_bytes, cudaMemcpyDeviceToHost);
58  cudaMemcpy(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost);
59 
60  cudaFree(d_in);
61  cudaFree(d_out);
62 }
63 
64 
65 template<typename Kernel, typename Input, typename Output>
66 void run_and_compare_to_cuda(const Kernel& ker, int n, const Input& in, Output& out)
67 {
68  Input in_ref, in_cuda;
69  Output out_ref, out_cuda;
70  #ifndef __CUDA_ARCH__
71  in_ref = in_cuda = in;
72  out_ref = out_cuda = out;
73  #endif
74  run_on_cpu (ker, n, in_ref, out_ref);
75  run_on_cuda(ker, n, in_cuda, out_cuda);
76  #ifndef __CUDA_ARCH__
77  VERIFY_IS_APPROX(in_ref, in_cuda);
78  VERIFY_IS_APPROX(out_ref, out_cuda);
79  #endif
80 }
81 
82 
84 {
85  int device = 0;
86  cudaDeviceProp deviceProp;
87  cudaGetDeviceProperties(&deviceProp, device);
88  std::cout << "CUDA device info:\n";
89  std::cout << " name: " << deviceProp.name << "\n";
90  std::cout << " capability: " << deviceProp.major << "." << deviceProp.minor << "\n";
91  std::cout << " multiProcessorCount: " << deviceProp.multiProcessorCount << "\n";
92  std::cout << " maxThreadsPerMultiProcessor: " << deviceProp.maxThreadsPerMultiProcessor << "\n";
93  std::cout << " warpSize: " << deviceProp.warpSize << "\n";
94  std::cout << " regsPerBlock: " << deviceProp.regsPerBlock << "\n";
95  std::cout << " concurrentKernels: " << deviceProp.concurrentKernels << "\n";
96  std::cout << " clockRate: " << deviceProp.clockRate << "\n";
97  std::cout << " canMapHostMemory: " << deviceProp.canMapHostMemory << "\n";
98  std::cout << " computeMode: " << deviceProp.computeMode << "\n";
99 }
100 
101 #endif // EIGEN_TEST_CUDA_COMMON_H
SCALAR Scalar
Definition: bench_gemm.cpp:33
__global__ void run_on_cuda_meta_kernel(const Kernel ker, int n, const Input *in, Output *out)
Definition: cuda_common.h:24
int n
dim3 threadIdx
Definition: cuda_common.h:11
void run_on_cuda(const Kernel &ker, int n, const Input &in, Output &out)
Definition: cuda_common.h:34
#define VERIFY_IS_APPROX(a, b)
void ei_test_init_cuda()
Definition: cuda_common.h:83
dim3 blockIdx
Definition: cuda_common.h:11
void run_on_cpu(const Kernel &ker, int n, const Input &in, Output &out)
Definition: cuda_common.h:15
dim3 blockDim
Definition: cuda_common.h:11
void run_and_compare_to_cuda(const Kernel &ker, int n, const Input &in, Output &out)
Definition: cuda_common.h:66


gtsam
Author(s):
autogenerated on Sat May 8 2021 02:41:54