OpenCL mode is working!
This commit is contained in:
parent
fddb092a56
commit
a4d0a5dca8
45
gol.cl
45
gol.cl
@ -1,37 +1,30 @@
|
|||||||
void kernel generation(global const int *previous, global int *new, global const int *size) {
|
void kernel generation(global const bool *previous, global bool *new, global const int *size) {
|
||||||
int ID, Nthreads, n, ratio, start, stop, x, y, left, right, up, down, neighbors, i;
|
int ID, Nthreads, n, ratio, start, stop, x, y, left, right, up, down, neighbors, i;
|
||||||
|
|
||||||
ID = get_global_id(0);
|
ID = get_global_id(0);
|
||||||
Nthreads = get_global_size(0);
|
|
||||||
n = size[0] * size[1];
|
|
||||||
|
|
||||||
ratio = (n / Nthreads); // number of elements for each thread
|
x = ID % size[0];
|
||||||
start = ratio * ID;
|
y = ID / size[0];
|
||||||
stop = ratio * (ID + 1);
|
|
||||||
|
|
||||||
for (i = start; i < stop; i++)
|
left = (x - 1 + size[0]) % size[0];
|
||||||
x = i % size[1];
|
right = (x + 1) % size[0];
|
||||||
y = i / size[0];
|
|
||||||
|
|
||||||
left = x - 1;
|
up = (y - 1 + size[1]) % size[1];
|
||||||
right = (x + 1) % size[0];
|
down = (y + 1) % size[1];
|
||||||
|
|
||||||
up = (y - 1 + size[1]) % size[1];
|
// Get the number of neighbors
|
||||||
down = (y + 1) % size[1];
|
neighbors =
|
||||||
|
previous[size[0] * up + left]
|
||||||
|
+ previous[size[0] * up + x]
|
||||||
|
+ previous[size[0] * up + right]
|
||||||
|
|
||||||
// Get the number of neighbors
|
+ previous[size[0] * y + left]
|
||||||
neighbors =
|
+ previous[size[0] * y + right]
|
||||||
previous[size[0] * up + left]
|
|
||||||
+ previous[size[0] * up + x]
|
|
||||||
+ previous[size[0] * up + right]
|
|
||||||
|
|
||||||
+ previous[size[0] * y + left]
|
+ previous[size[0] * down + left]
|
||||||
+ previous[size[0] * y + right]
|
+ previous[size[0] * down + x]
|
||||||
|
+ previous[size[0] * down + right];
|
||||||
|
|
||||||
+ previous[size[0] * down + left]
|
// Update cell
|
||||||
+ previous[size[0] * down + x]
|
new[size[0] * y + x] = (neighbors == 3) + previous[size[0] * y + x] * (neighbors == 2);
|
||||||
+ previous[size[0] * down + right];
|
|
||||||
|
|
||||||
// Update cell
|
|
||||||
new[size[0] * y + x] = (neighbors == 3) + previous[size[0] * y + x] * (neighbors == 2);
|
|
||||||
}
|
}
|
36
main.cpp
36
main.cpp
@ -212,6 +212,8 @@ void main_opencl(std::string infile, std::string outfile, int num_generations) {
|
|||||||
std::ostringstream ss;
|
std::ostringstream ss;
|
||||||
ss << file.rdbuf();
|
ss << file.rdbuf();
|
||||||
kernel_code = ss.str();
|
kernel_code = ss.str();
|
||||||
|
} else {
|
||||||
|
std::cout << "Error: Couldn't read Kernel source!" << std::endl;
|
||||||
}
|
}
|
||||||
sources.push_back({kernel_code.c_str(), kernel_code.length()});
|
sources.push_back({kernel_code.c_str(), kernel_code.length()});
|
||||||
|
|
||||||
@ -236,6 +238,7 @@ void main_opencl(std::string infile, std::string outfile, int num_generations) {
|
|||||||
int size_y = std::stoi(y_str);
|
int size_y = std::stoi(y_str);
|
||||||
|
|
||||||
bool *world = new bool[size_x * size_y];
|
bool *world = new bool[size_x * size_y];
|
||||||
|
bool *result = new bool[size_x * size_y];
|
||||||
|
|
||||||
// Set the data
|
// Set the data
|
||||||
for (int y = 0; y < size_y; y++) {
|
for (int y = 0; y < size_y; y++) {
|
||||||
@ -266,19 +269,31 @@ void main_opencl(std::string infile, std::string outfile, int num_generations) {
|
|||||||
|
|
||||||
// push write commands to queue
|
// push write commands to queue
|
||||||
queue.enqueueWriteBuffer(buffer_previous, CL_TRUE, 0, sizeof(bool) * n, world);
|
queue.enqueueWriteBuffer(buffer_previous, CL_TRUE, 0, sizeof(bool) * n, world);
|
||||||
queue.enqueueWriteBuffer(buffer_new, CL_TRUE, 0, sizeof(bool) * n, world); // TODO: pass empty array
|
queue.enqueueWriteBuffer(buffer_new, CL_TRUE, 0, sizeof(bool) * n, result);
|
||||||
queue.enqueueWriteBuffer(buffer_size, CL_TRUE, 0, sizeof(int) * 2, size); // TODO: pass empty array
|
queue.enqueueWriteBuffer(buffer_size, CL_TRUE, 0, sizeof(int) * 2, size);
|
||||||
|
|
||||||
// RUN ZE KERNEL
|
// RUN ZE KERNEL
|
||||||
cl::Kernel gol_kernel(program, "gol");
|
cl::Kernel gol_kernel(program, "generation");
|
||||||
gol_kernel.setArg(0, buffer_previous);
|
|
||||||
gol_kernel.setArg(1, buffer_new);
|
for (int i = 0; i < num_generations; i++) {
|
||||||
gol_kernel.setArg(2, buffer_size);
|
gol_kernel.setArg(0, buffer_previous);
|
||||||
queue.enqueueNDRangeKernel(gol_kernel, cl::NullRange, cl::NDRange(10), cl::NullRange);
|
gol_kernel.setArg(1, buffer_new);
|
||||||
|
gol_kernel.setArg(2, buffer_size);
|
||||||
|
|
||||||
|
queue.enqueueNDRangeKernel(gol_kernel, cl::NullRange, cl::NDRange(n), cl::NullRange);
|
||||||
|
queue.finish();
|
||||||
|
|
||||||
|
std::swap(buffer_previous, buffer_new);
|
||||||
|
}
|
||||||
queue.finish();
|
queue.finish();
|
||||||
|
|
||||||
// read result from GPU to here
|
// Since we swap after every generation, we need to proceed differently depending on
|
||||||
queue.enqueueReadBuffer(buffer_new, CL_TRUE, 0, sizeof(bool) * n, world); // TODO: pass different empty?
|
// whether we're in swapped mode or not at the moment
|
||||||
|
if (num_generations % 2 == 0) {
|
||||||
|
queue.enqueueReadBuffer(buffer_previous, CL_TRUE, 0, sizeof(bool) * n, result);
|
||||||
|
} else {
|
||||||
|
queue.enqueueReadBuffer(buffer_new, CL_TRUE, 0, sizeof(bool) * n, result);
|
||||||
|
}
|
||||||
|
|
||||||
// Write the result
|
// Write the result
|
||||||
std::ofstream result_file;
|
std::ofstream result_file;
|
||||||
@ -288,11 +303,10 @@ void main_opencl(std::string infile, std::string outfile, int num_generations) {
|
|||||||
|
|
||||||
for (int y = 0; y < size_y; y++) {
|
for (int y = 0; y < size_y; y++) {
|
||||||
std::string line;
|
std::string line;
|
||||||
getline(world_file, line);
|
|
||||||
|
|
||||||
for (int x = 0; x < size_x; x++) {
|
for (int x = 0; x < size_x; x++) {
|
||||||
// Convert 1 and 0 to 'x' and '.' again
|
// Convert 1 and 0 to 'x' and '.' again
|
||||||
line += world[y * size_x + x] ? 'x' : '.';
|
line += result[y * size_x + x] ? 'x' : '.';
|
||||||
}
|
}
|
||||||
|
|
||||||
result_file << line << '\n';
|
result_file << line << '\n';
|
||||||
|
Loading…
x
Reference in New Issue
Block a user