C/C++教程

使用CUDA C进行并行编程

本文主要是介绍使用CUDA C进行并行编程,对大家解决编程问题具有一定的参考价值,需要的程序猿们随着小编来一起学习吧!

1、CUDA程序结构

  CUDA程序是在主机或者设备上执行的函数的组合。不显示并行性的函数在CPU上执行,显示数据并行性的函数在GPU上执行,GPU在编译期间要隔离这些函数。CUDA代码基本上与C代码相同,只是添加了一些开发数据并行性所需的关键字

2、CUDA C中的双变量加法程序

  编写一个将两个变量相加的内核(kernel)函数

#include<iostream>
#include<cuda.h>
#include<cuda_runtime.h>
__global__ void gpuAdd(int d_a, int d_b, int* d_c) {
    *d_c = d_a + d_b;
}

  main函数代码

int main(void) {
    //定义主机变量,统一以h_开头
    int h_c;
    //定义设备变量,统一以d_开头,作为参数传递给内核的指针应该仅指向设备显存
    int* d_c;
    //在设备上分配d_c变量的内存,cudaMalloc类似于C中的malloc函数
    //分配了一个整数变量大小的内存块,并返回指向该内存位置的指针d_c
    cudaMalloc((void**)&d_c, sizeof(int));
    //调用gpuAdd,其中1和4是两个输入变量,d_c是一个作为输出指针变量的设备显存指针
    //
    gpuAdd << <1, 1 >> > (1, 4, d_c);
    //将结果从设备内存复制到主机内存中
    cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
    //释放设备上通过cudaMalloc使用的所有内存
    cudaFree(d_c);
    return 0;
}

  需要注意的是,作为参数传递给内核的指针应该仅指向设备显存,如果它指向主机内存,会导致程序崩溃。其次,当通过对调用的内核传递参数时i,所有的指针仅指向设备显存。

3、内核调用

  使用ASNSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码(Host Code)通过内核调用的方式来启动的。简单地说,内核调用的含义是我们从主机代码启动设备代码。内核调用通常会产生大量的块(Block)和线程(Thread)来在GPU上并行地处理数据。内核使用__global__关键字定义,使用kernel<< <希望执行的块数, 每个块将具有的线程数, 内核使用的共享内存大小>> >来启动内核。在变量相加程序中,内核启动的语法如下:

gpuAdd << <1, 1 >> > (1, 4, d_c)

4、CUDA中的向量运算

       上面的双变量加法没有利用到GPU设备的并行处理能力,下面利用GPU的并行结构修改CPU上的向量加法程序,首先在CPU上编写向量加法程序

4.1、CPU向量加法

#include "stdio.h"
#include<iostream>
//定义向量元素个数
#define N    5
//定义CPU向量加法
void cpuAdd(int *h_a, int *h_b, int *h_c) {
    int tid = 0;    
    while (tid < N)
    {
        h_c[tid] = h_a[tid] + h_b[tid];
        tid += 1;
    }
}

int main(void) {
    int h_a[N], h_b[N], h_c[N];
        //向量初始化
    for (int i = 0; i < N; i++) {
        h_a[i] = 2 * i*i;
        h_b[i] = i;
    }
    //调用CPU向量加法函数
    cpuAdd (h_a, h_b, h_c);
    //输出结果
    printf("Vector addition on CPU\n");
    for (int i = 0; i < N; i++) {
        printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    return 0;
}

4.2、GPU向量加法

       在GPU上添加修改后的内核函数

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//定义向量元素个数
#define N 5
//定义向量加法内核函数
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
    //获取当前内核的块ID
    int tid = blockIdx.x;    //在当前块ID来初始化tid
    //因为是同时启动的,每个块必须直到自己要处理的数据是什么
    if (tid < N)
        d_c[tid] = d_a[tid] + d_b[tid];
}

