-
Notifications
You must be signed in to change notification settings - Fork 1
/
Compute.cu
140 lines (99 loc) · 4 KB
/
Compute.cu
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
#include "globals.h"
#include "timing.h"
#include "Compute.h"
#include <sstream>
void write_kspace_data(const char*, complex<float>*);
__global__ void d_multiplyComplex(cufftComplex*, cufftComplex*,
cufftComplex*, int);
__global__ void d_prepareDensity(int, float*, cufftComplex*, int);
Compute::Compute() {};
Compute::~Compute() {};
void Compute::Initialize(string command) {
this->input_command = command;
this->compute_id = n_computes - 1;
istringstream iss(command);
string word;
// Put the first word, which should be "compute"
iss >> word;
// Detect compute style
iss >> word;
if (word == "avg_sk") {
// For avg_sk, will calculate the average S(k) for one type of particle.
// iparams[0] = particle type
// iparams[1] will track the number of calls to doCompute()
this->style = "avg_sk";
iss >> word;
this->iparams[0] = stoi(word) - 1; // Store relevant type, shifting by one to zero index.
if (this->iparams[0] >= ntypes)
die("Invalid type to calculate avg_sk");
this->iparams[1] = 0;
// Set defaults for optional arguments
this->compute_wait = 0;
this->compute_freq = 100;
// Check for optional arguments
while (!iss.eof()) {
iss >> word;
if (word == "freq") {
iss >> word;
this->compute_freq = stoi(word);
}
else if (word == "wait") {
iss >> word;
this->compute_wait = stoi(word);
}
}
cout << " Calculating <S(k)> for component " << iparams[0] + 1 << " every " << this->compute_freq <<
" steps after " << this->compute_wait << " steps have passed." << endl;
}// word == avg_sk
else {
die("Undefined compute style!");
}
}
void Compute::allocStorage() {
if ( this->style == "avg_sk" ) {
this->cpx.resize(M);
for (int i = 0; i < M; i++)
this->cpx.at(i) = 0.0f;
cout << " this->cpx has initial size " << this->cpx.size() << " and capacity " << this->cpx.capacity() << endl;
}
}
void Compute::doCompute() {
compute_t_in = int(time(0));
if (this->style == "avg_sk") {
// Extract the density of the relevant type
d_prepareDensity<<<M_Grid, M_Block>>>(this->iparams[0], d_all_rho, d_cpx1, M);
check_cudaError("Compute->doCompute.style = avg_sk prepare density");
// fourier from d_cpx1 to d_cpx2 forward
cufftExecC2C(fftplan, d_cpx1, d_cpx2, CUFFT_FORWARD);
check_cudaError("Compute->doCompute.style = avg_sk cufftExec");
// Multiply by the complex conjugate and scale by 1/M
// Store it in d_cpx1 as the values inside are not needed at this point
d_multiplyComplex<<<M_Grid, M_Block>>> (d_cpx2, d_cpx2, d_cpx1, M);
check_cudaError("Compute->doCompute.style = avg_sk multiplyComplex");
// Copy data to host and store
// NOTE: this should probably be stored on the device and only
// communicated when writing, but may be OK for now.
cudaMemcpy(cpx1, d_cpx1, M * sizeof(cufftComplex), cudaMemcpyDeviceToHost);
for (int i = 0; i < M; i++)
this->cpx.at(i) += cpx1[i].x + I * cpx1[i].y;
this->iparams[1] += 1;
}
compute_t_out = int(time(0));
compute_tot_time += compute_t_out - compute_t_in;
}
// Results are going to be written at log_freq intervals
void Compute::writeResults(int compute_id) {
if (this->style == "avg_sk") {
for (int i = 0; i < M; i++) {
if (this->iparams[1] > 0)
k_tmp[i] = this->cpx[i] / float(this->iparams[1]);
else
k_tmp[i] = 0.0f;
}
char nm[50];
// compute_id is used in the name instead of "type" in case multiple
// computes operate on the same type
sprintf(nm, "avg_sk_%d.dat", compute_id);
write_kspace_data(nm, k_tmp);
}
}