GPU Mandelbrot Set with CUDA C++

Build a parallel Mandelbrot fractal generator using CUDA on your GPU.

Introduction

30 Second Summary

Every time you play a video game, train an AI model, or watch a movie with CGI effects, your computer's GPU is crunching millions of calculations at the same time. That parallel processing power is what makes modern computing fast.

In this project, you will write a program in CUDA C++ that generates a fractal image by running thousands of calculations simultaneously across your GPU's cores. You will build both a slow CPU version and a fast GPU version, then compare their speeds to see parallelism in action.

What You'll Build

A command-line CUDA C++ program that computes the Mandelbrot set fractal and outputs it as a PPM image file, with both sequential CPU and parallel GPU implementations timed for comparison.

By the end of this project, you'll have:

  • A working CUDA development environment with Visual Studio 2022 and the NVIDIA CUDA Toolkit installed on your Windows machine.
  • Two versions of a Mandelbrot generator: a sequential CPU version and a parallel GPU version that demonstrates real speedup across your GPU's 1280 CUDA cores.
  • A solid understanding of CUDA's thread/block/grid hierarchy and host-device memory flow that transfers to any GPU programming task.
  • Secret Mission: Add smooth color gradient mapping and implement a zoom function that renders magnified regions of the fractal.

Are there any prerequisites?

You need basic programming experience in any language. No prior C++ or CUDA experience is required. You also need a Windows computer with an NVIDIA GPU (this project uses a GTX 1650 Super, but any NVIDIA GPU with CUDA support will work).

Install Visual Studio 2022 and CUDA Toolkit

Before you can write GPU code, you need two things: a C++ compiler and the CUDA Toolkit.

Visual Studio 2022 provides the MSVC compiler that CUDA's compiler (nvcc) depends on. The CUDA Toolkit provides nvcc itself along with the runtime libraries that let your code talk to the GPU.

In this step, get ready to:
  • Install Visual Studio 2022 with the C++ compiler toolchain.
  • Install the NVIDIA CUDA Toolkit 13.3.
  • Verify both tools are working and create your project folder.
Install Visual Studio 2022 Community Edition

CUDA's compiler (nvcc) does not compile code on its own. It delegates the host-side C++ compilation to MSVC, the compiler that ships with Visual Studio.

  • Open https://visualstudio.microsoft.com/vs/community/ in your browser.
  • Click Free download to download the installer.
  • Run the downloaded installer and wait for it to load the workload selection screen.
  • Check the box for Desktop development with C++.
  • Click Install and wait for the download and installation to complete. This may take 10-15 minutes depending on your internet speed.

Why Visual Studio if we're coding in VS Code?

You are not using Visual Studio as your code editor. You only need it for the MSVC compiler it bundles.

Think of Visual Studio as the supplier of a tool (the compiler) and VS Code as the workshop where you write your code.

Install the CUDA Toolkit

The CUDA Toolkit gives you nvcc (the CUDA compiler), GPU runtime libraries, and header files for writing kernel code.

  • Open https://developer.nvidia.com/cuda-downloads in your browser.
  • Select Windows as the operating system.
  • Select x86_64 as the architecture.
  • Select your Windows version, then choose the exe (local) installer type.
  • Click Download and run the installer once it finishes.
  • Accept the default installation options and wait for installation to complete.

What about my GPU driver?

Starting with CUDA 13, the NVIDIA driver is no longer bundled with the CUDA Toolkit installer. Since you have a GTX 1650 Super (a gaming GPU), your driver is almost certainly up to date from Game Ready driver updates.

You will confirm your driver is working in the next substep.

Verify your installations

Both tools are installed. Now you need to open a special command prompt that has the MSVC compiler paths loaded, and confirm CUDA is ready.

  • Press the Windows key to open your search bar.
  • Type x64 Native Tools Command Prompt for VS 2022 and press Enter to open it.

Why this specific command prompt?

This command prompt runs a script that sets up environment variables for the 64-bit MSVC compiler. Without it, nvcc cannot find the C++ compiler it needs.

  • Verify the CUDA compiler is installed by running this command:
nvcc -V

✔️ I see CUDA compilation tools and a version number

You should see output mentioning Cuda compilation tools and a release version like 13.3. Your CUDA compiler is ready.

ⓧ I see a different version

If you see an older CUDA version, you may have a previous toolkit installed. Uninstall the older version from Add or remove programs in Windows Settings, then re-run the CUDA 13.3 installer.

ⓧ Command not recognized

Make sure you are using the x64 Native Tools Command Prompt for VS 2022 and not a regular Command Prompt or PowerShell window.

If it still fails, the CUDA Toolkit installer may not have added itself to your PATH. Try closing and reopening the x64 Native Tools Command Prompt. If that does not help, re-run the CUDA Toolkit installer.

  • Verify your GPU is detected by running this command:
nvidia-smi

✔️ I see a table with my GPU name

You should see a table showing your GPU name (GeForce GTX 1650 SUPER), driver version, and memory usage. This confirms your GPU driver is working and your card is ready for CUDA programs.

ⓧ Command not recognized or no GPU listed

Your NVIDIA driver may be missing or outdated. Download the latest Game Ready driver from the NVIDIA website for your GTX 1650 Super, install it, restart your computer, then try nvidia-smi again.

Still having trouble?

Make sure your NVIDIA GPU driver is updated separately from the CUDA Toolkit. Since CUDA 13, the driver is no longer included in the toolkit installer.

help me troubleshoot why nvcc or nvidia-smi is not working on my Windows machine with a GTX 1650 Super

Create your project folder

With your toolchain confirmed, you need a dedicated folder for the CUDA files you will write in the next steps.

  • In the x64 Native Tools Command Prompt, create and navigate to your project folder by running these commands:
