vecadd.cu

#include <cstdio>
#include <cstdlib>
#include <chrono>
#include <cmath>

__global__ void vecAdd(const float *a, const float *b, float *c, int n) {
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid < n) c[gid] = a[gid] + b[gid];
}

int main() {
    int N = 1<<20; // 1M elements (tweak as needed)
    size_t bytes = N * sizeof(float);

    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);

    for (int i = 0; i < N; ++i) {
        h_a[i] = 1.0f * i;
        h_b[i] = 2.0f * i;
    }

    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;

    // warm-up
    vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);
    cudaDeviceSynchronize();

    // timed launch
    auto t0 = std::chrono::high_resolution_clock::now();
    vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);
    cudaDeviceSynchronize();
    auto t1 = std::chrono::high_resolution_clock::now();

    double ms = std::chrono::duration<double, std::milli>(t1 - t0).count();
    printf("Vector add: N=%d time=%.3f ms\n", N, ms);

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // quick correctness check
    for (int i = 0; i < 10; ++i) {
        if (fabs(h_c[i] - (h_a[i] + h_b[i])) > 1e-6f) {
            printf("Mismatch at %d\n", i); break;
        }
    }
    printf("Sample c[0..9]: ");
    for (int i = 0; i < 10; ++i) printf("%.1f ", h_c[i]);
    printf("\n");

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}