Skip to content

Commit b9ce683

Browse files
kunkunblueberrykunkunblueberry
committed
fix:dngvd.hip.cu run properly in DCU
1 parent 34f564e commit b9ce683

1 file changed

Lines changed: 37 additions & 255 deletions

File tree

‎source/module_hsolver/kernels/rocm/dngvd_op.hip.cu‎

Lines changed: 37 additions & 255 deletions
Original file line numberDiff line numberDiff line change
@@ -10,18 +10,11 @@ namespace hsolver {
1010
static hipsolverHandle_t hipsolver_H = nullptr;
1111

1212
void createGpuSolverHandle() {
13-
if (hipsolver_H == nullptr)
14-
{
15-
hipsolverErrcheck(hipsolverCreate(&hipsolver_H));
16-
}
13+
return;
1714
}
1815

1916
void destroyGpuSolverHandle() {
20-
if (hipsolver_H != nullptr)
21-
{
22-
hipsolverErrcheck(hipsolverDestroy(hipsolver_H));
23-
hipsolver_H = nullptr;
24-
}
17+
return;
2518
}
2619

2720
#ifdef __LCAO
@@ -35,222 +28,50 @@ void dngvd_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DE
3528
double* _vcc,
3629
int* fail_info)
3730
{
38-
// copied from ../cuda/dngvd_op.cu, "dngvd_op"
39-
assert(nstart == ldh);
40-
41-
hipErrcheck(hipMemcpy(_vcc, _hcc, sizeof(double) * ldh * nstart, hipMemcpyDeviceToDevice));
42-
// now vcc contains hcc
43-
44-
// prepare some values for hipsolverDnZhegvd_bufferSize
45-
int * devInfo = nullptr;
46-
int lwork = 0, info_gpu = 0;
47-
double * work = nullptr;
48-
hipErrcheck(hipMalloc((void**)&devInfo, sizeof(int)));
49-
hipsolverFillMode_t uplo = HIPSOLVER_FILL_MODE_UPPER;
50-
51-
// calculate the sizes needed for pre-allocated buffer.
52-
hipsolverErrcheck(hipsolverDnDsygvd_bufferSize(
53-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
54-
nstart,
55-
_vcc, ldh,
56-
_scc, ldh,
57-
_eigenvalue,
58-
&lwork));
59-
60-
// allocate memery
61-
hipErrcheck(hipMalloc((void**)&work, sizeof(double) * lwork));
62-
63-
// compute eigenvalues and eigenvectors.
64-
hipsolverErrcheck(hipsolverDnDsygvd(
65-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
66-
nstart,
67-
_vcc, ldh,
68-
const_cast<double *>(_scc), ldh,
69-
_eigenvalue,
70-
work, lwork, devInfo));
71-
72-
hipErrcheck(hipMemcpy(&info_gpu, devInfo, sizeof(int), hipMemcpyDeviceToHost));
73-
74-
// free the buffer
75-
hipErrcheck(hipFree(work));
76-
hipErrcheck(hipFree(devInfo));
77-
if(fail_info != nullptr) *fail_info = info_gpu;
78-
79-
80-
//std::vector<double> hcc(nstart * nstart, 0.0);
81-
//std::vector<double> scc(nstart * nstart, 0.0);
82-
//std::vector<double> vcc(nstart * nstart, 0.0);
83-
//std::vector<double> eigenvalue(nstart, 0);
84-
//hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(double) * hcc.size(), hipMemcpyDeviceToHost));
85-
//hipErrcheck(hipMemcpy(scc.data(), _scc, sizeof(double) * scc.size(), hipMemcpyDeviceToHost));
86-
//base_device::DEVICE_CPU* cpu_ctx = {};
87-
//dngvd_op<double, base_device::DEVICE_CPU>()(cpu_ctx,
88-
// nstart,
89-
// ldh,
90-
// hcc.data(),
91-
// scc.data(),
92-
// eigenvalue.data(),
93-
// vcc.data(),
94-
// fail_info);
95-
//hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(double) * vcc.size(), hipMemcpyHostToDevice));
96-
//hipErrcheck(hipMemcpy(_eigenvalue, eigenvalue.data(), sizeof(double) * eigenvalue.size(), hipMemcpyHostToDevice));
97-
}
98-
#endif // __LCAO
99-
100-
template <>
101-
void dngvd_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* ctx,
102-
const int nstart,
103-
const int ldh,
104-
const std::complex<float>* _hcc,
31+
std::vector<double> hcc(ldh * nstart, 0.0);
32+
std::vector<double> scc(ldh * nstart, 0.0);
33+
std::vector<double> vcc(ldh * nstart, 0.0);
34+
std::vector<double> eigenvalue(nstart, 0);
35+
hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(double) * hcc.size(), hipMemcpyDeviceToHost));
36+
hipErrcheck(hipMemcpy(scc.data(), _scc, sizeof(double) * scc.size(), hipMemcpyDeviceToHost));
37+
base_device::DEVICE_CPU* cpu_ctx = {};
38+
dngvd_op<double, base_device::DEVICE_CPU>()(cpu_ctx,
39+
nstart,
40+
ldh,
41+
hcc.data(),
42+
scc.data(),
43+
eigenvalue.data(),
44+
vcc.data(),
45+
fail_info);
46+
hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(double) * vcc.size(), hipMemcpyHostToDevice));
47+
hip const std::complex<float>* _hcc,
10548
const std::complex<float>* _scc,
10649
float* _eigenvalue,
10750
std::complex<float>* _vcc,
10851
int* fail_info)
10952
{
110-
// copied from ../cuda/dngvd_op.cu, "dngvd_op"
111-
// assert(nstart == ldh);
112-
113-
hipErrcheck(hipMemcpy(_vcc, _hcc, sizeof(std::complex<float>) * ldh * nstart, hipMemcpyDeviceToDevice));
114-
// now vcc contains hcc
115-
116-
// prepare some values for hipsolverDnZhegvd_bufferSize
117-
int * devInfo = nullptr;
118-
int lwork = 0, info_gpu = 0;
119-
float2 * work = nullptr;
120-
hipErrcheck(hipMalloc((void**)&devInfo, sizeof(int)));
121-
hipsolverFillMode_t uplo = HIPSOLVER_FILL_MODE_UPPER;
122-
123-
// calculate the sizes needed for pre-allocated buffer.
124-
hipsolverErrcheck(hipsolverDnChegvd_bufferSize(
125-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
126-
nstart,
127-
reinterpret_cast<const float2 *>(_vcc), ldh,
128-
reinterpret_cast<const float2 *>(_scc), ldh,
129-
_eigenvalue,
130-
&lwork));
131-
132-
// allocate memery
133-
hipErrcheck(hipMalloc((void**)&work, sizeof(float2) * lwork));
134-
135-
// compute eigenvalues and eigenvectors.
136-
hipsolverErrcheck(hipsolverDnChegvd(
137-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
138-
nstart,
139-
reinterpret_cast<float2 *>(_vcc), ldh,
140-
const_cast<float2 *>(reinterpret_cast<const float2 *>(_scc)), ldh,
141-
_eigenvalue,
142-
work, lwork, devInfo));
143-
144-
hipErrcheck(hipMemcpy(&info_gpu, devInfo, sizeof(int), hipMemcpyDeviceToHost));
145-
// free the buffer
146-
hipErrcheck(hipFree(work));
147-
hipErrcheck(hipFree(devInfo));
148-
if(fail_info != nullptr) *fail_info = info_gpu;
149-
150-
151-
//std::vector<std::complex<float>> hcc(nstart * nstart, {0, 0});
152-
//std::vector<std::complex<float>> scc(nstart * nstart, {0, 0});
153-
//std::vector<std::complex<float>> vcc(nstart * nstart, {0, 0});
154-
//std::vector<float> eigenvalue(nstart, 0);
155-
//hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(std::complex<float>) * hcc.size(), hipMemcpyDeviceToHost));
156-
//hipErrcheck(hipMemcpy(scc.data(), _scc, sizeof(std::complex<float>) * scc.size(), hipMemcpyDeviceToHost));
157-
//base_device::DEVICE_CPU* cpu_ctx = {};
158-
//dngvd_op<std::complex<float>, base_device::DEVICE_CPU>()(cpu_ctx,
159-
// nstart,
160-
// ldh,
161-
// hcc.data(),
162-
// scc.data(),
163-
// eigenvalue.data(),
164-
// vcc.data(),
165-
// fail_info);
166-
//hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(std::complex<float>) * vcc.size(), hipMemcpyHostToDevice));
167-
//hipErrcheck(hipMemcpy(_eigenvalue, eigenvalue.data(), sizeof(float) * eigenvalue.size(), hipMemcpyHostToDevice));
168-
}
169-
170-
template <>
171-
void dngvd_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* ctx,
172-
const int nstart,
173-
const int ldh,
53+
54+
std::vector<std::complex<float>> hcc(ldh * nstart, {0, 0});
55+
std::vector<s nstart,
56+
ldh,
57+
hcc.data(),
58+
scc.data(),
59+
eigenvalue.data(),
60+
vcc.data(),
61+
fail_info);
62+
hipErrchec const int ldh,
17463
const std::complex<double>* _hcc,
17564
const std::complex<double>* _scc,
17665
double* _eigenvalue,
17766
std::complex<double>* _vcc,
178-
int* fail_info)
179-
{
180-
// copied from ../cuda/dngvd_op.cu, "dngvd_op"
181-
// assert(nstart == ldh);
182-
183-
// save a copy of scc in case the diagonalization fails
184-
std::vector<std::complex<double>> scc(nstart * nstart, {0, 0});
185-
hipErrcheck(hipMemcpy(scc.data(), _scc, sizeof(std::complex<double>) * scc.size(), hipMemcpyDeviceToHost));
186-
187-
hipErrcheck(hipMemcpy(_vcc, _hcc, sizeof(std::complex<double>) * ldh * nstart, hipMemcpyDeviceToDevice));
188-
189-
// now vcc contains hcc
190-
191-
// prepare some values for hipsolverDnZhegvd_bufferSize
192-
int * devInfo = nullptr;
193-
int lwork = 0, info_gpu = 0;
194-
double2 * work = nullptr;
195-
hipErrcheck(hipMalloc((void**)&devInfo, sizeof(int)));
196-
hipsolverFillMode_t uplo = HIPSOLVER_FILL_MODE_UPPER;
197-
198-
// calculate the sizes needed for pre-allocated buffer.
199-
hipsolverErrcheck(hipsolverDnZhegvd_bufferSize(
200-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
201-
nstart,
202-
reinterpret_cast<const double2 *>(_vcc), ldh,
203-
reinterpret_cast<const double2 *>(_scc), ldh,
204-
_eigenvalue,
205-
&lwork));
206-
207-
// allocate memery
208-
hipErrcheck(hipMalloc((void**)&work, sizeof(double2) * lwork));
209-
210-
// compute eigenvalues and eigenvectors.
211-
hipsolverErrcheck(hipsolverDnZhegvd(
212-
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
213-
nstart,
214-
reinterpret_cast<double2 *>(_vcc), ldh,
215-
const_cast<double2 *>(reinterpret_cast<const double2 *>(_scc)), ldh,
216-
_eigenvalue,
217-
work, lwork, devInfo));
218-
219-
hipErrcheck(hipMemcpy(&info_gpu, devInfo, sizeof(int), hipMemcpyDeviceToHost));
220-
// free the buffer
221-
hipErrcheck(hipFree(work));
222-
hipErrcheck(hipFree(devInfo));
223-
if(fail_info != nullptr) *fail_info = info_gpu;
224-
225-
226-
227-
228-
229-
230-
231-
/*std::vector<std::complex<double>> hcc(nstart * nstart, {0, 0});
232-
std::vector<std::complex<double>> scc(nstart * nstart, {0, 0});
233-
std::vector<std::complex<double>> vcc(nstart * nstart, {0, 0});
234-
std::vector<double> eigenvalue(nstart, 0);
235-
hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(std::complex<double>) * hcc.size(), hipMemcpyDeviceToHost));
236-
hipErrcheck(hipMemcpy(scc.data(), _scc, sizeof(std::complex<double>) * scc.size(), hipMemcpyDeviceToHost));
237-
base_device::DEVICE_CPU* cpu_ctx = {};
67+
int* f base_device::DEVICE_CPU* cpu_ctx = {};
23868
dngvd_op<std::complex<double>, base_device::DEVICE_CPU>()(cpu_ctx,
23969
nstart,
24070
ldh,
24171
hcc.data(),
24272
scc.data(),
24373
eigenvalue.data(),
244-
vcc.data(),
245-
fail_info);
246-
hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(std::complex<double>) * vcc.size(), hipMemcpyHostToDevice));
247-
hipErrcheck(hipMemcpy(_eigenvalue, eigenvalue.data(), sizeof(double) * eigenvalue.size(), hipMemcpyHostToDevice));*/
248-
}
249-
250-
#ifdef __LCAO
251-
template <>
252-
void dnevx_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* ctx,
253-
const int nstart,
74+
const int nstart,
25475
const int ldh,
25576
const double* _hcc,
25677
const int m,
@@ -280,75 +101,36 @@ void dnevx_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const ba
280101
std::vector<std::complex<float>> hcc(ldh * ldh, {0, 0});
281102
std::vector<std::complex<float>> vcc(ldh * ldh, {0, 0});
282103
std::vector<float> eigenvalue(ldh, 0);
283-
hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(std::complex<float>) * hcc.size(), hipMemcpyDeviceToHost));
284-
base_device::DEVICE_CPU* cpu_ctx = {};
285-
dnevx_op<std::complex<float>, base_device::DEVICE_CPU>()(cpu_ctx,
286-
nstart,
287-
ldh,
288-
hcc.data(),
289-
m,
290-
eigenvalue.data(),
291-
vcc.data());
104+
hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(std::complex<float>) * hcc.size(), hipM vcc.data());
292105
hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(std::complex<float>) * vcc.size(), hipMemcpyHostToDevice));
293106
hipErrcheck(hipMemcpy(_eigenvalue, eigenvalue.data(), sizeof(float) * eigenvalue.size(), hipMemcpyHostToDevice));
294107
}
295108