mkdir C:\cuda-mandelbrot
cd C:\cuda-mandelbrot

What does this do?

mkdir creates a new folder. cd moves your command prompt into that folder so all future commands run from there.

  • Confirm the folder was created by running:
cd

You should see C:\cuda-mandelbrot printed as your current directory.

  • Switch to VS Code and open your new project folder.
  • Click File in the top menu bar.
  • Click Open Folder.
  • Navigate to C:\cuda-mandelbrot, select it, and click Select Folder.

Your VS Code Explorer sidebar should now show the empty cuda-mandelbrot folder.

Don't see the folder in the file picker?

Navigate to your C: drive root in the file picker dialog. The cuda-mandelbrot folder should be listed there.

help me find my cuda-mandelbrot folder in VS Code's Open Folder dialog on Windows

Your development environment is ready. You have a C++ compiler, the CUDA toolkit, and a confirmed GPU. Next up, you will write your first C++ program that computes the Mandelbrot set on the CPU.

Write the CPU Mandelbrot Generator

Your development environment is ready. Now it is time to write your first program.

The Mandelbrot set is a fractal shape defined by a simple math formula. For each pixel in the image, you repeat a calculation (squaring and adding complex numbers) until the value "escapes" past a threshold or you hit a maximum iteration count. Every pixel is independent, which makes it a perfect candidate for parallel GPU computation later. In this step, you will write the sequential CPU version first so you have a correct reference and a baseline time to compare against.

In this step, get ready to:
  • Write the Mandelbrot escape-time algorithm in C++.
  • Output the result as a PPM image file.
  • Measure the CPU execution time in milliseconds.
Create the file and write the Mandelbrot computation

The escape-time algorithm works by mapping each pixel to a point on the complex number plane. It then repeatedly squares that number and adds the original point. If the value grows larger than 2 (we check the squared magnitude against 4 to avoid a square root), the point has "escaped" and is not in the Mandelbrot set. The number of iterations before escape determines the pixel's color.

  • In VS Code, create a new file called mandelbrot_cpu.cu inside your C:\cuda-mandelbrot folder.

Why .cu and not .cpp?

The .cu file extension tells nvcc (the CUDA compiler) that this file may contain GPU device code. Even though this first file is CPU-only, using .cu lets you compile with nvcc right away without extra configuration.

  • Add the includes, image constants, and the computeMandelbrot function by pasting the following code into mandelbrot_cpu.cu:
#include <stdio.h>
#include <stdlib.h>
#include <chrono>

const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;

void computeMandelbrot(unsigned char* image, int width, int height, int maxIter) {
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            // Map pixel to complex plane
            float cx = (float)x / width * 3.5f - 2.5f;
            float cy = (float)y / height * 2.0f - 1.0f;

            float zx = 0.0f;
            float zy = 0.0f;
            int iter = 0;

            while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
                float tmp = zx * zx - zy * zy + cx;
                zy = 2.0f * zx * zy + cy;
                zx = tmp;
                iter++;
            }

            int idx = (y * width + x) * 3;
            unsigned char color = (unsigned char)(iter % 256);
            image[idx] = color;
            image[idx + 1] = color;
            image[idx + 2] = color;
        }
    }
}

What does this code do?

The constants define a 1920x1080 image with up to 1000 iterations per pixel. More iterations produce finer detail at the fractal boundary.

Inside computeMandelbrot, the nested for loops visit every pixel one at a time (sequentially). For each pixel, cx and cy map the pixel coordinates into the complex plane (real range -2.5 to 1.0, imaginary range -1.0 to 1.0). The while loop runs the escape-time formula. The final iteration count is mapped to a grayscale color using modulo 256.

  • You will not see a change yet. The next chunk adds the output and timing code that completes the file.
Add image output and timing

The computation alone does not produce anything visible. You need a function to write the pixel data to an image file, and a main function that allocates memory, runs the computation with a timer, and saves the result.

  • Add the following code below the computeMandelbrot function in mandelbrot_cpu.cu:
void writePPM(const char* filename, unsigned char* image, int width, int height) {
    FILE* fp = fopen(filename, "w");
    fprintf(fp, "P3\n%d %d\n255\n", width, height);
    for (int i = 0; i < width * height * 3; i += 3) {
        fprintf(fp, "%d %d %d\n", image[i], image[i + 1], image[i + 2]);
    }
    fclose(fp);
}

int main() {
    int imageSize = WIDTH * HEIGHT * 3;
    unsigned char* image = new unsigned char[imageSize];

    printf("Computing Mandelbrot set on CPU (%dx%d, max %d iterations)...\n", WIDTH, HEIGHT, MAX_ITER);

    auto start = std::chrono::high_resolution_clock::now();
    computeMandelbrot(image, WIDTH, HEIGHT, MAX_ITER);
    auto end = std::chrono::high_resolution_clock::now();

    float cpuTime = std::chrono::duration<float, std::milli>(end - start).count();
    printf("CPU Time: %.2f ms\n", cpuTime);

    writePPM("mandelbrot.ppm", image, WIDTH, HEIGHT);
    printf("Image saved to mandelbrot.ppm\n");

    delete[] image;
    return 0;
}

What does this code do?

writePPM writes the pixel data in PPM format. PPM is a plain-text image format where the header (P3, width, height, max color value) is followed by RGB triplets for every pixel. It requires zero external libraries.

In main, new unsigned char[imageSize] allocates the pixel array on the heap (the stack would overflow at 6 million bytes). The <chrono> timer wraps the computation call, and delete[] frees the heap memory when the program finishes.

  • Save mandelbrot_cpu.cu.

✔️ Awesome, I've got everything!

Great. Double check you have saved your file before moving on to compile.

ⓧ I'd like to double check the full code

