CUDA Testing Example: How to Verify Your GPU Code

Short answer: A CUDA test should prove two things: the kernel returns the correct result, and the CUDA runtime reported no error. The pattern is to prepare small input, run the kernel, call cudaGetLastError() and cudaDeviceSynchronize(), copy the result back, and compare it with a CPU-computed reference. Test awkward sizes (1, one less than a block, one more than a block) and run compute-sanitizer to catch memory errors that a correct-looking output can hide.

A CUDA kernel can compile cleanly, launch without an obvious error, and still return the wrong answer. That is why a useful CUDA testing example does more than confirm that code runs. It checks output against a known result, catches CUDA API failures at the right point, and makes debugging repeatable.

CUDA is NVIDIA's parallel computing platform and programming model for running general-purpose code on NVIDIA GPUs. It only runs on NVIDIA hardware, so a Mac, a lightweight laptop or a standard office PC without an NVIDIA GPU cannot run these tests locally. A cloud Windows workstation with a dedicated NVIDIA GPU gives you a practical place to build, run and test CUDA code without buying a workstation just for a project.

Check Your CUDA Setup First

Before testing a kernel, confirm that the machine can run CUDA at all.

These two numbers can differ, and that is normal. nvidia-smi reports what the driver can support. nvcc reports the toolkit you will compile with. If nvcc is not found, the CUDA Toolkit is not installed or not on your path.

On Windows, nvcc needs a C++ host compiler. Install Visual Studio with the C++ workload, then run your build from a Visual Studio developer command prompt (such as "x64 Native Tools Command Prompt"). For install details, follow NVIDIA's current CUDA installation guide.

What a CUDA Test Should Actually Prove

GPU testing has two jobs. First, it verifies that your kernel produces the expected result. Second, it verifies that the CUDA runtime did not report a failure that your application ignored.

Those are separate checks. A kernel launch is asynchronous. The CPU may continue before the GPU has completed the work. If you only call the kernel and copy results later, an invalid memory access or bad launch configuration can be harder to trace.

A good test therefore follows a simple pattern: prepare small input data, run the GPU kernel, synchronize, check for errors, copy the result back, and compare it with a CPU-generated reference result. Keep the first tests small. A vector with 16 or 64 elements is easier to inspect than a production-sized dataset.

CUDA Testing Example: A Small Vector Addition Test

Vector addition is simple, but it covers the core mechanics of CUDA testing: device allocation, memory transfer, a kernel launch, synchronization, cleanup and result validation.

#include <cuda_runtime.h>
#include <cmath>
#include <iostream>
#include <vector>

#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t error = (call);                                         \
        if (error != cudaSuccess) {                                         \
            std::cerr << "CUDA error: " << cudaGetErrorString(error)        \
                      << " at line " << __LINE__ << std::endl;              \
            return 1;                                                       \
        }                                                                   \
    } while (0)

__global__ void addVectors(const float* a, const float* b, float* result,
                           int count) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < count) {
        result[index] = a[index] + b[index];
    }
}

