API tools faq
paste
Advertisement
Guest User

Untitled

a guest
May 18th, 2017
583
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 13.00 KB | None | 0 0
raw download clone embed print report
  1. [----------] 26 tests from NetTest/1, where TypeParam = caffe::CPUDevice<double>
  2. [ RUN ] NetTest/1.TestGetBlob
  3. [ OK ] NetTest/1.TestGetBlob (0 ms)
  4. [ RUN ] NetTest/1.TestBottomNeedBackwardForce
  5. [ OK ] NetTest/1.TestBottomNeedBackwardForce (1 ms)
  6. [ RUN ] NetTest/1.TestReshape
  7. Build Status = -2 ( Err = -11 )
  8. Log: stringInput.cl:111:49: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  9. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  10. stringInput.cl:111:70: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  11. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  12. stringInput.cl:176:49: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  13. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  14. stringInput.cl:176:89: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  15. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  16. stringInput.cl:187:16: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  17. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  18. stringInput.cl:188:10: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  19. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  20. stringInput.cl:190:1: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  21. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  22. stringInput.cl:209:22: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  23. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  24. stringInput.cl:211:48: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  25. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  26. stringInput.cl:211:88: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  27. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  28. stringInput.cl:222:16: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  29. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  30. stringInput.cl:223:10: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  31. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  32. stringInput.cl:225:1: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  33. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  34. stringInput.cl:244:22: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  35. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  36. stringInput.cl:306:44: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  37. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  38. stringInput.cl:306:81: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  39. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  40. stringInput.cl:316:10: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  41. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  42. stringInput.cl:317:16: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  43. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  44. stringInput.cl:333:1: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
  45. stringInput.cl:20:15: note: expanded from macro 'Dtype'
  46. stringInput.cl:333:18: warning: double precision constant requires cl_khr_fp64, casting to single precision
  47.  
  48. Sources: #if defined(cl_khr_fp64)
  49. #pragma OPENCL EXTENSION cl_khr_fp64 : enable
  50. #define DOUBLE_SUPPORT_AVAILABLE
  51. #elif defined(cl_amd_fp64)
  52. #pragma OPENCL EXTENSION cl_amd_fp64 : enable
  53. #define DOUBLE_SUPPORT_AVAILABLE
  54. #endif
  55. #if defined(cl_khr_int32_base_atomics)
  56. #pragma OPENCL EXTENSION cl_khr_int32_base_atomics : enable
  57. #define ATOMICS_32_AVAILABLE
  58. #endif
  59. #if defined(cl_khr_global_int32_base_atomics)
  60. #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
  61. #define ATOMICS_32_AVAILABLE
  62. #endif
  63. #if defined(cl_khr_int64_base_atomics)
  64. #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
  65. #define ATOMICS_64_AVAILABLE
  66. #endif
  67. #define Dtype double
  68. #define Dtype1 double
  69. #define Dtype2 double2
  70. #define Dtype4 double4
  71. #define Dtype8 double8
  72. #define Dtype16 double16
  73. #define VEC_1_0(X) X
  74. #define VEC_2_0(X) X.x
  75. #define VEC_2_1(X) X.y
  76. #define VEC_4_0(X) X.x
  77. #define VEC_4_1(X) X.y
  78. #define VEC_4_2(X) X.z
  79. #define VEC_4_3(X) X.w
  80. #define VEC_8_0(X) X.s0
  81. #define VEC_8_1(X) X.s1
  82. #define VEC_8_2(X) X.s2
  83. #define VEC_8_3(X) X.s3
  84. #define VEC_8_4(X) X.s4
  85. #define VEC_8_5(X) X.s5
  86. #define VEC_8_6(X) X.s6
  87. #define VEC_8_7(X) X.s7
  88. #define VEC_16_0(X) X.s0
  89. #define VEC_16_1(X) X.s1
  90. #define VEC_16_2(X) X.s2
  91. #define VEC_16_3(X) X.s3
  92. #define VEC_16_4(X) X.s4
  93. #define VEC_16_5(X) X.s5
  94. #define VEC_16_6(X) X.s6
  95. #define VEC_16_7(X) X.s7
  96. #define VEC_16_8(X) X.s8
  97. #define VEC_16_9(X) X.s9
  98. #define VEC_16_10(X) X.sA
  99. #define VEC_16_11(X) X.sB
  100. #define VEC_16_12(X) X.sC
  101. #define VEC_16_13(X) X.sD
  102. #define VEC_16_14(X) X.sE
  103. #define VEC_16_15(X) X.sF
  104. #define int_tp int
  105. #define uint_tp unsigned int
  106. #define int_tpc int
  107. #define uint_tpc unsigned int
  108. #ifdef ATOMICS_64_AVAILABLE
  109. inline void atomicAdd(volatile __global Dtype* source, const Dtype operand) {
  110. union {
  111. unsigned long intVal;
  112. Dtype floatVal;
  113. } next, expected, current;
  114. current.floatVal = *source;
  115. do {
  116. expected.floatVal = current.floatVal;
  117. next.floatVal = expected.floatVal + operand;
  118. current.intVal = atom_cmpxchg((volatile __global unsigned long *)source, expected.intVal, next.intVal);
  119. } while (current.intVal != expected.intVal);
  120. }
  121. inline void atomicSub(volatile __global Dtype* source, const Dtype operand) {
  122. union {
  123. unsigned long intVal;
  124. Dtype floatVal;
  125. } next, expected, current;
  126. current.floatVal = *source;
  127. do {
  128. expected.floatVal = current.floatVal;
  129. next.floatVal = expected.floatVal - operand;
  130. current.intVal = atom_cmpxchg((volatile __global unsigned long *)source, expected.intVal, next.intVal);
  131. } while (current.intVal != expected.intVal);
  132. }
  133. inline void atomicMul(volatile __global Dtype* source, const Dtype operand) {
  134. union {
  135. unsigned long intVal;
  136. Dtype floatVal;
  137. } next, expected, current;
  138. current.floatVal = *source;
  139. do {
  140. expected.floatVal = current.floatVal;
  141. next.floatVal = expected.floatVal * operand;
  142. current.intVal = atom_cmpxchg((volatile __global unsigned long *)source, expected.intVal, next.intVal);
  143. } while (current.intVal != expected.intVal);
  144. }
  145. inline void atomicDiv(volatile __global Dtype* source, const Dtype operand) {
  146. union {
  147. unsigned long intVal;
  148. Dtype floatVal;
  149. } next, expected, current;
  150. current.floatVal = *source;
  151. do {
  152. expected.floatVal = current.floatVal;
  153. next.floatVal = expected.floatVal / operand;
  154. current.intVal = atom_cmpxchg((volatile __global unsigned long *)source, expected.intVal, next.intVal);
  155. } while (current.intVal != expected.intVal);
  156. }
  157. #endif
  158. __kernel void fill_memory(const int_tp n, const Dtype alpha,__global Dtype* x, const int_tp offx) {
  159. for (int_tp index = get_global_id(0); index < n; index += get_global_size(0)) {
  160. x[index + offx] = alpha;
  161. }
  162. }
  163. #ifdef v_nax
  164. #undef v_nax
  165. #endif
  166. #define v_nax 2
  167. #ifdef v_k_0
  168. #undef v_k_0
  169. #endif
  170. #define v_k_0 2
  171. #ifdef v_k_1
  172. #undef v_k_1
  173. #endif
  174. #define v_k_1 2
  175. #ifdef v_p_0
  176. #undef v_p_0
  177. #endif
  178. #define v_p_0 0
  179. #ifdef v_p_1
  180. #undef v_p_1
  181. #endif
  182. #define v_p_1 0
  183. #ifdef v_s_0
  184. #undef v_s_0
  185. #endif
  186. #define v_s_0 2
  187. #ifdef v_s_1
  188. #undef v_s_1
  189. #endif
  190. #define v_s_1 2
  191. #ifdef v_d_0
  192. #undef v_d_0
  193. #endif
  194. #define v_d_0 1
  195. #ifdef v_d_1
  196. #undef v_d_1
  197. #endif
  198. #define v_d_1 1
  199. #ifdef v_imsi_0
  200. #undef v_imsi_0
  201. #endif
  202. #define v_imsi_0 49
  203. #ifdef v_imso_0
  204. #undef v_imso_0
  205. #endif
  206. #define v_imso_0 25
  207. #ifdef v_imsi_1
  208. #undef v_imsi_1
  209. #endif
  210. #define v_imsi_1 49
  211. #ifdef v_imso_1
  212. #undef v_imso_1
  213. #endif
  214. #define v_imso_1 25
  215. #ifdef v_imsi
  216. #undef v_imsi
  217. #endif
  218. #define v_imsi 2401
  219. #ifdef v_imso
  220. #undef v_imso
  221. #endif
  222. #define v_imso 625
  223. __kernel void pool_forward_train(__global const Dtype* __restrict bottom_data, __global Dtype* __restrict top_data, __global int_tp* __restrict mask, int_tp channels, int_tp batch_size) {
  224. int_tp out_idx = get_global_id(0);
  225. if (get_global_id(1) >= channels * batch_size) {return;}
  226. int_tp idx_0 = get_global_id(0);
  227. int_tp idx_1 = (idx_0 % v_imso_1);
  228. idx_1 = idx_1 * v_s_1 - v_p_1;
  229. idx_0 /= v_imso_1;
  230. if (idx_0 >= v_imso_0) {return;}
  231. idx_0 = idx_0 * v_s_0 - v_p_0;
  232. int_tp in_idx = idx_0;
  233. in_idx = in_idx * v_imsi_1 + idx_1;
  234. __global const Dtype* in_ptr = bottom_data + get_global_id(1) * v_imsi + in_idx;
  235. __global Dtype* out_ptr = top_data + get_global_id(1) * v_imso;
  236. __global int_tp* mask_ptr = mask + get_global_id(1) * v_imso;
  237. Dtype val = -FLT_MAX;
  238. int_tp maxidx = -1;
  239. if (in_ptr[0] > val) {
  240. maxidx = in_idx + 0;
  241. val = in_ptr[0];
  242. }
  243. if (idx_1 < v_imsi_1 - 1 && in_ptr[1] > val) {
  244. maxidx = in_idx + 1;
  245. val = in_ptr[1];
  246. }
  247. if (idx_0 < v_imsi_0 - 1 && in_ptr[49] > val) {
  248. maxidx = in_idx + 49;
  249. val = in_ptr[49];
  250. }
  251. if (idx_0 < v_imsi_0 - 1 && idx_1 < v_imsi_1 - 1 && in_ptr[50] > val) {
  252. maxidx = in_idx + 50;
  253. val = in_ptr[50];
  254. }
  255. out_ptr[out_idx] = val;
  256. mask_ptr[out_idx] = (Dtype)maxidx;
  257. }
  258. __kernel void pool_forward_test(__global const Dtype* __restrict bottom_data, __global Dtype* __restrict top_data, __global int_tp* __restrict mask, int_tp channels, int_tp batch_size) {
  259. int_tp out_idx = get_global_id(0);
  260. if (get_global_id(1) >= channels * batch_size) {return;}
  261. int_tp idx_0 = get_global_id(0);
  262. int_tp idx_1 = (idx_0 % v_imso_1);
  263. idx_1 = idx_1 * v_s_1 - v_p_1;
  264. idx_0 /= v_imso_1;
  265. if (idx_0 >= v_imso_0) {return;}
  266. idx_0 = idx_0 * v_s_0 - v_p_0;
  267. int_tp in_idx = idx_0;
  268. in_idx = in_idx * v_imsi_1 + idx_1;
  269. __global const Dtype* in_ptr = bottom_data + get_global_id(1) * v_imsi + in_idx;
  270. __global Dtype* out_ptr = top_data + get_global_id(1) * v_imso;
  271. __global int_tp* mask_ptr = mask + get_global_id(1) * v_imso;
  272. Dtype val = -FLT_MAX;
  273. int_tp maxidx = -1;
  274. if (in_ptr[0] > val) {
  275. maxidx = in_idx + 0;
  276. val = in_ptr[0];
  277. }
  278. if (idx_1 < v_imsi_1 - 1 && in_ptr[1] > val) {
  279. maxidx = in_idx + 1;
  280. val = in_ptr[1];
  281. }
  282. if (idx_0 < v_imsi_0 - 1 && in_ptr[49] > val) {
  283. maxidx = in_idx + 49;
  284. val = in_ptr[49];
  285. }
  286. if (idx_0 < v_imsi_0 - 1 && idx_1 < v_imsi_1 - 1 && in_ptr[50] > val) {
  287. maxidx = in_idx + 50;
  288. val = in_ptr[50];
  289. }
  290. out_ptr[out_idx] = val;
  291. mask_ptr[out_idx] = (Dtype)maxidx;
  292. }
  293. #ifdef v_nax
  294. #undef v_nax
  295. #endif
  296. #define v_nax 2
  297. #ifdef v_k_0
  298. #undef v_k_0
  299. #endif
  300. #define v_k_0 2
  301. #ifdef v_k_1
  302. #undef v_k_1
  303. #endif
  304. #define v_k_1 2
  305. #ifdef v_p_0
  306. #undef v_p_0
  307. #endif
  308. #define v_p_0 0
  309. #ifdef v_p_1
  310. #undef v_p_1
  311. #endif
  312. #define v_p_1 0
  313. #ifdef v_s_0
  314. #undef v_s_0
  315. #endif
  316. #define v_s_0 2
  317. #ifdef v_s_1
  318. #undef v_s_1
  319. #endif
  320. #define v_s_1 2
  321. #ifdef v_d_0
  322. #undef v_d_0
  323. #endif
  324. #define v_d_0 1
  325. #ifdef v_d_1
  326. #undef v_d_1
  327. #endif
  328. #define v_d_1 1
  329. #ifdef v_imsi_0
  330. #undef v_imsi_0
  331. #endif
  332. #define v_imsi_0 49
  333. #ifdef v_imso_0
  334. #undef v_imso_0
  335. #endif
  336. #define v_imso_0 25
  337. #ifdef v_imsi_1
  338. #undef v_imsi_1
  339. #endif
  340. #define v_imsi_1 49
  341. #ifdef v_imso_1
  342. #undef v_imso_1
  343. #endif
  344. #define v_imso_1 25
  345. #ifdef v_imsi
  346. #undef v_imsi
  347. #endif
  348. #define v_imsi 2401
  349. #ifdef v_imso
  350. #undef v_imso
  351. #endif
  352. #define v_imso 625
  353. __kernel void pool_backward(__global const Dtype* __restrict top_diff, __global Dtype* __restrict bottom_diff, __global const int_tp* __restrict mask, int_tp channels, int_tp batch_size) {
  354. int_tp d_start[2];
  355. int_tp d_end[2];
  356. int_tp d_iter[2];
  357. int_tp out_idx = get_global_id(0);
  358. int_tp idx_0 = get_global_id(0);
  359. if (get_global_id(1) >= channels * batch_size) {return;}
  360. int_tp idx_1 = (idx_0 % v_imsi_1);
  361. idx_0 /= v_imsi_1;
  362. if (idx_0 >= v_imsi_0) {return;}
  363. __global Dtype* out_ptr = bottom_diff + get_global_id(1) * v_imsi + out_idx;
  364. __global const Dtype* in_ptr = top_diff + get_global_id(1) * v_imso;
  365. __global const int_tp* mask_ptr = mask + get_global_id(1) * v_imso;
  366. d_start[0] = (idx_0 + v_p_0 < ((v_k_0 - 1) * v_d_0 + 1)) ? 0 : (idx_0 + v_p_0 - ((v_k_0 - 1) * v_d_0 + 1)) / v_s_0 + 1;
  367. d_end[0] = min(v_imso_0 - 1, (idx_0 + v_p_0) / v_s_0);
  368. d_iter[0] = d_start[0];
  369. if (d_start[0] > d_end[0]) {
  370. out_ptr[0] = 0;
  371. return;
  372. }
  373. d_start[1] = (idx_1 + v_p_1 < ((v_k_1 - 1) * v_d_1 + 1)) ? 0 : (idx_1 + v_p_1 - ((v_k_1 - 1) * v_d_1 + 1)) / v_s_1 + 1;
  374. d_end[1] = min(v_imso_1 - 1, (idx_1 + v_p_1) / v_s_1);
  375. d_iter[1] = d_start[1];
  376. if (d_start[1] > d_end[1]) {
  377. out_ptr[0] = 0;
  378. return;
  379. }
  380. Dtype gradient = 0.0;
  381. bool incremented;
  382. do {
  383. int_tp offset = 0;
  384. offset += d_iter[0];
  385. offset *= v_imso_1;
  386. offset += d_iter[1];
  387. if ((int_tp)mask_ptr[offset] == out_idx) {
  388. gradient += in_ptr[offset];
  389. }
  390. incremented = false;
  391. for (int_tp i = v_nax - 1; i >= 0; --i) {
  392. if (d_iter[i] >= d_end[i]) {
  393. d_iter[i] = d_start[i];
  394. } else {
  395. ++d_iter[i];
  396. incremented = true;
  397. break;
  398. }}} while (incremented);
  399. out_ptr[0] = gradient;
  400. }
  401.  
  402. unknown file: Failure
  403. C++ exception with description "ViennaCL: FATAL ERROR: CL_INVALID_PROGRAM_EXECUTABLE.
  404. If you think that this is a bug in ViennaCL, please report it at viennacl-support@lists.sourceforge.net and supply at least the following information:
  405. * Operating System
  406. * Which OpenCL implementation (AMD, NVIDIA, etc.)
  407. * ViennaCL version
  408. Many thanks in advance!" thrown in the test body.
  409. [ FAILED ] NetTest/1.TestReshape, where TypeParam = caffe::CPUDevice<double> (159 ms)
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!