int main(void) {
    //定义主机变量
    int h_a[N], h_b[N], h_c[N];
    //定义设备指针,指向cudaMalloc分配的显存
    int *d_a, *d_b, *d_c;
    //给设备变量分配内存
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));
    //主机向量初始化
    for (int i = 0; i < N; i++) {
        h_a[i] = 2*i*i;
        h_b[i] = i ;
    }
    //将主机向量复制到设备向量中
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
    //使用内核启动符号<<<>>>调用内核函数,启动N个块(和向量长度一样,使用不同的块并行处理向量N个位置的加法),每个块1个线程,传递参数为设备指针
    gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
    //将结果从设备内存复制到主机内存中
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    printf("Vector addition on GPU \n");
    //在控制台中输出结果
    for (int i = 0; i < N; i++) {
        printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    //释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    return 0;
}

       在main函数中调用内核函数时,并行启动了N个块,这意味着我们同时启动了N个执行该内核代码的线程副本。因为N个块是同时启动的,只有每个块分工合作效率才会高,main函数调用gpuAdd内核函数时,启动的块是随机的,但可以通过blockIdx.x查看启动的块的ID,用该ID作为块处理的数据对应的ID,这样就可以让N个块分别处理向量中N个数据的加法。GPU的并行计算的方式比CPU的串行计算提高了吞吐量。

4.3、GPU对向量每个元素进行平方

       上面是并行启动N个块,每个块只有一个线程。也可以在1个块里面并行启动N个线程。

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//定义向量元素个数
#define N    5
//平方内核函数
__global__ void gpuSquare(float *d_in, float *d_out) {
    //获取当前调用内核函数启动的线程ID
    int tid = threadIdx.x;    //使用内核函数计算线程ID对应的向量中的元素
    float temp = d_in[tid];
    d_out[tid] = temp*temp;
}

int main(void) {
    //定义主机向量
    float h_in[N], h_out[N];
    //定义设备指针
    float *d_in, *d_out;

    //在CPU上分配内存
    cudaMalloc((void**)&d_in, N * sizeof(float));
    cudaMalloc((void**)&d_out, N * sizeof(float));
    //Initializing Array
    for (int i = 0; i < N; i++) {
        h_in[i] = i;
    }
    //将主机上的数组复制到设备上
    cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
    //调用1个块,N个线程的内核函数
    gpuSquare << <1, N >> >(d_in, d_out);
    //将设备上的结果复制到主机内存上
    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    //控制台输出结果
    printf("Square of Number on GPU \n");
    for (int i = 0; i < N; i++) {
        printf("The square of %f is %f\n", h_in[i], h_out[i]);
    }
    //释放内存
    cudaFree(d_in);
    cudaFree(d_out);
    return 0;
}

       当使用这种方式启动N个线程时,每个块的最大线程数不超过512或1024,如果N=2000,而GPU的线程最大数量为512,那么不能写成<<<1, 2000>>>,而应该用<<<4, 500>>>。应当选择合适的块和线程的数量

5、并行通信模式

       当多个线程并行执行时,他们遵循一定的通信模式,直到它们在显存哪里输入,哪里输出,通讯模式包括映射、收集、分散式、蒙版、转置。

5.1、映射(Map)

       单一输入,单一输出。Map的代码模式看起来如下:

d_out[i] = d_in[i] * 2;

5.2、收集(Gather)

      多个输入,单个输出。Gather模式的代码看起来如下:

out[i] = (in[i - 1] + in[i] + in[i + 1]) / 3;

5.3、分散式(Scatter)

      单一输入,多个输出。Scatter模式的代码看起来如下:

out[i - 1] += 2 * in[i] and out[i + 1] += 3 * in[i];

5.4、蒙版(Stencil)

      从数组中读取固定形状的相邻元素时,称为stencil模式。比如3X3或者5X5卷积就是Stencil,它是Gather操作的一种模式,代码和Gather相似。

5.5、转置(Transpose)

      对矩阵进行转换。

       

这篇关于使用CUDA C进行并行编程的文章就介绍到这儿,希望我们推荐的文章对大家有所帮助,也希望大家多多支持为之网!