CUDA ripple effect on the GPU, with measured frames and BLOCK explanation comment for help with indexing.

ripple.cu

#include <chrono>
#include "../common/book.h"
#include "../common/cpu_anim.h"

#define DIM 1024

using namespace std;
using namespace std::chrono;

// 2D array of blocks and dimensions
// dim3 grid(3, 3);
// dim3 threads(3, 3);
// kernel<<<grid, threads>>>
// 3 x 3 blocks, each with 3 x 3 threads
// blocks
// [[0, 0], [0, 1], [0, 2]]
// [[1, 0], [1, 1], [1, 2]]
// [[2, 0], [2, 1], [2, 2]]
//
// threads in each block
// [[0, 0], [0, 1], [0, 2]]
// [[1, 0], [1, 1], [1, 2]]
// [[2, 0], [2, 1], [2, 2]]
//
// representation
// ╔═══════════════╦═══════════════╦═══════════════╗
// ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║
// ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║
// ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║
// ╠═══════════════╬═══════════════╬═══════════════╣
// ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║
// ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║
// ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║
// ╠═══════════════╬═══════════════╬═══════════════╣
// ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║ ┌───┬───┬───┐ ║
// ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║ │ 0 │ 1 │ 2 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║ │ 3 │ 4 │ 5 │ ║
// ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║ ├───┼───┼───┤ ║
// ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║ │ 6 │ 7 │ 8 │ ║
// ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║ └───┴───┴───┘ ║
// ╚═══════════════╩═══════════════╩═══════════════╝
//
__global__ void kernel(unsigned char* ptr, int ticks)
{
    // map from threadIdx/blockIdx to pixel position
    unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int offset = x + y * blockDim.x * gridDim.x;

    // now calculate the value at that position
    float fx = (float)x - DIM / 2;
    float fy = (float)y - DIM / 2;
    float d = sqrtf(fx * fx + fy * fy);

    auto gray = (unsigned char)(
            128.0f + 127.0f
            * cos(d / 10.0f - (float)ticks / 7.0f)
            / (d / 10.0f + 1.0f));

    ptr[offset * 4 + 0] = gray;
    ptr[offset * 4 + 1] = gray;
    ptr[offset * 4 + 2] = gray;
    ptr[offset * 4 + 3] = 255;
}

struct DataBlock {
    unsigned char* dev_bitmap;
    CPUAnimBitmap* bitmap;
};

/// cleanup memory on the GPU
void cleanup(DataBlock* db) {
    cudaFree(db->dev_bitmap);
}


void generate_frame(DataBlock* db, int ticks) {
    auto ts_start = high_resolution_clock::now();

    dim3 grid(DIM / 16, DIM / 16);
    dim3 threads(16, 16);
    kernel<<<grid, threads>>>(db->dev_bitmap, ticks);
    cudaMemcpy(db->bitmap->get_ptr(), db->dev_bitmap, db->bitmap->image_size(), cudaMemcpyDeviceToHost);

    auto end = duration<double, milli>(high_resolution_clock::now() - ts_start).count();
    printf("frame took: %f ms\n", end);
}

int main()
{
    DataBlock data;
    CPUAnimBitmap bitmap(DIM, DIM, &data);
    data.bitmap = &bitmap;

    HANDLE_ERROR(cudaMalloc((void**)&data.dev_bitmap, bitmap.image_size()));

    bitmap.anim_and_exit((void (*)(void*, int))generate_frame, (void (*)(void*))cleanup);

    return EXIT_SUCCESS;
}

frame measuring

the ripple effect:

b8fcaaac9ba3202e6ad97ae35e9d1d5b.mp4