#include <stdio.h>
#include <stdlib.h>
#include <chrono>

const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;

void computeMandelbrot(unsigned char* image, int width, int height, int maxIter) {
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            // Map pixel to complex plane
            float cx = (float)x / width * 3.5f - 2.5f;
            float cy = (float)y / height * 2.0f - 1.0f;

            float zx = 0.0f;
            float zy = 0.0f;
            int iter = 0;

            while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
                float tmp = zx * zx - zy * zy + cx;
                zy = 2.0f * zx * zy + cy;
                zx = tmp;
                iter++;
            }

            int idx = (y * width + x) * 3;
            unsigned char color = (unsigned char)(iter % 256);
            image[idx] = color;
            image[idx + 1] = color;
            image[idx + 2] = color;
        }
    }
}

void writePPM(const char* filename, unsigned char* image, int width, int height) {
    FILE* fp = fopen(filename, "w");
    fprintf(fp, "P3\n%d %d\n255\n", width, height);
    for (int i = 0; i < width * height * 3; i += 3) {
        fprintf(fp, "%d %d %d\n", image[i], image[i + 1], image[i + 2]);
    }
    fclose(fp);
}

int main() {
    int imageSize = WIDTH * HEIGHT * 3;
    unsigned char* image = new unsigned char[imageSize];

    printf("Computing Mandelbrot set on CPU (%dx%d, max %d iterations)...\n", WIDTH, HEIGHT, MAX_ITER);

    auto start = std::chrono::high_resolution_clock::now();
    computeMandelbrot(image, WIDTH, HEIGHT, MAX_ITER);
    auto end = std::chrono::high_resolution_clock::now();

    float cpuTime = std::chrono::duration<float, std::milli>(end - start).count();
    printf("CPU Time: %.2f ms\n", cpuTime);

    writePPM("mandelbrot.ppm", image, WIDTH, HEIGHT);
    printf("Image saved to mandelbrot.ppm\n");

    delete[] image;
    return 0;
}
Compile and run the program

You will compile using nvcc, the CUDA compiler, from the x64 Native Tools Command Prompt you opened in the previous step.

  • Switch back to the x64 Native Tools Command Prompt for VS 2022 from earlier.
  • Navigate to your project folder by running:
cd C:\cuda-mandelbrot
  • Compile mandelbrot_cpu.cu by running:
nvcc mandelbrot_cpu.cu -o mandelbrot_cpu.exe
  • Run the compiled program:
mandelbrot_cpu.exe
  • You should see output similar to this:

What should I see?

The console prints three lines. The first announces the computation parameters. The second shows the CPU time in milliseconds (expect a value in the thousands, depending on your processor). The third confirms the image file was saved.

Computing Mandelbrot set on CPU (1920x1080, max 1000 iterations)...
CPU Time: 3500.42 ms
Image saved to mandelbrot.ppm
  • Open the generated mandelbrot.ppm file. The easiest method is to drag the file from File Explorer into your web browser window. You can also open it with GIMP or IrfanView if you have either installed.
  • You should see the classic Mandelbrot shape: a black cardioid bulb on the left with a circular bud attached, surrounded by grayscale fringing.

Don't see the fractal image?

  • If the browser shows garbled text instead of an image, try GIMP (free at gimp.org) which handles PPM natively.
  • If you see a compile error mentioning chrono, make sure you are compiling from the x64 Native Tools Command Prompt (not plain Command Prompt or PowerShell). The MSVC headers must be on the include path.
  • If nvcc reports it cannot find cl.exe, you are in the wrong terminal. Close it and open x64 Native Tools Command Prompt for VS 2022 from the Start menu.

Help me troubleshoot my Mandelbrot CPU program compile or runtime error.

Note the CPU time value printed in the console. You will compare this against the GPU version in the next step to measure the speedup from parallel computation.

Your CPU baseline is established. Next up, you will write a CUDA kernel that computes every pixel in parallel across your GPU's 1280 cores.

Write the CUDA GPU Kernel

Your CPU version computes pixels one at a time, looping through 2 million pixels sequentially. A CUDA kernel flips this model: instead of one thread doing all the work, you launch millions of threads that each compute one pixel simultaneously.

In this step, you will write a __global__ kernel function, allocate memory on the GPU, launch the kernel across a 2D grid of threads, and copy the results back to the CPU. You will also time both versions and see the speedup your GTX 1650 Super delivers.

In this step, get ready to:
  • Write a CUDA kernel that computes Mandelbrot pixels in parallel.
  • Set up host-device memory transfers with cudaMalloc and cudaMemcpy.
  • Measure GPU performance with CUDA events and compare against the CPU.
Create the file and write the kernel function

The kernel is the function that runs on the GPU. Each GPU thread executes the same kernel code, but on a different pixel.

  • In VS Code, create a new file called mandelbrot_gpu.cu inside your C:\cuda-mandelbrot folder.
  • Add the file header with includes, constants, and an error-checking macro by pasting this code:
#include <stdio.h>
#include <stdlib.h>
#include <chrono>

const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;

// Error-checking macro for CUDA API calls
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while (0)

What does this code do?

The constants are the same as your CPU version: a 1920x1080 image with up to 1000 iterations per pixel.

The CUDA_CHECK macro wraps every CUDA API call. If any call fails, it prints the exact file, line number, and error message, then exits. Without this, CUDA errors fail silently and produce corrupted output with no explanation.

  • Below the macro, add the CPU comparison function (same escape-time logic you wrote before, renamed for clarity):
