Re: [Starpu-devel] GPU issue with r12137

[Samuel Thibault <samuel.thibault@ens-lyon.org>]

Xavier Lacoste, le Mon 31 Mar 2014 15:31:27 +0200, a écrit :
> I applied the patch on top of starpu-1.1 (r12511) and this seems to work.

Ok. I discussed a bit with Cedric about it, he thinks it could be a mere
threadsafety issue in CUDA. Could you try the attached patch instead?

Samuel
Index: src/drivers/cuda/driver_cuda.c
===================================================================
--- src/drivers/cuda/driver_cuda.c      (révision 12509)
+++ src/drivers/cuda/driver_cuda.c      (copie de travail)
@@ -42,8 +42,8 @@
 static size_t global_mem[STARPU_NMAXWORKERS];
 #ifdef STARPU_USE_CUDA
 static cudaStream_t streams[STARPU_NMAXWORKERS];
-static cudaStream_t out_transfer_streams[STARPU_MAXCUDADEVS];
-static cudaStream_t in_transfer_streams[STARPU_MAXCUDADEVS];
+static cudaStream_t 
out_transfer_streams[STARPU_NMAXWORKERS][STARPU_MAXCUDADEVS];
+static cudaStream_t 
in_transfer_streams[STARPU_NMAXWORKERS][STARPU_MAXCUDADEVS];
 static cudaStream_t 
peer_transfer_streams[STARPU_MAXCUDADEVS][STARPU_MAXCUDADEVS];
 static struct cudaDeviceProp props[STARPU_MAXCUDADEVS];
 #endif /* STARPU_USE_CUDA */
@@ -118,15 +118,17 @@
 cudaStream_t starpu_cuda_get_in_transfer_stream(unsigned node)
 {
        int devid = _starpu_memory_node_get_devid(node);
+       int worker = starpu_worker_get_id();
 
-       return in_transfer_streams[devid];
+       return in_transfer_streams[worker][devid];
 }
 
 cudaStream_t starpu_cuda_get_out_transfer_stream(unsigned node)
 {
        int devid = _starpu_memory_node_get_devid(node);
+       int worker = starpu_worker_get_id();
 
-       return out_transfer_streams[devid];
+       return out_transfer_streams[worker][devid];
 }
 
 cudaStream_t starpu_cuda_get_peer_transfer_stream(unsigned src_node, 
unsigned dst_node)
@@ -199,6 +201,7 @@
 {
        cudaError_t cures;
        int workerid;
+       int nworkers = starpu_worker_get_count();
        int i;
 
        /* TODO: cudaSetDeviceFlag(cudaDeviceMapHost) */
@@ -208,7 +211,6 @@
 #ifdef HAVE_CUDA_MEMCPY_PEER
        if (starpu_get_env_number("STARPU_ENABLE_CUDA_GPU_GPU_DIRECT") != 0)
        {
-               int nworkers = starpu_worker_get_count();
                for (workerid = 0; workerid < nworkers; workerid++)
                {
                        struct _starpu_worker *worker = 
_starpu_get_worker_struct(workerid);
@@ -256,13 +258,16 @@
        if (STARPU_UNLIKELY(cures))
                STARPU_CUDA_REPORT_ERROR(cures);
 
-       cures = cudaStreamCreate(&in_transfer_streams[devid]);
-       if (STARPU_UNLIKELY(cures))
-               STARPU_CUDA_REPORT_ERROR(cures);
+       for (workerid = 0; workerid < nworkers; workerid++)
+       {
+               cures = 
cudaStreamCreate(&in_transfer_streams[workerid][devid]);
+               if (STARPU_UNLIKELY(cures))
+                       STARPU_CUDA_REPORT_ERROR(cures);
 
-       cures = cudaStreamCreate(&out_transfer_streams[devid]);
-       if (STARPU_UNLIKELY(cures))
-               STARPU_CUDA_REPORT_ERROR(cures);
+               cures = 
cudaStreamCreate(&out_transfer_streams[workerid][devid]);
+               if (STARPU_UNLIKELY(cures))
+                       STARPU_CUDA_REPORT_ERROR(cures);
+       }
 
        for (i = 0; i < ncudagpus; i++)
        {
@@ -277,10 +282,14 @@
        cudaError_t cures;
        int devid = starpu_worker_get_devid(workerid);
        int i;
+       int nworkers = starpu_worker_get_count();
 
        cudaStreamDestroy(streams[workerid]);
-       cudaStreamDestroy(in_transfer_streams[devid]);
-       cudaStreamDestroy(out_transfer_streams[devid]);
+       for (i = 0; i < nworkers; i++)
+       {
+               cudaStreamDestroy(in_transfer_streams[i][devid]);
+               cudaStreamDestroy(out_transfer_streams[i][devid]);
+       }
        for (i = 0; i < ncudagpus; i++)
                cudaStreamDestroy(peer_transfer_streams[i][devid]);