int main() {
    const int count = 64;
    const size_t bytes = count * sizeof(float);

    std::vector<float> hostA(count);
    std::vector<float> hostB(count);
    std::vector<float> hostResult(count);
    std::vector<float> expected(count);

    for (int i = 0; i < count; ++i) {
        hostA[i] = static_cast<float>(i);
        hostB[i] = static_cast<float>(count - i);
        expected[i] = hostA[i] + hostB[i];
    }

    float *deviceA, *deviceB, *deviceResult;
    CUDA_CHECK(cudaMalloc(&deviceA, bytes));
    CUDA_CHECK(cudaMalloc(&deviceB, bytes));
    CUDA_CHECK(cudaMalloc(&deviceResult, bytes));

    CUDA_CHECK(cudaMemcpy(deviceA, hostA.data(), bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(deviceB, hostB.data(), bytes, cudaMemcpyHostToDevice));

    const int threadsPerBlock = 32;
    const int blocks = (count + threadsPerBlock - 1) / threadsPerBlock;
    addVectors<<<blocks, threadsPerBlock>>>(deviceA, deviceB, deviceResult, count);

    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    CUDA_CHECK(cudaMemcpy(hostResult.data(), deviceResult, bytes,
                          cudaMemcpyDeviceToHost));

    for (int i = 0; i < count; ++i) {
        if (std::fabs(hostResult[i] - expected[i]) > 0.0001f) {
            std::cerr << "Test failed at index " << i
                      << ": expected " << expected[i]
                      << ", got " << hostResult[i] << std::endl;
            return 1;
        }
    }

    cudaFree(deviceA);
    cudaFree(deviceB);
    cudaFree(deviceResult);

    std::cout << "CUDA vector addition test passed." << std::endl;
    return 0;
}

Save the file as vector_test.cu, then compile it with the CUDA compiler:

nvcc vector_test.cu -o vector_test

On Windows, run vector_test.exe from a terminal configured for your CUDA toolchain. A passing run prints CUDA vector addition test passed. That tells you this specific kernel produced the expected output on the GPU available to that machine. It does not prove every future input or kernel configuration is correct. That is where a test strategy matters.

Why the Error Checks Are in Two Places

cudaGetLastError() checks whether the launch itself was accepted. Common launch-side problems include an invalid configuration, invalid arguments, or a kernel that could not be started.

cudaDeviceSynchronize() waits for the GPU and reports errors that happen while the kernel executes. This is where out-of-bounds memory access and some execution failures may surface. Without synchronization, a test may report success before the GPU has finished its work.

Use both checks in early development and automated tests. Once you understand the behavior of a larger application, you may use streams and more targeted synchronization for better throughput. The tradeoff is simple: less synchronization can improve performance, but it also makes failures more difficult to locate.

Use a CPU Reference, Not a Guess

The strongest part of the example is not the CUDA code. It is the expected vector. The CPU computes the same operation using straightforward code, then the test compares every GPU result to that reference.

For integer operations, exact equality is often appropriate. Floating-point work needs more care. Different instruction ordering, fused operations and reduction patterns can create small numerical differences. Comparing floats with == can create false failures.

Choose a tolerance that matches the algorithm. For a basic float addition test, a small absolute tolerance can be enough. For matrix operations, reductions or mixed-precision code, use an absolute and a relative tolerance. Test the result against what the workload actually requires, not against an unrealistically perfect decimal match.

Add Cases That Expose Real Kernel Bugs

A single happy-path test is not enough. CUDA indexing mistakes often appear at boundaries, especially when the input size is not evenly divisible by the block size.

Build tests around small and awkward sizes: one element, one less than a block, exactly one block, one more than a block, and a size that spans several blocks. Also test negative values, repeated values and large values when the algorithm handles them. A launch with zero elements is a special case, because a grid with zero blocks is an invalid launch configuration, so handle empty input in your own code before launching.

Here is the same test turned into a size sweep. It also frees GPU memory automatically if a check fails partway through, which the first version does not.

#include <cuda_runtime.h>
#include <cmath>
#include <iostream>
#include <vector>

#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t error = (call);                                         \
        if (error != cudaSuccess) {                                         \
            std::cerr << "CUDA error: " << cudaGetErrorString(error)        \
                      << " at line " << __LINE__ << std::endl;              \
            return false;                                                   \
        }                                                                   \
    } while (0)

// Frees GPU memory automatically, even when a check returns early.
// cudaFree(nullptr) is a harmless no-op.
struct DeviceBuffer {
    float* ptr = nullptr;
    ~DeviceBuffer() { cudaFree(ptr); }
};

__global__ void addVectors(const float* a, const float* b, float* result,
                           int count) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < count) {
        result[index] = a[index] + b[index];
    }
}

