CUDA编程学习(2)

P4 向量加法程序解析

#include <stdio.h>
#include <cuda.h>

#include "uax.h"

typedef float FLOAT;

/* host, add */ 
// CPU 上的加法
void vec_add_host(FLOAT *x, FLOAT *y, FLOAT *z, int N);  

/* device function */
// GPU 上的加法
__global__ void vec_add(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
    /* 1D block */
    int idx = get_tid();

    if (idx < N) z[idx] = z[idx] + y[idx] + x[idx];  // N 以内对应IDx的相加
}

void vec_add_host(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
    int i;

    for (i = 0; i < N; i++) z[i] = z[i] + y[i] + x[i];  // 顺序相加
}


#define get_tid() (blockDim.x*(blockIdx.x+blockIdx.y*blockDim.x)+threadIdx.x)

#define get_bid() (blockIdx.x + blockIdx.y*blockDim.x)

#include<sys/time.h>
#include<time.h>

double get_time(void){
    struct timeval tv;
    double t;
    gettimeofday(&tv, (struct timezone *)0);
	t = tv.tv_sec + (double)tv.tv_usec * 1e-6;

	return t;
}


__global__ void warmpup_knl() {  // 热身无意义
	int i, j;
	i = 1;
	j = 2;
	i = i + j;
}

void warmup() {
	int i;
	for (i = 0; i < 8; i++) {  // 热身的 kernel 调用了8次
		warmpup_knl <<<1, 256 >>>();
	}

}

int main()
{
    int N = 20000000;
    int nbytes = N * sizeof(FLOAT);  // sizeof(float) 输出的是float的字节数

    /* 1D block */
    int bs = 256;  // 块的大小

    /* 2D grid */
    int s = ceil(sqrt((N + bs - 1.) / bs));  // 得到网格的大小
    dim3 grid = dim3(s, s);  // s*s个block

    FLOAT *dx = NULL, *hx = NULL;  // 三维向量相加
    FLOAT *dy = NULL, *hy = NULL;
    FLOAT *dz = NULL, *hz = NULL;

    int itr = 30;
    int i;
    double th, td;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);
    cudaMalloc((void **)&dy, nbytes);
    cudaMalloc((void **)&dz, nbytes);

    if (dx == NULL || dy == NULL || dz == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    printf("allocated %.2f MB on GPU\n", nbytes / (1024.f * 1024.f));

    /* alllocate CPU mem */
    hx = (FLOAT *) malloc(nbytes);
    hy = (FLOAT *) malloc(nbytes);
    hz = (FLOAT *) malloc(nbytes);

    if (hx == NULL || hy == NULL || hz == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }
    printf("allocated %.2f MB on CPU\n", nbytes / (1024.f * 1024.f));

    /* init */
    // CPU 初始化
    for (i = 0; i < N; i++) {
        hx[i] = 1;
        hy[i] = 1;
        hz[i] = 1;
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dz, hz, nbytes, cudaMemcpyHostToDevice);

    // warm up
    warmup();

    /* call GPU */
    cudaDeviceSynchronize();  // 等所有的GPU动作执行完

    td = get_time();  // 获取GPU起始时间
    
    for (i = 0; i < itr; i++) vec_add<<<grid, bs>>>(dx, dy, dz, N);  // 调用了30个核函数，由于cuda程序是异构的，所以这里CPU调用是很快的，只是GPU还不一定算完

    cudaDeviceSynchronize();
    td = get_time() - td;  // 获取GPU结束时间

    /* CPU */
    th = get_time();  // 获取CPU起始时间
    for (i = 0; i < itr; i++) vec_add_host(hx, hy, hz, N);
    th = get_time() - th;  // 获取CPU结束时间

    printf("GPU time: %e, CPU time: %e, speedup: %g\n", td, th, th / td);

    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);

    free(hx);
    free(hy);
    free(hz);

    return 0;
}
 

• nvcc 编译指令：

nvcc -O3 -arch=sm_52 -Xcompiler "-Wall -Wextra" -o my-vec-add vec-add.cu 
# arch=sm_52 这个参数是与具体显卡的型号有关系，不同显卡包含的sm数量不一样
# 不同的显卡可以参考： https://blog.youkuaiyun.com/lb1244206405/article/details/90718040  
# -Xcompiler 表示设置编译器相关的信息
# -Wall 表示打印所有的警告
# -Wextra 表示打印警告之外的相关信息

./my-vec-add