1	/*
     2	cuda_memory.cu
     3	
     4	CUDA memory utilities
     5	*/
     6	
     7	#include <stdlib.h>
     8	#include <string.h>
     9	
    10	// malloc and clear
    11	void cuda_malloc(int gpu, int um, void **ptr, size_t size)
    12	{
    13		if (gpu) {
    14			if (um) {
    15				cudaMallocManaged(ptr, size);
    16			}
    17			else {
    18				cudaMalloc(ptr, size);
    19			}
    20			cudaMemset(*ptr, 0, size);
    21		}
    22		else {
    23			*ptr = malloc(size);
    24			memset(*ptr, 0, size);
    25		}
    26	}
    27	
    28	// free
    29	void cuda_free(int gpu, void *ptr)
    30	{
    31		if (gpu) {
    32			cudaFree(ptr);
    33		}
    34		else {
    35			free(ptr);
    36		}
    37	}
    38	
    39	// memset
    40	void cuda_memset(int gpu, void *ptr, int c, size_t size)
    41	{
    42		if (gpu) {
    43			cudaMemset(ptr, c, size);
    44		}
    45		else {
    46			memset(ptr, c, size);
    47		}
    48	}
    49	
    50	// memcpy
    51	void cuda_memcpy(int gpu, void *dst, const void *src, size_t size, cudaMemcpyKind kind)
    52	{
    53		if (gpu) {
    54			cudaMemcpy(dst, src, size, kind);
    55		}
    56		else {
    57			memcpy(dst, src, size);
    58		}
    59	}

     1	/*
     2	add two vectors (CUDA)
     3	
     4	Compile,Link:
     5	> nvcc -O3 -o vadd_cuda vadd_cuda.cu cuda_memory.cu
     6	
     7	Usage:
     8	> vadd_cuda [-gpu|-cpu] [-hdm|-um] [-device <device>] <num> <loop>
     9	*/
    10	
    11	// GPU/CPU
    12	__host__ __device__
    13	static void vadd_calc(float a, float b, float *c)
    14	{
    15		*c = a + b;
    16	}
    17	
    18	// GPU
    19	__global__
    20	static void vadd_gpu(int n, const float *a, const float *b, float *c)
    21	{
    22		int tid = threadIdx.x + (blockIdx.x * blockDim.x);
    23		if (tid < n) {
    24			vadd_calc(a[tid], b[tid], &c[tid]);
    25		}
    26	}
    27	
    28	// CPU
    29	static void vadd_cpu(int n, const float *a, const float *b, float *c)
    30	{
    31		for (int i = 0; i < n; i++) {
    32			vadd_calc(a[i], b[i], &c[i]);
    33		}
    34	}
    35	
    36	// GPU/CPU
    37	static void vadd(int gpu, int n, const float *a, const float *b, float *c)
    38	{
    39		if (gpu) {
    40			int block = 256;
    41			int grid = (n + block - 1) / block;
    42			vadd_gpu<<<grid, block>>>(n, a, b, c);
    43		}
    44		else {
    45			vadd_cpu(n, a, b, c);
    46		}
    47	}
    48	
    49	#include <stdlib.h>
    50	#include <stdio.h>
    51	#include <string.h>
    52	#include <time.h>
    53	
    54	extern void cuda_malloc(int, int, void **, size_t);
    55	extern void cuda_free(int, void *);
    56	extern void cuda_memcpy(int, void *, const void *, size_t, cudaMemcpyKind);
    57	
    58	int main(int argc, char **argv)
    59	{
    60		int    gpu = 1;
    61		int    um = 0;
    62		int    device = 0;
    63		int    n = 1000;
    64		int    nloop = 1000;
    65		float  *a, *b, *c;
    66		clock_t t0 = 0, t1 = 0;
    67	
    68		// arguments
    69		while (--argc) {
    70			argv++;
    71			if      (!strcmp(*argv, "-gpu")) {
    72				gpu = 1;
    73			}
    74			else if (!strcmp(*argv, "-cpu")) {
    75				gpu = 0;
    76			}
    77			else if (!strcmp(*argv, "-hdm")) {
    78				um = 0;
    79			}
    80			else if (!strcmp(*argv, "-um")) {
    81				um = 1;
    82			}
    83			else if (!strcmp(*argv, "-device")) {
    84				if (--argc) {
    85					device = atoi(*++argv);
    86				}
    87			}
    88			else if (argc == 2) {
    89				n = atoi(*argv);
    90			}
    91			else if (argc == 1) {
    92				nloop = atoi(*argv);
    93			}
    94		}
    95	
    96		// GPU info and set device
    97		if (gpu) {
    98			int ndevice;
    99			cudaDeviceProp prop;
   100			cudaGetDeviceCount(&ndevice);
   101			if (device >= ndevice) device = ndevice - 1;
   102			cudaGetDeviceProperties(&prop, device);
   103			printf("GPU-%d : %s, C.C.%d.%d, U.M.%s\n", device, prop.name, prop.major, prop.minor, (um ? "ON" : "OFF"));
   104			cudaSetDevice(device);
   105		}
   106	
   107		// alloc device or managed memory
   108		size_t size = n * sizeof(float);
   109		cuda_malloc(gpu, um, (void **)&a, size);
   110		cuda_malloc(gpu, um, (void **)&b, size);
   111		cuda_malloc(gpu, um, (void **)&c, size);
   112	
   113		// alloc host memory
   114		float *h_a = (float *)malloc(size);
   115		float *h_b = (float *)malloc(size);
   116	
   117		// setup problem
   118		for (int i = 0; i < n; i++) {
   119			h_a[i] = i;
   120			h_b[i] = i + 1;
   121		}
   122	
   123		// copy host to device
   124		cuda_memcpy(gpu, a, h_a, size, cudaMemcpyHostToDevice);
   125		cuda_memcpy(gpu, b, h_b, size, cudaMemcpyHostToDevice);
   126	
   127		// timer
   128		t0 = clock();
   129	
   130		// calculation
   131		for (int loop = 0; loop < nloop; loop++) {
   132			vadd(gpu, n, a, b, c);
   133		}
   134		if (gpu) cudaDeviceSynchronize();
   135	
   136		// timer
   137		t1 = clock();
   138	
   139		// copy device to host
   140		float *h_c = (float *)malloc(size);
   141		cuda_memcpy(gpu, h_c, c, size, cudaMemcpyDeviceToHost);
   142	
   143		// sum
   144		double sum = 0;
   145		for (int i = 0; i < n; i++) {
   146			sum += h_c[i];
   147		}
   148	
   149		// output
   150		double exact = (double)n * n;
   151		double sec = (double)(t1 - t0) / CLOCKS_PER_SEC;
   152		printf("nvector=%d nloop=%d %e(%e) %s[sec]=%.3f\n",
   153			n, nloop, sum, exact, (gpu ? "GPU" : "CPU"), sec);
   154	
   155		// free
   156		free(h_a);
   157		free(h_b);
   158		free(h_c);
   159		cuda_free(gpu, a);
   160		cuda_free(gpu, b);
   161		cuda_free(gpu, c);
   162	
   163		return 0;
   164	}