// CPU version for comparison
void computeMandelbrotCPU(unsigned char* image, int width, int height, int maxIter) {
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float cx = (float)x / width * 3.5f - 2.5f;
            float cy = (float)y / height * 2.0f - 1.0f;

            float zx = 0.0f;
            float zy = 0.0f;
            int iter = 0;

            while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
                float tmp = zx * zx - zy * zy + cx;
                zy = 2.0f * zx * zy + cy;
                zx = tmp;
                iter++;
            }

            int idx = (y * width + x) * 3;
            unsigned char color = (unsigned char)(iter % 256);
            image[idx] = color;
            image[idx + 1] = color;
            image[idx + 2] = color;
        }
    }
}

Why include the CPU version here?

Keeping the CPU function in the same file lets you time both versions in one run and print the speedup ratio directly. This is the same logic from mandelbrot_cpu.cu, just renamed to computeMandelbrotCPU to distinguish it from the kernel.

  • Now add the GPU kernel function below the CPU function. This is the code that runs on every GPU thread in parallel:
// GPU kernel - each thread computes one pixel
__global__ void mandelbrotKernel(unsigned char* image, int width, int height, int maxIter) {
    // Calculate this thread's pixel position from its ID
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // Bounds check - some threads may be outside the image
    if (x >= width || y >= height) return;

    // Map pixel to complex plane coordinates
    float cx = (float)x / width * 3.5f - 2.5f;
    float cy = (float)y / height * 2.0f - 1.0f;

    float zx = 0.0f;
    float zy = 0.0f;
    int iter = 0;

    // Escape-time algorithm (same math as CPU)
    while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
        float tmp = zx * zx - zy * zy + cx;
        zy = 2.0f * zx * zy + cy;
        zx = tmp;
        iter++;
    }

    // Write grayscale color to the image array
    int idx = (y * width + x) * 3;
    unsigned char color = (unsigned char)(iter % 256);
    image[idx] = color;
    image[idx + 1] = color;
    image[idx + 2] = color;
}

What does this kernel do?

The __global__ keyword tells the CUDA compiler this function runs on the GPU but is launched from the CPU.

  • Each thread calculates its unique pixel position using blockIdx, blockDim, and threadIdx. These are built-in CUDA variables that identify which thread is running.
  • The bounds check (if (x >= width || y >= height) return;) handles the case where the grid of threads is slightly larger than the image. Extra threads just exit early.
  • The escape-time math is identical to the CPU version. The only difference is that millions of threads execute it simultaneously, each for a different pixel.

What is the thread/block/grid hierarchy?

CUDA organizes threads into blocks, and blocks into a grid. A block is a group of threads that can cooperate (share memory, synchronize). The grid is the full collection of blocks covering your entire image.

For a 1920x1080 image with 16x16 thread blocks, the grid is 120x68 blocks, totaling over 130,000 threads launched simultaneously. Your GPU schedules them across its 1280 CUDA cores automatically.

  • Save mandelbrot_gpu.cu.

You will not see any output yet. The kernel needs a main() function to allocate GPU memory and launch it.

Seeing red squiggles in VS Code?

VS Code may not recognize __global__ or CUDA-specific keywords. This is a VS Code IntelliSense limitation, not a code error. The code compiles fine with nvcc regardless of editor warnings.

help me fix CUDA syntax highlighting in VS Code

Write the main function with host-device memory flow

The main function orchestrates everything: it allocates GPU memory, configures the thread grid, launches the kernel, then copies the result back to CPU memory. This host-to-device-to-host flow is the core pattern of every CUDA program.

  • Below the kernel, add the writePPM function (same as your CPU file) and then the main() function by pasting this code:
void writePPM(const char* filename, unsigned char* image, int width, int height) {
    FILE* fp = fopen(filename, "w");
    fprintf(fp, "P3\n%d %d\n255\n", width, height);
    for (int i = 0; i < width * height * 3; i += 3) {
        fprintf(fp, "%d %d %d\n", image[i], image[i + 1], image[i + 2]);
    }
    fclose(fp);
}

int main() {
    int imageSize = WIDTH * HEIGHT * 3;
    unsigned char* h_image = new unsigned char[imageSize];

    // Print GPU info
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("GPU: %s\n", prop.name);
    printf("SMs: %d, Max threads/block: %d, Total memory: %zu MB\n",
           prop.multiProcessorCount, prop.maxThreadsPerBlock,
           prop.totalGlobalMem / (1024 * 1024));
    printf("Image: %dx%d, Max iterations: %d\n\n", WIDTH, HEIGHT, MAX_ITER);

    // --- CPU Benchmark ---
    printf("Running CPU version...\n");
    auto cpuStart = std::chrono::high_resolution_clock::now();
    computeMandelbrotCPU(h_image, WIDTH, HEIGHT, MAX_ITER);
    auto cpuEnd = std::chrono::high_resolution_clock::now();
    float cpuTime = std::chrono::duration<float, std::milli>(cpuEnd - cpuStart).count();
    printf("CPU Time: %.2f ms\n\n", cpuTime);

    // --- GPU Benchmark ---
    printf("Running GPU version...\n");

    // Allocate memory on the GPU (device)
    unsigned char* d_image;
    CUDA_CHECK(cudaMalloc(&d_image, imageSize));

    // Configure 2D grid of 16x16 thread blocks
    dim3 blockSize(16, 16);
    dim3 gridSize((WIDTH + blockSize.x - 1) / blockSize.x,
                  (HEIGHT + blockSize.y - 1) / blockSize.y);

    // CUDA event timing (more accurate than chrono for GPU)
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    // Launch the kernel
    CUDA_CHECK(cudaEventRecord(start));
    mandelbrotKernel<<<gridSize, blockSize>>>(d_image, WIDTH, HEIGHT, MAX_ITER);
    CUDA_CHECK(cudaEventRecord(stop));

    CUDA_CHECK(cudaEventSynchronize(stop));

    // Check for kernel launch errors
    CUDA_CHECK(cudaGetLastError());

    float gpuTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&gpuTime, start, stop));
    printf("GPU Time: %.2f ms\n", gpuTime);
    printf("Speedup: %.1fx\n\n", cpuTime / gpuTime);

    // Copy result back to host
    CUDA_CHECK(cudaMemcpy(h_image, d_image, imageSize, cudaMemcpyDeviceToHost));

    // Save image
    writePPM("mandelbrot_gpu.ppm", h_image, WIDTH, HEIGHT);
    printf("Image saved to mandelbrot_gpu.ppm\n");

    // Cleanup
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    CUDA_CHECK(cudaFree(d_image));
    delete[] h_image;

    return 0;
}

