API tools faq
paste
rajeswar

Untitled

rajeswar
Aug 9th, 2014
313
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
Python 47.34 KB | None | 0 0
raw download clone embed print report
  1. ... loading data
  2. ... building the model
  3. 1 #include <Python.h>
  4. 2 #include <iostream>
  5. 3 #include "cuda_ndarray.cuh"
  6. 4 #include <math.h>
  7. 5 #include <numpy/arrayobject.h>
  8. 6 #include <numpy/arrayscalars.h>
  9. 7 //////////////////////
  10. 8 ////  Support Code
  11. 9 //////////////////////
  12. 10
  13. 11
  14. 12 __global__ void
  15. 13 //thread block size = out_dur
  16. 14 //grid block size =(out_len*out_wid, nb kern *nb batch)
  17. 15 //
  18. 16 conv_rows_stack( float* img, float* kern, float* bias, float* out,
  19. 17                  int img_len, int img_wid, int img_dur,
  20. 18                  int kern_height, int kern_wid, int kern_dur,
  21. 19                  int nkern, int input_channels,
  22. 20                  int dr, int dc, int dt,
  23. 21                  int img_stride_frame, int img_stride_col, int img_stride_row,
  24. 22                  int img_stride_ochannel, int img_stride_batch,
  25. 23                  int kern_stride_frame, int kern_stride_col, int kern_stride_row,
  26. 24                  int kern_stride_stack, int kern_stride_okern)
  27. 25 {
  28. 26   int __shared__ out_len, out_wid, out_dur, batch_id, kern_id;
  29. 27   float  __shared__ *d_img, *d_kern;
  30. 28   out_len = int( (img_len - kern_height) / dr )+1;
  31. 29   out_wid = int( (img_wid - kern_wid) / dc )+1;
  32. 30   out_dur = int( (img_dur - kern_dur) / dt )+1;
  33. 31
  34. 32   batch_id= blockIdx.y/nkern;
  35. 33   kern_id = blockIdx.y - batch_id*nkern;
  36. 34
  37. 35   const int out_row = blockIdx.x%out_len;
  38. 36   const int out_col = blockIdx.x/out_len;
  39. 37   const int out_frame=threadIdx.x;
  40. 38
  41. 39   img += batch_id*img_stride_batch + out_row*dr*img_stride_row + out_col*dc*img_stride_col+out_frame*dt*img_stride_frame;
  42. 40   kern += kern_id*kern_stride_okern;
  43. 41     float sum = 0.0f;
  44. 42     for (int z = 0; z < input_channels; z++) {//1 for first layer
  45. 43         for (int k =0; k < kern_height; k++) {
  46. 44           for (int l = 0; l < kern_wid; l++) {
  47. 45             for (int m = 0; m < kern_dur; m++) {
  48. 46               sum += img[img_stride_ochannel*z+img_stride_row*k+img_stride_col*l+img_stride_frame*m] *
  49. 47                          kern[kern_stride_stack*z+kern_stride_row*k+kern_stride_col*l+kern_stride_frame*m];
  50. 48             }
  51. 49           }
  52. 50         }
  53. 51
  54. 52       out[batch_id*nkern*out_len*out_wid*out_dur+//the good batch
  55. 53           out_frame*nkern+//the output frame
  56. 54           out_row*out_wid*out_dur*nkern+//the output row
  57. 55           out_col*out_dur*nkern + //the output_col
  58. 56           kern_id //the output image (channel)
  59. 57 ] = sum + bias[kern_id];
  60. 58     }
  61. 59 }
  62. 60
  63. 61
  64. 62            
  65. 63
  66. 64     namespace {
  67. 65     struct __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae {
  68. 66         PyObject* __ERROR;
  69. 67
  70. 68         PyObject* storage_V3;
  71. 69 PyObject* storage_V5;
  72. 70 PyObject* storage_V7;
  73. 71 PyObject* storage_V9;
  74. 72 PyObject* storage_V1;
  75. 73        
  76. 74
  77. 75         __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae() {}
  78. 76         ~__struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae(void) {
  79. 77             cleanup();
  80. 78         }
  81. 79
  82. 80         int init(PyObject* __ERROR, PyObject* storage_V3, PyObject* storage_V5, PyObject* storage_V7, PyObject* storage_V9, PyObject* storage_V1) {
  83. 81             Py_XINCREF(storage_V3);
  84. 82 Py_XINCREF(storage_V5);
  85. 83 Py_XINCREF(storage_V7);
  86. 84 Py_XINCREF(storage_V9);
  87. 85 Py_XINCREF(storage_V1);
  88. 86             this->storage_V3 = storage_V3;
  89. 87 this->storage_V5 = storage_V5;
  90. 88 this->storage_V7 = storage_V7;
  91. 89 this->storage_V9 = storage_V9;
  92. 90 this->storage_V1 = storage_V1;
  93. 91             int __failure = 0;
  94. 92            
  95. 93 {
  96. 94
  97. 95 {
  98. 96
  99. 97 {
  100. 98
  101. 99 {
  102. 100
  103. 101 {
  104. 102
  105. 103             this->__ERROR = __ERROR;
  106. 104             return 0;
  107. 105             __label_9:
  108. 106
  109. 107 double __DUMMY_9;
  110. 108
  111. 109 }
  112. 110 __label_7:
  113. 111
  114. 112 double __DUMMY_7;
  115. 113
  116. 114 }
  117. 115 __label_5:
  118. 116
  119. 117 double __DUMMY_5;
  120. 118
  121. 119 }
  122. 120 __label_3:
  123. 121
  124. 122 double __DUMMY_3;
  125. 123
  126. 124 }
  127. 125 __label_1:
  128. 126
  129. 127 double __DUMMY_1;
  130. 128
  131. 129 }
  132. 130
  133. 131             Py_XDECREF(this->storage_V3);
  134. 132 Py_XDECREF(this->storage_V5);
  135. 133 Py_XDECREF(this->storage_V7);
  136. 134 Py_XDECREF(this->storage_V9);
  137. 135 Py_XDECREF(this->storage_V1);
  138. 136            
  139. 137         if (__failure) {
  140. 138             // When there is a failure, this code puts the exception
  141. 139             // in __ERROR.
  142. 140             PyObject* err_type = NULL;
  143. 141             PyObject* err_msg = NULL;
  144. 142             PyObject* err_traceback = NULL;
  145. 143             PyErr_Fetch(&err_type, &err_msg, &err_traceback);
  146. 144             if (!err_type) {err_type = Py_None;Py_INCREF(Py_None);}
  147. 145             if (!err_msg) {err_msg = Py_None; Py_INCREF(Py_None);}
  148. 146             if (!err_traceback) {err_traceback = Py_None; Py_INCREF(Py_None);}
  149. 147             PyObject* old_err_type = PyList_GET_ITEM(__ERROR, 0);
  150. 148             PyObject* old_err_msg = PyList_GET_ITEM(__ERROR, 1);
  151. 149             PyObject* old_err_traceback = PyList_GET_ITEM(__ERROR, 2);
  152. 150             PyList_SET_ITEM(__ERROR, 0, err_type);
  153. 151             PyList_SET_ITEM(__ERROR, 1, err_msg);
  154. 152             PyList_SET_ITEM(__ERROR, 2, err_traceback);
  155. 153             {Py_XDECREF(old_err_type);}
  156. 154             {Py_XDECREF(old_err_msg);}
  157. 155             {Py_XDECREF(old_err_traceback);}
  158. 156         }
  159. 157         // The failure code is returned to index what code block failed.
  160. 158         return __failure;
  161. 159        
  162. 160         }
  163. 161         void cleanup(void) {
  164. 162             __label_1:
  165. 163
  166. 164 double __DUMMY_1;
  167. 165 __label_3:
  168. 166
  169. 167 double __DUMMY_3;
  170. 168 __label_5:
  171. 169
  172. 170 double __DUMMY_5;
  173. 171 __label_7:
  174. 172
  175. 173 double __DUMMY_7;
  176. 174 __label_9:
  177. 175
  178. 176 double __DUMMY_9;
  179. 177
  180. 178             Py_XDECREF(this->storage_V3);
  181. 179 Py_XDECREF(this->storage_V5);
  182. 180 Py_XDECREF(this->storage_V7);
  183. 181 Py_XDECREF(this->storage_V9);
  184. 182 Py_XDECREF(this->storage_V1);
  185. 183         }
  186. 184         int run(void) {
  187. 185             int __failure = 0;
  188. 186            
  189. 187     PyObject* py_V1;
  190. 188      CudaNdarray * V1;
  191. 189     PyObject* py_V3;
  192. 190      CudaNdarray * V3;
  193. 191     PyObject* py_V5;
  194. 192      CudaNdarray * V5;
  195. 193     PyObject* py_V7;
  196. 194      CudaNdarray * V7;
  197. 195     PyObject* py_V9;
  198. 196    
  199. 197         PyArrayObject* V9;
  200. 198        
  201. 199             typedef npy_int64 dtype_V9;
  202. 200            
  203. 201 {
  204. 202
  205. 203     py_V1 = PyList_GET_ITEM(storage_V1, 0);
  206. 204     {Py_XINCREF(py_V1);}
  207. 205    
  208. 206         if (py_V1 == Py_None)
  209. 207         {
  210. 208             V1 = NULL;
  211. 209         }
  212. 210         else
  213. 211         {
  214. 212            
  215. 213         assert(py_V1->ob_refcnt >= 2); // There should be at least one ref from the container object,
  216. 214         // and one ref from the local scope.
  217. 215
  218. 216         if (CudaNdarray_Check(py_V1))
  219. 217         {
  220. 218             //fprintf(stderr, "c_extract CNDA object w refcnt %p %i\n", py_V1, (py_V1->ob_refcnt));
  221. 219             V1 = (CudaNdarray*)py_V1;
  222. 220             //std::cerr << "c_extract " << V1 << '\n';
  223. 221        
  224. 222
  225. 223                 if (V1->nd != 5)
  226. 224                 {
  227. 225                     PyErr_Format(PyExc_RuntimeError,
  228. 226                                  "c_extract: Some CudaNdarray has rank %i, it was supposed to have rank 5",
  229. 227                                  V1->nd);
  230. 228                     V1 = NULL;
  231. 229                     {
  232. 230         __failure = 2;
  233. 231         if (!PyErr_Occurred()) {
  234. 232             PyErr_SetString(PyExc_RuntimeError,
  235. 233                 "Unexpected error in an Op's C code. "
  236. 234                 "No Python exception was set.");
  237. 235             }
  238. 236         goto __label_2;};
  239. 237                 }
  240. 238                 //std::cerr << "c_extract " << V1 << " nd check passed\n";
  241. 239            
  242. 240
  243. 241                 assert(V1);
  244. 242                 Py_INCREF(py_V1);
  245. 243             }
  246. 244             else if (py_V1 == Py_None)
  247. 245             {
  248. 246                 PyErr_SetString(PyExc_TypeError,
  249. 247                                 "expected a CudaNdarray, not None");
  250. 248                 V1 = NULL;
  251. 249                 {
  252. 250         __failure = 2;
  253. 251         if (!PyErr_Occurred()) {
  254. 252             PyErr_SetString(PyExc_RuntimeError,
  255. 253                 "Unexpected error in an Op's C code. "
  256. 254                 "No Python exception was set.");
  257. 255             }
  258. 256         goto __label_2;};
  259. 257             }
  260. 258             else
  261. 259             {
  262. 260                 //fprintf(stderr, "FAILING c_extract CNDA object w refcnt %p %i\n", py_V1, (py_V1->ob_refcnt));
  263. 261                 PyErr_SetString(PyExc_TypeError, "Argument not a CudaNdarray");
  264. 262                 V1 = NULL;
  265. 263                 {
  266. 264         __failure = 2;
  267. 265         if (!PyErr_Occurred()) {
  268. 266             PyErr_SetString(PyExc_RuntimeError,
  269. 267                 "Unexpected error in an Op's C code. "
  270. 268                 "No Python exception was set.");
  271. 269             }
  272. 270         goto __label_2;};
  273. 271             }
  274. 272             //std::cerr << "c_extract done " << V1 << '\n';
  275. 273            
  276. 274
  277. 275         }
  278. 276        
  279. 277 {
  280. 278
  281. 279     py_V3 = PyList_GET_ITEM(storage_V3, 0);
  282. 280     {Py_XINCREF(py_V3);}
  283. 281    
  284. 282         assert(py_V3->ob_refcnt >= 2); // There should be at least one ref from the container object,
  285. 283         // and one ref from the local scope.
  286. 284
  287. 285         if (CudaNdarray_Check(py_V3))
  288. 286         {
  289. 287             //fprintf(stderr, "c_extract CNDA object w refcnt %p %i\n", py_V3, (py_V3->ob_refcnt));
  290. 288             V3 = (CudaNdarray*)py_V3;
  291. 289             //std::cerr << "c_extract " << V3 << '\n';
  292. 290        
  293. 291
  294. 292                 if (V3->nd != 5)
  295. 293                 {
  296. 294                     PyErr_Format(PyExc_RuntimeError,
  297. 295                                  "c_extract: Some CudaNdarray has rank %i, it was supposed to have rank 5",
  298. 296                                  V3->nd);
  299. 297                     V3 = NULL;
  300. 298                     {
  301. 299         __failure = 4;
  302. 300         if (!PyErr_Occurred()) {
  303. 301             PyErr_SetString(PyExc_RuntimeError,
  304. 302                 "Unexpected error in an Op's C code. "
  305. 303                 "No Python exception was set.");
  306. 304             }
  307. 305         goto __label_4;};
  308. 306                 }
  309. 307                 //std::cerr << "c_extract " << V3 << " nd check passed\n";
  310. 308            
  311. 309
  312. 310                 if (CudaNdarray_HOST_DIMS(V3)[2] != 1)
  313. 311                 {
  314. 312                     PyErr_Format(PyExc_RuntimeError,
  315. 313                                  "c_extract: Some CudaNdarray has dim %i on broadcastable dimension %i",
  316. 314                                  CudaNdarray_HOST_DIMS(V3)[2], 2);
  317. 315                     V3 = NULL;
  318. 316                     {
  319. 317         __failure = 4;
  320. 318         if (!PyErr_Occurred()) {
  321. 319             PyErr_SetString(PyExc_RuntimeError,
  322. 320                 "Unexpected error in an Op's C code. "
  323. 321                 "No Python exception was set.");
  324. 322             }
  325. 323         goto __label_4;};
  326. 324                 }
  327. 325                 //std::cerr << "c_extract " << V3 << "dim check 2 passed\n";
  328. 326                 //std::cerr << "c_extract " << V3 << "checking bcast 2 <" << V3->str<< ">\n";
  329. 327                 //std::cerr << "c_extract " << V3->str[2] << "\n";
  330. 328                 if (CudaNdarray_HOST_STRIDES(V3)[2])
  331. 329                 {
  332. 330                     //std::cerr << "c_extract bad stride detected...\n";
  333. 331                     PyErr_Format(PyExc_RuntimeError,
  334. 332                                  "c_extract: Some CudaNdarray has a nonzero stride %i on a broadcastable dimension %i",
  335. 333                                  CudaNdarray_HOST_STRIDES(V3)[2], 2);
  336. 334                     V3 = NULL;
  337. 335                     {
  338. 336         __failure = 4;
  339. 337         if (!PyErr_Occurred()) {
  340. 338             PyErr_SetString(PyExc_RuntimeError,
  341. 339                 "Unexpected error in an Op's C code. "
  342. 340                 "No Python exception was set.");
  343. 341             }
  344. 342         goto __label_4;};
  345. 343                 }
  346. 344                 //std::cerr << "c_extract " << V3 << "bcast check 2 passed\n";
  347. 345                    
  348. 346
  349. 347                 assert(V3);
  350. 348                 Py_INCREF(py_V3);
  351. 349             }
  352. 350             else if (py_V3 == Py_None)
  353. 351             {
  354. 352                 PyErr_SetString(PyExc_TypeError,
  355. 353                                 "expected a CudaNdarray, not None");
  356. 354                 V3 = NULL;
  357. 355                 {
  358. 356         __failure = 4;
  359. 357         if (!PyErr_Occurred()) {
  360. 358             PyErr_SetString(PyExc_RuntimeError,
  361. 359                 "Unexpected error in an Op's C code. "
  362. 360                 "No Python exception was set.");
  363. 361             }
  364. 362         goto __label_4;};
  365. 363             }
  366. 364             else
  367. 365             {
  368. 366                 //fprintf(stderr, "FAILING c_extract CNDA object w refcnt %p %i\n", py_V3, (py_V3->ob_refcnt));
  369. 367                 PyErr_SetString(PyExc_TypeError, "Argument not a CudaNdarray");
  370. 368                 V3 = NULL;
  371. 369                 {
  372. 370         __failure = 4;
  373. 371         if (!PyErr_Occurred()) {
  374. 372             PyErr_SetString(PyExc_RuntimeError,
  375. 373                 "Unexpected error in an Op's C code. "
  376. 374                 "No Python exception was set.");
  377. 375             }
  378. 376         goto __label_4;};
  379. 377             }
  380. 378             //std::cerr << "c_extract done " << V3 << '\n';
  381. 379            
  382. 380
  383. 381 {
  384. 382
  385. 383     py_V5 = PyList_GET_ITEM(storage_V5, 0);
  386. 384     {Py_XINCREF(py_V5);}
  387. 385    
  388. 386         assert(py_V5->ob_refcnt >= 2); // There should be at least one ref from the container object,
  389. 387         // and one ref from the local scope.
  390. 388
  391. 389         if (CudaNdarray_Check(py_V5))
  392. 390         {
  393. 391             //fprintf(stderr, "c_extract CNDA object w refcnt %p %i\n", py_V5, (py_V5->ob_refcnt));
  394. 392             V5 = (CudaNdarray*)py_V5;
  395. 393             //std::cerr << "c_extract " << V5 << '\n';
  396. 394        
  397. 395
  398. 396                 if (V5->nd != 5)
  399. 397                 {
  400. 398                     PyErr_Format(PyExc_RuntimeError,
  401. 399                                  "c_extract: Some CudaNdarray has rank %i, it was supposed to have rank 5",
  402. 400                                  V5->nd);
  403. 401                     V5 = NULL;
  404. 402                     {
  405. 403         __failure = 6;
  406. 404         if (!PyErr_Occurred()) {
  407. 405             PyErr_SetString(PyExc_RuntimeError,
  408. 406                 "Unexpected error in an Op's C code. "
  409. 407                 "No Python exception was set.");
  410. 408             }
  411. 409         goto __label_6;};
  412. 410                 }
  413. 411                 //std::cerr << "c_extract " << V5 << " nd check passed\n";
  414. 412            
  415. 413
  416. 414                 assert(V5);
  417. 415                 Py_INCREF(py_V5);
  418. 416             }
  419. 417             else if (py_V5 == Py_None)
  420. 418             {
  421. 419                 PyErr_SetString(PyExc_TypeError,
  422. 420                                 "expected a CudaNdarray, not None");
  423. 421                 V5 = NULL;
  424. 422                 {
  425. 423         __failure = 6;
  426. 424         if (!PyErr_Occurred()) {
  427. 425             PyErr_SetString(PyExc_RuntimeError,
  428. 426                 "Unexpected error in an Op's C code. "
  429. 427                 "No Python exception was set.");
  430. 428             }
  431. 429         goto __label_6;};
  432. 430             }
  433. 431             else
  434. 432             {
  435. 433                 //fprintf(stderr, "FAILING c_extract CNDA object w refcnt %p %i\n", py_V5, (py_V5->ob_refcnt));
  436. 434                 PyErr_SetString(PyExc_TypeError, "Argument not a CudaNdarray");
  437. 435                 V5 = NULL;
  438. 436                 {
  439. 437         __failure = 6;
  440. 438         if (!PyErr_Occurred()) {
  441. 439             PyErr_SetString(PyExc_RuntimeError,
  442. 440                 "Unexpected error in an Op's C code. "
  443. 441                 "No Python exception was set.");
  444. 442             }
  445. 443         goto __label_6;};
  446. 444             }
  447. 445             //std::cerr << "c_extract done " << V5 << '\n';
  448. 446            
  449. 447
  450. 448 {
  451. 449
  452. 450     py_V7 = PyList_GET_ITEM(storage_V7, 0);
  453. 451     {Py_XINCREF(py_V7);}
  454. 452    
  455. 453         assert(py_V7->ob_refcnt >= 2); // There should be at least one ref from the container object,
  456. 454         // and one ref from the local scope.
  457. 455
  458. 456         if (CudaNdarray_Check(py_V7))
  459. 457         {
  460. 458             //fprintf(stderr, "c_extract CNDA object w refcnt %p %i\n", py_V7, (py_V7->ob_refcnt));
  461. 459             V7 = (CudaNdarray*)py_V7;
  462. 460             //std::cerr << "c_extract " << V7 << '\n';
  463. 461        
  464. 462
  465. 463                 if (V7->nd != 1)
  466. 464                 {
  467. 465                     PyErr_Format(PyExc_RuntimeError,
  468. 466                                  "c_extract: Some CudaNdarray has rank %i, it was supposed to have rank 1",
  469. 467                                  V7->nd);
  470. 468                     V7 = NULL;
  471. 469                     {
  472. 470         __failure = 8;
  473. 471         if (!PyErr_Occurred()) {
  474. 472             PyErr_SetString(PyExc_RuntimeError,
  475. 473                 "Unexpected error in an Op's C code. "
  476. 474                 "No Python exception was set.");
  477. 475             }
  478. 476         goto __label_8;};
  479. 477                 }
  480. 478                 //std::cerr << "c_extract " << V7 << " nd check passed\n";
  481. 479            
  482. 480
  483. 481                 assert(V7);
  484. 482                 Py_INCREF(py_V7);
  485. 483             }
  486. 484             else if (py_V7 == Py_None)
  487. 485             {
  488. 486                 PyErr_SetString(PyExc_TypeError,
  489. 487                                 "expected a CudaNdarray, not None");
  490. 488                 V7 = NULL;
  491. 489                 {
  492. 490         __failure = 8;
  493. 491         if (!PyErr_Occurred()) {
  494. 492             PyErr_SetString(PyExc_RuntimeError,
  495. 493                 "Unexpected error in an Op's C code. "
  496. 494                 "No Python exception was set.");
  497. 495             }
  498. 496         goto __label_8;};
  499. 497             }
  500. 498             else
  501. 499             {
  502. 500                 //fprintf(stderr, "FAILING c_extract CNDA object w refcnt %p %i\n", py_V7, (py_V7->ob_refcnt));
  503. 501                 PyErr_SetString(PyExc_TypeError, "Argument not a CudaNdarray");
  504. 502                 V7 = NULL;
  505. 503                 {
  506. 504         __failure = 8;
  507. 505         if (!PyErr_Occurred()) {
  508. 506             PyErr_SetString(PyExc_RuntimeError,
  509. 507                 "Unexpected error in an Op's C code. "
  510. 508                 "No Python exception was set.");
  511. 509             }
  512. 510         goto __label_8;};
  513. 511             }
  514. 512             //std::cerr << "c_extract done " << V7 << '\n';
  515. 513            
  516. 514
  517. 515 {
  518. 516
  519. 517     py_V9 = PyList_GET_ITEM(storage_V9, 0);
  520. 518     {Py_XINCREF(py_V9);}
  521. 519    
  522. 520             V9 = NULL;
  523. 521             if (py_V9 == Py_None) {
  524. 522                 // We can either fail here or set V9 to NULL and rely on Ops
  525. 523                 // using tensors to handle the NULL case, but if they fail to do so
  526. 524                 // they'll end up with nasty segfaults, so this is public service.
  527. 525                 PyErr_SetString(PyExc_ValueError, "expected an ndarray, not None");
  528. 526                 {
  529. 527         __failure = 10;
  530. 528         if (!PyErr_Occurred()) {
  531. 529             PyErr_SetString(PyExc_RuntimeError,
  532. 530                 "Unexpected error in an Op's C code. "
  533. 531                 "No Python exception was set.");
  534. 532             }
  535. 533         goto __label_10;}
  536. 534             }
  537. 535             if (!PyArray_Check(py_V9)) {
  538. 536                 PyErr_SetString(PyExc_ValueError, "expected an ndarray");
  539. 537                 {
  540. 538         __failure = 10;
  541. 539         if (!PyErr_Occurred()) {
  542. 540             PyErr_SetString(PyExc_RuntimeError,
  543. 541                 "Unexpected error in an Op's C code. "
  544. 542                 "No Python exception was set.");
  545. 543             }
  546. 544         goto __label_10;}
  547. 545             }
  548. 546             // We expect NPY_INT64
  549. 547             if (!PyArray_ISALIGNED((PyArrayObject*) py_V9)) {
  550. 548                 PyArrayObject * tmp = (PyArrayObject*) py_V9;
  551. 549                 PyErr_Format(PyExc_NotImplementedError,
  552. 550                              "expected an aligned array of type %ld "
  553. 551                              "(NPY_INT64), got non-aligned array of type %ld"
  554. 552                              " with %ld dimensions, with 3 last dims "
  555. 553                              "%ld, %ld, %ld"
  556. 554                              " and 3 last strides %ld %ld, %ld.",
  557. 555                              (long int) NPY_INT64,
  558. 556                              (long int) PyArray_TYPE((PyArrayObject*) py_V9),
  559. 557                              (long int) PyArray_NDIM(tmp),
  560. 558                              (long int) PyArray_NDIM(tmp) >= 3 ?
  561. 559             PyArray_DIMS(tmp)[PyArray_NDIM(tmp)-3] : -1,
  562. 560                              (long int) PyArray_NDIM(tmp) >= 2 ?
  563. 561             PyArray_DIMS(tmp)[PyArray_NDIM(tmp)-2] : -1,
  564. 562                              (long int) PyArray_NDIM(tmp) >= 1 ?
  565. 563             PyArray_DIMS(tmp)[PyArray_NDIM(tmp)-1] : -1,
  566. 564                              (long int) PyArray_NDIM(tmp) >= 3 ?
  567. 565             PyArray_STRIDES(tmp)[PyArray_NDIM(tmp)-3] : -1,
  568. 566                              (long int) PyArray_NDIM(tmp) >= 2 ?
  569. 567             PyArray_STRIDES(tmp)[PyArray_NDIM(tmp)-2] : -1,
  570. 568                              (long int) PyArray_NDIM(tmp) >= 1 ?
  571. 569             PyArray_STRIDES(tmp)[PyArray_NDIM(tmp)-1] : -1
  572. 570             );
  573. 571                 {
  574. 572         __failure = 10;
  575. 573         if (!PyErr_Occurred()) {
  576. 574             PyErr_SetString(PyExc_RuntimeError,
  577. 575                 "Unexpected error in an Op's C code. "
  578. 576                 "No Python exception was set.");
  579. 577             }
  580. 578         goto __label_10;}
  581. 579             }
  582. 580             // This is a TypeError to be consistent with DEBUG_MODE
  583. 581             // Note: DEBUG_MODE also tells the name of the container
  584. 582             if (PyArray_TYPE((PyArrayObject*) py_V9) != NPY_INT64) {
  585. 583                 PyErr_Format(PyExc_TypeError,
  586. 584                              "expected type_num %d (NPY_INT64) got %d",
  587. 585                              NPY_INT64, PyArray_TYPE((PyArrayObject*) py_V9));
  588. 586                 {
  589. 587         __failure = 10;
  590. 588         if (!PyErr_Occurred()) {
  591. 589             PyErr_SetString(PyExc_RuntimeError,
  592. 590                 "Unexpected error in an Op's C code. "
  593. 591                 "No Python exception was set.");
  594. 592             }
  595. 593         goto __label_10;}
  596. 594             }
  597. 595            
  598. 596         V9 = (PyArrayObject*)(py_V9);
  599. 597         Py_XINCREF(V9);
  600. 598        
  601. 599 {
  602. 600 // Op class GpuConv3D
  603. 601
  604. 602                         ///////////// < code generated by GpuConv3D >
  605. 603
  606. 604                         //printf("              Conv3DGPU c code\n");
  607. 605
  608. 606                         //Check dimensionality of inputs
  609. 607                         if (CudaNdarray_NDIM(V5) != 5)
  610. 608                         {
  611. 609                 PyErr_Format(PyExc_ValueError, "GpuConv3D: W must be a 5 dimensional CudaNdarray");
  612. 610                             {
  613. 611         __failure = 11;
  614. 612         if (!PyErr_Occurred()) {
  615. 613             PyErr_SetString(PyExc_RuntimeError,
  616. 614                 "Unexpected error in an Op's C code. "
  617. 615                 "No Python exception was set.");
  618. 616             }
  619. 617         goto __label_11;}
  620. 618                         }
  621. 619
  622. 620                         if (CudaNdarray_NDIM(V3) != 5)
  623. 621                         {
  624. 622                 PyErr_Format(PyExc_ValueError, "GpuConv3D: V must be a 5 dimensional CudaNdarray");
  625. 623                             {
  626. 624         __failure = 11;
  627. 625         if (!PyErr_Occurred()) {
  628. 626             PyErr_SetString(PyExc_RuntimeError,
  629. 627                 "Unexpected error in an Op's C code. "
  630. 628                 "No Python exception was set.");
  631. 629             }
  632. 630         goto __label_11;}
  633. 631                         }
  634. 632
  635. 633                         if (CudaNdarray_NDIM(V7) != 1)
  636. 634                         {
  637. 635                 PyErr_Format(PyExc_ValueError, "GpuConv3D: b must be a vector CudaNdarray");
  638. 636                             {
  639. 637         __failure = 11;
  640. 638         if (!PyErr_Occurred()) {
  641. 639             PyErr_SetString(PyExc_RuntimeError,
  642. 640                 "Unexpected error in an Op's C code. "
  643. 641                 "No Python exception was set.");
  644. 642             }
  645. 643         goto __label_11;}
  646. 644                         }
  647. 645
  648. 646                         if (CudaNdarray_NDIM(V9) != 1)
  649. 647                         {
  650. 648 PyErr_Format(PyExc_ValueError, "GpuConv3D: d must be a vector CudaNdarray");
  651. 649                             {
  652. 650         __failure = 11;
  653. 651         if (!PyErr_Occurred()) {
  654. 652             PyErr_SetString(PyExc_RuntimeError,
  655. 653                 "Unexpected error in an Op's C code. "
  656. 654                 "No Python exception was set.");
  657. 655             }
  658. 656         goto __label_11;}
  659. 657
  660. 658                         }
  661. 659                         if (PyArray_DIMS(V9)[0] != 3)
  662. 660                         {
  663. 661                 PyErr_Format(PyExc_ValueError, "GpuConv3D: 3 stride length arguments expected (row, col, time) but %li were given", PyArray_DIMS(V9)[0]);
  664. 662                             {
  665. 663         __failure = 11;
  666. 664         if (!PyErr_Occurred()) {
  667. 665             PyErr_SetString(PyExc_RuntimeError,
  668. 666                 "Unexpected error in an Op's C code. "
  669. 667                 "No Python exception was set.");
  670. 668             }
  671. 669         goto __label_11;}
  672. 670
  673. 671                         }
  674. 672
  675. 673 { //extra scope so fail doesn't jump over declarations
  676. 674                         //Read and check sizes of inputs
  677. 675                         const int batchSize = CudaNdarray_HOST_DIMS(V3)[0];
  678. 676                         const int outputChannels =  CudaNdarray_HOST_DIMS(V5)[0];
  679. 677                         const int inputChannels = CudaNdarray_HOST_DIMS(V3)[4];
  680. 678                         if (CudaNdarray_HOST_DIMS(V5)[4] != inputChannels)
  681. 679                         {
  682. 680                             PyErr_Format(PyExc_ValueError, "GpuConv3D: W operates on a %i channel image but the image has %i channels",CudaNdarray_HOST_DIMS(V5)[4],inputChannels);
  683. 681                             {
  684. 682         __failure = 11;
  685. 683         if (!PyErr_Occurred()) {
  686. 684             PyErr_SetString(PyExc_RuntimeError,
  687. 685                 "Unexpected error in an Op's C code. "
  688. 686                 "No Python exception was set.");
  689. 687             }
  690. 688         goto __label_11;}
  691. 689                         }
  692. 690 {  //extra scope so error handler jumps don't cause errors
  693. 691                         const int filterHeight = CudaNdarray_HOST_DIMS(V5)[1];
  694. 692                         const int filterWidth = CudaNdarray_HOST_DIMS(V5)[2];
  695. 693                         const int filterDur = CudaNdarray_HOST_DIMS(V5)[3];
  696. 694                         const int vidHeight = CudaNdarray_HOST_DIMS(V3)[1];
  697. 695                         const int vidWidth = CudaNdarray_HOST_DIMS(V3)[2];
  698. 696                         const int vidDur = CudaNdarray_HOST_DIMS(V3)[3];
  699. 697             if (vidHeight < filterHeight)
  700. 698             {
  701. 699                 PyErr_Format(PyExc_ValueError, "GpuConv3D: W has a height of %i but V is only %i pixels tall",filterHeight,vidHeight);
  702. 700                 {
  703. 701         __failure = 11;
  704. 702         if (!PyErr_Occurred()) {
  705. 703             PyErr_SetString(PyExc_RuntimeError,
  706. 704                 "Unexpected error in an Op's C code. "
  707. 705                 "No Python exception was set.");
  708. 706             }
  709. 707         goto __label_11;}
  710. 708             }
  711. 709 { // extra scope so fail works
  712. 710             if (vidWidth < filterWidth)
  713. 711             {
  714. 712                 PyErr_Format(PyExc_ValueError, "GpuConv3D: W has a width of %i but V is only %i pixels wide",filterWidth,vidWidth);
  715. 713                 {
  716. 714         __failure = 11;
  717. 715         if (!PyErr_Occurred()) {
  718. 716             PyErr_SetString(PyExc_RuntimeError,
  719. 717                 "Unexpected error in an Op's C code. "
  720. 718                 "No Python exception was set.");
  721. 719             }
  722. 720         goto __label_11;}
  723. 721             }
  724. 722 { // extra scope so fail works
  725. 723             if (vidDur < filterDur)
  726. 724             {
  727. 725                 PyErr_Format(PyExc_ValueError, "GpuConv3D: W has a duration of %i but V is only %i pixels long",filterDur,vidDur);
  728. 726                 {
  729. 727         __failure = 11;
  730. 728         if (!PyErr_Occurred()) {
  731. 729             PyErr_SetString(PyExc_RuntimeError,
  732. 730                 "Unexpected error in an Op's C code. "
  733. 731                 "No Python exception was set.");
  734. 732             }
  735. 733         goto __label_11;}
  736. 734             }
  737. 735 { // extra scope so fail works
  738. 736
  739. 737                         //Read and check stride arguments
  740. 738                         const int dr = *(dtype_V9*)PyArray_GETPTR1(V9,0);
  741. 739                         const int dc = *(dtype_V9*)PyArray_GETPTR1(V9,1);
  742. 740                         const int dt = *(dtype_V9*)PyArray_GETPTR1(V9,2);
  743. 741                         if (dr <= 0 || dc <= 0 || dt <= 0)
  744. 742                         {
  745. 743                 PyErr_Format(PyExc_ValueError, "GpuConv3D: Strides must all be positive but are %i, %i, %i", dr, dc, dt);
  746. 744                 {
  747. 745         __failure = 11;
  748. 746         if (!PyErr_Occurred()) {
  749. 747             PyErr_SetString(PyExc_RuntimeError,
  750. 748                 "Unexpected error in an Op's C code. "
  751. 749                 "No Python exception was set.");
  752. 750             }
  753. 751         goto __label_11;}
  754. 752                         }
  755. 753 { // extra scope so fail works
  756. 754
  757. 755                         //Make correctly sized output
  758. 756                         const int outputHeight = int( (vidHeight - filterHeight) / dr )+1;
  759. 757                         const int outputWidth = int( (vidWidth - filterWidth) / dc )+1;
  760. 758                         const int outputDur = int( (vidDur - filterDur) / dt ) +1;
  761. 759
  762. 760                         npy_intp dims[5];
  763. 761                         dims[0] = batchSize;
  764. 762                         dims[4] = outputChannels;
  765. 763                         dims[1] = outputHeight;
  766. 764                         dims[2] = outputWidth;
  767. 765                         dims[3] = outputDur;
  768. 766
  769. 767                         if(!(V1) || CudaNdarray_HOST_DIMS(V1)[0]!=dims[0] ||
  770. 768                         CudaNdarray_HOST_DIMS(V1)[1]!=dims[1] ||
  771. 769                         CudaNdarray_HOST_DIMS(V1)[2]!=dims[2] ||
  772. 770                         CudaNdarray_HOST_DIMS(V1)[3]!=dims[3] ||
  773. 771                         CudaNdarray_HOST_DIMS(V1)[4]!=dims[4]){
  774. 772                                 Py_XDECREF(V1);
  775. 773                                 V1 = (CudaNdarray*)CudaNdarray_NewDims(5,dims);
  776. 774                                 if (!(V1)) {
  777. 775
  778. 776                     PyErr_Format(PyExc_MemoryError, "GpuConv3D: could not allocate output");
  779. 777                             {
  780. 778         __failure = 11;
  781. 779         if (!PyErr_Occurred()) {
  782. 780             PyErr_SetString(PyExc_RuntimeError,
  783. 781                 "Unexpected error in an Op's C code. "
  784. 782                 "No Python exception was set.");
  785. 783             }
  786. 784         goto __label_11;}
  787. 785                                 }
  788. 786                         }
  789. 787 { // extra scope so fail will not cross declarations
  790. 788                         //#define ELEM_AT(x, i) * ( dtype_ ## x *) ( x->data + (i) )####################
  791. 789
  792. 790                         const int ws4 = CudaNdarray_HOST_STRIDES(V5)[4];
  793. 791                         const int vs4 = CudaNdarray_HOST_STRIDES(V3)[4];
  794. 792                         const int ws3 = CudaNdarray_HOST_STRIDES(V5)[3];
  795. 793                         const int vs3 = CudaNdarray_HOST_STRIDES(V3)[3];
  796. 794                         const int ws2 = CudaNdarray_HOST_STRIDES(V5)[2];
  797. 795                         const int vs2 = CudaNdarray_HOST_STRIDES(V3)[2];
  798. 796                         const int ws1 = CudaNdarray_HOST_STRIDES(V5)[1];
  799. 797                         const int vs1 = CudaNdarray_HOST_STRIDES(V3)[1];
  800. 798                         const int ws0 = CudaNdarray_HOST_STRIDES(V5)[0];
  801. 799                         const int vs0 = CudaNdarray_HOST_STRIDES(V3)[0];
  802. 800
  803. 801                         // Compute H
  804. 802                         //H[i,x,y,t,j] = b_j + sum_k sum_l sum_m sum_z W[j,k,l,m,z] V[i, dr*r+k,dc*c+l,dt*t+m,z]
  805. 803
  806. 804 bool out_contiguous = CudaNdarray_is_c_contiguous(V1);
  807. 805 int version = -1;
  808. 806 int verbose = 0;
  809. 807 bool subsample =(dr>1)||(dc>1)||(dt>1);
  810. 808 bool b_strided = (CudaNdarray_HOST_STRIDES(V7)[0]!=1) && !(CudaNdarray_HOST_STRIDES(V7)[0]==0 && outputChannels==1);
  811. 809 bool work_complete = false;
  812. 810
  813. 811 if(out_contiguous && !b_strided && (version==0||version==-1) && outputDur<=512 && !work_complete){
  814. 812     //conv_rows_stack
  815. 813     dim3 grid(outputHeight*outputWidth,batchSize*outputChannels);
  816. 814     dim3 threads(outputDur);
  817. 815
  818. 816     int shared_size=0;
  819. 817         conv_rows_stack<<<grid, threads, shared_size>>>(
  820. 818         CudaNdarray_DEV_DATA(V3), CudaNdarray_DEV_DATA(V5), CudaNdarray_DEV_DATA(V7), CudaNdarray_DEV_DATA(V1),
  821. 819         vidHeight, vidWidth, vidDur,
  822. 820         filterHeight, filterWidth, filterDur,
  823. 821         outputChannels, inputChannels,
  824. 822         dr,dc,dt,
  825. 823         vs3,vs2,vs1,vs4,vs0,
  826. 824         ws3,ws2,ws1,ws4,ws0);
  827. 825
  828. 826         CNDA_THREAD_SYNC;
  829. 827         cudaError_t sts = cudaGetLastError();
  830. 828         if (cudaSuccess == sts)
  831. 829         {
  832. 830             work_complete = true;
  833. 831             if (verbose>1) printf("threads.x=%i, threads.y=%i, grid.x=%i, grid.y=%i, shared_size=%i, nb_threads=%i\n", threads.x, threads.y, grid.x, grid.y, shared_size, threads.x * threads.y);
  834. 832             if (verbose) printf("INFO: used 'conv_rows_stack' version\n");
  835. 833         }
  836. 834         else
  837. 835         {
  838. 836             if (verbose) printf("threads.x=%i, threads.y=%i, grid.x=%i, grid.y=%i, shared_size=%i, nb_threads=%i\n", threads.x, threads.y, grid.x, grid.y, shared_size, threads.x * threads.y);
  839. 837             if (verbose) printf("ERROR: all implementations failed for GpuConv3D! (%s)",cudaGetErrorString(sts));
  840. 838             PyErr_Format(PyExc_RuntimeError, "ERROR: all implementations failed for GpuConv3D! (%s)",
  841. 839                     cudaGetErrorString(sts));
  842. 840             {
  843. 841         __failure = 11;
  844. 842         if (!PyErr_Occurred()) {
  845. 843             PyErr_SetString(PyExc_RuntimeError,
  846. 844                 "Unexpected error in an Op's C code. "
  847. 845                 "No Python exception was set.");
  848. 846             }
  849. 847         goto __label_11;}
  850. 848         }
  851. 849
  852. 850
  853. 851 }
  854. 852
  855. 853 if(!work_complete){
  856. 854             PyErr_Format(PyExc_RuntimeError, "ERROR: no implementations executed for this GpuConv3D!");
  857. 855             {
  858. 856         __failure = 11;
  859. 857         if (!PyErr_Occurred()) {
  860. 858             PyErr_SetString(PyExc_RuntimeError,
  861. 859                 "Unexpected error in an Op's C code. "
  862. 860                 "No Python exception was set.");
  863. 861             }
  864. 862         goto __label_11;}
  865. 863 }
  866. 864
  867. 865 }}}}}}} //extra scope so error handler jumps don't cross declarations
  868. 866                         ///////////// < /code generated by GpuConv3D >
  869. 867         __label_11:
  870. 868
  871. 869 double __DUMMY_11;
  872. 870
  873. 871 }
  874. 872 __label_10:
  875. 873
  876. 874         if (V9) {
  877. 875             Py_XDECREF(V9);
  878. 876         }
  879. 877        
  880. 878     {Py_XDECREF(py_V9);}
  881. 879    
  882. 880 double __DUMMY_10;
  883. 881
  884. 882 }
  885. 883 __label_8:
  886. 884
  887. 885         //std::cerr << "cleanup " << py_V7 << " " << V7 << "\n";
  888. 886         //fprintf(stderr, "c_cleanup CNDA py_object w refcnt %p %i\n", py_V7, (py_V7->ob_refcnt));
  889. 887         if (V7)
  890. 888         {
  891. 889             //fprintf(stderr, "c_cleanup CNDA cn_object w refcnt %p %i\n", V7, (V7->ob_refcnt));
  892. 890             Py_XDECREF(V7);
  893. 891         }
  894. 892         //std::cerr << "cleanup done" << py_V7 << "\n";
  895. 893        
  896. 894     {Py_XDECREF(py_V7);}
  897. 895    
  898. 896 double __DUMMY_8;
  899. 897
  900. 898 }
  901. 899 __label_6:
  902. 900
  903. 901         //std::cerr << "cleanup " << py_V5 << " " << V5 << "\n";
  904. 902         //fprintf(stderr, "c_cleanup CNDA py_object w refcnt %p %i\n", py_V5, (py_V5->ob_refcnt));
  905. 903         if (V5)
  906. 904         {
  907. 905             //fprintf(stderr, "c_cleanup CNDA cn_object w refcnt %p %i\n", V5, (V5->ob_refcnt));
  908. 906             Py_XDECREF(V5);
  909. 907         }
  910. 908         //std::cerr << "cleanup done" << py_V5 << "\n";
  911. 909        
  912. 910     {Py_XDECREF(py_V5);}
  913. 911    
  914. 912 double __DUMMY_6;
  915. 913
  916. 914 }
  917. 915 __label_4:
  918. 916
  919. 917         //std::cerr << "cleanup " << py_V3 << " " << V3 << "\n";
  920. 918         //fprintf(stderr, "c_cleanup CNDA py_object w refcnt %p %i\n", py_V3, (py_V3->ob_refcnt));
  921. 919         if (V3)
  922. 920         {
  923. 921             //fprintf(stderr, "c_cleanup CNDA cn_object w refcnt %p %i\n", V3, (V3->ob_refcnt));
  924. 922             Py_XDECREF(V3);
  925. 923         }
  926. 924         //std::cerr << "cleanup done" << py_V3 << "\n";
  927. 925        
  928. 926     {Py_XDECREF(py_V3);}
  929. 927    
  930. 928 double __DUMMY_4;
  931. 929
  932. 930 }
  933. 931 __label_2:
  934. 932
  935. 933     if (!__failure) {
  936. 934      
  937. 935         //std::cerr << "sync\n";
  938. 936         if (NULL == V1) {
  939. 937             // failure: sync None to storage
  940. 938             Py_XDECREF(py_V1);
  941. 939             py_V1 = Py_None;
  942. 940             Py_INCREF(py_V1);
  943. 941         }
  944. 942         else
  945. 943         {
  946. 944             if (py_V1 != (PyObject*)V1)
  947. 945             {
  948. 946                 Py_XDECREF(py_V1);
  949. 947                 py_V1 = (PyObject*)V1;
  950. 948                 Py_INCREF(py_V1);
  951. 949             }
  952. 950             assert(py_V1->ob_refcnt);
  953. 951         }
  954. 952        
  955. 953       PyObject* old = PyList_GET_ITEM(storage_V1, 0);
  956. 954       {Py_XINCREF(py_V1);}
  957. 955       PyList_SET_ITEM(storage_V1, 0, py_V1);
  958. 956       {Py_XDECREF(old);}
  959. 957     }
  960. 958    
  961. 959         //std::cerr << "cleanup " << py_V1 << " " << V1 << "\n";
  962. 960         //fprintf(stderr, "c_cleanup CNDA py_object w refcnt %p %i\n", py_V1, (py_V1->ob_refcnt));
  963. 961         if (V1)
  964. 962         {
  965. 963             //fprintf(stderr, "c_cleanup CNDA cn_object w refcnt %p %i\n", V1, (V1->ob_refcnt));
  966. 964             Py_XDECREF(V1);
  967. 965         }
  968. 966         //std::cerr << "cleanup done" << py_V1 << "\n";
  969. 967        
  970. 968     {Py_XDECREF(py_V1);}
  971. 969    
  972. 970 double __DUMMY_2;
  973. 971
  974. 972 }
  975. 973
  976. 974            
  977. 975         if (__failure) {
  978. 976             // When there is a failure, this code puts the exception
  979. 977             // in __ERROR.
  980. 978             PyObject* err_type = NULL;
  981. 979             PyObject* err_msg = NULL;
  982. 980             PyObject* err_traceback = NULL;
  983. 981             PyErr_Fetch(&err_type, &err_msg, &err_traceback);
  984. 982             if (!err_type) {err_type = Py_None;Py_INCREF(Py_None);}
  985. 983             if (!err_msg) {err_msg = Py_None; Py_INCREF(Py_None);}
  986. 984             if (!err_traceback) {err_traceback = Py_None; Py_INCREF(Py_None);}
  987. 985             PyObject* old_err_type = PyList_GET_ITEM(__ERROR, 0);
  988. 986             PyObject* old_err_msg = PyList_GET_ITEM(__ERROR, 1);
  989. 987             PyObject* old_err_traceback = PyList_GET_ITEM(__ERROR, 2);
  990. 988             PyList_SET_ITEM(__ERROR, 0, err_type);
  991. 989             PyList_SET_ITEM(__ERROR, 1, err_msg);
  992. 990             PyList_SET_ITEM(__ERROR, 2, err_traceback);
  993. 991             {Py_XDECREF(old_err_type);}
  994. 992             {Py_XDECREF(old_err_msg);}
  995. 993             {Py_XDECREF(old_err_traceback);}
  996. 994         }
  997. 995         // The failure code is returned to index what code block failed.
  998. 996         return __failure;
  999. 997        
  1000. 998         }
  1001. 999     };
  1002. 1000     }
  1003. 1001    
  1004. 1002
  1005. 1003         static int __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae_executor(__struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae* self) {
  1006. 1004             return self->run();
  1007. 1005         }
  1008. 1006
  1009. 1007         static void __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae_destructor(void* executor, void* self) {
  1010. 1008             delete ((__struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae*)self);
  1011. 1009         }
  1012. 1010        
  1013. 1011 //////////////////////
  1014. 1012 ////  Functions
  1015. 1013 //////////////////////
  1016. 1014 static PyObject * instantiate(PyObject * self, PyObject *argtuple) {
  1017. 1015   assert(PyTuple_Check(argtuple));
  1018. 1016   if (6 != PyTuple_Size(argtuple)){
  1019. 1017      PyErr_Format(PyExc_TypeError, "Wrong number of arguments, expected 6, got %i", (int)PyTuple_Size(argtuple));
  1020. 1018      return NULL;
  1021. 1019   }
  1022. 1020   __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae* struct_ptr = new __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae();
  1023. 1021   struct_ptr->init( PyTuple_GET_ITEM(argtuple, 0),PyTuple_GET_ITEM(argtuple, 1),PyTuple_GET_ITEM(argtuple, 2),PyTuple_GET_ITEM(argtuple, 3),PyTuple_GET_ITEM(argtuple, 4),PyTuple_GET_ITEM(argtuple, 5) );
  1024. 1022   PyObject* thunk = PyCObject_FromVoidPtrAndDesc((void*)(&__struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae_executor), struct_ptr, __struct_compiled_op_fb52f4ac7a6c97b676900c51ebdcc3ae_destructor);
  1025. 1023   return thunk; }
  1026. 1024
  1027. 1025 //////////////////////
  1028. 1026 ////  Module init
  1029. 1027 //////////////////////
  1030. 1028 static PyMethodDef MyMethods[] = {
  1031. 1029    {"instantiate", instantiate, METH_VARARGS, "undocumented"} ,
  1032. 1030    {NULL, NULL, 0, NULL}
  1033. 1031 };
  1034. 1032 PyMODINIT_FUNC initfb52f4ac7a6c97b676900c51ebdcc3ae(void){
  1035. 1033    import_array();
  1036. 1034    (void) Py_InitModule("fb52f4ac7a6c97b676900c51ebdcc3ae", MyMethods);
  1037. 1035 }
  1038. 1036
  1039. ===============================
  1040. In file included from /opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7/Python.h:8,
  1041.                 from mod.cu:1:
  1042. /opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7/pyconfig.h:1179:1: warning: "_POSIX_C_SOURCE" redefined
  1043. In file included from /opt/apps/cuda/5.5/bin/..//include/host_config.h:119,
  1044.                 from /opt/apps/cuda/5.5/bin/..//include/cuda_runtime.h:59,
  1045.                 from <command-line>:0:
  1046. /usr/include/features.h:162:1: warning: this is the location of the previous definition
  1047. In file included from /opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7/Python.h:8,
  1048.                 from mod.cu:1:
  1049. /opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7/pyconfig.h:1201:1: warning: "_XOPEN_SOURCE" redefined
  1050. In file included from /opt/apps/cuda/5.5/bin/..//include/host_config.h:119,
  1051.                 from /opt/apps/cuda/5.5/bin/..//include/cuda_runtime.h:59,
  1052.                 from <command-line>:0:
  1053. /usr/include/features.h:164:1: warning: this is the location of the previous definition
  1054. mod.cu(646): error: class "tagPyArrayObject" has no member "nd"
  1055.  
  1056. ['nvcc', '-shared', '-g1 error detected in the compilation of "/tmp/tmpxft_0000c10d_00000000-8_mod.cpp1.ii".
  1057. ', '-O3', '-use_fast_math', '-arch=sm_35', '-m64', '-Xcompiler', '-fno-math-errno,-Wno-unused-label,-Wno-unused-variable,-Wno-write-strings,-DCUDA_NDARRAY_CUH=aea3043c359bf8caa2336b3dc5baec95,-D NPY_NO_DEPRECATED_API=NPY_1_7_API_VERSION,-fPIC', '-Xlinker', '-rpath,/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray', '-I/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray', '-I=/opt/apps/cuda/5.5/bin/include', '-I/opt/apps/intel14/mvapich2_2_0/python/2.7.6/lib/python2.7/site-packages/numpy/core/include', '-I/opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7', '-I/home1/02246/srm89/Theano-3dfftconv/theano/sandbox/cuda', '-o', '/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/tmpGsO6tr/fb52f4ac7a6c97b676900c51ebdcc3ae.so', 'mod.cu', '-L/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray', '-L/opt/apps/intel14/mvapich2_2_0/python/2.7.6/lib', '-lpython2.7', '-lcudart', '-lcublas', '-lcuda_ndarray']
  1058. Traceback (most recent call last):
  1059.  File "lenet_final.py", line 304, in <module>
  1060.    evaluate_lenet5()
  1061.  File "lenet_final.py", line 224, in evaluate_lenet5
  1062.    y: test_set_y[index * batch_size: (index + 1) * batch_size]},mode=mode_fft)
  1063.  File "/home1/02246/srm89/Theano-3dfftconv/theano/compile/function.py", line 261, in function
  1064.    profile=profile)
  1065.  File "/home1/02246/srm89/Theano-3dfftconv/theano/compile/pfunc.py", line 511, in pfunc
  1066.    on_unused_input=on_unused_input)
  1067.  File "/home1/02246/srm89/Theano-3dfftconv/theano/compile/function_module.py", line 1332, in orig_function
  1068.    defaults)
  1069.  File "/home1/02246/srm89/Theano-3dfftconv/theano/compile/function_module.py", line 1198, in create
  1070.    _fn, _i, _o = self.linker.make_thunk(input_storage=input_storage_lists)
  1071.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/link.py", line 489, in make_thunk
  1072.    output_storage=output_storage)[:3]
  1073.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/vm.py", line 888, in make_all
  1074.    no_recycling))
  1075.  File "/home1/02246/srm89/Theano-3dfftconv/theano/sandbox/cuda/__init__.py", line 236, in make_thunk
  1076.    compute_map, no_recycling)
  1077.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/op.py", line 618, in make_thunk
  1078.    output_storage=node_output_storage)
  1079.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/cc.py", line 992, in make_thunk
  1080.    keep_lock=keep_lock)
  1081.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/cc.py", line 935, in __compile__
  1082.    keep_lock=keep_lock)
  1083.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/cc.py", line 1372, in cthunk_factory
  1084.    key=key, fn=self.compile_cmodule_by_step, keep_lock=keep_lock)
  1085.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/cmodule.py", line 1005, in module_from_key
  1086.    module = next(compile_steps)
  1087.  File "/home1/02246/srm89/Theano-3dfftconv/theano/gof/cc.py", line 1287, in compile_cmodule_by_step
  1088.    preargs=preargs)
  1089.  File "/home1/02246/srm89/Theano-3dfftconv/theano/sandbox/cuda/nvcc_compiler.py", line 414, in compile_str
  1090.    'for cmd', ' '.join(cmd))
  1091. Exception: ('The following error happened while compiling the node', GpuConv3D(GpuReshape{5}.0, <CudaNdarrayType(float32, 5D)>, <CudaNdarrayType(float32, vector)>, <TensorType(int64, vector)>), '\n', 'nvcc return status', 2, 'for cmd', 'nvcc -shared -g -O3 -use_fast_math -arch=sm_35 -m64 -Xcompiler -fno-math-errno,-Wno-unused-label,-Wno-unused-variable,-Wno-write-strings,-DCUDA_NDARRAY_CUH=aea3043c359bf8caa2336b3dc5baec95,-D NPY_NO_DEPRECATED_API=NPY_1_7_API_VERSION,-fPIC -Xlinker -rpath,/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray -I/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray -I=/opt/apps/cuda/5.5/bin/include -I/opt/apps/intel14/mvapich2_2_0/python/2.7.6/lib/python2.7/site-packages/numpy/core/include -I/opt/apps/intel14/mvapich2_2_0/python/2.7.6/include/python2.7 -I/home1/02246/srm89/Theano-3dfftconv/theano/sandbox/cuda -o /home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/tmpGsO6tr/fb52f4ac7a6c97b676900c51ebdcc3ae.so mod.cu -L/home1/02246/srm89/.theano/compiledir_Linux-2.6.32-431.17.1.el6.x86_64-x86_64-with-centos-6.5-Final-x86_64-2.7.6-64/cuda_ndarray -L/opt/apps/intel14/mvapich2_2_0/python/2.7.6/lib -lpython2.7 -lcudart -lcublas -lcuda_ndarray', '[GpuConv3D(<CudaNdarrayType(float32, (False, False, True, False, False))>, <CudaNdarrayType(float32, 5D)>, <CudaNdarrayType(float32, vector)>, <TensorType(int64, vector)>)]')
  1092. Skipping empty Profile
Advertisement
Add Comment
Please, Sign In to add comment
Public Pastes
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!