1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
|
/*************************************************************************
* Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/
#ifndef COPY_KERNEL_H_
#define COPY_KERNEL_H_
#include "common_kernel.h"
template<typename T>
struct FuncPassA {
__device__ T operator()(const T x, const T y) const {
return x;
}
};
#ifdef CUDA_HAS_HALF
template <>
struct FuncPassA<half> {
__device__ half2 operator()(const half2 x, const half2 y) const {
return x;
}
__device__ half operator()(const half x, const half y) const {
return x;
}
};
#endif
// Assumptions:
// - there is exactly 1 block
// - THREADS is the number of producer threads
// - this function is called by all producer threads
template<int UNROLL, int THREADS, typename T>
__device__ void Copy(volatile T * __restrict__ const dest,
const volatile T * __restrict__ const src, const int N) {
ReduceOrCopy<UNROLL, THREADS, FuncPassA<T>, T, false, false>(threadIdx.x,
dest, nullptr, src, nullptr, N);
}
// Assumptions:
// - there is exactly 1 block
// - THREADS is the number of producer threads
// - this function is called by all producer threads
template<int UNROLL, int THREADS, typename T>
__device__ void DoubleCopy(volatile T * __restrict__ const dest0,
volatile T * __restrict__ const dest1,
const volatile T * __restrict__ const src, const int N) {
ReduceOrCopy<UNROLL, THREADS, FuncPassA<T>, T, true, false>(threadIdx.x,
dest0, dest1, src, nullptr, N);
}
#endif // COPY_KERNEL_H_
|
