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.cuvoid 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(){    // Loads initial input    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(){    // Loads initial input    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?
CUDA program that emulates a kind of cellular automata slower on the GPU than on the CPU
performance;beginner;c;cuda;cellular automata