-
Notifications
You must be signed in to change notification settings - Fork 2
/
Copy pathcuda_fusedR.h
176 lines (166 loc) · 6.33 KB
/
cuda_fusedR.h
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
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
/**
*
* OHIO STATE UNIVERSITY SOFTWARE DISTRIBUTION LICENSE
*
* Parallel CCD++ on GPU (the “Software”) Copyright (c) 2017, The Ohio State
* University. All rights reserved.
*
* The Software is available for download and use subject to the terms and
* conditions of this License. Access or use of the Software constitutes acceptance
* and agreement to the terms and conditions of this License. Redistribution and
* use of the Software in source and binary forms, with or without modification,
* are permitted provided that the following conditions are met:
*
* 1. Redistributions of source code must retain the above copyright notice, this
* list of conditions and the capitalized paragraph below.
*
* 2. Redistributions in binary form must reproduce the above copyright notice,
* this list of conditions and the capitalized paragraph below in the documentation
* and/or other materials provided with the distribution.
*
* 3. The names of Ohio State University, or its faculty, staff or students may not
* be used to endorse or promote products derived from the Software without
* specific prior written permission.
*
* This software was produced with support from the National Science Foundation
* (NSF) through Award 1629548. Nothing in this work should be construed as
* reflecting the official policy or position of the Defense Department, the United
* States government, Ohio State University.
*
* THIS SOFTWARE HAS BEEN APPROVED FOR PUBLIC RELEASE, UNLIMITED DISTRIBUTION. THE
* SOFTWARE IS PROVIDED “AS IS” AND WITHOUT ANY EXPRESS, IMPLIED OR STATUTORY
* WARRANTIES, INCLUDING, BUT NOT LIMITED TO, WARRANTIES OF ACCURACY, COMPLETENESS,
* NONINFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. ACCESS OR USE OF THE SOFTWARE IS ENTIRELY AT THE USER’S RISK. IN
* NO EVENT SHALL OHIO STATE UNIVERSITY OR ITS FACULTY, STAFF OR STUDENTS BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
* DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
* SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR
* TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
* THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. THE SOFTWARE
* USER SHALL INDEMNIFY, DEFEND AND HOLD HARMLESS OHIO STATE UNIVERSITY AND ITS
* FACULTY, STAFF AND STUDENTS FROM ANY AND ALL CLAIMS, ACTIONS, DAMAGES, LOSSES,
* LIABILITIES, COSTS AND EXPENSES, INCLUDING ATTORNEYS’ FEES AND COURT COSTS,
* DIRECTLY OR INDIRECTLY ARISING OUT OF OR IN CONNECTION WITH ACCESS OR USE OF THE
* SOFTWARE.
*
*/
/**
*
* Author:
* Israt ([email protected])
*
* Contacts:
* Israt ([email protected])
* Aravind Sukumaran-Rajam ([email protected])
* P. (Saday) Sadayappan ([email protected])
*
*/
template<bool UseCache, unsigned POWER, unsigned Nth_POWER>
__global__ void updateR_gen(int const* __restrict__ R_colPtr,
int const* __restrict__ R_rowLim, unsigned * R_rowIdx, DTYPE *R_val,
const DTYPE * __restrict__ u, const DTYPE * __restrict__ v, int m,
int n, bool add, int * __restrict__ rowGroupPtr, int numRowsPerGroup,
DTYPE lambda, DTYPE * __restrict__ g_arr, DTYPE * __restrict__ h_arr,
const DTYPE * __restrict__ u_new, const DTYPE * __restrict__ Wt_p,
const DTYPE * __restrict__ Ht_p, int t) {
unsigned int tId = threadIdx.x;
unsigned int laneId = tId & (POWER - 1);
unsigned int c = (blockIdx.x * blockDim.x + tId) >> (Nth_POWER);
if (c < numRowsPerGroup) {
c = rowGroupPtr[c];
DTYPE g = 0, h = 0;
DTYPE vc = v[c];
DTYPE vc_p = Ht_p[c];
unsigned int colPtr = R_colPtr[c];
unsigned int nnz_row = R_rowLim[c] - R_colPtr[c]; //nnz_row = R_colPtr[c+1] - colPtr;
if (add) {
for (unsigned short i = laneId; i < nnz_row; i += POWER) {
int ii = R_rowIdx[colPtr + i];
DTYPE &Rval = R_val[colPtr + i];
DTYPE u_val = u[ii];
if (t > 0)
Rval -= Wt_p[ii] * vc_p;
Rval += u_val * vc;
DTYPE ul_new = u_new[ii];
g += ul_new * R_val[colPtr + i];
h += ul_new * ul_new;
}
#pragma unroll Nth_POWER
for (unsigned i = (POWER) >> 1; i >= 1; i = i >> 1) {
g += __shfl_down(g, i);
h += __shfl_down(h, i);
}
if (laneId == 0) {
h += lambda * (nnz_row);
g_arr[c] += g;
h_arr[c] += h;
}
} else {
for (unsigned short i = laneId; i < nnz_row; i += POWER)
R_val[colPtr + i] -= u[R_rowIdx[colPtr + i]] * vc;
}
}
}
template<bool UseCache>
__global__ void updateR_7(int const* __restrict__ R_colPtr,
int const* __restrict__ R_rowLim, unsigned * R_rowIdx, DTYPE *R_val,
const DTYPE * __restrict__ u, const DTYPE * __restrict__ v, int m,
int n, bool add, int * __restrict__ rowGroupPtr, int numRowsPerGroup,
DTYPE lambda, DTYPE * __restrict__ g_arr, DTYPE * __restrict__ h_arr,
const DTYPE * __restrict__ u_new, const DTYPE * __restrict__ Wt_p,
const DTYPE * __restrict__ Ht_p, int t) {
unsigned int tId = threadIdx.x;
unsigned int laneId = tId & 63;
unsigned int c = (blockIdx.x * blockDim.x + tId) >> 6;
extern __shared__ volatile DTYPE SD[];
if (c < numRowsPerGroup) {
c = rowGroupPtr[c];
DTYPE g = 0, h = 0;
DTYPE vc = v[c];
DTYPE vc_p = Ht_p[c];
unsigned int colPtr = R_colPtr[c], nnz_row = R_rowLim[c] - colPtr; //nnz_row = R_colPtr[c+1] - colPtr;
if (add) {
for (long i = laneId; i < nnz_row; i += 64) {
int ii = R_rowIdx[colPtr + i];
DTYPE &Rval = R_val[colPtr + i];
DTYPE u_val = u[ii];
if (t > 0)
Rval -= Wt_p[ii] * vc_p;
Rval += u_val * vc;
DTYPE ul_new = u_new[ii];
g += ul_new * Rval;
h += ul_new * ul_new;
}
DTYPE newvj = 0;
g += __shfl_down(g, 16);
g += __shfl_down(g, 8);
g += __shfl_down(g, 4);
g += __shfl_down(g, 2);
g += __shfl_down(g, 1);
h += __shfl_down(h, 16);
h += __shfl_down(h, 8);
h += __shfl_down(h, 4);
h += __shfl_down(h, 2);
h += __shfl_down(h, 1);
if ((tId & 31) == 0) {
SD[tId >> 5] = g;
SD[blockDim.x / 32 + (tId >> 5)] = h;
}
__syncthreads();
if (laneId == 0) {
g += SD[(tId >> 5) + 1];
h += SD[(blockDim.x / 32 + (tId >> 5)) + 1];
h += lambda * (nnz_row);
g_arr[c] += g;
h_arr[c] += h;
}
} else {
for (long i = laneId; i < nnz_row; i += 64) {
int ii = R_rowIdx[colPtr + i];
R_val[colPtr + i] -= u[ii] * vc;
}
}
}
}