1/*
2 * se_class1_cuda.hpp
3 *
4 * Created on: Jan 13, 2019
5 * Author: i-bird
6 */
7
8#ifndef SE_CLASS1_CUDA_HPP_
9#define SE_CLASS1_CUDA_HPP_
10
11#include "util/se_util.hpp"
12#include <type_traits>
13#include <string>
14
15//! Void structure
16template<typename> struct Void_dev
17{
18 //! define void type
19 typedef void type;
20};
21
22template<typename T, typename Sfinae = void>
23struct has_check_device_pointer: std::false_type {};
24
25/*! \brief has_check_device_pointer check if a type has defined a member yes_has_check_device_pointer
26 *
27 * This mean that the class support a way to check if it is the owner od a particular device pointer
28 *
29 *
30 * return true if T::yes_has_check_device_pointer is a valid type
31 *
32 */
33template<typename T>
34struct has_check_device_pointer<T, typename Void_dev< typename T::yes_has_check_device_pointer >::type> : std::true_type
35{};
36
37struct pointer_check
38{
39 //! Indicate if the pointer match
40 bool match;
41
42 //! match string
43 std::string match_str;
44};
45
46template<typename T, int type_of_t=has_check_device_pointer<T>::value>
47struct check_type
48{
49 static pointer_check check(void * ptr, int prp, T & arg)
50 {
51 pointer_check pc;
52
53 pc.match = false;
54
55 return pc;
56 }
57};
58
59
60
61template<typename T>
62struct check_type<T,1>
63{
64 static pointer_check check(void * ptr, int prp, T & arg)
65 {
66 return arg.check_device_pointer(ptr);
67 }
68};
69
70struct pos_pc
71{
72 int pos;
73 pointer_check pc;
74};
75
76template<typename ArgL>
77pos_pc error_args_impl(void * ptr, int prp, ArgL argl)
78{
79 pos_pc pp;
80 pointer_check pc = check_type<ArgL>::check(ptr,prp,argl);
81 if (pc.match == true)
82 {
83 pp.pos = 0;
84 pp.pc = pc;
85 return pp;
86 }
87
88 pp.pos = -1;
89
90 return pp;
91}
92
93template<typename ArgL, typename ... Args>
94pos_pc error_args_impl(void * ptr, int prp, ArgL argl, Args ... args)
95{
96 pos_pc pp;
97 pointer_check pc = check_type<ArgL>::check(ptr,prp,argl);
98 if (pc.match == true)
99 {
100 pp.pos = sizeof...(args);
101 pp.pc = pc;
102 return pp;
103 }
104 return error_args_impl(ptr, prp, args ...);
105}
106
107template<typename ... Args>pos_pc error_arg(void * ptr, int prp, Args ... args)
108{
109 pos_pc pp;
110 pp = error_args_impl(ptr, prp, args ... );
111 pp.pos = sizeof...(args) - pp.pos - 1;
112 return pp;
113}
114
115#include <boost/algorithm/string.hpp>
116
117#if defined(SE_CLASS1) && !defined(__clang__)
118#define CUDA_LAUNCH_ERROR_OBJECT std::runtime_error("Runtime vector error");
119#define CHECK_SE_CLASS1_PRE int dev_mem[] = {0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
120
121#if !defined(CUDA_ON_CPU)
122
123 #define CHECK_SE_CLASS1_POST(kernel_call,...) \
124 cudaError_t e1 = cudaMemcpyFromSymbol(dev_mem,global_cuda_error_array,sizeof(dev_mem)); \
125 if (e1 != cudaSuccess)\
126 {\
127 std::string error = cudaGetErrorString(e1);\
128 std::cout << "Cuda Error in cudaMemcpyFromSymbol: " << __FILE__ << ":" << __LINE__ << " " << error << std::endl;\
129 }\
130 if (dev_mem[0] != 0)\
131 {\
132 void * ptr = (void *)*(size_t *)&dev_mem[1]; \
133 int prp_err = dev_mem[3];\
134 pos_pc ea = error_arg(ptr,prp_err,__VA_ARGS__);\
135 std::string args_s( #__VA_ARGS__ );\
136 std::vector<std::string> results;\
137 boost::split(results, args_s, [](char c){return c == ',';});\
138 std::string data_s;\
139 if (ea.pos >= results.size())\
140 {data_s = "Internal";}\
141 else\
142 {data_s = results[ea.pos];}\
143 std::cout << __FILE__ << ":" << __LINE__ << " Overflow detected in Kernel: " << kernel_call << " from the structure: " << data_s << " property: " << prp_err << " index:(" ;\
144 int i = 0; \
145 for ( ; i < dev_mem[4]-1 ; i++)\
146 {\
147 std::cout << dev_mem[5+i] << ",";\
148 }\
149 std::cout << dev_mem[5+i];\
150 std::cout << ")";\
151 std::cout << " thread: " << "(" << dev_mem[6+i] << "," << dev_mem[7+i] << "," << dev_mem[8+i] << ")*(" << dev_mem[9+i] << "," << dev_mem[10+i] << "," << dev_mem[11+i] << ")+(" << dev_mem[12+i] << "," << dev_mem[13+i] << "," << dev_mem[14+i] << ")" << std::endl;\
152 std::cout << "Internal error report: " << ea.pc.match_str << std::endl;\
153 int dev_mem_null[] = {0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};\
154 cudaError_t e2 = cudaMemcpyToSymbol(global_cuda_error_array,dev_mem_null,sizeof(dev_mem_null),0,cudaMemcpyHostToDevice);\
155 if (e2 != cudaSuccess)\
156 {\
157 std::string error = cudaGetErrorString(e2);\
158 std::cout << "Cuda Error in cudaMemcpyToSymbol: " << __FILE__ << ":" << __LINE__ << " " << error << std::endl;\
159 }\
160 ACTION_ON_ERROR(CUDA_LAUNCH_ERROR_OBJECT);\
161 }\
162
163#else
164
165 #define CHECK_SE_CLASS1_POST(kernel_call,...) \
166 memcpy(dev_mem,global_cuda_error_array,sizeof(dev_mem)); \
167 if (dev_mem[0] != 0)\
168 {\
169 void * ptr = (void *)*(size_t *)&dev_mem[1]; \
170 int prp_err = dev_mem[3];\
171 pos_pc ea = error_arg(ptr,prp_err,__VA_ARGS__);\
172 std::string args_s( #__VA_ARGS__ );\
173 std::vector<std::string> results;\
174 boost::split(results, args_s, [](char c){return c == ',';});\
175 std::string data_s;\
176 if (ea.pos >= results.size())\
177 {data_s = "Internal";}\
178 else\
179 {data_s = results[ea.pos];}\
180 std::cout << __FILE__ << ":" << __LINE__ << " Overflow detected in Kernel: " << kernel_call << " from the structure: " << data_s << " property: " << prp_err << " index:(" ;\
181 int i = 0; \
182 for ( ; i < dev_mem[4]-1 ; i++)\
183 {\
184 std::cout << dev_mem[5+i] << ",";\
185 }\
186 std::cout << dev_mem[5+i];\
187 std::cout << ")";\
188 std::cout << " thread: " << "(" << dev_mem[6+i] << "," << dev_mem[7+i] << "," << dev_mem[8+i] << ")*(" << dev_mem[9+i] << "," << dev_mem[10+i] << "," << dev_mem[11+i] << ")+(" << dev_mem[12+i] << "," << dev_mem[13+i] << "," << dev_mem[14+i] << ")" << std::endl;\
189 std::cout << "Internal error report: " << ea.pc.match_str << std::endl;\
190 int dev_mem_null[] = {0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};\
191 memcpy(global_cuda_error_array,dev_mem_null,sizeof(dev_mem_null));\
192 ACTION_ON_ERROR(CUDA_LAUNCH_ERROR_OBJECT);\
193 }\
194
195#endif
196
197#else
198#define CHECK_SE_CLASS1_PRE
199#define CHECK_SE_CLASS1_POST(kernel_call,...)
200#endif
201
202
203#endif /* SE_CLASS1_CUDA_HPP_ */
204