csr_interface.c 24 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735
  1. /* StarPU --- Runtime system for heterogeneous multicore architectures.
  2. *
  3. * Copyright (C) 2009-2012 Université de Bordeaux 1
  4. * Copyright (C) 2010 Mehdi Juhoor <mjuhoor@gmail.com>
  5. * Copyright (C) 2010, 2011, 2012 Centre National de la Recherche Scientifique
  6. *
  7. * StarPU is free software; you can redistribute it and/or modify
  8. * it under the terms of the GNU Lesser General Public License as published by
  9. * the Free Software Foundation; either version 2.1 of the License, or (at
  10. * your option) any later version.
  11. *
  12. * StarPU is distributed in the hope that it will be useful, but
  13. * WITHOUT ANY WARRANTY; without even the implied warranty of
  14. * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
  15. *
  16. * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  17. */
  18. #include <starpu.h>
  19. #include <common/config.h>
  20. #include <datawizard/coherency.h>
  21. #include <datawizard/copy_driver.h>
  22. #include <datawizard/filters.h>
  23. #include <starpu_hash.h>
  24. #include <starpu_cuda.h>
  25. #include <starpu_opencl.h>
  26. #include <drivers/opencl/driver_opencl.h>
  27. static int copy_ram_to_ram(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  28. #ifdef STARPU_USE_CUDA
  29. static int copy_ram_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  30. static int copy_cuda_to_ram(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  31. static int copy_cuda_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  32. static int copy_cuda_to_ram_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  33. static int copy_ram_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  34. static int copy_cuda_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  35. #endif
  36. #ifdef STARPU_USE_OPENCL
  37. static int copy_ram_to_opencl(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  38. static int copy_opencl_to_ram(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node);
  39. #endif
  40. static struct starpu_data_copy_methods csr_copy_data_methods_s =
  41. {
  42. .ram_to_ram = copy_ram_to_ram,
  43. .ram_to_spu = NULL,
  44. #ifdef STARPU_USE_CUDA
  45. .ram_to_cuda = copy_ram_to_cuda,
  46. .cuda_to_ram = copy_cuda_to_ram,
  47. .cuda_to_cuda = copy_cuda_to_cuda,
  48. .ram_to_cuda_async = copy_ram_to_cuda_async,
  49. .cuda_to_ram_async = copy_cuda_to_ram_async,
  50. .cuda_to_cuda_async = copy_cuda_to_cuda_async,
  51. #endif
  52. #ifdef STARPU_USE_OPENCL
  53. .ram_to_opencl = copy_ram_to_opencl,
  54. .opencl_to_ram = copy_opencl_to_ram,
  55. #endif
  56. .cuda_to_spu = NULL,
  57. .spu_to_ram = NULL,
  58. .spu_to_cuda = NULL,
  59. .spu_to_spu = NULL
  60. };
  61. static void register_csr_handle(starpu_data_handle_t handle, uint32_t home_node, void *data_interface);
  62. static ssize_t allocate_csr_buffer_on_node(void *data_interface_, uint32_t dst_node);
  63. static void free_csr_buffer_on_node(void *data_interface, uint32_t node);
  64. static size_t csr_interface_get_size(starpu_data_handle_t handle);
  65. static int csr_compare(void *data_interface_a, void *data_interface_b);
  66. static uint32_t footprint_csr_interface_crc32(starpu_data_handle_t handle);
  67. static struct starpu_data_interface_ops interface_csr_ops =
  68. {
  69. .register_data_handle = register_csr_handle,
  70. .allocate_data_on_node = allocate_csr_buffer_on_node,
  71. .free_data_on_node = free_csr_buffer_on_node,
  72. .copy_methods = &csr_copy_data_methods_s,
  73. .get_size = csr_interface_get_size,
  74. .interfaceid = STARPU_CSR_INTERFACE_ID,
  75. .interface_size = sizeof(struct starpu_csr_interface),
  76. .footprint = footprint_csr_interface_crc32,
  77. .compare = csr_compare,
  78. };
  79. static void register_csr_handle(starpu_data_handle_t handle, uint32_t home_node, void *data_interface)
  80. {
  81. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *) data_interface;
  82. unsigned node;
  83. for (node = 0; node < STARPU_MAXNODES; node++)
  84. {
  85. struct starpu_csr_interface *local_interface = (struct starpu_csr_interface *)
  86. starpu_data_get_interface_on_node(handle, node);
  87. if (node == home_node)
  88. {
  89. local_interface->nzval = csr_interface->nzval;
  90. local_interface->colind = csr_interface->colind;
  91. }
  92. else
  93. {
  94. local_interface->nzval = 0;
  95. local_interface->colind = NULL;
  96. }
  97. local_interface->rowptr = csr_interface->rowptr;
  98. local_interface->nnz = csr_interface->nnz;
  99. local_interface->nrow = csr_interface->nrow;
  100. local_interface->firstentry = csr_interface->firstentry;
  101. local_interface->elemsize = csr_interface->elemsize;
  102. }
  103. }
  104. /* declare a new data with the BLAS interface */
  105. void starpu_csr_data_register(starpu_data_handle_t *handleptr, uint32_t home_node,
  106. uint32_t nnz, uint32_t nrow, uintptr_t nzval, uint32_t *colind, uint32_t *rowptr, uint32_t firstentry, size_t elemsize)
  107. {
  108. struct starpu_csr_interface csr_interface =
  109. {
  110. .nnz = nnz,
  111. .nrow = nrow,
  112. .nzval = nzval,
  113. .colind = colind,
  114. .rowptr = rowptr,
  115. .firstentry = firstentry,
  116. .elemsize = elemsize
  117. };
  118. starpu_data_register(handleptr, home_node, &csr_interface, &interface_csr_ops);
  119. }
  120. static uint32_t footprint_csr_interface_crc32(starpu_data_handle_t handle)
  121. {
  122. return starpu_crc32_be(starpu_csr_get_nnz(handle), 0);
  123. }
  124. static int csr_compare(void *data_interface_a, void *data_interface_b)
  125. {
  126. struct starpu_csr_interface *csr_a = (struct starpu_csr_interface *) data_interface_a;
  127. struct starpu_csr_interface *csr_b = (struct starpu_csr_interface *) data_interface_b;
  128. /* Two matricess are considered compatible if they have the same size */
  129. return ((csr_a->nnz == csr_b->nnz)
  130. && (csr_a->nrow == csr_b->nrow)
  131. && (csr_a->elemsize == csr_b->elemsize));
  132. }
  133. /* offer an access to the data parameters */
  134. uint32_t starpu_csr_get_nnz(starpu_data_handle_t handle)
  135. {
  136. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  137. starpu_data_get_interface_on_node(handle, 0);
  138. return csr_interface->nnz;
  139. }
  140. uint32_t starpu_csr_get_nrow(starpu_data_handle_t handle)
  141. {
  142. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  143. starpu_data_get_interface_on_node(handle, 0);
  144. return csr_interface->nrow;
  145. }
  146. uint32_t starpu_csr_get_firstentry(starpu_data_handle_t handle)
  147. {
  148. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  149. starpu_data_get_interface_on_node(handle, 0);
  150. return csr_interface->firstentry;
  151. }
  152. size_t starpu_csr_get_elemsize(starpu_data_handle_t handle)
  153. {
  154. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  155. starpu_data_get_interface_on_node(handle, 0);
  156. return csr_interface->elemsize;
  157. }
  158. uintptr_t starpu_csr_get_local_nzval(starpu_data_handle_t handle)
  159. {
  160. unsigned node;
  161. node = _starpu_get_local_memory_node();
  162. STARPU_ASSERT(starpu_data_test_if_allocated_on_node(handle, node));
  163. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  164. starpu_data_get_interface_on_node(handle, node);
  165. return csr_interface->nzval;
  166. }
  167. uint32_t *starpu_csr_get_local_colind(starpu_data_handle_t handle)
  168. {
  169. unsigned node;
  170. node = _starpu_get_local_memory_node();
  171. STARPU_ASSERT(starpu_data_test_if_allocated_on_node(handle, node));
  172. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  173. starpu_data_get_interface_on_node(handle, node);
  174. return csr_interface->colind;
  175. }
  176. uint32_t *starpu_csr_get_local_rowptr(starpu_data_handle_t handle)
  177. {
  178. unsigned node;
  179. node = _starpu_get_local_memory_node();
  180. STARPU_ASSERT(starpu_data_test_if_allocated_on_node(handle, node));
  181. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *)
  182. starpu_data_get_interface_on_node(handle, node);
  183. return csr_interface->rowptr;
  184. }
  185. static size_t csr_interface_get_size(starpu_data_handle_t handle)
  186. {
  187. size_t size;
  188. uint32_t nnz = starpu_csr_get_nnz(handle);
  189. uint32_t nrow = starpu_csr_get_nrow(handle);
  190. size_t elemsize = starpu_csr_get_elemsize(handle);
  191. size = nnz*elemsize + nnz*sizeof(uint32_t) + (nrow+1)*sizeof(uint32_t);
  192. return size;
  193. }
  194. /* memory allocation/deallocation primitives for the BLAS interface */
  195. /* returns the size of the allocated area */
  196. static ssize_t allocate_csr_buffer_on_node(void *data_interface_, uint32_t dst_node)
  197. {
  198. uintptr_t addr_nzval = 0;
  199. uint32_t *addr_colind = NULL, *addr_rowptr = NULL;
  200. ssize_t allocated_memory;
  201. /* we need the 3 arrays to be allocated */
  202. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *) data_interface_;
  203. uint32_t nnz = csr_interface->nnz;
  204. uint32_t nrow = csr_interface->nrow;
  205. size_t elemsize = csr_interface->elemsize;
  206. enum starpu_node_kind kind = starpu_node_get_kind(dst_node);
  207. switch(kind)
  208. {
  209. case STARPU_CPU_RAM:
  210. addr_nzval = (uintptr_t)malloc(nnz*elemsize);
  211. if (!addr_nzval)
  212. goto fail_nzval;
  213. addr_colind = (uint32_t *) malloc(nnz*sizeof(uint32_t));
  214. if (!addr_colind)
  215. goto fail_colind;
  216. addr_rowptr = (uint32_t *) malloc((nrow+1)*sizeof(uint32_t));
  217. if (!addr_rowptr)
  218. goto fail_rowptr;
  219. break;
  220. #ifdef STARPU_USE_CUDA
  221. case STARPU_CUDA_RAM:
  222. cudaMalloc((void **)&addr_nzval, nnz*elemsize);
  223. if (!addr_nzval)
  224. goto fail_nzval;
  225. cudaMalloc((void **)&addr_colind, nnz*sizeof(uint32_t));
  226. if (!addr_colind)
  227. goto fail_colind;
  228. cudaMalloc((void **)&addr_rowptr, (nrow+1)*sizeof(uint32_t));
  229. if (!addr_rowptr)
  230. goto fail_rowptr;
  231. break;
  232. #endif
  233. #ifdef STARPU_USE_OPENCL
  234. case STARPU_OPENCL_RAM:
  235. {
  236. int ret;
  237. cl_mem ptr;
  238. ret = starpu_opencl_allocate_memory(&ptr, nnz*elemsize, CL_MEM_READ_WRITE);
  239. addr_nzval = (uintptr_t)ptr;
  240. if (ret) goto fail_nzval;
  241. ret = starpu_opencl_allocate_memory(&ptr, nnz*sizeof(uint32_t), CL_MEM_READ_WRITE);
  242. addr_colind = (void*) ptr;
  243. if (ret) goto fail_colind;
  244. ret = starpu_opencl_allocate_memory(&ptr, (nrow+1)*sizeof(uint32_t), CL_MEM_READ_WRITE);
  245. addr_rowptr = (void*) ptr;
  246. if (ret) goto fail_rowptr;
  247. break;
  248. }
  249. #endif
  250. default:
  251. STARPU_ABORT();
  252. }
  253. /* allocation succeeded */
  254. allocated_memory =
  255. nnz*elemsize + nnz*sizeof(uint32_t) + (nrow+1)*sizeof(uint32_t);
  256. /* update the data properly in consequence */
  257. csr_interface->nzval = addr_nzval;
  258. csr_interface->colind = addr_colind;
  259. csr_interface->rowptr = addr_rowptr;
  260. return allocated_memory;
  261. fail_rowptr:
  262. switch(kind)
  263. {
  264. case STARPU_CPU_RAM:
  265. free((void *)addr_colind);
  266. #ifdef STARPU_USE_CUDA
  267. case STARPU_CUDA_RAM:
  268. {
  269. cudaError_t err;
  270. err = cudaFree((void*)addr_colind);
  271. if (STARPU_UNLIKELY(err != cudaSuccess))
  272. STARPU_CUDA_REPORT_ERROR(err);
  273. break;
  274. }
  275. #endif
  276. #ifdef STARPU_USE_OPENCL
  277. case STARPU_OPENCL_RAM:
  278. clReleaseMemObject((void*)addr_colind);
  279. break;
  280. #endif
  281. default:
  282. STARPU_ABORT();
  283. }
  284. fail_colind:
  285. switch(kind)
  286. {
  287. case STARPU_CPU_RAM:
  288. free((void *)addr_nzval);
  289. #ifdef STARPU_USE_CUDA
  290. case STARPU_CUDA_RAM:
  291. {
  292. cudaError_t err;
  293. err = cudaFree((void*)addr_nzval);
  294. if (STARPU_UNLIKELY(err != cudaSuccess))
  295. STARPU_CUDA_REPORT_ERROR(err);
  296. break;
  297. }
  298. #endif
  299. #ifdef STARPU_USE_OPENCL
  300. case STARPU_OPENCL_RAM:
  301. clReleaseMemObject((void*)addr_nzval);
  302. break;
  303. #endif
  304. default:
  305. STARPU_ABORT();
  306. }
  307. fail_nzval:
  308. /* allocation failed */
  309. return -ENOMEM;
  310. }
  311. static void free_csr_buffer_on_node(void *data_interface, uint32_t node)
  312. {
  313. struct starpu_csr_interface *csr_interface = (struct starpu_csr_interface *) data_interface;
  314. enum starpu_node_kind kind = starpu_node_get_kind(node);
  315. switch(kind)
  316. {
  317. case STARPU_CPU_RAM:
  318. free((void*)csr_interface->nzval);
  319. free((void*)csr_interface->colind);
  320. free((void*)csr_interface->rowptr);
  321. break;
  322. #ifdef STARPU_USE_CUDA
  323. case STARPU_CUDA_RAM:
  324. {
  325. cudaError_t err;
  326. err = cudaFree((void*)csr_interface->nzval);
  327. if (STARPU_UNLIKELY(err != cudaSuccess))
  328. STARPU_CUDA_REPORT_ERROR(err);
  329. err = cudaFree((void*)csr_interface->colind);
  330. if (STARPU_UNLIKELY(err != cudaSuccess))
  331. STARPU_CUDA_REPORT_ERROR(err);
  332. err = cudaFree((void*)csr_interface->rowptr);
  333. if (STARPU_UNLIKELY(err != cudaSuccess))
  334. STARPU_CUDA_REPORT_ERROR(err);
  335. break;
  336. }
  337. #endif
  338. #ifdef STARPU_USE_OPENCL
  339. case STARPU_OPENCL_RAM:
  340. clReleaseMemObject((void*)csr_interface->nzval);
  341. clReleaseMemObject((void*)csr_interface->colind);
  342. clReleaseMemObject((void*)csr_interface->rowptr);
  343. break;
  344. #endif
  345. default:
  346. STARPU_ABORT();
  347. }
  348. }
  349. #ifdef STARPU_USE_CUDA
  350. static int copy_cuda_common(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED, enum cudaMemcpyKind kind)
  351. {
  352. struct starpu_csr_interface *src_csr = src_interface;
  353. struct starpu_csr_interface *dst_csr = dst_interface;
  354. uint32_t nnz = src_csr->nnz;
  355. uint32_t nrow = src_csr->nrow;
  356. size_t elemsize = src_csr->elemsize;
  357. cudaError_t cures;
  358. cures = cudaMemcpy((char *)dst_csr->nzval, (char *)src_csr->nzval, nnz*elemsize, kind);
  359. if (STARPU_UNLIKELY(cures))
  360. STARPU_CUDA_REPORT_ERROR(cures);
  361. cures = cudaMemcpy((char *)dst_csr->colind, (char *)src_csr->colind, nnz*sizeof(uint32_t), kind);
  362. if (STARPU_UNLIKELY(cures))
  363. STARPU_CUDA_REPORT_ERROR(cures);
  364. cures = cudaMemcpy((char *)dst_csr->rowptr, (char *)src_csr->rowptr, (nrow+1)*sizeof(uint32_t), kind);
  365. if (STARPU_UNLIKELY(cures))
  366. STARPU_CUDA_REPORT_ERROR(cures);
  367. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  368. return 0;
  369. }
  370. static int copy_cuda_common_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED, enum cudaMemcpyKind kind, cudaStream_t stream)
  371. {
  372. struct starpu_csr_interface *src_csr = src_interface;
  373. struct starpu_csr_interface *dst_csr = dst_interface;
  374. uint32_t nnz = src_csr->nnz;
  375. uint32_t nrow = src_csr->nrow;
  376. size_t elemsize = src_csr->elemsize;
  377. cudaError_t cures;
  378. int synchronous_fallback = 0;
  379. _STARPU_TRACE_START_DRIVER_COPY_ASYNC(src_node, dst_node);
  380. cures = cudaMemcpyAsync((char *)dst_csr->nzval, (char *)src_csr->nzval, nnz*elemsize, kind, stream);
  381. if (cures)
  382. {
  383. synchronous_fallback = 1;
  384. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  385. cures = cudaMemcpy((char *)dst_csr->nzval, (char *)src_csr->nzval, nnz*elemsize, kind);
  386. if (STARPU_UNLIKELY(cures))
  387. STARPU_CUDA_REPORT_ERROR(cures);
  388. }
  389. if (!synchronous_fallback)
  390. {
  391. cures = cudaMemcpyAsync((char *)dst_csr->colind, (char *)src_csr->colind, nnz*sizeof(uint32_t), kind, stream);
  392. }
  393. if (synchronous_fallback || cures != cudaSuccess)
  394. {
  395. synchronous_fallback = 1;
  396. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  397. cures = cudaMemcpy((char *)dst_csr->colind, (char *)src_csr->colind, nnz*sizeof(uint32_t), kind);
  398. if (STARPU_UNLIKELY(cures))
  399. STARPU_CUDA_REPORT_ERROR(cures);
  400. }
  401. if (!synchronous_fallback)
  402. {
  403. cures = cudaMemcpyAsync((char *)dst_csr->rowptr, (char *)src_csr->rowptr, (nrow+1)*sizeof(uint32_t), kind, stream);
  404. }
  405. if (synchronous_fallback || cures != cudaSuccess)
  406. {
  407. synchronous_fallback = 1;
  408. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  409. cures = cudaMemcpy((char *)dst_csr->rowptr, (char *)src_csr->rowptr, (nrow+1)*sizeof(uint32_t), kind);
  410. if (STARPU_UNLIKELY(cures))
  411. STARPU_CUDA_REPORT_ERROR(cures);
  412. }
  413. if (synchronous_fallback)
  414. {
  415. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  416. return 0;
  417. }
  418. else
  419. {
  420. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  421. return -EAGAIN;
  422. }
  423. }
  424. static int copy_cuda_peer(void *src_interface STARPU_ATTRIBUTE_UNUSED, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface STARPU_ATTRIBUTE_UNUSED, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  425. {
  426. #ifdef HAVE_CUDA_MEMCPY_PEER
  427. struct starpu_csr_interface *src_csr = src_interface;
  428. struct starpu_csr_interface *dst_csr = dst_interface;
  429. uint32_t nnz = src_csr->nnz;
  430. uint32_t nrow = src_csr->nrow;
  431. size_t elemsize = src_csr->elemsize;
  432. int src_dev = _starpu_memory_node_to_devid(src_node);
  433. int dst_dev = _starpu_memory_node_to_devid(dst_node);
  434. cudaError_t cures;
  435. cures = cudaMemcpyPeer((char *)dst_csr->nzval, dst_dev, (char *)src_csr->nzval, src_dev, nnz*elemsize);
  436. if (STARPU_UNLIKELY(cures))
  437. STARPU_CUDA_REPORT_ERROR(cures);
  438. cures = cudaMemcpyPeer((char *)dst_csr->colind, dst_dev, (char *)src_csr->colind, src_dev, nnz*sizeof(uint32_t));
  439. if (STARPU_UNLIKELY(cures))
  440. STARPU_CUDA_REPORT_ERROR(cures);
  441. cures = cudaMemcpyPeer((char *)dst_csr->rowptr, dst_dev, (char *)src_csr->rowptr, src_dev, (nrow+1)*sizeof(uint32_t));
  442. if (STARPU_UNLIKELY(cures))
  443. STARPU_CUDA_REPORT_ERROR(cures);
  444. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  445. return 0;
  446. #else
  447. STARPU_ABORT();
  448. return 0;
  449. #endif
  450. }
  451. static int copy_cuda_peer_async(void *src_interface STARPU_ATTRIBUTE_UNUSED, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  452. void *dst_interface STARPU_ATTRIBUTE_UNUSED, unsigned dst_node STARPU_ATTRIBUTE_UNUSED, cudaStream_t stream STARPU_ATTRIBUTE_UNUSED)
  453. {
  454. #ifdef HAVE_CUDA_MEMCPY_PEER
  455. struct starpu_csr_interface *src_csr = src_interface;
  456. struct starpu_csr_interface *dst_csr = dst_interface;
  457. uint32_t nnz = src_csr->nnz;
  458. uint32_t nrow = src_csr->nrow;
  459. size_t elemsize = src_csr->elemsize;
  460. cudaError_t cures;
  461. int src_dev = _starpu_memory_node_to_devid(src_node);
  462. int dst_dev = _starpu_memory_node_to_devid(dst_node);
  463. int synchronous_fallback = 0;
  464. _STARPU_TRACE_START_DRIVER_COPY_ASYNC(src_node, dst_node);
  465. cures = cudaMemcpyPeerAsync((char *)dst_csr->nzval, dst_dev, (char *)src_csr->nzval, src_dev, nnz*elemsize, stream);
  466. if (cures)
  467. {
  468. synchronous_fallback = 1;
  469. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  470. cures = cudaMemcpyPeer((char *)dst_csr->nzval, dst_dev, (char *)src_csr->nzval, src_dev, nnz*elemsize);
  471. if (STARPU_UNLIKELY(cures))
  472. STARPU_CUDA_REPORT_ERROR(cures);
  473. }
  474. if (!synchronous_fallback)
  475. {
  476. cures = cudaMemcpyPeerAsync((char *)dst_csr->colind, dst_dev, (char *)src_csr->colind, src_dev, nnz*sizeof(uint32_t), stream);
  477. }
  478. if (synchronous_fallback || cures != cudaSuccess)
  479. {
  480. synchronous_fallback = 1;
  481. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  482. cures = cudaMemcpyPeer((char *)dst_csr->colind, dst_dev, (char *)src_csr->colind, src_dev, nnz*sizeof(uint32_t));
  483. if (STARPU_UNLIKELY(cures))
  484. STARPU_CUDA_REPORT_ERROR(cures);
  485. }
  486. if (!synchronous_fallback)
  487. {
  488. cures = cudaMemcpyPeerAsync((char *)dst_csr->rowptr, dst_dev, (char *)src_csr->rowptr, src_dev, (nrow+1)*sizeof(uint32_t), stream);
  489. }
  490. if (synchronous_fallback || cures != cudaSuccess)
  491. {
  492. synchronous_fallback = 1;
  493. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  494. cures = cudaMemcpyPeer((char *)dst_csr->rowptr, dst_dev, (char *)src_csr->rowptr, src_dev, (nrow+1)*sizeof(uint32_t));
  495. if (STARPU_UNLIKELY(cures))
  496. STARPU_CUDA_REPORT_ERROR(cures);
  497. }
  498. if (synchronous_fallback)
  499. {
  500. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  501. return 0;
  502. }
  503. else
  504. {
  505. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  506. return -EAGAIN;
  507. }
  508. #else
  509. /* Illegal without Peer tranfers */
  510. STARPU_ABORT();
  511. return 0;
  512. #endif
  513. }
  514. static int copy_cuda_to_ram(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  515. {
  516. return copy_cuda_common(src_interface, src_node, dst_interface, dst_node, cudaMemcpyDeviceToHost);
  517. }
  518. static int copy_ram_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  519. {
  520. return copy_cuda_common(src_interface, src_node, dst_interface, dst_node, cudaMemcpyHostToDevice);
  521. }
  522. static int copy_cuda_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  523. {
  524. if (src_node == dst_node)
  525. return copy_cuda_common(src_interface, src_node, dst_interface, dst_node, cudaMemcpyDeviceToDevice);
  526. else
  527. return copy_cuda_peer(src_interface, src_node, dst_interface, dst_node);
  528. }
  529. static int copy_cuda_to_ram_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  530. {
  531. return copy_cuda_common_async(src_interface, src_node, dst_interface, dst_node, cudaMemcpyDeviceToHost, stream);
  532. }
  533. static int copy_ram_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  534. {
  535. return copy_cuda_common_async(src_interface, src_node, dst_interface, dst_node, cudaMemcpyHostToDevice, stream);
  536. }
  537. static int copy_cuda_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  538. {
  539. if (src_node == dst_node)
  540. return copy_cuda_common_async(src_interface, src_node, dst_interface, dst_node, cudaMemcpyDeviceToDevice, stream);
  541. else
  542. return copy_cuda_peer_async(src_interface, src_node, dst_interface, dst_node, stream);
  543. }
  544. #endif // STARPU_USE_CUDA
  545. #ifdef STARPU_USE_OPENCL
  546. static int copy_opencl_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  547. {
  548. struct starpu_csr_interface *src_csr = src_interface;
  549. struct starpu_csr_interface *dst_csr = dst_interface;
  550. uint32_t nnz = src_csr->nnz;
  551. uint32_t nrow = src_csr->nrow;
  552. size_t elemsize = src_csr->elemsize;
  553. int err;
  554. err = starpu_opencl_copy_opencl_to_ram((cl_mem)src_csr->nzval, src_node, (void *)dst_csr->nzval, dst_node, nnz*elemsize, 0, NULL);
  555. if (STARPU_UNLIKELY(err))
  556. STARPU_OPENCL_REPORT_ERROR(err);
  557. err = starpu_opencl_copy_opencl_to_ram((cl_mem)src_csr->colind, src_node, (void *)dst_csr->colind, dst_node, nnz*sizeof(uint32_t), 0, NULL);
  558. if (STARPU_UNLIKELY(err))
  559. STARPU_OPENCL_REPORT_ERROR(err);
  560. err = starpu_opencl_copy_opencl_to_ram((cl_mem)src_csr->rowptr, src_node, (void *)dst_csr->rowptr, dst_node, (nrow+1)*sizeof(uint32_t), 0, NULL);
  561. if (STARPU_UNLIKELY(err))
  562. STARPU_OPENCL_REPORT_ERROR(err);
  563. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  564. return 0;
  565. }
  566. static int copy_ram_to_opencl(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  567. {
  568. struct starpu_csr_interface *src_csr = src_interface;
  569. struct starpu_csr_interface *dst_csr = dst_interface;
  570. uint32_t nnz = src_csr->nnz;
  571. uint32_t nrow = src_csr->nrow;
  572. size_t elemsize = src_csr->elemsize;
  573. int err;
  574. err = starpu_opencl_copy_ram_to_opencl((void *)src_csr->nzval, src_node, (cl_mem)dst_csr->nzval, dst_node, nnz*elemsize, 0, NULL);
  575. if (STARPU_UNLIKELY(err))
  576. STARPU_OPENCL_REPORT_ERROR(err);
  577. err = starpu_opencl_copy_ram_to_opencl((void *)src_csr->colind, src_node, (cl_mem)dst_csr->colind, dst_node, nnz*sizeof(uint32_t), 0, NULL);
  578. if (STARPU_UNLIKELY(err))
  579. STARPU_OPENCL_REPORT_ERROR(err);
  580. err = starpu_opencl_copy_ram_to_opencl((void *)src_csr->rowptr, src_node, (cl_mem)dst_csr->rowptr, dst_node, (nrow+1)*sizeof(uint32_t), 0, NULL);
  581. if (STARPU_UNLIKELY(err))
  582. STARPU_OPENCL_REPORT_ERROR(err);
  583. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  584. return 0;
  585. }
  586. #endif // STARPU_USE_OPENCL
  587. /* as not all platform easily have a BLAS lib installed ... */
  588. static int copy_ram_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  589. {
  590. struct starpu_csr_interface *src_csr = (struct starpu_csr_interface *) src_interface;
  591. struct starpu_csr_interface *dst_csr = (struct starpu_csr_interface *) dst_interface;
  592. uint32_t nnz = src_csr->nnz;
  593. uint32_t nrow = src_csr->nrow;
  594. size_t elemsize = src_csr->elemsize;
  595. memcpy((void *)dst_csr->nzval, (void *)src_csr->nzval, nnz*elemsize);
  596. memcpy((void *)dst_csr->colind, (void *)src_csr->colind, nnz*sizeof(uint32_t));
  597. memcpy((void *)dst_csr->rowptr, (void *)src_csr->rowptr, (nrow+1)*sizeof(uint32_t));
  598. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nnz*elemsize + (nnz+nrow+1)*sizeof(uint32_t));
  599. return 0;
  600. }