3
\$\begingroup\$

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(){
    // 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?

\$\endgroup\$

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.