What does this code do?

This is the full host-device memory flow that every CUDA program follows:

  • cudaGetDeviceProperties queries your GPU's hardware specs (SM count, max threads, memory) so you can see what you are working with.
  • cudaMalloc allocates memory on the GPU. The CPU cannot read this memory directly. The d_ prefix (device) and h_ prefix (host) are CUDA naming conventions.
  • dim3 blockSize(16, 16) creates blocks of 256 threads (16 times 16). gridSize calculates how many blocks are needed to cover the full image.
  • The <<<gridSize, blockSize>>> syntax launches the kernel across the entire grid. This is how you tell the GPU "run this function on thousands of threads at once."
  • cudaMemcpy with cudaMemcpyDeviceToHost copies the finished image from GPU memory back to CPU memory so you can save it to disk.
  • cudaFree releases GPU memory. Always pair it with cudaMalloc to avoid memory leaks on the device.

Why use CUDA events instead of chrono for GPU timing?

Kernel launches are asynchronous. The CPU queues work on the GPU and immediately continues executing the next line. chrono would measure the time to queue the kernel, not the time to execute it.

CUDA events are timestamps on the GPU's own timeline. cudaEventElapsedTime measures the actual GPU execution between two events, giving you an accurate measurement.

  • Save mandelbrot_gpu.cu.

Unsure about the memory flow?

Think of it as a three-step relay: allocate GPU memory (cudaMalloc), run the kernel (which writes into that GPU memory), then copy results back (cudaMemcpy). The CPU never touches GPU memory directly.

explain the CUDA host-device memory model for my Mandelbrot program

✔️ Awesome, I've got everything!

Great. Double check you have saved mandelbrot_gpu.cu with all the code above.

ⓧ I'd like to double check the full code

Here is the complete mandelbrot_gpu.cu file. Make sure yours matches exactly:

#include <stdio.h>
#include <stdlib.h>
#include <chrono>

const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while (0)

void computeMandelbrotCPU(unsigned char* image, int width, int height, int maxIter) {
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float cx = (float)x / width * 3.5f - 2.5f;
            float cy = (float)y / height * 2.0f - 1.0f;

            float zx = 0.0f;
            float zy = 0.0f;
            int iter = 0;

            while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
                float tmp = zx * zx - zy * zy + cx;
                zy = 2.0f * zx * zy + cy;
                zx = tmp;
                iter++;
            }

            int idx = (y * width + x) * 3;
            unsigned char color = (unsigned char)(iter % 256);
            image[idx] = color;
            image[idx + 1] = color;
            image[idx + 2] = color;
        }
    }
}

__global__ void mandelbrotKernel(unsigned char* image, int width, int height, int maxIter) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x >= width || y >= height) return;

    float cx = (float)x / width * 3.5f - 2.5f;
    float cy = (float)y / height * 2.0f - 1.0f;

    float zx = 0.0f;
    float zy = 0.0f;
    int iter = 0;

    while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
        float tmp = zx * zx - zy * zy + cx;
        zy = 2.0f * zx * zy + cy;
        zx = tmp;
        iter++;
    }

    int idx = (y * width + x) * 3;
    unsigned char color = (unsigned char)(iter % 256);
    image[idx] = color;
    image[idx + 1] = color;
    image[idx + 2] = color;
}

void writePPM(const char* filename, unsigned char* image, int width, int height) {
    FILE* fp = fopen(filename, "w");
    fprintf(fp, "P3\n%d %d\n255\n", width, height);
    for (int i = 0; i < width * height * 3; i += 3) {
        fprintf(fp, "%d %d %d\n", image[i], image[i + 1], image[i + 2]);
    }
    fclose(fp);
}

int main() {
    int imageSize = WIDTH * HEIGHT * 3;
    unsigned char* h_image = new unsigned char[imageSize];

    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("GPU: %s\n", prop.name);
    printf("SMs: %d, Max threads/block: %d, Total memory: %zu MB\n",
           prop.multiProcessorCount, prop.maxThreadsPerBlock,
           prop.totalGlobalMem / (1024 * 1024));
    printf("Image: %dx%d, Max iterations: %d\n\n", WIDTH, HEIGHT, MAX_ITER);

    printf("Running CPU version...\n");
    auto cpuStart = std::chrono::high_resolution_clock::now();
    computeMandelbrotCPU(h_image, WIDTH, HEIGHT, MAX_ITER);
    auto cpuEnd = std::chrono::high_resolution_clock::now();
    float cpuTime = std::chrono::duration<float, std::milli>(cpuEnd - cpuStart).count();
    printf("CPU Time: %.2f ms\n\n", cpuTime);

    printf("Running GPU version...\n");

    unsigned char* d_image;
    CUDA_CHECK(cudaMalloc(&d_image, imageSize));

    dim3 blockSize(16, 16);
    dim3 gridSize((WIDTH + blockSize.x - 1) / blockSize.x,
                  (HEIGHT + blockSize.y - 1) / blockSize.y);

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    CUDA_CHECK(cudaEventRecord(start));
    mandelbrotKernel<<<gridSize, blockSize>>>(d_image, WIDTH, HEIGHT, MAX_ITER);
    CUDA_CHECK(cudaEventRecord(stop));

    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaGetLastError());

    float gpuTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&gpuTime, start, stop));
    printf("GPU Time: %.2f ms\n", gpuTime);
    printf("Speedup: %.1fx\n\n", cpuTime / gpuTime);

    CUDA_CHECK(cudaMemcpy(h_image, d_image, imageSize, cudaMemcpyDeviceToHost));

    writePPM("mandelbrot_gpu.ppm", h_image, WIDTH, HEIGHT);
    printf("Image saved to mandelbrot_gpu.ppm\n");

    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    CUDA_CHECK(cudaFree(d_image));
    delete[] h_image;

    return 0;
}
Compile and compare performance

