OpenACC
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
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
/* Model of a forest fire - a 2D rectangular grid of trees is initialized
with one random tree caught on fire. At each time step, trees that are not
on fire yet check their neighbors to the north, east, south, and west, and
if any of them are on fire, the tree catches fire with some percent chance.
The model runs for a certain number of time steps, which can be controlled
by the user. At the end of the simulation, the program outputs the total
percentage of trees burned. Tree data can also be output at each time step
if a filename is provided by the user.
There is one GPGPU thread for each tree in the
forest, including the boundaries. Each part of the simulation is fed to the
GPGPU as a kernel. If each time step is to be printed to a file, the GPGPU
copies the forest data back to the host at each time step first. Otherwise,
the GPGPU does not copy any data back until the end, when it sends back the
total number of trees burned to the host for printing.
*/
/* Author: Aaron Weeden, Shodor, 2015 */
/* Naming convention:
ALL_CAPS for constants
CamelCase for globals and functions
lowerCase for locals
*/
#include <stdbool.h> /* bool type */
#include <stdio.h> /* printf() */
#include <stdlib.h> /* atoi(), exit(), EXIT_FAILURE, malloc(), free(),
random() */
#include <string.h> /* strcpy() */
#include <unistd.h> /* getopt() */
#include "fire-curand.h"
/* Define descriptions of command line options */
#define N_ROWS_DESCR \
"The forest has this many rows of trees (positive integer)"
#define N_COLS_DESCR \
"The forest has this many columns of trees (positive integer)"
#define BURN_PROB_DESCR \
"Chance of catching fire if next to burning tree (positive integer [0..100])"
#define N_MAX_BURN_STEPS_DESCR \
"A burning tree stops burning after this many time steps (positive integer bigger than 1)"
#define N_STEPS_DESCR \
"Run for this many time steps (positive integer)"
#define RAND_SEED_DESCR \
"Seed value for the random number generator (positive integer)"
#define OUTPUT_FILENAME_DESCR \
"Filename to output tree data at each time step (file must not already exist)"
#define IS_RAND_FIRST_TREE_DESCR \
"Start the first on a random first tree as opposed to the middle tree"
/* Define default values for simulation parameters - each of these parameters
can also be changed later via user input */
#define N_ROWS_DEFAULT 21
#define N_COLS_DEFAULT N_ROWS_DEFAULT
#define BURN_PROB_DEFAULT 100
#define N_MAX_BURN_STEPS_DEFAULT 2
#define N_STEPS_DEFAULT N_ROWS_DEFAULT
#define RAND_SEED_DEFAULT 1
#define DEFAULT_IS_OUTPUTTING_EACH_STEP false
#define DEFAULT_IS_RAND_FIRST_TREE false
/* Define characters used on the command line to change the values of input
parameters */
#define N_ROWS_CHAR 'r'
#define N_COLS_CHAR 'c'
#define BURN_PROB_CHAR 'b'
#define N_MAX_BURN_STEPS_CHAR 'm'
#define N_STEPS_CHAR 't'
#define RAND_SEED_CHAR 's'
#define OUTPUT_FILENAME_CHAR 'o'
#define IS_RAND_FIRST_TREE_CHAR 'f'
/* Define options string used by getopt() - a colon after the character means
the parameter's value is specified by the user */
const char GETOPT_STRING[] = {
N_ROWS_CHAR, ':',
N_COLS_CHAR, ':',
BURN_PROB_CHAR, ':',
N_MAX_BURN_STEPS_CHAR, ':',
N_STEPS_CHAR, ':',
RAND_SEED_CHAR, ':',
OUTPUT_FILENAME_CHAR, ':',
IS_RAND_FIRST_TREE_CHAR
};
/* Define a mapping from the row and column of a given tree in a forest with
boundaries to the index of that tree in a 1D array that includes
boundaries */
#define TREE_MAP(row, col, nColsPlusBounds) ((row) * (nColsPlusBounds) + (col))
/* Define a mapping from the row and column of a given tree in a forest with
boundaries to the index of that tree in a 1D array that does not include
boundaries */
#define NEW_TREE_MAP(row, col, nCols) ((row - 1) * (nCols) + (col - 1))
/* Declare global parameters */
int NRows = N_ROWS_DEFAULT;
int NCols = N_COLS_DEFAULT;
int BurnProb = BURN_PROB_DEFAULT;
int NMaxBurnSteps = N_MAX_BURN_STEPS_DEFAULT;
int NSteps = N_STEPS_DEFAULT;
int RandSeed = RAND_SEED_DEFAULT;
bool IsOutputtingEachStep = DEFAULT_IS_OUTPUTTING_EACH_STEP;
bool IsRandFirstTree = DEFAULT_IS_RAND_FIRST_TREE;
char *OutputFilename;
/* Declare other needed global variables */
bool AreParamsValid = true; /* Do the model parameters have valid values? */
int NTrees; /* Total number of trees in the forest */
int NRowsPlusBounds; /* Number of rows of trees plus the boundary rows */
int NColsPlusBounds; /* Number of columns of trees plus the boundary columns */
int NTreesPlusBounds; /* Total number of trees plus the boundaries */
int MiddleRow; /* The tree in the middle is here. If an even number of rows,
this tree is just below the middle */
int MiddleCol; /* The tree in the middle is here. If an even number of cols,
this tree is just to the right of the middle */
int CurStep; /* The current time step */
int NBurnedTrees; /* The total number of burned trees */
char ExeName[32]; /* The name of the program executable */
int NMaxBurnStepsDigits; /* The number of digits in the max burn steps; used for
outputting tree data */
int *Trees; /* 1D tree array, contains a boundary around the outside of the
forest so the same neighbor checking algorithm can be used on
all cells */
int *NewTrees; /* Copy of 1D tree array - used so that we don't update the
forest too soon as we are deciding which new trees
should burn -- does not contain boundary */
FILE *OutputFile; /* For outputting tree data to a file */
float *RandNums; /* Random numbers, one per tree */
/* DECLARE FUNCTIONS */
/* Prints out a description of an integer command line option
@param optChar The character used to specify the option
@param optDescr The description of the option
@param optDefault The default value of the option
*/
void DescribeOptionInt(const char optChar, const char *optDescr,
const int optDefault) {
fprintf(stderr, "-%c : \n\t%s\n\tdefault: %d\n", optChar, optDescr,
optDefault);
}
/* Prints out a description of a string command line option
@param optChar The character used to specify the option
@param optDescr The description of the option
*/
void DescribeOptionNoDefault(const char optChar, const char *optDescr) {
fprintf(stderr, "-%c : \n\t%s\n", optChar, optDescr);
}
/* Print an error message
@param errorMsg Buffer containing the message
*/
void PrintError(const char *errorMsg) {
fprintf(stderr, "%s", errorMsg);
AreParamsValid = false;
}
/* Display to the user what options are available for running the program and
exit the program in failure
@param errorMsg The error message to print
*/
void PrintUsageAndExit() {
fprintf(stderr, "Usage: ");
fprintf(stderr, "%s [OPTIONS]\n", ExeName);
fprintf(stderr, "Where OPTIONS can be any of the following:\n");
DescribeOptionInt(N_ROWS_CHAR, N_ROWS_DESCR, N_ROWS_DEFAULT);
DescribeOptionInt(N_COLS_CHAR, N_COLS_DESCR, N_COLS_DEFAULT);
DescribeOptionInt(BURN_PROB_CHAR, BURN_PROB_DESCR, BURN_PROB_DEFAULT);
DescribeOptionInt(N_MAX_BURN_STEPS_CHAR, N_MAX_BURN_STEPS_DESCR,
N_MAX_BURN_STEPS_DEFAULT);
DescribeOptionInt(N_STEPS_CHAR, N_STEPS_DESCR, N_STEPS_DEFAULT);
DescribeOptionInt(RAND_SEED_CHAR, RAND_SEED_DESCR, RAND_SEED_DEFAULT);
DescribeOptionNoDefault(OUTPUT_FILENAME_CHAR, OUTPUT_FILENAME_DESCR);
DescribeOptionNoDefault(IS_RAND_FIRST_TREE_CHAR,
IS_RAND_FIRST_TREE_DESCR);
exit(EXIT_FAILURE);
}
/* Assert that a user's input value is an integer. If it is not, print
a usage message and an error message and exit the program.
@param param The user's input value
@param optChar The character used to specify the user's input value
*/
void AssertInteger(int param, const char optChar) {
char errorStr[64];
/* Get the user's input value, assume floating point */
const float floatParam = atof(optarg);
/* Make sure positive and integer */
if (floatParam != param) {
sprintf(errorStr, "ERROR: value for -%c must be an integer\n",
optChar);
PrintError(errorStr);
}
}
/* Assert that a user's input value is a positive integer. If it is not, print
a usage message and an error message and exit the program.
@param param The user's input value
@param optChar The character used the specify the user's input value
*/
void AssertPositiveInteger(int param, const char optChar) {
char errorStr[64];
/* Get the user's input value, assume floating point */
const float floatParam = atof(optarg);
/* Make sure positive and integer */
if (param < 1 || floatParam != param) {
sprintf(errorStr, "ERROR: value for -%c must be positive integer\n",
optChar);
PrintError(errorStr);
}
}
/* Assert that a user's input value is bigger than a value. If it is
not, print a usage message and an error message and exit the program.
@param param The user's input value
@param low The value the parameter needs to be bigger than
@param optChar The character used the specify the user's input value
*/
void AssertBigger(int param, const int val, const char optChar) {
char errorStr[64];
if (param <= val) {
sprintf(errorStr,
"ERROR: value for -%c must be bigger than %d\n", optChar, val);
PrintError(errorStr);
}
}
/* Assert that a user's input value is between two values, inclusive. If it is
not, print a usage message and an error message and exit the program.
@param param The user's input value
@param low The lowest value the parameter can be
@param high The highest value the parameter can be
@param optChar The character used the specify the user's input value
*/
void AssertBetweenInclusive(int param, const int low, const int high,
const char optChar) {
char errorStr[64];
if (param < low || param > high) {
sprintf(errorStr,
"ERROR: value for -%c must be between %d and %d, inclusive\n",
optChar, low, high);
PrintError(errorStr);
}
}
/* Exit if a file already exists */
void AssertFileDNE(const char *filename) {
char errorStr[64];
if (access(filename, F_OK) != -1) {
sprintf(errorStr,
"ERROR: File '%s' already exists\n", filename);
PrintError(errorStr);
}
}
/* Allow the user to change simulation parameters via the command line
@param argc The number of command line arguments to parse
@param argv The array of command line arguments to parse
*/
void GetUserOptions(const int argc, char **argv) {
char c; /* Loop control variable */
/* Loop over argv, setting parameter values given */
while ((c = getopt(argc, argv, GETOPT_STRING)) != -1) {
switch(c) {
case N_ROWS_CHAR:
NRows = atoi(optarg);
AssertPositiveInteger(NRows, N_ROWS_CHAR);
break;
case N_COLS_CHAR:
NCols = atoi(optarg);
AssertPositiveInteger(NCols, N_COLS_CHAR);
break;
case BURN_PROB_CHAR:
BurnProb = atoi(optarg);
AssertInteger(BurnProb, BURN_PROB_CHAR);
AssertBetweenInclusive(BurnProb, 0, 100, BURN_PROB_CHAR);
break;
case N_MAX_BURN_STEPS_CHAR:
NMaxBurnSteps = atoi(optarg);
AssertPositiveInteger(NMaxBurnSteps, N_MAX_BURN_STEPS_CHAR);
AssertBigger(NMaxBurnSteps, 1, N_MAX_BURN_STEPS_CHAR);
break;
case N_STEPS_CHAR:
NSteps = atoi(optarg);
AssertPositiveInteger(NSteps, N_STEPS_CHAR);
break;
case RAND_SEED_CHAR:
RandSeed = atoi(optarg);
AssertPositiveInteger(RandSeed, RAND_SEED_CHAR);
break;
case OUTPUT_FILENAME_CHAR:
IsOutputtingEachStep = true;
OutputFilename = optarg;
break;
case IS_RAND_FIRST_TREE_CHAR:
IsRandFirstTree = true;
break;
case '?':
default:
PrintError("ERROR: illegal option\n");
PrintUsageAndExit();
}
}
if (IsOutputtingEachStep) {
/* Make sure the output file does not exist (DNE) */
AssertFileDNE(OutputFilename);
}
}
/* Allocate dynamic memory */
void AllocateMemory() {
Trees = (int*)malloc(NTreesPlusBounds * sizeof(int));
NewTrees = (int*)malloc( NTrees * sizeof(int));
RandNums = (float*)malloc( NTrees * sizeof(float));
}
/* Generate a random integer between [min..max]
@param min Smallest integer to generate
@param max 1 more than the biggest integer to generate
@param id ID of tree for which to get a random number
@return random integer
*/
int RandBetween(const int min, const int max, const int id) {
return min + (RandNums[id] * (max - min));
}
/* Light a random tree on fire, set all other trees to be not burning */
void InitData(const int randRow, const int randCol) {
int row;
int col;
/* Set all trees as having burned for 0 time steps */
#pragma acc parallel loop present(Trees[0:NTreesPlusBounds], NewTrees[0:NTrees])
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 0;
}
}
/* Set the boundaries as burnt out */
#pragma acc parallel loop present(Trees[0:NTreesPlusBounds])
for (row = 0; row < NRowsPlusBounds; row++) {
/* Left */
Trees[TREE_MAP(row, 0, NColsPlusBounds)] = NMaxBurnSteps;
/* Top/Bottom */
if ((row == 0) || (row == NRows + 1)) {
for (col = 1; col < NCols + 1; col++) {
Trees[TREE_MAP(row, col, NColsPlusBounds)] = NMaxBurnSteps;
}
}
/* Right */
Trees[TREE_MAP(row, NCols + 1, NColsPlusBounds)] = NMaxBurnSteps;
}
#pragma acc update host(Trees[0:NTreesPlusBounds], NewTrees[0:NTrees])
if (IsRandFirstTree) {
/* Light a random tree on fire */
row = randRow;
col = randCol;
}
else {
/* Light the middle tree on fire */
row = MiddleRow + 1;
col = MiddleCol + 1;
}
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 1;
NBurnedTrees++;
#pragma acc update device(Trees[0:NTreesPlusBounds], NewTrees[0:NTrees], \
NBurnedTrees)
}
/* Output tree data for the current time step */
void OutputData() {
int row;
int col;
char buf[64];
/* Write the header for the current time step */
sprintf(buf, "Time step %d\n", CurStep);
fprintf(OutputFile, "%s", buf);
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
sprintf(buf, "%*d ",
NMaxBurnStepsDigits, Trees[TREE_MAP(row, col, NColsPlusBounds)]);
fprintf(OutputFile, "%s", buf);
}
fprintf(OutputFile, "\n");
}
/* Write the newline between time steps */
fprintf(OutputFile, "\n");
}
/* Return whether a given tree has burnt out
@param row The row index of the tree
@param col The column index of the tree
@return Whether the tree in the given row and column has burnt out
*/
bool IsBurntOut(const int row, const int col) {
return Trees[TREE_MAP(row, col, NColsPlusBounds)] >= NMaxBurnSteps;
}
/* Return whether a given tree is on fire
@param row The row index of the tree
@param col The column index of the tree
@return Whether the tree in the given row and column is on fire
*/
bool IsOnFire(const int row, const int col) {
return Trees[TREE_MAP(row, col, NColsPlusBounds)] > 0 &&
!IsBurntOut(row, col);
}
/* For trees already burning, increment the number of time steps they have
burned
*/
void ContinueBurning() {
int row;
int col;
#pragma acc parallel loop present(Trees[0:NTreesPlusBounds], NewTrees[0:NTrees])
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
if (IsOnFire(row, col)) {
NewTrees[NEW_TREE_MAP(row, col, NCols)] =
Trees[ TREE_MAP(row, col, NColsPlusBounds)] + 1;
}
}
}
}
/* Find trees that are not on fire yet and try to catch them on fire from
burning neighbor trees
*/
void BurnNew() {
int row;
int col;
/* Use curand to generate an array of random numbers, one per tree */
#pragma acc host_data use_device(RandNums[0:NTrees])
{
GenerateUniform(RandNums, NTrees);
}
#pragma acc parallel loop present(Trees[0:NTreesPlusBounds], \
NewTrees[0:NTrees], RandNums[0:NTrees], NBurnedTrees) \
reduction(+:NBurnedTrees)
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
if (!IsOnFire(row, col) && !IsBurntOut(row, col)) {
/* Check neighbors */
/* Top */
if ((IsOnFire(row-1, col) ||
/* Left */
IsOnFire(row, col-1) ||
/* Bottom */
IsOnFire(row+1, col) ||
/* Right */
IsOnFire(row, col+1)) &&
/* Apply random chance */
(RandBetween(0, 100, NEW_TREE_MAP(row, col, NCols)) <
BurnProb)) {
/* Catch the tree on fire */
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 1;
NBurnedTrees++;
}
}
}
}
}
/* Copy new tree data into old tree data */
void AdvanceTime() {
int row;
int col;
#pragma acc parallel loop present(Trees[0:NTreesWithBounds], NewTrees[0:NTrees])
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)];
}
}
}
/* Free allocated memory */
void FreeMemory() {
free(RandNums);
free(NewTrees);
free(Trees);
}
/* @param argc The number of command line arguments
@param argv String of command line arguments
*/
int main(int argc, char **argv) {
/* Set the program executable name */
strcpy(ExeName, argv[0]);
/* Allow the user to change simulation parameters via the command line */
GetUserOptions(argc, argv);
if (!AreParamsValid) {
/* Model parameters are not valid; exit early */
PrintUsageAndExit();
}
if (IsOutputtingEachStep) {
/* Open the output file */
OutputFile = fopen(OutputFilename, "w");
}
/* Do some calculations before splitting up the rows */
NTrees = NRows * NCols;
NRowsPlusBounds = NRows + 2;
NColsPlusBounds = NCols + 2;
NTreesPlusBounds = NRowsPlusBounds * NColsPlusBounds;
MiddleRow = NRows / 2;
MiddleCol = NCols / 2;
/* Allocate dynamic memory for the 1D tree arrays */
AllocateMemory();
/* Seed the random number generator */
InitCurandGenerator(RandSeed);
srandom(RandSeed);
/* Initialize number of burned trees */
NBurnedTrees = 0;
#pragma acc data create(Trees[0:NTreesPlusBounds], NewTrees[0:NTrees], \
RandNums[0:NTrees]) copy(NBurnedTrees)
{
/* Light a random tree on fire, set all other trees to be not burning */
InitData(random() % NRowsPlusBounds, random() % NColsPlusBounds);
/* Start the simulation looping for the specified number of time steps */
for (CurStep = 0; CurStep < NSteps; CurStep++) {
if (IsOutputtingEachStep) {
/* Output tree data for the current time step */
#pragma acc update host(Trees[0:NTreesPlusBounds])
OutputData();
}
/* For trees already burning, increment the number of time steps they have
burned */
ContinueBurning();
/* Find trees that are not on fire yet and try to catch them on fire from
burning neighbor trees */
BurnNew();
/* Copy new tree data into old tree data */
AdvanceTime();
}
}
/* Print the total percentage of trees burned */
printf("%.2f%% of the trees were burned\n",
(100.0 * NBurnedTrees) / NTrees);
if (IsOutputtingEachStep) {
/* Close the output file */
fclose(OutputFile);
}
/* Free allocated memory */
FreeMemory();
return 0;
}
CUDA
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
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
/* Model of a forest fire - a 2D rectangular grid of trees is initialized
with one random tree caught on fire. At each time step, trees that are not
on fire yet check their neighbors to the north, east, south, and west, and
if any of them are on fire, the tree catches fire with some percent chance.
The model runs for a certain number of time steps, which can be controlled
by the user. At the end of the simulation, the program outputs the total
percentage of trees burned. Tree data can also be output at each time step
if a filename is provided by the user.
There is one GPGPU thread for each tree in the
forest, including the boundaries. Each part of the simulation is fed to the
GPGPU as a kernel. If each time step is to be printed to a file, the GPGPU
copies the forest data back to the host at each time step first. Otherwise,
the GPGPU does not copy any data back until the end, when it sends back the
total number of trees burned to the host for printing.
*/
/* Author: Aaron Weeden, Shodor, 2015 */
/* Naming convention:
ALL_CAPS for constants
CamelCase for globals and functions
lowerCase for locals
*/
#include <curand.h>
#include <curand_kernel.h>
#include <stdbool.h> /* bool type */
#include <stdio.h> /* printf() */
#include <stdlib.h> /* atoi(), exit(), EXIT_FAILURE, malloc(), free(),
random() */
#include <string.h> /* strcpy() */
#include <unistd.h> /* getopt() */
/* CUDA threads per block = threads per warp * max warps per block */
#define THREADS_PER_BLOCK 1024
/* Define descriptions of command line options */
#define N_ROWS_DESCR \
"The forest has this many rows of trees (positive integer)"
#define N_COLS_DESCR \
"The forest has this many columns of trees (positive integer)"
#define BURN_PROB_DESCR \
"Chance of catching fire if next to burning tree (positive integer [0..100])"
#define N_MAX_BURN_STEPS_DESCR \
"A burning tree stops burning after this many time steps (positive integer bigger than 1)"
#define N_STEPS_DESCR \
"Run for this many time steps (positive integer)"
#define RAND_SEED_DESCR \
"Seed value for the random number generator (positive integer)"
#define OUTPUT_FILENAME_DESCR \
"Filename to output tree data at each time step (file must not already exist)"
#define IS_RAND_FIRST_TREE_DESCR \
"Start the first on a random first tree as opposed to the middle tree"
/* Define default values for simulation parameters - each of these parameters
can also be changed later via user input */
#define N_ROWS_DEFAULT 21
#define N_COLS_DEFAULT N_ROWS_DEFAULT
#define BURN_PROB_DEFAULT 100
#define N_MAX_BURN_STEPS_DEFAULT 2
#define N_STEPS_DEFAULT N_ROWS_DEFAULT
#define RAND_SEED_DEFAULT 1
#define DEFAULT_IS_OUTPUTTING_EACH_STEP false
#define DEFAULT_IS_RAND_FIRST_TREE false
/* Define characters used on the command line to change the values of input
parameters */
#define N_ROWS_CHAR 'r'
#define N_COLS_CHAR 'c'
#define BURN_PROB_CHAR 'b'
#define N_MAX_BURN_STEPS_CHAR 'm'
#define N_STEPS_CHAR 't'
#define RAND_SEED_CHAR 's'
#define OUTPUT_FILENAME_CHAR 'o'
#define IS_RAND_FIRST_TREE_CHAR 'f'
/* Define options string used by getopt() - a colon after the character means
the parameter's value is specified by the user */
const char GETOPT_STRING[] = {
N_ROWS_CHAR, ':',
N_COLS_CHAR, ':',
BURN_PROB_CHAR, ':',
N_MAX_BURN_STEPS_CHAR, ':',
N_STEPS_CHAR, ':',
RAND_SEED_CHAR, ':',
OUTPUT_FILENAME_CHAR, ':',
IS_RAND_FIRST_TREE_CHAR
};
/* Define a mapping from the row and column of a given tree in a forest with
boundaries to the index of that tree in a 1D array that includes
boundaries */
#define TREE_MAP(row, col, nColsPlusBounds) ((row) * (nColsPlusBounds) + (col))
/* Define a mapping from the row and column of a given tree in a forest with
boundaries to the index of that tree in a 1D array that does not include
boundaries */
#define NEW_TREE_MAP(row, col, nCols) ((row - 1) * (nCols) + (col - 1))
/* Declare global parameters */
int NRows = N_ROWS_DEFAULT;
int NCols = N_COLS_DEFAULT;
int BurnProb = BURN_PROB_DEFAULT;
int NMaxBurnSteps = N_MAX_BURN_STEPS_DEFAULT;
int NSteps = N_STEPS_DEFAULT;
int RandSeed = RAND_SEED_DEFAULT;
bool IsOutputtingEachStep = DEFAULT_IS_OUTPUTTING_EACH_STEP;
bool IsRandFirstTree = DEFAULT_IS_RAND_FIRST_TREE;
char *OutputFilename;
/* Declare other needed global variables */
bool AreParamsValid = true; /* Do the model parameters have valid values? */
int NTrees; /* Total number of trees in the forest */
int NRowsPlusBounds; /* Number of rows of trees plus the boundary rows */
int NColsPlusBounds; /* Number of columns of trees plus the boundary columns */
int NTreesPlusBounds; /* Total number of trees plus the boundaries */
int MiddleRow; /* The tree in the middle is here. If an even number of rows,
this tree is just below the middle */
int MiddleCol; /* The tree in the middle is here. If an even number of cols,
this tree is just to the right of the middle */
int CurStep; /* The current time step */
int NBurnedTrees; /* The total number of burned trees */
char ExeName[32]; /* The name of the program executable */
int NMaxBurnStepsDigits; /* The number of digits in the max burn steps; used for
outputting tree data */
int *Trees; /* 1D tree array, contains a boundary around the outside of the
forest so the same neighbor checking algorithm can be used on
all cells */
int *NewTrees; /* Copy of 1D tree array - used so that we don't update the
forest too soon as we are deciding which new trees
should burn -- does not contain boundary */
FILE *OutputFile; /* For outputting tree data to a file */
int NBlocks; /* Number of CUDA blocks to invoke for each kernel */
curandState *D_RandStates; /* Array of random states for curand */
/* Copies of arrays for the CUDA device */
int *D_Trees;
int *D_NewTrees;
int *D_NBurnedTrees;
/* DECLARE FUNCTIONS */
/* Set the random states array for each CUDA thread
@param Seed value for the random number generator
@param states The random states array
*/
__global__ void SetRandStates(const int RandSeed, curandState *states) {
curand_init(RandSeed, /* seed */
threadIdx.x, /* sequence */
0, /* offset */
&states[threadIdx.x]); /* state */
}
/* Prints out a description of an integer command line option
@param optChar The character used to specify the option
@param optDescr The description of the option
@param optDefault The default value of the option
*/
void DescribeOptionInt(const char optChar, const char *optDescr,
const int optDefault) {
fprintf(stderr, "-%c : \n\t%s\n\tdefault: %d\n", optChar, optDescr,
optDefault);
}
/* Prints out a description of a string command line option
@param optChar The character used to specify the option
@param optDescr The description of the option
*/
void DescribeOptionNoDefault(const char optChar, const char *optDescr) {
fprintf(stderr, "-%c : \n\t%s\n", optChar, optDescr);
}
/* Print an error message
@param errorMsg Buffer containing the message
*/
void PrintError(const char *errorMsg) {
fprintf(stderr, "%s", errorMsg);
AreParamsValid = false;
}
/* Display to the user what options are available for running the program and
exit the program in failure
@param errorMsg The error message to print
*/
void PrintUsageAndExit() {
fprintf(stderr, "Usage: ");
fprintf(stderr, "%s [OPTIONS]\n", ExeName);
fprintf(stderr, "Where OPTIONS can be any of the following:\n");
DescribeOptionInt(N_ROWS_CHAR, N_ROWS_DESCR, N_ROWS_DEFAULT);
DescribeOptionInt(N_COLS_CHAR, N_COLS_DESCR, N_COLS_DEFAULT);
DescribeOptionInt(BURN_PROB_CHAR, BURN_PROB_DESCR, BURN_PROB_DEFAULT);
DescribeOptionInt(N_MAX_BURN_STEPS_CHAR, N_MAX_BURN_STEPS_DESCR,
N_MAX_BURN_STEPS_DEFAULT);
DescribeOptionInt(N_STEPS_CHAR, N_STEPS_DESCR, N_STEPS_DEFAULT);
DescribeOptionInt(RAND_SEED_CHAR, RAND_SEED_DESCR, RAND_SEED_DEFAULT);
DescribeOptionNoDefault(OUTPUT_FILENAME_CHAR, OUTPUT_FILENAME_DESCR);
DescribeOptionNoDefault(IS_RAND_FIRST_TREE_CHAR,
IS_RAND_FIRST_TREE_DESCR);
exit(EXIT_FAILURE);
}
/* Assert that a user's input value is an integer. If it is not, print
a usage message and an error message and exit the program.
@param param The user's input value
@param optChar The character used to specify the user's input value
*/
void AssertInteger(int param, const char optChar) {
char errorStr[64];
/* Get the user's input value, assume floating point */
const float floatParam = atof(optarg);
/* Make sure positive and integer */
if (floatParam != param) {
sprintf(errorStr, "ERROR: value for -%c must be an integer\n",
optChar);
PrintError(errorStr);
}
}
/* Assert that a user's input value is a positive integer. If it is not, print
a usage message and an error message and exit the program.
@param param The user's input value
@param optChar The character used the specify the user's input value
*/
void AssertPositiveInteger(int param, const char optChar) {
char errorStr[64];
/* Get the user's input value, assume floating point */
const float floatParam = atof(optarg);
/* Make sure positive and integer */
if (param < 1 || floatParam != param) {
sprintf(errorStr, "ERROR: value for -%c must be positive integer\n",
optChar);
PrintError(errorStr);
}
}
/* Assert that a user's input value is bigger than a value. If it is
not, print a usage message and an error message and exit the program.
@param param The user's input value
@param low The value the parameter needs to be bigger than
@param optChar The character used the specify the user's input value
*/
void AssertBigger(int param, const int val, const char optChar) {
char errorStr[64];
if (param <= val) {
sprintf(errorStr,
"ERROR: value for -%c must be bigger than %d\n", optChar, val);
PrintError(errorStr);
}
}
/* Assert that a user's input value is between two values, inclusive. If it is
not, print a usage message and an error message and exit the program.
@param param The user's input value
@param low The lowest value the parameter can be
@param high The highest value the parameter can be
@param optChar The character used the specify the user's input value
*/
void AssertBetweenInclusive(int param, const int low, const int high,
const char optChar) {
char errorStr[64];
if (param < low || param > high) {
sprintf(errorStr,
"ERROR: value for -%c must be between %d and %d, inclusive\n",
optChar, low, high);
PrintError(errorStr);
}
}
/* Exit if a file already exists */
void AssertFileDNE(const char *filename) {
char errorStr[64];
if (access(filename, F_OK) != -1) {
sprintf(errorStr,
"ERROR: File '%s' already exists\n", filename);
PrintError(errorStr);
}
}
/* Allow the user to change simulation parameters via the command line
@param argc The number of command line arguments to parse
@param argv The array of command line arguments to parse
*/
void GetUserOptions(const int argc, char **argv) {
char c; /* Loop control variable */
/* Loop over argv, setting parameter values given */
while ((c = getopt(argc, argv, GETOPT_STRING)) != -1) {
switch(c) {
case N_ROWS_CHAR:
NRows = atoi(optarg);
AssertPositiveInteger(NRows, N_ROWS_CHAR);
break;
case N_COLS_CHAR:
NCols = atoi(optarg);
AssertPositiveInteger(NCols, N_COLS_CHAR);
break;
case BURN_PROB_CHAR:
BurnProb = atoi(optarg);
AssertInteger(BurnProb, BURN_PROB_CHAR);
AssertBetweenInclusive(BurnProb, 0, 100, BURN_PROB_CHAR);
break;
case N_MAX_BURN_STEPS_CHAR:
NMaxBurnSteps = atoi(optarg);
AssertPositiveInteger(NMaxBurnSteps, N_MAX_BURN_STEPS_CHAR);
AssertBigger(NMaxBurnSteps, 1, N_MAX_BURN_STEPS_CHAR);
break;
case N_STEPS_CHAR:
NSteps = atoi(optarg);
AssertPositiveInteger(NSteps, N_STEPS_CHAR);
break;
case RAND_SEED_CHAR:
RandSeed = atoi(optarg);
AssertPositiveInteger(RandSeed, RAND_SEED_CHAR);
break;
case OUTPUT_FILENAME_CHAR:
IsOutputtingEachStep = true;
OutputFilename = optarg;
break;
case IS_RAND_FIRST_TREE_CHAR:
IsRandFirstTree = true;
break;
case '?':
default:
PrintError("ERROR: illegal option\n");
PrintUsageAndExit();
}
}
if (IsOutputtingEachStep) {
/* Make sure the output file does not exist (DNE) */
AssertFileDNE(OutputFilename);
}
}
/* Allocate dynamic memory */
void AllocateMemory() {
Trees = (int*)malloc(NTreesPlusBounds * sizeof(int));
NewTrees = (int*)malloc( NTrees * sizeof(int));
cudaMalloc((void**)&D_Trees, NTreesPlusBounds * sizeof(int));
cudaMalloc((void**)&D_NewTrees, NTrees * sizeof(int));
cudaMalloc((void**)&D_NBurnedTrees, 1 * sizeof(int));
cudaMalloc((void**)&D_RandStates, NRows * sizeof(curandState));
}
/* Generate a random integer between [min..max]
@param min Smallest integer to generate
@param max 1 more than the biggest integer to generate
@param states Array of random states, 1 per thread
@return random integer
*/
__device__ int RandBetween(const int min, const int max, curandState *states) {
return min + (curand_uniform(&states[threadIdx.x]) * (max - min));
}
/* Determine row and column for a given CUDA thread
@param threadIdx CUDA index for the thread
@param blockIdx CUDA index for the thread's block
@param blockDim CUDA dimensions for the thread's block
@param nColsPlusBounds Number of columns in the forest plus boundary columns
@param row The resulting row for the given thread
@param col The resulting column for the given thread
*/
__device__ void getRowCol(const dim3 threadIdx, const dim3 blockIdx,
const dim3 blockDim, const int nColsPlusBounds, int *row, int *col) {
const int id = blockIdx.x * blockDim.x + threadIdx.x;
*row = id / nColsPlusBounds;
*col = id % nColsPlusBounds;
}
/* Light a random tree on fire, set all other trees to be not burning */
__global__ void InitData(const int NRows, const int NRowsPlusBounds,
const int NCols, const int NColsPlusBounds, const int NMaxBurnSteps,
const bool IsRandFirstTree, const int MiddleRow, const int MiddleCol,
int *Trees, int *NewTrees, int *NBurnedTrees,
const int randRow, const int randCol) {
int row;
int col;
int myRow;
int myCol;
/* Set all trees as having burned for 0 time steps */
getRowCol(threadIdx, blockDim, blockIdx, NColsPlusBounds, &row, &col);
if (row >= 1 && row < NRows + 1 &&
col >= 1 && col < NCols + 1) {
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 0;
}
/* Set the boundaries as burnt out */
if (row >= 0 && row < NRowsPlusBounds) {
if (col == 0) {
/* Left */
Trees[TREE_MAP(row, 0, NColsPlusBounds)] = NMaxBurnSteps;
}
/* Top/Bottom */
if ((row == 0) || (row == NRows + 1)) {
if (col >= 1 && col < NCols + 1) {
Trees[TREE_MAP(row, col, NColsPlusBounds)] = NMaxBurnSteps;
}
}
/* Right */
if (col == NCols + 1) {
Trees[TREE_MAP(row, NCols + 1, NColsPlusBounds)] = NMaxBurnSteps;
}
}
myRow = row;
myCol = col;
if (IsRandFirstTree) {
/* Light a random tree on fire */
row = randRow;
col = randCol;
}
else {
/* Light the middle tree on fire */
row = MiddleRow + 1;
col = MiddleCol + 1;
}
if (row == myRow && col == myCol) {
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 1;
atomicAdd(NBurnedTrees, 1);
}
}
/* Output tree data for the current time step */
void OutputData() {
int row;
int col;
char buf[64];
/* Write the header for the current time step */
sprintf(buf, "Time step %d\n", CurStep);
fprintf(OutputFile, "%s", buf);
for (row = 1; row < NRows + 1; row++) {
for (col = 1; col < NCols + 1; col++) {
sprintf(buf, "%*d ",
NMaxBurnStepsDigits, Trees[TREE_MAP(row, col, NColsPlusBounds)]);
fprintf(OutputFile, "%s", buf);
}
fprintf(OutputFile, "\n");
}
/* Write the newline between time steps */
fprintf(OutputFile, "\n");
}
/* Return whether a given tree has burnt out
@param row The row index of the tree
@param col The column index of the tree
@return Whether the tree in the given row and column has burnt out
*/
__device__ bool IsBurntOut(const int row, const int col,
const int NColsPlusBounds, const int NMaxBurnSteps, const int *Trees) {
return Trees[TREE_MAP(row, col, NColsPlusBounds)] >= NMaxBurnSteps;
}
/* Return whether a given tree is on fire
@param row The row index of the tree
@param col The column index of the tree
@return Whether the tree in the given row and column is on fire
*/
__device__ bool IsOnFire(const int row, const int col,
const int NColsPlusBounds, const int NMaxBurnSteps, const int *Trees) {
return Trees[TREE_MAP(row, col, NColsPlusBounds)] > 0 &&
!IsBurntOut(row, col, NColsPlusBounds, NMaxBurnSteps, Trees);
}
/* For trees already burning, increment the number of time steps they have
burned
*/
__global__ void ContinueBurning(const int NRows, const int NCols,
const int NColsPlusBounds, const int NMaxBurnSteps, const int *Trees,
int *NewTrees) {
int row;
int col;
getRowCol(threadIdx, blockDim, blockIdx, NColsPlusBounds, &row, &col);
if (row >= 1 && row < NRows + 1 &&
col >= 1 && col < NCols + 1) {
if (IsOnFire(row, col, NColsPlusBounds, NMaxBurnSteps, Trees)) {
NewTrees[NEW_TREE_MAP(row, col, NCols)] =
Trees[ TREE_MAP(row, col, NColsPlusBounds)] + 1;
}
}
}
/* Find trees that are not on fire yet and try to catch them on fire from
burning neighbor trees
*/
__global__ void BurnNew(const int NRows, const int NCols,
const int NColsPlusBounds, const int BurnProb, const int NMaxBurnSteps,
const int *Trees, int *NewTrees, int *NBurnedTrees,
curandState *randStates) {
int row;
int col;
getRowCol(threadIdx, blockDim, blockIdx, NColsPlusBounds, &row, &col);
if (row >= 1 && row < NRows + 1 &&
col >= 1 && col < NCols + 1) {
if (!IsOnFire(row, col, NColsPlusBounds, NMaxBurnSteps, Trees) &&
!IsBurntOut(row, col, NColsPlusBounds, NMaxBurnSteps, Trees)) {
/* Check neighbors */
/* Top */
if ((IsOnFire(row-1, col, NColsPlusBounds, NMaxBurnSteps, Trees) ||
/* Left */
IsOnFire(row, col-1, NColsPlusBounds, NMaxBurnSteps, Trees) ||
/* Bottom */
IsOnFire(row+1, col, NColsPlusBounds, NMaxBurnSteps, Trees) ||
/* Right */
IsOnFire(row, col+1, NColsPlusBounds, NMaxBurnSteps, Trees)) &&
/* Apply random chance */
(RandBetween(0, 100, randStates) < BurnProb)) {
/* Catch the tree on fire */
NewTrees[NEW_TREE_MAP(row, col, NCols)] = 1;
/* One thread at a time increments the burned tree count */
atomicAdd(NBurnedTrees, 1);
}
}
}
}
/* Copy new tree data into old tree data */
__global__ void AdvanceTime(const int NRows, const int NCols,
const int NColsPlusBounds, int *Trees, int *NewTrees) {
int row;
int col;
getRowCol(threadIdx, blockDim, blockIdx, NColsPlusBounds, &row, &col);
if (row >= 1 && row < NRows + 1 &&
col >= 1 && col < NCols + 1) {
Trees[ TREE_MAP(row, col, NColsPlusBounds)] =
NewTrees[NEW_TREE_MAP(row, col, NCols)];
}
}
/* Free allocated memory */
void FreeMemory() {
cudaFree(D_RandStates);
cudaFree(D_NBurnedTrees);
cudaFree(D_NewTrees);
cudaFree(D_Trees);
free(NewTrees);
free(Trees);
}
/* @param argc The number of command line arguments
@param argv String of command line arguments
*/
int main(int argc, char **argv) {
/* Set the program executable name */
strcpy(ExeName, argv[0]);
/* Allow the user to change simulation parameters via the command line */
GetUserOptions(argc, argv);
if (!AreParamsValid) {
/* Model parameters are not valid; exit early */
PrintUsageAndExit();
}
if (IsOutputtingEachStep) {
/* Open the output file */
OutputFile = fopen(OutputFilename, "w");
}
/* Do some calculations before splitting up the rows */
NTrees = NRows * NCols;
NRowsPlusBounds = NRows + 2;
NColsPlusBounds = NCols + 2;
NTreesPlusBounds = NRowsPlusBounds * NColsPlusBounds;
MiddleRow = NRows / 2;
MiddleCol = NCols / 2;
/* Calculate the number of CUDA blocks needed for each tree to be assigned to
a thread, including the boundaries
*/
NBlocks = ceil((double)NTreesPlusBounds / THREADS_PER_BLOCK);
/* Allocate dynamic memory for the 1D tree arrays */
AllocateMemory();
/* Seed the random number generator */
SetRandStates<<<NBlocks, THREADS_PER_BLOCK>>>(RandSeed,
D_RandStates);
srandom(RandSeed);
/* Initialize number of burned trees */
NBurnedTrees = 0;
cudaMemcpy(D_NBurnedTrees, &NBurnedTrees, 1 * sizeof(int),
cudaMemcpyHostToDevice);
/* Light a random tree on fire, set all other trees to be not burning */
InitData<<<NBlocks, THREADS_PER_BLOCK>>>(NRows, NRowsPlusBounds,
NCols, NColsPlusBounds, NMaxBurnSteps, IsRandFirstTree, MiddleRow,
MiddleCol, D_Trees, D_NewTrees, D_NBurnedTrees,
random() % NRowsPlusBounds, random() % NColsPlusBounds);
/* Start the simulation looping for the specified number of time steps */
for (CurStep = 0; CurStep < NSteps; CurStep++) {
if (IsOutputtingEachStep) {
/* Output tree data for the current time step */
cudaMemcpy(Trees, D_Trees, NTreesPlusBounds * sizeof(int),
cudaMemcpyDeviceToHost);
OutputData();
}
/* For trees already burning, increment the number of time steps they have
burned */
ContinueBurning<<<NBlocks, THREADS_PER_BLOCK>>>(NRows, NCols,
NColsPlusBounds, NMaxBurnSteps, D_Trees, D_NewTrees);
/* Find trees that are not on fire yet and try to catch them on fire from
burning neighbor trees */
BurnNew<<<NBlocks, THREADS_PER_BLOCK>>>(NRows, NCols,
NColsPlusBounds, BurnProb, NMaxBurnSteps, D_Trees, D_NewTrees,
D_NBurnedTrees, D_RandStates);
/* Copy new tree data into old tree data */
AdvanceTime<<<NBlocks, THREADS_PER_BLOCK>>>(NRows, NCols,
NColsPlusBounds, D_Trees, D_NewTrees);
}
/* Print the total percentage of trees burned */
cudaMemcpy(&NBurnedTrees, D_NBurnedTrees, 1 * sizeof(int),
cudaMemcpyDeviceToHost);
printf("%.2f%% of the trees were burned\n",
(100.0 * NBurnedTrees) / NTrees);
if (IsOutputtingEachStep) {
/* Close the output file */
fclose(OutputFile);
}
/* Free allocated memory */
FreeMemory();
return 0;
}
45 differences: 140 lines, 237 inline differences in 113 changed lines
Added(24,117)
Deleted(3,47)
Changed(113)
Changed in changed(73)
Ignored
Generated on October 21, 2015, 2:31 PM by ExamDiff Pro 8.0.0.3.