// Runs one test for a given element count. Returns true if the GPU result
// matches the CPU reference. Sizes start at 1: a launch with zero blocks
// is an invalid configuration.
bool runVectorAddTest(int count) {
    const size_t bytes = count * sizeof(float);

    std::vector<float> hostA(count), hostB(count);
    std::vector<float> hostResult(count), expected(count);
    for (int i = 0; i < count; ++i) {
        hostA[i] = static_cast<float>(i);
        hostB[i] = static_cast<float>(count - i);
        expected[i] = hostA[i] + hostB[i];  // CPU reference
    }

    DeviceBuffer a, b, result;
    CUDA_CHECK(cudaMalloc(&a.ptr, bytes));
    CUDA_CHECK(cudaMalloc(&b.ptr, bytes));
    CUDA_CHECK(cudaMalloc(&result.ptr, bytes));
    CUDA_CHECK(cudaMemcpy(a.ptr, hostA.data(), bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(b.ptr, hostB.data(), bytes, cudaMemcpyHostToDevice));

    const int threadsPerBlock = 32;
    const int blocks = (count + threadsPerBlock - 1) / threadsPerBlock;
    addVectors<<<blocks, threadsPerBlock>>>(a.ptr, b.ptr, result.ptr, count);

    CUDA_CHECK(cudaGetLastError());       // was the launch accepted?
    CUDA_CHECK(cudaDeviceSynchronize());  // did the kernel finish cleanly?
    CUDA_CHECK(cudaMemcpy(hostResult.data(), result.ptr, bytes,
                          cudaMemcpyDeviceToHost));

    for (int i = 0; i < count; ++i) {
        if (std::fabs(hostResult[i] - expected[i]) > 0.0001f) {
            std::cerr << "Test failed (count=" << count << ") at index " << i
                      << ": expected " << expected[i]
                      << ", got " << hostResult[i] << std::endl;
            return false;
        }
    }
    return true;
}

int main() {
    // One element, just under / exactly at / just over a block, several
    // blocks, and a larger size.
    const int sizes[] = {1, 31, 32, 33, 64, 1000, 1 << 20};
    bool allPassed = true;
    for (int count : sizes) {
        bool passed = runVectorAddTest(count);
        std::cout << "count=" << count << (passed ? " passed" : " FAILED")
                  << std::endl;
        allPassed = allPassed && passed;
    }
    return allPassed ? 0 : 1;
}

Every size should print passed. Try removing the if (index < count) check and re-running the sweep. The sizes that are not a multiple of 32 are where out-of-range threads write past the buffer, which is the bug this sweep is designed to expose.

For image, tensor or matrix kernels, vary the dimensions. A square 1024-by-1024 input may hide an indexing error that appears immediately with a 17-by-31 input. If your kernel uses shared memory, test dimensions near the tile boundaries. If it uses multiple streams, test synchronization and output ordering explicitly.

Catch Memory Problems Before They Reach Production

Correct output on one input does not rule out unsafe memory behavior. An out-of-range thread may write beyond the output buffer and happen not to affect the result you inspect.

Use NVIDIA's compute-sanitizer, which ships with the CUDA Toolkit, as part of development:

compute-sanitizer vector_test

By default it runs the memory checker, which reports invalid global memory accesses. Other tools in the same program cover race conditions (--tool racecheck), uninitialized memory (--tool initcheck) and synchronization errors (--tool synccheck). Compiling with -lineinfo lets the reports point to source lines. These tools add overhead, so use them in a debug workflow rather than as the only run for large workloads. NVIDIA documents the options in the Compute Sanitizer docs.

Also check every allocation, copy and cleanup call. A failed cudaMalloc is easy to miss if your code assumes the device has enough memory. The CUDA_CHECK macro keeps failures close to the line that caused them, which is much more useful than discovering a generic failure later.

Turn the Example Into Repeatable Tests

As a project grows, move the setup into helper functions and run the cases through a test framework such as GoogleTest or Catch2. The CUDA pattern stays the same: generate input, compute a reference result, launch, synchronize, copy back, compare.

Keep unit tests focused on one kernel or operation. Then add integration tests for pipelines that involve preprocessing, several kernels, and file or model inputs. Unit tests find indexing and math errors fast. Integration tests catch mismatched assumptions between stages.

Where to Learn More CUDA

For a hands-on path, start with NVIDIA's own material: the CUDA documentation (including the programming guide) and the official CUDA samples repository for reference applications. Write and test small kernels first, as in this article, before moving to shared memory, streams and performance tuning.

Where a Cloud GPU Machine Fits

A dedicated cloud GPU machine is useful when your local device cannot run the test environment comfortably, such as a Mac, a thin laptop or an office PC with no NVIDIA GPU. SensePC provides a Windows 11 cloud PC with dedicated NVIDIA GPU configurations and persistent SSD storage, so source code, tools and test data remain available between sessions. You install your own toolchain (CUDA Toolkit and a C++ compiler), then return to the same configured workspace instead of rebuilding a local setup every time. Access is hourly, daily or monthly.

Cloud access is not automatically the best choice for every workflow. If you run long GPU jobs daily and need consistent local access, owning hardware may be the better financial fit. If your CUDA work is project-based, experimental, or limited by a laptop or Mac, hourly cloud GPU access can offer more control over cost and hardware commitment.

The practical next step is simple: make your first CUDA test intentionally small, verify both output and runtime errors, then add the edge cases that your real data will eventually hit. A kernel that passes those tests is far easier to trust when the workload gets bigger.

SensePC at a Glance

What is a good CUDA testing example for beginners?

A vector addition test. Allocate small arrays, run the kernel, call cudaGetLastError() and cudaDeviceSynchronize(), copy the result back and compare it with a CPU-computed reference. It covers allocation, transfer, launch, synchronization and validation in under 100 lines.

How do I check my CUDA version?

Run nvcc --version to see the installed CUDA Toolkit version, and nvidia-smi to see the driver and the highest CUDA version it supports. The two can differ, and that is normal.

Why do I need both cudaGetLastError() and cudaDeviceSynchronize()?

cudaGetLastError() reports whether the kernel launch was accepted. cudaDeviceSynchronize() waits for the GPU and reports errors that occur during execution, such as invalid memory access. Without both, a test can pass before the kernel has finished, or miss a failure.

How do I test a CUDA kernel's output is correct?

Compute the same result on the CPU with straightforward code, then compare every GPU value against it. Use exact equality for integers and a tolerance for floating-point results.

What input sizes should I test?

One element, one less than a block size, exactly one block, one more than a block, and a size that spans several blocks. Boundary sizes expose indexing mistakes that a round number like 1,024 hides.

How do I find memory errors in CUDA code?

Run your program under compute-sanitizer, which ships with the CUDA Toolkit. Its default memory checker reports invalid global memory accesses. Other tools check for race conditions, uninitialized memory and synchronization problems. Compile with -lineinfo so reports show source lines.

Can I run CUDA code on a Mac?

Not on the Mac itself, because CUDA needs an NVIDIA GPU. Mac users can run CUDA on a remote NVIDIA GPU machine, such as a cloud Windows PC, and work from the Mac over remote desktop.

What do I need to compile CUDA on Windows?

An NVIDIA GPU with a current driver, the CUDA Toolkit, and Visual Studio with the C++ workload for the host compiler. Build from a Visual Studio developer command prompt. Follow NVIDIA's current installation guide for versions.

How do I install CUDA?

Install the NVIDIA driver, then the CUDA Toolkit from NVIDIA, then confirm with nvcc --version and nvidia-smi. The exact steps depend on your operating system and GPU, so use NVIDIA's installation guide for your platform.

What are the best resources to learn CUDA programming?

Start with NVIDIA's CUDA documentation and programming guide, and the official CUDA samples on GitHub. Write small kernels with tests first, then move on to shared memory and performance tuning.