Stokhos Package Browser (Single Doxygen Collection) Version of the Day
Loading...
Searching...
No Matches
Stokhos_DynArrayTraits.hpp
Go to the documentation of this file.
1// @HEADER
2// ***********************************************************************
3//
4// Stokhos Package
5// Copyright (2009) Sandia Corporation
6//
7// Under terms of Contract DE-AC04-94AL85000, there is a non-exclusive
8// license for use of this work by or on behalf of the U.S. Government.
9//
10// Redistribution and use in source and binary forms, with or without
11// modification, are permitted provided that the following conditions are
12// met:
13//
14// 1. Redistributions of source code must retain the above copyright
15// notice, this list of conditions and the following disclaimer.
16//
17// 2. Redistributions in binary form must reproduce the above copyright
18// notice, this list of conditions and the following disclaimer in the
19// documentation and/or other materials provided with the distribution.
20//
21// 3. Neither the name of the Corporation nor the names of the
22// contributors may be used to endorse or promote products derived from
23// this software without specific prior written permission.
24//
25// THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
26// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
27// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
28// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
29// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
30// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
31// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
32// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
33// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
34// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
35// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
36//
37// Questions? Contact Eric T. Phipps (etphipp@sandia.gov).
38//
39// ***********************************************************************
40// @HEADER
41
42#ifndef STOKHOS_DYN_ARRAY_TRAITS_HPP
43#define STOKHOS_DYN_ARRAY_TRAITS_HPP
44
45#include <new>
46#include <cstring>
47
48#include "Kokkos_Core_fwd.hpp"
49
50namespace Stokhos {
51
53
57 template <typename T> struct IsScalarType2 {
58 static const bool value = false;
59 };
60
62#define STOKHOS_BUILTIN_SPECIALIZATION(t) \
63 template <> struct IsScalarType2< t > { \
64 static const bool value = true; \
65 };
66
71
72#undef STOKHOS_BUILTIN_SPECIALIZATION
73
78 template <typename T, typename device_t,
79 bool isScalar = IsScalarType2<T>::value>
81
82 typedef T value_type;
83 typedef device_t execution_space;
84
86 static
87 KOKKOS_INLINE_FUNCTION
88 void copy(const volatile T* src, volatile T* dest, std::size_t sz) {
89 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
90 for (std::size_t i=0; i<sz; ++i)
91 *(dest++) = *(src++);
92 }
93
95 static
96 KOKKOS_INLINE_FUNCTION
97 void copy(const volatile T* src, T* dest, std::size_t sz) {
98 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
99 for (std::size_t i=0; i<sz; ++i)
100 *(dest++) = *(src++);
101 }
102
104 static
105 KOKKOS_INLINE_FUNCTION
106 void copy(const T* src, volatile T* dest, std::size_t sz) {
107 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
108 for (std::size_t i=0; i<sz; ++i)
109 *(dest++) = *(src++);
110 }
111
113 static
114 KOKKOS_INLINE_FUNCTION
115 void copy(const T* src, T* dest, std::size_t sz) {
116 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
117 for (std::size_t i=0; i<sz; ++i)
118 *(dest++) = *(src++);
119 }
120
122 static
123 KOKKOS_INLINE_FUNCTION
124 void zero(T* dest, std::size_t sz) {
125 if (sz > 0) std::memset(dest,0,sz*sizeof(T));
126 }
127
129 static
130 KOKKOS_INLINE_FUNCTION
131 void zero(volatile T* dest, std::size_t sz) {
132 //if (sz > 0) std::memset(dest,0,sz*sizeof(T));
133 for (std::size_t i=0; i<sz; ++i)
134 *(dest++) = T(0.);
135 }
136
138 static
139 KOKKOS_INLINE_FUNCTION
140 void fill(T* dest, std::size_t sz, const T& v) {
141 //if (sz > 0) std::memset(dest,v,sz*sizeof(T));
142 for (std::size_t i=0; i<sz; ++i)
143 *(dest++) = v;
144 }
145
147 static
148 KOKKOS_INLINE_FUNCTION
149 void fill(volatile T* dest, std::size_t sz, const T& v) {
150 //if (sz > 0) std::memset(dest,v,sz*sizeof(T));
151 for (std::size_t i=0; i<sz; ++i)
152 *(dest++) = v;
153 }
154
156 static
157 KOKKOS_INLINE_FUNCTION
158 T* get_and_fill(std::size_t sz, const T& x = T(0.0)) {
159 T* m = 0;
160 if (sz > 0) {
161 m = static_cast<T* >(operator new(sz*sizeof(T)));
162 //std::memset(m,x,sz*sizeof(T));
163 for (std::size_t i=0; i<sz; ++i)
164 m[i] = x;
165 }
166 return m;
167 }
168
173 static
174 KOKKOS_INLINE_FUNCTION
175 T* get_and_fill(const T* src, std::size_t sz) {
176 T* m = 0;
177 if (sz > 0) {
178 m = static_cast<T* >(operator new(sz*sizeof(T)));
179 for (std::size_t i=0; i<sz; ++i)
180 m[i] = src[i];
181 }
182 return m;
183 }
184
189 static
190 KOKKOS_INLINE_FUNCTION
191 T* get_and_fill(const volatile T* src, std::size_t sz) {
192 T* m = 0;
193 if (sz > 0) {
194 m = static_cast<T* >(operator new(sz*sizeof(T)));
195 for (std::size_t i=0; i<sz; ++i)
196 m[i] = src[i];
197 }
198 return m;
199 }
200
202 static
203 KOKKOS_INLINE_FUNCTION
204 void destroy_and_release(T* m, std::size_t sz) {
205 if (sz > 0) operator delete((void*) m);
206 }
207
209 static
210 KOKKOS_INLINE_FUNCTION
211 void destroy_and_release(volatile T* m, std::size_t sz) {
212 if (sz > 0) operator delete((void*) m);
213 }
214 };
215
219 template <typename T, typename device_t>
220 struct DynArrayTraits<T, device_t, false> {
221
222 typedef T value_type;
223 typedef device_t execution_space;
224
226 static
227 KOKKOS_INLINE_FUNCTION
228 void fill(T* dest, std::size_t sz, const T& v) {
229 for (std::size_t i=0; i<sz; ++i)
230 *(dest++) = v;
231 }
232
234 static
235 KOKKOS_INLINE_FUNCTION
236 void fill(volatile T* dest, std::size_t sz, const T& v) {
237 for (std::size_t i=0; i<sz; ++i)
238 *(dest++) = v;
239 }
240
242 static
243 KOKKOS_INLINE_FUNCTION
244 void copy(const volatile T* src, volatile T* dest, std::size_t sz) {
245 for (std::size_t i=0; i<sz; ++i)
246 *(dest++) = *(src++);
247 }
248
250 static
251 KOKKOS_INLINE_FUNCTION
252 void copy(const volatile T* src, T* dest, std::size_t sz) {
253 for (std::size_t i=0; i<sz; ++i)
254 *(dest++) = *(src++);
255 }
256
258 static
259 KOKKOS_INLINE_FUNCTION
260 void copy(const T* src, volatile T* dest, std::size_t sz) {
261 for (std::size_t i=0; i<sz; ++i)
262 *(dest++) = *(src++);
263 }
264
266 static
267 KOKKOS_INLINE_FUNCTION
268 void copy(const T* src, T* dest, std::size_t sz) {
269 for (std::size_t i=0; i<sz; ++i)
270 *(dest++) = *(src++);
271 }
272
274 static
275 KOKKOS_INLINE_FUNCTION
276 void zero(T* dest, std::size_t sz) {
277 for (std::size_t i=0; i<sz; ++i)
278 *(dest++) = T(0.);
279 }
280
282 static
283 KOKKOS_INLINE_FUNCTION
284 void zero(volatile T* dest, std::size_t sz) {
285 for (std::size_t i=0; i<sz; ++i)
286 *(dest++) = T(0.);
287 }
288
290 static
291 KOKKOS_INLINE_FUNCTION
292 T* get_and_fill(std::size_t sz, const T& x = T(0.0)) {
293 T* m = 0;
294 if (sz > 0) {
295 m = static_cast<T* >(operator new(sz*sizeof(T)));
296 T* p = m;
297 for (std::size_t i=0; i<sz; ++i)
298 new (p++) T(x);
299 }
300 return m;
301 }
302
307 static
308 KOKKOS_INLINE_FUNCTION
309 T* get_and_fill(const T* src, std::size_t sz) {
310 T* m = 0;
311 if (sz > 0) {
312 m = static_cast<T* >(operator new(sz*sizeof(T)));
313 T* p = m;
314 for (std::size_t i=0; i<sz; ++i)
315 new (p++) T(*(src++));
316 }
317 return m;
318 }
319
324 static
325 KOKKOS_INLINE_FUNCTION
326 T* get_and_fill(const volatile T* src, std::size_t sz) {
327 T* m = 0;
328 if (sz > 0) {
329 m = static_cast<T* >(operator new(sz*sizeof(T)));
330 T* p = m;
331 for (std::size_t i=0; i<sz; ++i)
332 new (p++) T(*(src++));
333 }
334 return m;
335 }
336
338 static
339 KOKKOS_INLINE_FUNCTION
340 void destroy_and_release(T* m, std::size_t sz) {
341 T* e = m+sz;
342 for (T* b = m; b!=e; b++)
343 b->~T();
344 operator delete((void*) m);
345 }
346
348 static
349 KOKKOS_INLINE_FUNCTION
350 void destroy_and_release(volatile T* m, std::size_t sz) {
351 T* e = m+sz;
352 for (T* b = m; b!=e; b++)
353 b->~T();
354 operator delete((void*) m);
355 }
356 };
357
358#if defined(KOKKOS_ENABLE_CUDA)
359
366 template <typename T>
367 struct DynArrayTraits<T,Kokkos::Cuda,true> {
368
369 typedef T value_type;
370 typedef Kokkos::Cuda execution_space;
371
373 static
374 KOKKOS_INLINE_FUNCTION
375 void copy(const volatile T* src, volatile T* dest, std::size_t sz) {
376 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
377 for (std::size_t i=0; i<sz; ++i)
378 *(dest++) = *(src++);
379 }
380
382 static
383 KOKKOS_INLINE_FUNCTION
384 void copy(const volatile T* src, T* dest, std::size_t sz) {
385 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
386 for (std::size_t i=0; i<sz; ++i)
387 *(dest++) = *(src++);
388 }
389
391 static
392 KOKKOS_INLINE_FUNCTION
393 void copy(const T* src, volatile T* dest, std::size_t sz) {
394 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
395 for (std::size_t i=0; i<sz; ++i)
396 *(dest++) = *(src++);
397 }
398
400 static
401 KOKKOS_INLINE_FUNCTION
402 void copy(const T* src, T* dest, std::size_t sz) {
403 //if (sz > 0) std::memcpy(dest,src,sz*sizeof(T));
404 for (std::size_t i=0; i<sz; ++i)
405 *(dest++) = *(src++);
406 }
407
409 static
410 KOKKOS_INLINE_FUNCTION
411 void zero(T* dest, std::size_t sz) {
412 if (sz > 0) std::memset(dest,0,sz*sizeof(T));
413 }
414
416 static
417 KOKKOS_INLINE_FUNCTION
418 void zero(volatile T* dest, std::size_t sz) {
419 //if (sz > 0) std::memset(dest,0,sz*sizeof(T));
420 for (std::size_t i=0; i<sz; ++i)
421 *(dest++) = T(0.);
422 }
423
425 static
426 KOKKOS_INLINE_FUNCTION
427 void fill(T* dest, std::size_t sz, const T& v) {
428 //if (sz > 0) std::memset(dest,v,sz*sizeof(T));
429 for (std::size_t i=0; i<sz; ++i)
430 *(dest++) = v;
431 }
432
434 static
435 KOKKOS_INLINE_FUNCTION
436 void fill(volatile T* dest, std::size_t sz, const T& v) {
437 //if (sz > 0) std::memset(dest,v,sz*sizeof(T));
438 for (std::size_t i=0; i<sz; ++i)
439 *(dest++) = v;
440 }
441
443 static
444 KOKKOS_INLINE_FUNCTION
445 T* get_and_fill(std::size_t sz, const T& x = T(0.0)) {
446 T* m = 0;
447 if (sz > 0) {
448#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
449 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
450#else
451 m = static_cast<T* >(operator new(sz*sizeof(T)));
452#endif
453 //std::memset(m,x,sz*sizeof(T));
454 for (std::size_t i=0; i<sz; ++i)
455 m[i] = x;
456 }
457 return m;
458 }
459
464 static
465 KOKKOS_INLINE_FUNCTION
466 T* get_and_fill(const T* src, std::size_t sz) {
467 T* m = 0;
468 if (sz > 0) {
469#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
470 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
471#else
472 m = static_cast<T* >(operator new(sz*sizeof(T)));
473#endif
474 for (std::size_t i=0; i<sz; ++i)
475 m[i] = src[i];
476 }
477 return m;
478 }
479
484 static
485 KOKKOS_INLINE_FUNCTION
486 T* get_and_fill(const volatile T* src, std::size_t sz) {
487 T* m = 0;
488 if (sz > 0) {
489#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
490 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
491#else
492 m = static_cast<T* >(operator new(sz*sizeof(T)));
493#endif
494 for (std::size_t i=0; i<sz; ++i)
495 m[i] = src[i];
496 }
497 return m;
498 }
499
501 static
502 KOKKOS_INLINE_FUNCTION
503 void destroy_and_release(T* m, std::size_t sz) {
504#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
505 cudaFree(m);
506#else
507 if (sz > 0) operator delete((void*) m);
508#endif
509 }
510
512 static
513 KOKKOS_INLINE_FUNCTION
514 void destroy_and_release(volatile T* m, std::size_t sz) {
515#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
516 cudaFree(m);
517#else
518 if (sz > 0) operator delete((void*) m);
519#endif
520 }
521 };
522
528 template <typename T>
529 struct DynArrayTraits<T, Kokkos::Cuda, false> {
530
531 typedef T value_type;
532 typedef Kokkos::Cuda execution_space;
533
535 static
536 KOKKOS_INLINE_FUNCTION
537 void fill(T* dest, std::size_t sz, const T& v) {
538 for (std::size_t i=0; i<sz; ++i)
539 *(dest++) = v;
540 }
541
543 static
544 KOKKOS_INLINE_FUNCTION
545 void fill(volatile T* dest, std::size_t sz, const T& v) {
546 for (std::size_t i=0; i<sz; ++i)
547 *(dest++) = v;
548 }
549
551 static
552 KOKKOS_INLINE_FUNCTION
553 void copy(const volatile T* src, volatile T* dest, std::size_t sz) {
554 for (std::size_t i=0; i<sz; ++i)
555 *(dest++) = *(src++);
556 }
557
559 static
560 KOKKOS_INLINE_FUNCTION
561 void copy(const volatile T* src, T* dest, std::size_t sz) {
562 for (std::size_t i=0; i<sz; ++i)
563 *(dest++) = *(src++);
564 }
565
567 static
568 KOKKOS_INLINE_FUNCTION
569 void copy(const T* src, volatile T* dest, std::size_t sz) {
570 for (std::size_t i=0; i<sz; ++i)
571 *(dest++) = *(src++);
572 }
573
575 static
576 KOKKOS_INLINE_FUNCTION
577 void copy(const T* src, T* dest, std::size_t sz) {
578 for (std::size_t i=0; i<sz; ++i)
579 *(dest++) = *(src++);
580 }
581
583 static
584 KOKKOS_INLINE_FUNCTION
585 void zero(T* dest, std::size_t sz) {
586 for (std::size_t i=0; i<sz; ++i)
587 *(dest++) = T(0.);
588 }
589
591 static
592 KOKKOS_INLINE_FUNCTION
593 void zero(volatile T* dest, std::size_t sz) {
594 for (std::size_t i=0; i<sz; ++i)
595 *(dest++) = T(0.);
596 }
597
599 static
600 KOKKOS_INLINE_FUNCTION
601 T* get_and_fill(std::size_t sz, const T& x = T(0.0)) {
602 T* m = 0;
603 if (sz > 0) {
604#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
605 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
606#else
607 m = static_cast<T* >(operator new(sz*sizeof(T)));
608#endif
609 T* p = m;
610 for (std::size_t i=0; i<sz; ++i)
611 new (p++) T(x);
612 }
613 return m;
614 }
615
620 static
621 KOKKOS_INLINE_FUNCTION
622 T* get_and_fill(const T* src, std::size_t sz) {
623 T* m = 0;
624 if (sz > 0) {
625#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
626 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
627#else
628 m = static_cast<T* >(operator new(sz*sizeof(T)));
629#endif
630 T* p = m;
631 for (std::size_t i=0; i<sz; ++i)
632 new (p++) T(*(src++));
633 }
634 return m;
635 }
636
641 static
642 KOKKOS_INLINE_FUNCTION
643 T* get_and_fill(const volatile T* src, std::size_t sz) {
644 T* m = 0;
645 if (sz > 0) {
646#if defined( CUDA_VERSION ) && ( 6000 <= CUDA_VERSION ) && defined(KOKKOS_ENABLE_CUDA_UVM) && !defined( __CUDA_ARCH__ )
647 cudaMallocManaged( (void**) &m, sz*sizeof(T), cudaMemAttachGlobal );
648#else
649 m = static_cast<T* >(operator new(sz*sizeof(T)));
650#endif
651 T* p = m;
652 for (std::size_t i=0; i<sz; ++i)
653 new (p++) T(*(src++));
654 }
655 return m;
656 }
657
659 static
660 KOKKOS_INLINE_FUNCTION
661 void destroy_and_release(T* m, std::size_t sz) {
662 T* e = m+sz;
663 for (T* b = m; b!=e; b++)
664 b->~T();
665 operator delete((void*) m);
666 }
667
669 static
670 KOKKOS_INLINE_FUNCTION
671 void destroy_and_release(volatile T* m, std::size_t sz) {
672 T* e = m+sz;
673 for (T* b = m; b!=e; b++)
674 b->~T();
675 operator delete((void*) m);
676 }
677 };
678
679#endif
680
681} // namespace Stokhos
682
683#endif // STOKHOS_DYN_ARRAY_TRAITS_HPP
#define STOKHOS_BUILTIN_SPECIALIZATION(t)
Specialization of above classes to built-in types.
Top-level namespace for Stokhos classes and functions.
static KOKKOS_INLINE_FUNCTION void destroy_and_release(volatile T *m, std::size_t sz)
Destroy array elements and release memory.
static KOKKOS_INLINE_FUNCTION void fill(T *dest, std::size_t sz, const T &v)
Fill array dest of length sz with value v.
static KOKKOS_INLINE_FUNCTION void copy(const volatile T *src, volatile T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(const T *src, std::size_t sz)
Get memory for new array of length sz and fill with entries from src.
static KOKKOS_INLINE_FUNCTION void copy(const T *src, volatile T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION void zero(T *dest, std::size_t sz)
Zero out array dest of length sz.
static KOKKOS_INLINE_FUNCTION void copy(const T *src, T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(const volatile T *src, std::size_t sz)
Get memory for new array of length sz and fill with entries from src.
static KOKKOS_INLINE_FUNCTION void zero(volatile T *dest, std::size_t sz)
Zero out array dest of length sz.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(std::size_t sz, const T &x=T(0.0))
Get memory for new array of length sz and fill with zeros.
static KOKKOS_INLINE_FUNCTION void copy(const volatile T *src, T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION void fill(volatile T *dest, std::size_t sz, const T &v)
Fill array dest of length sz with value v.
static KOKKOS_INLINE_FUNCTION void destroy_and_release(T *m, std::size_t sz)
Destroy array elements and release memory.
Dynamic array allocation class that is specialized for scalar i.e., fundamental or built-in types (fl...
static KOKKOS_INLINE_FUNCTION void copy(const T *src, T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(std::size_t sz, const T &x=T(0.0))
Get memory for new array of length sz and fill with zeros.
static KOKKOS_INLINE_FUNCTION void destroy_and_release(T *m, std::size_t sz)
Destroy array elements and release memory.
static KOKKOS_INLINE_FUNCTION void fill(T *dest, std::size_t sz, const T &v)
Fill array dest of length sz with value v.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(const volatile T *src, std::size_t sz)
Get memory for new array of length sz and fill with entries from src.
static KOKKOS_INLINE_FUNCTION void copy(const volatile T *src, T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION T * get_and_fill(const T *src, std::size_t sz)
Get memory for new array of length sz and fill with entries from src.
static KOKKOS_INLINE_FUNCTION void destroy_and_release(volatile T *m, std::size_t sz)
Destroy array elements and release memory.
static KOKKOS_INLINE_FUNCTION void fill(volatile T *dest, std::size_t sz, const T &v)
Fill array dest of length sz with value v.
static KOKKOS_INLINE_FUNCTION void zero(T *dest, std::size_t sz)
Zero out array dest of length sz.
static KOKKOS_INLINE_FUNCTION void zero(volatile T *dest, std::size_t sz)
Zero out array dest of length sz.
static KOKKOS_INLINE_FUNCTION void copy(const T *src, volatile T *dest, std::size_t sz)
Copy array from src to dest of length sz.
static KOKKOS_INLINE_FUNCTION void copy(const volatile T *src, volatile T *dest, std::size_t sz)
Copy array from src to dest of length sz.
Base template specification for IsScalarType.