david-grs / geiger Goto Github PK

#include <iostream>
#include <cstdlib>
#include <cmath>
#include <chrono>
#include <algorithm>
#include <geiger/geiger.h>

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

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

int main()
{
    geiger::init();
    geiger::suite<> suite;

    size_t size = 1000;

    int block_size = 64;
    int grid_size = (int)ceil((float)size/block_size);

    float* a = (float*)malloc(sizeof(float)*size);
    float* b = (float*)malloc(sizeof(float)*size);
    float* c = (float*)malloc(sizeof(float)*size);

    std::fill(a, a + size, 5);
    std::fill(b, b + size, 3);

    float* d_a;
    float* d_b;
    float* d_c;

    cudaMalloc(&d_a, sizeof(float)*size);
    cudaMalloc(&d_b, sizeof(float)*size);
    cudaMalloc(&d_c, sizeof(float)*size);

    cudaMemcpy(d_a, a, sizeof(float) * size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(float) * size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_c, c, sizeof(float) * size, cudaMemcpyHostToDevice);

    suite.add("hadamard standard", [d_a, d_b, d_c,
				    size, block_size, grid_size]()
	      {
		  hadamard<<<grid_size, block_size>>>(size, d_a, d_b, d_c);
	      })
	.add("hadamard fast", [d_a, d_b, d_c,
			       size, block_size, grid_size]()
	      {
		  hadamard_fast<<<grid_size, block_size>>>(size, d_a, d_b, d_c);
	      })
	.set_printer<geiger::printer::console<>>()
	.run(std::chrono::milliseconds(1));

    cudaMemcpy(c, d_c, sizeof(float) * size, cudaMemcpyDeviceToHost);

    std::cout << "result: " << c[0] << std::endl;

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(a);
    free(b);
    free(c);

    return 0;
}