NAMD
CudaUtils.C
Go to the documentation of this file.
1 
2 #include <stdio.h>
3 #include "common.h"
4 #include "charm++.h"
5 #include "CudaUtils.h"
6 
7 #ifdef NAMD_CUDA
8 
9 void cudaDie(const char *msg, cudaError_t err) {
10  char host[128];
11  gethostname(host, 128); host[127] = 0;
12  char devstr[128] = "";
13  int devnum;
14  if ( cudaGetDevice(&devnum) == cudaSuccess ) {
15  sprintf(devstr, " device %d", devnum);
16  }
17  cudaDeviceProp deviceProp;
18  if ( cudaGetDeviceProperties(&deviceProp, devnum) == cudaSuccess ) {
19  sprintf(devstr, " device %d pci %x:%x:%x", devnum,
20  deviceProp.pciDomainID, deviceProp.pciBusID, deviceProp.pciDeviceID);
21  }
22  char errmsg[1024];
23  if (err == cudaSuccess) {
24  sprintf(errmsg,"CUDA error %s on Pe %d (%s%s)", msg, CkMyPe(), host, devstr);
25  } else {
26  sprintf(errmsg,"CUDA error %s on Pe %d (%s%s): %s", msg, CkMyPe(), host, devstr, cudaGetErrorString(err));
27  }
28  NAMD_die(errmsg);
29 }
30 
31 void cudaNAMD_bug(const char *msg) {NAMD_bug(msg);}
32 
34  int devcnt = 0;
35  cudaError_t err = cudaGetDeviceCount(&devcnt);
36  if ( devcnt == 1 ) { // only one device so it must be ours
37  int *dummy;
38  if ( err == cudaSuccess ) err = cudaSetDevice(0);
39  if ( err == cudaSuccess ) err = cudaSetDeviceFlags(cudaDeviceMapHost);
40  if ( err == cudaSuccess ) err = cudaMalloc(&dummy, 4);
41  }
42  if ( err != cudaSuccess ) {
43  char host[128];
44  gethostname(host, 128); host[127] = 0;
45  fprintf(stderr,"CUDA initialization error on %s: %s\n", host, cudaGetErrorString(err));
46  }
47 }
48 
49 //----------------------------------------------------------------------------------------
50 
51 void clear_device_array_async_T(void *data, const int ndata, cudaStream_t stream, const size_t sizeofT) {
52  cudaCheck(cudaMemsetAsync(data, 0, sizeofT*ndata, stream));
53 }
54 
55 void clear_device_array_T(void *data, const int ndata, const size_t sizeofT) {
56  cudaCheck(cudaMemset(data, 0, sizeofT*ndata));
57 }
58 
59 //----------------------------------------------------------------------------------------
60 //
61 // Allocate page-locked host memory
62 // pp = memory pointer
63 // len = length of the array
64 //
65 void allocate_host_T(void **pp, const int len, const size_t sizeofT) {
66  cudaCheck(cudaMallocHost(pp, sizeofT*len));
67 }
68 
69 //----------------------------------------------------------------------------------------
70 //
71 // Allocate gpu memory
72 // pp = memory pointer
73 // len = length of the array
74 //
75 void allocate_device_T(void **pp, const int len, const size_t sizeofT) {
76  cudaCheck(cudaMalloc(pp, sizeofT*len));
77 }
78 
79 //----------------------------------------------------------------------------------------
80 //
81 // Deallocate gpu memory
82 // pp = memory pointer
83 //
84 void deallocate_device_T(void **pp) {
85 
86  if (*pp != NULL) {
87  cudaCheck(cudaFree((void *)(*pp)));
88  *pp = NULL;
89  }
90 
91 }
92 
93 //----------------------------------------------------------------------------------------
94 //
95 // Deallocate page-locked host memory
96 // pp = memory pointer
97 //
98 void deallocate_host_T(void **pp) {
99 
100  if (*pp != NULL) {
101  cudaCheck(cudaFreeHost((void *)(*pp)));
102  *pp = NULL;
103  }
104 
105 }
106 
107 //----------------------------------------------------------------------------------------
108 //
109 // Allocate & re-allocate device memory
110 // pp = memory pointer
111 // curlen = current length of the array
112 // newlen = new required length of the array
113 // fac = extra space allocation factor: in case of re-allocation new length will be fac*newlen
114 //
115 // returns true if reallocation happened
116 //
117 bool reallocate_device_T(void **pp, int *curlen, const int newlen, const float fac, const size_t sizeofT) {
118 
119  if (*pp != NULL && *curlen < newlen) {
120  cudaCheck(cudaFree((void *)(*pp)));
121  *pp = NULL;
122  }
123 
124  if (*pp == NULL) {
125  if (fac > 1.0f) {
126  *curlen = (int)(((double)(newlen))*(double)fac);
127  } else {
128  *curlen = newlen;
129  }
130  cudaCheck(cudaMalloc(pp, sizeofT*(*curlen)));
131  return true;
132  }
133 
134  return false;
135 }
136 
137 //----------------------------------------------------------------------------------------
138 //
139 // Allocate & re-allocate page-locked host memory
140 // pp = memory pointer
141 // curlen = current length of the array
142 // newlen = new required length of the array
143 // fac = extra space allocation factor: in case of re-allocation new length will be fac*newlen
144 // flag = allocation type:
145 // cudaHostAllocDefault = default type, emulates cudaMallocHost
146 // cudaHostAllocMapped = maps allocation into CUDA address space
147 //
148 // returns true if reallocation happened
149 //
150 bool reallocate_host_T(void **pp, int *curlen, const int newlen,
151  const float fac, const unsigned int flag, const size_t sizeofT) {
152 
153  if (*pp != NULL && *curlen < newlen) {
154  cudaCheck(cudaFreeHost((void *)(*pp)));
155  *pp = NULL;
156  }
157 
158  if (*pp == NULL) {
159  if (fac > 1.0f) {
160  *curlen = (int)(((double)(newlen))*(double)fac);
161  } else {
162  *curlen = newlen;
163  }
164  cudaCheck(cudaHostAlloc(pp, sizeofT*(*curlen), flag));
165  return true;
166  }
167 
168  return false;
169 }
170 
171 //----------------------------------------------------------------------------------------
172 //
173 // Copies memory Host -> Device
174 //
175 void copy_HtoD_async_T(const void *h_array, void *d_array, int array_len, cudaStream_t stream,
176  const size_t sizeofT) {
177  cudaCheck(cudaMemcpyAsync(d_array, h_array, sizeofT*array_len, cudaMemcpyHostToDevice, stream));
178 }
179 
180 void copy_HtoD_T(const void *h_array, void *d_array, int array_len,
181  const size_t sizeofT) {
182  cudaCheck(cudaMemcpy(d_array, h_array, sizeofT*array_len, cudaMemcpyHostToDevice));
183 }
184 
185 //----------------------------------------------------------------------------------------
186 //
187 // Copies memory Device -> Host
188 //
189 void copy_DtoH_async_T(const void *d_array, void *h_array, const int array_len, cudaStream_t stream,
190  const size_t sizeofT) {
191  cudaCheck(cudaMemcpyAsync(h_array, d_array, sizeofT*array_len, cudaMemcpyDeviceToHost, stream));
192 }
193 
194 void copy_DtoH_T(const void *d_array, void *h_array, const int array_len, const size_t sizeofT) {
195  cudaCheck(cudaMemcpy(h_array, d_array, sizeofT*array_len, cudaMemcpyDeviceToHost));
196 }
197 
198 //----------------------------------------------------------------------------------------
199 //
200 // Copies memory Device -> Device
201 //
202 void copy_DtoD_async_T(const void *d_src, void *d_dst, const int array_len, cudaStream_t stream,
203  const size_t sizeofT) {
204  cudaCheck(cudaMemcpyAsync(d_dst, d_src, sizeofT*array_len, cudaMemcpyDeviceToDevice, stream));
205 }
206 
207 void copy_DtoD_T(const void *d_src, void *d_dst, const int array_len, const size_t sizeofT) {
208  cudaCheck(cudaMemcpy(d_dst, d_src, sizeofT*array_len, cudaMemcpyDeviceToDevice));
209 }
210 
211 //----------------------------------------------------------------------------------------
212 //
213 // Copies memory between two devices Device -> Device
214 //
215 void copy_PeerDtoD_async_T(const int src_dev, const int dst_dev,
216  const void *d_src, void *d_dst, const int array_len, cudaStream_t stream,
217  const size_t sizeofT) {
218  cudaCheck(cudaMemcpyPeerAsync(d_dst, dst_dev, d_src, src_dev, sizeofT*array_len, stream));
219 }
220 
221 //----------------------------------------------------------------------------------------
222 //
223 // Copies 3D memory block Host -> Device
224 //
225 void copy3D_HtoD_T(void* src_data, void* dst_data,
226  int src_x0, int src_y0, int src_z0,
227  size_t src_xsize, size_t src_ysize,
228  int dst_x0, int dst_y0, int dst_z0,
229  size_t dst_xsize, size_t dst_ysize,
230  size_t width, size_t height, size_t depth,
231  size_t sizeofT, cudaStream_t stream) {
232  cudaMemcpy3DParms parms = {0};
233 
234  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
235  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
236 
237  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
238  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
239 
240  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
241  parms.kind = cudaMemcpyHostToDevice;
242 
243  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
244 }
245 
246 //----------------------------------------------------------------------------------------
247 //
248 // Copies 3D memory block Device -> Host
249 //
250 void copy3D_DtoH_T(void* src_data, void* dst_data,
251  int src_x0, int src_y0, int src_z0,
252  size_t src_xsize, size_t src_ysize,
253  int dst_x0, int dst_y0, int dst_z0,
254  size_t dst_xsize, size_t dst_ysize,
255  size_t width, size_t height, size_t depth,
256  size_t sizeofT, cudaStream_t stream) {
257  cudaMemcpy3DParms parms = {0};
258 
259  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
260  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
261 
262  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
263  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
264 
265  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
266  parms.kind = cudaMemcpyDeviceToHost;
267 
268  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
269 }
270 
271 //----------------------------------------------------------------------------------------
272 //
273 // Copies 3D memory block Device -> Device
274 //
275 void copy3D_DtoD_T(void* src_data, void* dst_data,
276  int src_x0, int src_y0, int src_z0,
277  size_t src_xsize, size_t src_ysize,
278  int dst_x0, int dst_y0, int dst_z0,
279  size_t dst_xsize, size_t dst_ysize,
280  size_t width, size_t height, size_t depth,
281  size_t sizeofT, cudaStream_t stream) {
282  cudaMemcpy3DParms parms = {0};
283 
284  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
285  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
286 
287  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
288  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
289 
290  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
291  parms.kind = cudaMemcpyDeviceToDevice;
292 
293  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
294 }
295 
296 //----------------------------------------------------------------------------------------
297 //
298 // Copies 3D memory block between devices Device -> Device
299 //
300 void copy3D_PeerDtoD_T(int src_dev, int dst_dev,
301  void* src_data, void* dst_data,
302  int src_x0, int src_y0, int src_z0,
303  size_t src_xsize, size_t src_ysize,
304  int dst_x0, int dst_y0, int dst_z0,
305  size_t dst_xsize, size_t dst_ysize,
306  size_t width, size_t height, size_t depth,
307  size_t sizeofT, cudaStream_t stream) {
308  cudaMemcpy3DPeerParms parms = {0};
309 
310  parms.srcDevice = src_dev;
311  parms.dstDevice = dst_dev;
312 
313  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
314  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
315 
316  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
317  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
318 
319  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
320 
321  cudaCheck(cudaMemcpy3DPeerAsync(&parms, stream));
322 }
323 #endif // NAMD_CUDA
void deallocate_device_T(void **pp)
Definition: CudaUtils.C:84
void copy_DtoH_T(const void *d_array, void *h_array, const int array_len, const size_t sizeofT)
Definition: CudaUtils.C:194
void copy3D_HtoD_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:225
void allocate_device_T(void **pp, const int len, const size_t sizeofT)
Definition: CudaUtils.C:75
void copy_DtoD_T(const void *d_src, void *d_dst, const int array_len, const size_t sizeofT)
Definition: CudaUtils.C:207
void copy_PeerDtoD_async_T(const int src_dev, const int dst_dev, const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:215
__thread cudaStream_t stream
void copy_DtoH_async_T(const void *d_array, void *h_array, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:189
bool reallocate_host_T(void **pp, int *curlen, const int newlen, const float fac, const unsigned int flag, const size_t sizeofT)
Definition: CudaUtils.C:150
void NAMD_bug(const char *err_msg)
Definition: common.C:123
void copy3D_DtoD_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:275
void copy_HtoD_T(const void *h_array, void *d_array, int array_len, const size_t sizeofT)
Definition: CudaUtils.C:180
void cudaDie(const char *msg, cudaError_t err=cudaSuccess)
Definition: CudaUtils.C:9
void NAMD_die(const char *err_msg)
Definition: common.C:83
void clear_device_array_async_T(void *data, const int ndata, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:51
void cudaNAMD_bug(const char *msg)
Definition: CudaUtils.C:31
void copy_DtoD_async_T(const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:202
bool reallocate_device_T(void **pp, int *curlen, const int newlen, const float fac, const size_t sizeofT)
Definition: CudaUtils.C:117
void copy3D_PeerDtoD_T(int src_dev, int dst_dev, void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:300
void copy3D_DtoH_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:250
void clear_device_array_T(void *data, const int ndata, const size_t sizeofT)
Definition: CudaUtils.C:55
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void deallocate_host_T(void **pp)
Definition: CudaUtils.C:98
void cuda_affinity_initialize()
Definition: CudaUtils.C:33
void allocate_host_T(void **pp, const int len, const size_t sizeofT)
Definition: CudaUtils.C:65
void copy_HtoD_async_T(const void *h_array, void *d_array, int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:175