Time to compile the GPU version and see how it stacks up against the CPU. The -arch=sm_75 flag tells nvcc to optimize for your GTX 1650 Super's Turing architecture.

  • In your x64 Native Tools Command Prompt, compile the GPU version by running this command:
nvcc mandelbrot_gpu.cu -arch=sm_75 -o mandelbrot_gpu.exe
  • Run the compiled program:
mandelbrot_gpu.exe

You should see output showing your GPU name, SM count, total memory, the CPU execution time, the GPU execution time, and a speedup ratio. The GPU time should be dramatically faster than the CPU time.

  • Open mandelbrot_gpu.ppm in an image viewer. It should show the same Mandelbrot fractal as your CPU version.

What should I see?

Your console output will look something like this (exact numbers depend on your system):

  • GPU: NVIDIA GeForce GTX 1650 SUPER with SM count, max threads, and memory info.
  • A CPU time in the thousands of milliseconds (several seconds).
  • A GPU time in single-digit or low double-digit milliseconds.
  • A speedup ratio showing how many times faster the GPU is (expect 50x to 200x or more).

Seeing a CUDA error?

If you see a compilation error about sm_75, make sure you are using the x64 Native Tools Command Prompt (not regular Command Prompt or PowerShell).

If you see a runtime error from CUDA_CHECK, check that no other application is heavily using your GPU. Also verify your GPU driver is up to date by running nvidia-smi and checking the driver version.

help me debug my CUDA compilation or runtime error

Your GPU just computed 2 million pixels in parallel. Next up, you will experiment with different thread configurations to understand why block size matters for performance.

Experiment with Thread Configuration

You have a working GPU Mandelbrot generator that crushes the CPU version. But how do you know the 16x16 block size is actually optimal?

The GPU's performance depends heavily on how you organize threads into blocks. In this step, you will experiment with different configurations to build intuition for why block size matters and how your GPU scales with bigger workloads.

In this step, get ready to:
  • Test 8x8, 16x16, and 32x32 block sizes and compare GPU execution times.
  • Understand the error checking and GPU property reporting already in your code.
  • Push the GPU to 4K resolution and observe how it scales.
Test different block sizes

Your current code uses a 16x16 block, meaning 256 threads per block. Each block runs on one Streaming Multiprocessor (SM). Choosing the right block size affects how efficiently the GPU fills its SMs with work.

  • In VS Code, open mandelbrot_gpu.cu and find this line in main():
    dim3 blockSize(16, 16);
  • Change 16, 16 to 8, 8 so the line looks like this:
    dim3 blockSize(8, 8);
  • Save mandelbrot_gpu.cu.
  • Return to the x64 Native Tools Command Prompt and recompile by running:
nvcc mandelbrot_gpu.cu -arch=sm_75 -o mandelbrot_gpu.exe
  • Run the program by running:
mandelbrot_gpu.exe

Note the GPU time it prints. An 8x8 block means only 64 threads per block, which may underutilize each SM.

  • Now change the block size to 32x32 (1024 threads per block) by editing the same line:
    dim3 blockSize(32, 32);
  • Save, recompile, and run again:
nvcc mandelbrot_gpu.cu -arch=sm_75 -o mandelbrot_gpu.exe
mandelbrot_gpu.exe

32x32 = 1024 threads, which is the maximum your GPU allows per block. Compare all three times.

Why does block size matter?

Each SM on your GPU can run multiple blocks simultaneously, up to a limit (16 blocks on CC 7.5 hardware). Smaller blocks mean more blocks to schedule, but each block has fewer threads for the SM to work with.

Larger blocks pack more threads together but leave less room for multiple blocks per SM. The 16x16 sweet spot (256 threads) gives your GPU flexibility to fit several blocks on each SM while keeping warp scheduling efficient.

  • Set blockSize back to 16, 16 (the optimal default):
    dim3 blockSize(16, 16);
  • Save mandelbrot_gpu.cu.

Seeing a CUDA error with 32x32?

If you accidentally try a block size larger than 1024 total threads (for example, 64x64 = 4096), the CUDA_CHECK macro will catch the error and print a message. The max threads per block for your GPU is 1024. Stick to 8x8, 16x16, or 32x32.

Help me debug my CUDA block size error.

Examine error checking and GPU properties

Your code already includes two important defensive programming features. Understanding them now will save you hours of debugging in future CUDA projects.

  • In mandelbrot_gpu.cu, find the CUDA_CHECK macro near the top of the file:
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while (0)

What does CUDA_CHECK do?

Every CUDA API function returns a cudaError_t value. Without checking it, failures are silent. Your code could crash or produce garbage output with no indication of why.

