# CUDA program that emulates a kind of cellular automata slower on the GPU than on the CPU

I'm a very new to CUDA, so I'm still trying to understand how to make the best use of the GPU. I've ported a C algorithm to it. This algorithm works by loading the memory with an initial input, then calling a rewrite function, which rewrites 128 bits of memory, consecutively, until we arrive at an answer. Since rewrite is completely local, it is easy to parallelize, because you can call it at the same time in different memory positions.

This is how I tested it for the CPU:

#include <stdio.h>
#include "test_inputs.cu"
#include "rewriter.cu"

void process(int *mem, int len, int j){
for (int i=j; i<len-2; i+=3)
rewrite(mem+i*4, mem+i*4+4, mem+i*4+8);
};

int main(){
for (int i=0; i < memory_ints; ++i)
memory[i] = i < program_nodes*4 ? program[i] : 0;

// Prints it for visualization
print(memory, 40);

// Calls rewrite sequentially until the computation is done
for (int k=0; k<clocks*3; ++k)
for (int j=0; j<3; ++j) // for alternating alignments
process(memory, memory_nodes, j);

// Print the result (should have only 4 numbers)
print(memory, 40);
}


This is how I tested it for the GPU:

#include <stdio.h>
#include "test_inputs.cu"
#include "rewriter.cu"

__global__ void process(int *mem, int len, int j){
int i = (blockIdx.x*blockDim.x + threadIdx.x)*3 + j;
if (i >= len-2) return;
rewrite(mem+i*4, mem+i*4+4, mem+i*4+8);
};

int main(){

for (int i=0; i < memory_ints; ++i)
memory[i] = i < program_nodes*4 ? program[i] : 0;

// Prints it for visualization
print(memory, 40);

// Puts it on the GPU
int *device_memory;
cudaMalloc((void**)&device_memory, sizeof memory);
cudaMemcpy(device_memory, memory, memory_size, cudaMemcpyHostToDevice);

// Calls rewrite in parallel until the computation is done
dim3 block_size(16, 1);
dim3 grid_size(memory_nodes / block_size.x / 3, 1);
for (int k=0; k<clocks; ++k)
for (int j=0; j<3; ++j) // for alternating alignments
process<<<block_size, grid_size>>>(device_memory, memory_nodes, j);

// Gets data back from GPU
cudaMemcpy(memory, device_memory, memory_size, cudaMemcpyDeviceToHost);
cudaFree(device_memory);

// Print the result (should have only 4 numbers)
print(memory, 40);
}


Unfortunately, the parallel version is about 2 times slower than the sequential version. rewriter.cu and test_inputs.cu are on this repository. I'm aware this must be ridden of newbie mistakes, but that's why I'm asking here. What are the worst mistakes I made and how can I improve it?