vector_interface.c 18 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540
  1. /* StarPU --- Runtime system for heterogeneous multicore architectures.
  2. *
  3. * Copyright (C) 2009-2012 Université de Bordeaux 1
  4. * Copyright (C) 2010, 2011, 2012 Centre National de la Recherche Scientifique
  5. *
  6. * StarPU is free software; you can redistribute it and/or modify
  7. * it under the terms of the GNU Lesser General Public License as published by
  8. * the Free Software Foundation; either version 2.1 of the License, or (at
  9. * your option) any later version.
  10. *
  11. * StarPU is distributed in the hope that it will be useful, but
  12. * WITHOUT ANY WARRANTY; without even the implied warranty of
  13. * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
  14. *
  15. * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  16. */
  17. #include <starpu.h>
  18. #include <common/config.h>
  19. #include <datawizard/coherency.h>
  20. #include <datawizard/copy_driver.h>
  21. #include <datawizard/filters.h>
  22. #include <starpu_hash.h>
  23. #include <starpu_cuda.h>
  24. #include <starpu_opencl.h>
  25. #include <drivers/opencl/driver_opencl.h>
  26. static int copy_ram_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  27. #ifdef STARPU_USE_CUDA
  28. static int copy_ram_to_cuda(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  29. static int copy_cuda_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  30. static int copy_ram_to_cuda_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  31. static int copy_cuda_to_ram_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  32. static int copy_cuda_to_cuda(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED);
  33. static int copy_cuda_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream);
  34. #endif
  35. #ifdef STARPU_USE_OPENCL
  36. static int copy_ram_to_opencl(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  37. static int copy_opencl_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  38. static int copy_opencl_to_opencl(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node);
  39. static int copy_ram_to_opencl_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node, void *_event);
  40. static int copy_opencl_to_ram_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED, void *dst_interface, unsigned dst_node, void *_event);
  41. #endif
  42. static struct starpu_data_copy_methods vector_copy_data_methods_s =
  43. {
  44. .ram_to_ram = copy_ram_to_ram,
  45. .ram_to_spu = NULL,
  46. #ifdef STARPU_USE_CUDA
  47. .ram_to_cuda = copy_ram_to_cuda,
  48. .cuda_to_ram = copy_cuda_to_ram,
  49. .ram_to_cuda_async = copy_ram_to_cuda_async,
  50. .cuda_to_ram_async = copy_cuda_to_ram_async,
  51. .cuda_to_cuda = copy_cuda_to_cuda,
  52. .cuda_to_cuda_async = copy_cuda_to_cuda_async,
  53. #endif
  54. #ifdef STARPU_USE_OPENCL
  55. .ram_to_opencl = copy_ram_to_opencl,
  56. .opencl_to_ram = copy_opencl_to_ram,
  57. .opencl_to_opencl = copy_opencl_to_opencl,
  58. .ram_to_opencl_async = copy_ram_to_opencl_async,
  59. .opencl_to_ram_async = copy_opencl_to_ram_async,
  60. #endif
  61. .cuda_to_spu = NULL,
  62. .spu_to_ram = NULL,
  63. .spu_to_cuda = NULL,
  64. .spu_to_spu = NULL
  65. };
  66. static void register_vector_handle(starpu_data_handle_t handle, uint32_t home_node, void *data_interface);
  67. static ssize_t allocate_vector_buffer_on_node(void *data_interface_, uint32_t dst_node);
  68. static void *vector_handle_to_pointer(starpu_data_handle_t data_handle, uint32_t node);
  69. static void free_vector_buffer_on_node(void *data_interface, uint32_t node);
  70. static size_t vector_interface_get_size(starpu_data_handle_t handle);
  71. static uint32_t footprint_vector_interface_crc32(starpu_data_handle_t handle);
  72. static int vector_compare(void *data_interface_a, void *data_interface_b);
  73. static void display_vector_interface(starpu_data_handle_t handle, FILE *f);
  74. #ifdef STARPU_USE_GORDON
  75. static int convert_vector_to_gordon(void *data_interface, uint64_t *ptr, gordon_strideSize_t *ss);
  76. #endif
  77. static struct starpu_data_interface_ops interface_vector_ops =
  78. {
  79. .register_data_handle = register_vector_handle,
  80. .allocate_data_on_node = allocate_vector_buffer_on_node,
  81. .handle_to_pointer = vector_handle_to_pointer,
  82. .free_data_on_node = free_vector_buffer_on_node,
  83. .copy_methods = &vector_copy_data_methods_s,
  84. .get_size = vector_interface_get_size,
  85. .footprint = footprint_vector_interface_crc32,
  86. .compare = vector_compare,
  87. #ifdef STARPU_USE_GORDON
  88. .convert_to_gordon = convert_vector_to_gordon,
  89. #endif
  90. .interfaceid = STARPU_VECTOR_INTERFACE_ID,
  91. .interface_size = sizeof(struct starpu_vector_interface),
  92. .display = display_vector_interface,
  93. };
  94. static void *vector_handle_to_pointer(starpu_data_handle_t handle, uint32_t node)
  95. {
  96. STARPU_ASSERT(starpu_data_test_if_allocated_on_node(handle, node));
  97. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  98. starpu_data_get_interface_on_node(handle, node);
  99. return (void*) vector_interface->ptr;
  100. }
  101. static void register_vector_handle(starpu_data_handle_t handle, uint32_t home_node, void *data_interface)
  102. {
  103. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *) data_interface;
  104. unsigned node;
  105. for (node = 0; node < STARPU_MAXNODES; node++)
  106. {
  107. struct starpu_vector_interface *local_interface = (struct starpu_vector_interface *)
  108. starpu_data_get_interface_on_node(handle, node);
  109. if (node == home_node)
  110. {
  111. local_interface->ptr = vector_interface->ptr;
  112. local_interface->dev_handle = vector_interface->dev_handle;
  113. local_interface->offset = vector_interface->offset;
  114. }
  115. else
  116. {
  117. local_interface->ptr = 0;
  118. local_interface->dev_handle = 0;
  119. local_interface->offset = 0;
  120. }
  121. local_interface->nx = vector_interface->nx;
  122. local_interface->elemsize = vector_interface->elemsize;
  123. }
  124. }
  125. #ifdef STARPU_USE_GORDON
  126. int convert_vector_to_gordon(void *data_interface, uint64_t *ptr, gordon_strideSize_t *ss)
  127. {
  128. struct starpu_vector_interface *vector_interface = interface;
  129. *ptr = vector_interface->ptr;
  130. (*ss).size = vector_interface->nx * vector_interface->elemsize;
  131. return 0;
  132. }
  133. #endif
  134. /* declare a new data with the vector interface */
  135. void starpu_vector_data_register(starpu_data_handle_t *handleptr, uint32_t home_node,
  136. uintptr_t ptr, uint32_t nx, size_t elemsize)
  137. {
  138. struct starpu_vector_interface vector =
  139. {
  140. .ptr = ptr,
  141. .nx = nx,
  142. .elemsize = elemsize,
  143. .dev_handle = ptr,
  144. .offset = 0
  145. };
  146. starpu_data_register(handleptr, home_node, &vector, &interface_vector_ops);
  147. }
  148. static uint32_t footprint_vector_interface_crc32(starpu_data_handle_t handle)
  149. {
  150. return starpu_crc32_be(starpu_vector_get_nx(handle), 0);
  151. }
  152. static int vector_compare(void *data_interface_a, void *data_interface_b)
  153. {
  154. struct starpu_vector_interface *vector_a = (struct starpu_vector_interface *) data_interface_a;
  155. struct starpu_vector_interface *vector_b = (struct starpu_vector_interface *) data_interface_b;
  156. /* Two vectors are considered compatible if they have the same size */
  157. return ((vector_a->nx == vector_b->nx)
  158. && (vector_a->elemsize == vector_b->elemsize));
  159. }
  160. static void display_vector_interface(starpu_data_handle_t handle, FILE *f)
  161. {
  162. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  163. starpu_data_get_interface_on_node(handle, 0);
  164. fprintf(f, "%u\t", vector_interface->nx);
  165. }
  166. static size_t vector_interface_get_size(starpu_data_handle_t handle)
  167. {
  168. size_t size;
  169. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  170. starpu_data_get_interface_on_node(handle, 0);
  171. size = vector_interface->nx*vector_interface->elemsize;
  172. return size;
  173. }
  174. /* offer an access to the data parameters */
  175. uint32_t starpu_vector_get_nx(starpu_data_handle_t handle)
  176. {
  177. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  178. starpu_data_get_interface_on_node(handle, 0);
  179. return vector_interface->nx;
  180. }
  181. uintptr_t starpu_vector_get_local_ptr(starpu_data_handle_t handle)
  182. {
  183. unsigned node;
  184. node = _starpu_get_local_memory_node();
  185. STARPU_ASSERT(starpu_data_test_if_allocated_on_node(handle, node));
  186. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  187. starpu_data_get_interface_on_node(handle, node);
  188. return vector_interface->ptr;
  189. }
  190. size_t starpu_vector_get_elemsize(starpu_data_handle_t handle)
  191. {
  192. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *)
  193. starpu_data_get_interface_on_node(handle, 0);
  194. return vector_interface->elemsize;
  195. }
  196. /* memory allocation/deallocation primitives for the vector interface */
  197. /* returns the size of the allocated area */
  198. static ssize_t allocate_vector_buffer_on_node(void *data_interface_, uint32_t dst_node)
  199. {
  200. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *) data_interface_;
  201. unsigned fail = 0;
  202. uintptr_t addr = 0, handle = 0;
  203. ssize_t allocated_memory;
  204. uint32_t nx = vector_interface->nx;
  205. size_t elemsize = vector_interface->elemsize;
  206. enum starpu_node_kind kind = starpu_node_get_kind(dst_node);
  207. #ifdef STARPU_USE_CUDA
  208. cudaError_t status;
  209. #endif
  210. switch(kind)
  211. {
  212. case STARPU_CPU_RAM:
  213. addr = handle = (uintptr_t)malloc(nx*elemsize);
  214. if (!addr)
  215. fail = 1;
  216. break;
  217. #ifdef STARPU_USE_CUDA
  218. case STARPU_CUDA_RAM:
  219. status = cudaMalloc((void **)&addr, nx*elemsize);
  220. if (!addr || (status != cudaSuccess))
  221. {
  222. if (STARPU_UNLIKELY(status != cudaErrorMemoryAllocation))
  223. STARPU_CUDA_REPORT_ERROR(status);
  224. fail = 1;
  225. }
  226. handle = addr;
  227. break;
  228. #endif
  229. #ifdef STARPU_USE_OPENCL
  230. case STARPU_OPENCL_RAM:
  231. {
  232. int ret;
  233. cl_mem mem;
  234. ret = starpu_opencl_allocate_memory(&mem, nx*elemsize, CL_MEM_READ_WRITE);
  235. handle = (uintptr_t)mem;
  236. if (ret)
  237. {
  238. fail = 1;
  239. }
  240. break;
  241. }
  242. #endif
  243. default:
  244. STARPU_ABORT();
  245. }
  246. if (fail)
  247. return -ENOMEM;
  248. /* allocation succeeded */
  249. allocated_memory = nx*elemsize;
  250. /* update the data properly in consequence */
  251. vector_interface->ptr = addr;
  252. vector_interface->dev_handle = handle;
  253. vector_interface->offset = 0;
  254. return allocated_memory;
  255. }
  256. static void free_vector_buffer_on_node(void *data_interface, uint32_t node)
  257. {
  258. struct starpu_vector_interface *vector_interface = (struct starpu_vector_interface *) data_interface;
  259. #ifdef STARPU_USE_CUDA
  260. cudaError_t cures;
  261. #endif
  262. enum starpu_node_kind kind = starpu_node_get_kind(node);
  263. switch(kind)
  264. {
  265. case STARPU_CPU_RAM:
  266. free((void*)vector_interface->ptr);
  267. break;
  268. #ifdef STARPU_USE_CUDA
  269. case STARPU_CUDA_RAM:
  270. cures = cudaFree((void*)vector_interface->ptr);
  271. STARPU_ASSERT(cures == cudaSuccess);
  272. break;
  273. #endif
  274. #ifdef STARPU_USE_OPENCL
  275. case STARPU_OPENCL_RAM:
  276. clReleaseMemObject((cl_mem)vector_interface->dev_handle);
  277. break;
  278. #endif
  279. default:
  280. STARPU_ABORT();
  281. }
  282. }
  283. #ifdef STARPU_USE_CUDA
  284. static int copy_cuda_async_sync(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream, enum cudaMemcpyKind kind)
  285. {
  286. struct starpu_vector_interface *src_vector = src_interface;
  287. struct starpu_vector_interface *dst_vector = dst_interface;
  288. int ret;
  289. ret = starpu_cuda_copy_async_sync((void *)src_vector->ptr, src_node, (void *)dst_vector->ptr, dst_node, src_vector->nx*src_vector->elemsize, stream, kind);
  290. _STARPU_TRACE_DATA_COPY(src_node, dst_node, src_vector->nx*src_vector->elemsize);
  291. return ret;
  292. }
  293. #ifdef HAVE_CUDA_MEMCPY_PEER
  294. static int copy_cuda_peer_common(void *src_interface, unsigned src_node,
  295. void *dst_interface, unsigned dst_node,
  296. int is_async, cudaStream_t stream)
  297. {
  298. cudaError_t cures;
  299. struct starpu_vector_interface *src_vector = src_interface;
  300. struct starpu_vector_interface *dst_vector = dst_interface;
  301. size_t length = src_vector->nx*src_vector->elemsize;
  302. int src_dev = _starpu_memory_node_to_devid(src_node);
  303. int dst_dev = _starpu_memory_node_to_devid(dst_node);
  304. if (is_async)
  305. {
  306. _STARPU_TRACE_START_DRIVER_COPY_ASYNC(src_node, dst_node);
  307. cures = cudaMemcpyPeerAsync((char *)dst_vector->ptr, dst_dev,
  308. (char *)src_vector->ptr, src_dev,
  309. length, stream);
  310. _STARPU_TRACE_END_DRIVER_COPY_ASYNC(src_node, dst_node);
  311. if (!cures)
  312. return -EAGAIN;
  313. }
  314. cures = cudaMemcpyPeer((char *)dst_vector->ptr, dst_dev,
  315. (char *)src_vector->ptr, src_dev, length);
  316. if (STARPU_UNLIKELY(cures))
  317. STARPU_CUDA_REPORT_ERROR(cures);
  318. _STARPU_TRACE_DATA_COPY(src_node, dst_node, length);
  319. return 0;
  320. }
  321. #endif
  322. static int copy_cuda_to_ram(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  323. {
  324. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, NULL, cudaMemcpyDeviceToHost);
  325. }
  326. static int copy_ram_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  327. {
  328. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, NULL, cudaMemcpyHostToDevice);
  329. }
  330. static int copy_cuda_to_cuda(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node)
  331. {
  332. if (src_node == dst_node)
  333. {
  334. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, NULL, cudaMemcpyDeviceToDevice);
  335. }
  336. else
  337. {
  338. #ifdef HAVE_CUDA_MEMCPY_PEER
  339. return copy_cuda_peer_common(src_interface, src_node, dst_interface, dst_node, 0, 0);
  340. #else
  341. /* This is illegal without cudaMemcpyPeer */
  342. STARPU_ABORT();
  343. return 0;
  344. #endif
  345. }
  346. }
  347. static int copy_cuda_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  348. {
  349. if (src_node == dst_node)
  350. {
  351. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, stream, cudaMemcpyDeviceToDevice);
  352. }
  353. else
  354. {
  355. #ifdef HAVE_CUDA_MEMCPY_PEER
  356. return copy_cuda_peer_common(src_interface, src_node, dst_interface, dst_node, 1, stream);
  357. #else
  358. /* This is illegal without cudaMemcpyPeer */
  359. STARPU_ABORT();
  360. return 0;
  361. #endif
  362. }
  363. }
  364. static int copy_cuda_to_ram_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  365. {
  366. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, stream, cudaMemcpyDeviceToHost);
  367. }
  368. static int copy_ram_to_cuda_async(void *src_interface, unsigned src_node, void *dst_interface, unsigned dst_node, cudaStream_t stream)
  369. {
  370. return copy_cuda_async_sync(src_interface, src_node, dst_interface, dst_node, stream, cudaMemcpyHostToDevice);
  371. }
  372. #endif // STARPU_USE_CUDA
  373. #ifdef STARPU_USE_OPENCL
  374. static int copy_ram_to_opencl_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  375. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED, void *_event)
  376. {
  377. struct starpu_vector_interface *src_vector = src_interface;
  378. struct starpu_vector_interface *dst_vector = dst_interface;
  379. int err, ret;
  380. err = starpu_opencl_copy_ram_to_opencl((void*)src_vector->ptr, src_node, (cl_mem)dst_vector->dev_handle, dst_node,
  381. src_vector->nx*src_vector->elemsize,
  382. dst_vector->offset, (cl_event*)_event, &ret);
  383. if (STARPU_UNLIKELY(err))
  384. STARPU_OPENCL_REPORT_ERROR(err);
  385. _STARPU_TRACE_DATA_COPY(src_node, dst_node, src_vector->nx*src_vector->elemsize);
  386. return ret;
  387. }
  388. static int copy_opencl_to_ram_async(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  389. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED, void *_event)
  390. {
  391. struct starpu_vector_interface *src_vector = src_interface;
  392. struct starpu_vector_interface *dst_vector = dst_interface;
  393. int err, ret;
  394. err = starpu_opencl_copy_opencl_to_ram((cl_mem)src_vector->dev_handle, src_node, (void*)dst_vector->ptr, dst_node, src_vector->nx*src_vector->elemsize,
  395. src_vector->offset, (cl_event *)_event, &ret);
  396. if (STARPU_UNLIKELY(err))
  397. STARPU_OPENCL_REPORT_ERROR(err);
  398. _STARPU_TRACE_DATA_COPY(src_node, dst_node, src_vector->nx*src_vector->elemsize);
  399. return ret;
  400. }
  401. static int copy_ram_to_opencl(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  402. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  403. {
  404. return copy_ram_to_opencl_async(src_interface, src_node, dst_interface, dst_node, NULL);
  405. }
  406. static int copy_opencl_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  407. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  408. {
  409. return copy_opencl_to_ram_async(src_interface, src_node, dst_interface, dst_node, NULL);
  410. }
  411. static int copy_opencl_to_opencl(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  412. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  413. {
  414. int err;
  415. struct starpu_vector_interface *src_vector = src_interface;
  416. struct starpu_vector_interface *dst_vector = dst_interface;
  417. cl_command_queue cq;
  418. starpu_opencl_get_current_queue(&cq);
  419. size_t size = src_vector->nx*src_vector->elemsize;
  420. cl_event event;
  421. err = clEnqueueCopyBuffer(cq, (cl_mem)src_vector->dev_handle, (cl_mem)dst_vector->dev_handle, src_vector->offset, dst_vector->offset, size, 0, NULL, &event);
  422. if (STARPU_UNLIKELY(err))
  423. STARPU_OPENCL_REPORT_ERROR(err);
  424. err = clWaitForEvents(1, &event);
  425. if (STARPU_UNLIKELY(err))
  426. STARPU_OPENCL_REPORT_ERROR(err);
  427. err = clReleaseEvent(event);
  428. if (STARPU_UNLIKELY(err))
  429. STARPU_OPENCL_REPORT_ERROR(err);
  430. _STARPU_TRACE_DATA_COPY(src_node, dst_node, src_vector->nx*src_vector->elemsize);
  431. return 0;
  432. }
  433. #endif // STARPU_USE_OPENCL
  434. static int copy_ram_to_ram(void *src_interface, unsigned src_node STARPU_ATTRIBUTE_UNUSED,
  435. void *dst_interface, unsigned dst_node STARPU_ATTRIBUTE_UNUSED)
  436. {
  437. struct starpu_vector_interface *src_vector = (struct starpu_vector_interface *) src_interface;
  438. struct starpu_vector_interface *dst_vector = (struct starpu_vector_interface *) dst_interface;
  439. uint32_t nx = dst_vector->nx;
  440. size_t elemsize = dst_vector->elemsize;
  441. uintptr_t ptr_src = src_vector->ptr;
  442. uintptr_t ptr_dst = dst_vector->ptr;
  443. memcpy((void *)ptr_dst, (void *)ptr_src, nx*elemsize);
  444. _STARPU_TRACE_DATA_COPY(src_node, dst_node, nx*elemsize);
  445. return 0;
  446. }