This macro wraps every CUDA call and immediately halts the program with a helpful message (file, line number, error description) if anything goes wrong. It catches invalid block sizes, out-of-memory errors, and driver issues.

  • Now find the cudaGetDeviceProperties section at the top of main():
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("GPU: %s\n", prop.name);
    printf("SMs: %d, Max threads/block: %d, Total memory: %zu MB\n",
           prop.multiProcessorCount, prop.maxThreadsPerBlock,
           prop.totalGlobalMem / (1024 * 1024));
    printf("Image: %dx%d, Max iterations: %d\n\n", WIDTH, HEIGHT, MAX_ITER);

What does cudaGetDeviceProperties tell you?

This function queries your GPU's hardware specs at runtime. The output tells you exactly what your hardware can do:

  • prop.name prints the GPU model name (GTX 1650 Super).
  • prop.multiProcessorCount shows the number of Streaming Multiprocessors. More SMs means more blocks can run simultaneously.
  • prop.maxThreadsPerBlock confirms the 1024 thread limit you discovered in the block size experiment.
  • prop.totalGlobalMem shows total GPU memory (4 GB on your card).
  • Also find the cudaGetLastError() call after the kernel launch:
    CUDA_CHECK(cudaGetLastError());

Kernel launches are asynchronous. Errors during the kernel (like out-of-bounds memory access) only surface when you call cudaGetLastError() or a synchronizing call. This line catches those errors immediately after the kernel finishes.

Test at 4K resolution

The GPU's advantage grows with more pixels because there is more parallelism to exploit. A 4K image has four times the pixels of 1080p, meaning four times the independent work for the GPU to distribute across its cores.

  • In mandelbrot_gpu.cu, find the constants at the top of the file:
const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;
  • Change them to 4K resolution and a higher iteration count:
const int WIDTH = 3840;
const int HEIGHT = 2160;
const int MAX_ITER = 2000;

Why increase MAX_ITER too?

Higher resolution reveals finer detail at the fractal boundary. More iterations allow the algorithm to distinguish pixels that are "almost" in the set from those that escape slowly. Without increasing iterations, the extra pixels at the boundary would all appear black.

  • Save mandelbrot_gpu.cu, then recompile and run:
nvcc mandelbrot_gpu.cu -arch=sm_75 -o mandelbrot_gpu.exe
mandelbrot_gpu.exe

Compare the CPU and GPU times. The CPU time should increase dramatically (roughly 4x or more), while the GPU time increases much less because it distributes the extra pixels across all 1280 cores.

Why does the GPU scale better?

At 1920x1080, your GPU launches about 2 million threads. At 3840x2160, it launches about 8.3 million. Either way, the GPU processes them in large batches (limited by SM count and block scheduling). The CPU processes them one at a time. Doubling both dimensions quadruples the CPU's work but barely dents the GPU because it already had more threads than cores.

  • Open the generated mandelbrot_gpu.ppm file in an image viewer. The 4K image shows noticeably finer detail along the fractal boundary.

Program crashes or prints a CUDA error?

At 4K with 3 bytes per pixel, the image needs about 24 MB. Your GPU has 4 GB, so memory is not the issue. If you see an error, check that you did not accidentally remove the bounds check (if (x >= width || y >= height) return;) from the kernel.

Help me debug my 4K CUDA Mandelbrot render.

  • Once you have observed the 4K results, set WIDTH back to 1920, HEIGHT back to 1080, and MAX_ITER back to 1000:
const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;
  • Save mandelbrot_gpu.cu and recompile one final time:
nvcc mandelbrot_gpu.cu -arch=sm_75 -o mandelbrot_gpu.exe

✔️ Awesome, I've got everything!

Great. Double check you have saved your file with the original 1920x1080 resolution restored.

ⓧ I'd like to double check the full code

Here is the full mandelbrot_gpu.cu as it should look after reverting your experiments:

#include <stdio.h>
#include <stdlib.h>
#include <chrono>

const int WIDTH = 1920;
const int HEIGHT = 1080;
const int MAX_ITER = 1000;

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while (0)

// CPU version for comparison
void computeMandelbrotCPU(unsigned char* image, int width, int height, int maxIter) {
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float cx = (float)x / width * 3.5f - 2.5f;
            float cy = (float)y / height * 2.0f - 1.0f;

            float zx = 0.0f;
            float zy = 0.0f;
            int iter = 0;

            while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
                float tmp = zx * zx - zy * zy + cx;
                zy = 2.0f * zx * zy + cy;
                zx = tmp;
                iter++;
            }

            int idx = (y * width + x) * 3;
            unsigned char color = (unsigned char)(iter % 256);
            image[idx] = color;
            image[idx + 1] = color;
            image[idx + 2] = color;
        }
    }
}

// GPU kernel
__global__ void mandelbrotKernel(unsigned char* image, int width, int height, int maxIter) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x >= width || y >= height) return;

    float cx = (float)x / width * 3.5f - 2.5f;
    float cy = (float)y / height * 2.0f - 1.0f;

    float zx = 0.0f;
    float zy = 0.0f;
    int iter = 0;

    while (zx * zx + zy * zy < 4.0f && iter < maxIter) {
        float tmp = zx * zx - zy * zy + cx;
        zy = 2.0f * zx * zy + cy;
        zx = tmp;
        iter++;
    }

    int idx = (y * width + x) * 3;
    unsigned char color = (unsigned char)(iter % 256);
    image[idx] = color;
    image[idx + 1] = color;
    image[idx + 2] = color;
}

void writePPM(const char* filename, unsigned char* image, int width, int height) {
    FILE* fp = fopen(filename, "w");
    fprintf(fp, "P3\n%d %d\n255\n", width, height);
    for (int i = 0; i < width * height * 3; i += 3) {
        fprintf(fp, "%d %d %d\n", image[i], image[i + 1], image[i + 2]);
    }
    fclose(fp);
}

