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__
11 dim3 threadIdx, blockDim, blockIdx;
12 #endif
13
14 template<typename Kernel, typename Input, typename Output>
run_on_cpu(const Kernel & ker,int n,const Input & in,Output & out)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__
run_on_cuda_meta_kernel(const Kernel ker,int n,const Input * in,Output * out)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>
run_on_cuda(const Kernel & ker,int n,const Input & in,Output & out)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>
run_and_compare_to_cuda(const Kernel & ker,int n,const Input & in,Output & out)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
ei_test_init_cuda()83 void ei_test_init_cuda()
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
102