296109
template <>
297110
void dnevx_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* ctx,
298111
const int nstart,
299-
const int ldh,
300-
const std::complex<double>* _hcc,
301-
const int m,
302-
double* _eigenvalue,
303-
std::complex<double>* _vcc)
304-
{
305-
std::vector<std::complex<double>> hcc(ldh * ldh, {0, 0});
112+
td::complex<double>> hcc(ldh * ldh, {0, 0});
306113
std::vector<std::complex<double>> vcc(ldh * ldh, {0, 0});
307114
std::vector<double> eigenvalue(ldh, 0);
308115
hipErrcheck(hipMemcpy(hcc.data(), _hcc, sizeof(std::complex<double>) * hcc.size(), hipMemcpyDeviceToHost));
309116
base_device::DEVICE_CPU* cpu_ctx = {};
310117
dnevx_op<std::complex<double>, base_device::DEVICE_CPU>()(cpu_ctx,
311118
nstart,
312-
ldh,
313-
hcc.data(),
314-
m,
315-
eigenvalue.data(),
316-
vcc.data());
317-
hipErrcheck(hipMemcpy(_vcc, vcc.data(), sizeof(std::complex<double>) * vcc.size(), hipMemcpyHostToDevice));
318-
hipErrcheck(hipMemcpy(_eigenvalue, eigenvalue.data(), sizeof(double) * eigenvalue.size(), hipMemcpyHostToDevice));
119+
ldh,nvalue.size(), hipMemcpyHostToDevice));
319120
}
320121

321122
template <>
322123
void dngvx_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
323124
const int nbase,
324125
const int ldh,
325126
std::complex<float>* hcc,
326-
std::complex<float>* scc,
327-
const int m,
328-
float* eigenvalue,
329-
std::complex<float>* vcc,
330-
int* fail_info)
331-
{
332-
}
333-
334-
template <>
335-
void dngvx_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
127+
std::complexce::DEVICE_GPU* d,
336128
const int nbase,
337129
const int ldh,
338130
std::complex<double>* hcc,
339131
std::complex<double>* scc,
340132
const int m,
341-
double* eigenvalue,
342-
std::complex<double>* vcc,
343-
int* fail_info)
344-
{
345-
}
346-
347-
#ifdef __LCAO
348-
template <>
349-
void dngvx_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
350-
const int nbase,
351-
const int ldh,
133+
const int ldh,
352134
double* hcc,
353135
double* scc,
354136
const int m,
@@ -359,4 +141,4 @@ void dngvx_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DE
359141
}
360142
#endif // __LCAO
361143

362-
} // namespace hsolver
144+
} // namespace hsolver

0 commit comments

Comments
 (0)