API tools faq
paste
rajeswar

Untitled

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