int main() {
    int imageSize = WIDTH * HEIGHT * 3;
    unsigned char* h_image = new unsigned char[imageSize];

    // Print GPU info
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("GPU: %s\n", prop.name);
    printf("SMs: %d, Max threads/block: %d, Total memory: %zu MB\n",
           prop.multiProcessorCount, prop.maxThreadsPerBlock,
           prop.totalGlobalMem / (1024 * 1024));
    printf("Image: %dx%d, Max iterations: %d\n\n", WIDTH, HEIGHT, MAX_ITER);

    // --- CPU Benchmark ---
    printf("Running CPU version...\n");
    auto cpuStart = std::chrono::high_resolution_clock::now();
    computeMandelbrotCPU(h_image, WIDTH, HEIGHT, MAX_ITER);
    auto cpuEnd = std::chrono::high_resolution_clock::now();
    float cpuTime = std::chrono::duration<float, std::milli>(cpuEnd - cpuStart).count();
    printf("CPU Time: %.2f ms\n\n", cpuTime);

    // --- GPU Benchmark ---
    printf("Running GPU version...\n");

    unsigned char* d_image;
    CUDA_CHECK(cudaMalloc(&d_image, imageSize));

    dim3 blockSize(16, 16);
    dim3 gridSize((WIDTH + blockSize.x - 1) / blockSize.x,
                  (HEIGHT + blockSize.y - 1) / blockSize.y);

    // CUDA event timing
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    CUDA_CHECK(cudaEventRecord(start));
    mandelbrotKernel<<<gridSize, blockSize>>>(d_image, WIDTH, HEIGHT, MAX_ITER);
    CUDA_CHECK(cudaEventRecord(stop));

    CUDA_CHECK(cudaEventSynchronize(stop));

    // Check for kernel launch errors
    CUDA_CHECK(cudaGetLastError());

    float gpuTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&gpuTime, start, stop));
    printf("GPU Time: %.2f ms\n", gpuTime);
    printf("Speedup: %.1fx\n\n", cpuTime / gpuTime);

    // Copy result back to host
    CUDA_CHECK(cudaMemcpy(h_image, d_image, imageSize, cudaMemcpyDeviceToHost));

    // Save image
    writePPM("mandelbrot_gpu.ppm", h_image, WIDTH, HEIGHT);
    printf("Image saved to mandelbrot_gpu.ppm\n");

    // Cleanup
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    CUDA_CHECK(cudaFree(d_image));
    delete[] h_image;

    return 0;
}

You now have a solid understanding of how thread configuration affects GPU performance. Next up, you will take this further in the Secret Mission by adding smooth color gradients and fractal zoom rendering.

Secret mission

Color Gradients and Fractal Zoom

Your fractal images so far are grayscale and show the same view every time. What if you could add vibrant colors and zoom into the infinite detail hiding at the fractal's boundary? In this mission, you will build a new CUDA kernel that produces smooth color gradients and renders magnified views of the most beautiful regions of the Mandelbrot set.

Clean Up Your Resources

Clean Up Your Resources

This project runs entirely locally on your machine, so there are no ongoing costs. Decide whether to keep your resources running, pause them to come back later, or delete them entirely.

Resources you used:

  • The C:\cuda-mandelbrot project folder containing source files (.cu), executables (.exe), and generated PPM images.
  • Visual Studio 2022 Community Edition (MSVC compiler for building CUDA code).
  • NVIDIA CUDA Toolkit 13.3 (provides nvcc compiler and CUDA runtime).

Keep everything running

No action needed. Choose this if you plan to continue experimenting with CUDA programming or want to build more GPU-accelerated projects.

  • Your project folder stays at C:\cuda-mandelbrot with all source files, executables, and images intact.
  • Visual Studio 2022 and the CUDA Toolkit remain installed for future development. Together they use approximately 10 GB of disk space.
  • The generated .ppm image files are uncompressed and large. You can keep them for reference or delete them individually without affecting your source code.

Pause - I'll come back to this later

Free up disk space from generated output files while keeping your development tools and source code ready to pick back up.

  • Delete the large PPM image files by opening your x64 Native Tools Command Prompt and running:
cd C:\cuda-mandelbrot
del *.ppm
del *.exe
  • This removes the generated images and compiled executables. Your .cu source files remain so you can recompile any time.
  • Keep Visual Studio 2022 and the CUDA Toolkit installed. They cost nothing to leave on your machine.

Delete - I don't want to use this again

Remove all project files and optionally uninstall the development tools to reclaim disk space.

  • Delete the entire project folder by opening your x64 Native Tools Command Prompt and running:
rmdir /s /q C:\cuda-mandelbrot
  • To uninstall Visual Studio 2022, press the Windows key, type Add or remove programs, and press Enter.
  • Find Microsoft Visual Studio Community 2022 in the list, click the three-dot menu, and select Uninstall.
  • To uninstall the CUDA Toolkit, find NVIDIA CUDA Toolkit 13.3 in the same list, click the three-dot menu, and select Uninstall.
  • Removing both frees approximately 10 GB of disk space.

That's a Wrap!

That's a Wrap!

Nice work! You just built a GPU-accelerated Mandelbrot fractal generator from scratch using CUDA C++ and your GTX 1650 Super's 1280 cores.

You've learned how to:

  • Written a CUDA kernel that launches thousands of parallel threads across your GPU, computing every pixel of a 1920x1080 image simultaneously instead of one at a time.
  • Managed the host-device memory flow using cudaMalloc, cudaMemcpy, and cudaFree to transfer data between your CPU and GPU.
  • Measured real CPU vs GPU performance with CUDA event timing and experimented with thread/block configurations to understand how the GPU's execution model works.
  • (Secret Mission) Implemented smooth color gradient mapping and a fractal zoom function that renders magnified regions of the Mandelbrot set at multiple zoom levels.

Ready to quiz yourself?