一 CPU和GPU的区别

在这里插入图片描述

衡量处理器优劣的重要的两个指标:
延时性:同量的数据,所需要的处理时间
吞吐性:处理速度不快,但是每次处理量很大
GPU设计理念是最大化吞吐量,使用很小的控制单元对应很小的内存
cpu的设计理念:强大的控制单元和缓存单元

LocalCache将数据存储在内存中,因此可以快速访问缓存数据,提高应用程序的性能。LocalCache适用于许多不同的使用场景,特别是以下几种情况:
1.频繁访问的数据:如果应用程序中有一些频繁被访问的数据,将这些数据存储在LocalCache中可以大大提高访问速度。由于数据存储在内存中,相比于从磁盘或网络中读取数据,从本地缓存中获取数据的速度更快
在这里插入图片描述

CPU的理念:
![在这里插入图片描述](https://img-blog.csdnimg.cn/direct/94f03ea6b1734a60a460f5aa4b79e4e4.png在这里插入图片描述
在这里插入图片描述

在这里插入图片描述

二 cuda编程

1 cuda运行效率

h开头表示是在cpu上的数据,
d开头表示在GPU上的数据
在这里插入图片描述
线程块是GPU运行的最小单元,线程块里面的东西不好会被打散
在这里插入图片描述
在这里插入图片描述
数据读取速度的比较,基于此,如果读取的数据很少,但是处理的数据的计算量很大,这种情况的计算效率就是最高的
在这里插入图片描述在这里插入图片描述
合并全局内存:推荐的数据读取方式是连读位置读取数据,速度是最快的;数据的存储也是连续位置
在这里插入图片描述
避免线程发散:
(1)避免同一个线程块或者kernel中执行不同的代码,因为不同代码之间的运行所需时间是不一样的,但是线程块又要求同步机制,程序之间彼此会相互等待,会影响程序效率
(2)循环长度不一致
在这里插入图片描述在这里插入图片描述

2 理解同步性

前面提到,Block里面的线程运行是同步的,不同kernel各自运行是同步的,kenelA内的东西全部运行结束之后才会运行kernel B,Block的线程同步是需要使用者自己来控制,kernel上的同步运行是GPU自己来实现的
在这里插入图片描述
在这里插入图片描述

3 硬件和软件之间的对应关系

(1)SM是流处理器

kernel:核,可以理解为c/c++里面的一个函数function,紫色的就是一个个线程
一个GPU包含若干SM
在这里插入图片描述

4 kernel的不同加载方式

首先需要清楚自己电脑的性能,知道自己电脑的能承受的最大线程数,每个处理器的最大线程数,每个block的最大线程数,
也可以使用代码,自己去获取计算配置参数

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

5 cuda如何调用内存

(1)调用局部变量,也就是本地变量

在这里插入图片描述

(2)调用全局变量

都是传递指针变量
前半段是自定义
后半段是使用,使用的时候,host表示主机,CPU,device表示GPU,使用的时候需要在GPU上先分配呢内存,然后将主机上的数据拷贝到GPU上,然后加载核函数进行运算,最后再将GPU上运算完毕的数据拷贝回CPU
在这里插入图片描述

(3)共享变量的调用

共享变量是每一个线程都在往里面的指定位置写数据,存在一个时间差的问题,就需要加同步机制
在这里插入图片描述

共享数据,三个参数,的含义依次如下
在这里插入图片描述

(4)同步操作----原子操作,有些系统不支持,需要自己去实现

在这里插入图片描述

(4)同步操作---- 同步函数: _syncthreads()
(5)同步操作---- 同步函数: _threadfence() 以及——threadfence_block()

在这里插入图片描述

(5)同步操作---- 同步函数: CPU和GPU同步

主要在主代码里面使用
在这里插入图片描述

在这里插入图片描述
在这里插入图片描述

案例1 计算1-7的各自平方


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>



__global__ void square(float* d_out, float* d_in)
{
	int idx = threadIdx.x;
	float f = d_in[idx];
	d_out[idx] = f * f;
}


int main(int argc, char** argv)
{
	const int ARRAY_SIZE = 8;
	const int ARRAY_BITE = ARRAY_SIZE * sizeof(float);

	//**************************在cpu上输入数据 *********************
	float h_in[ARRAY_SIZE];  //h_in表示存在cpu上的输入数据, 
	for (int i = 0; i < ARRAY_SIZE; i++)
	{
		h_in[i] = float(i);
	}
	float h_out[ARRAY_SIZE]; //h_out用来存储GPU处理完之后传回cpu上的数据

	//**************************在GPU上分配内容空间,空间和CPU上数据空间大小一样 *********************
	float* d_in;      //用来存储从CPU传到GPU上的待处理数据
	float* d_out;     //用来存储GPU最终处理完的数据
	cudaMalloc((void**)&d_in, ARRAY_BITE);
	cudaMalloc((void**)&d_out, ARRAY_BITE);

	**************************把cpu上的待处理数据复制到GPU上 *********************
	cudaMemcpy(d_in, h_in, ARRAY_BITE, cudaMemcpyHostToDevice);

	//*****启动加载kernel到GPU,所谓kernel,在这里就是square函数,这个函数是要在GPU上进行实现的 *********************
	square << <1, ARRAY_SIZE >> > (d_out, d_in);  //使用1个线程块,每个线程块有 ARRAY_SIZE个线程,(d_out, d_in)是函数square的输入参数

	//**************************将在GPU上处理完的数据复制到CPU上 *********************
	cudaMemcpy(h_out, d_out, ARRAY_BITE, cudaMemcpyDeviceToHost);

	for (int i = 0; i < ARRAY_SIZE; i++)
	{
		printf("%f", h_out[i]);
		printf("%\n");
	}

	cudaFree(d_in);
	cudaFree(d_out);

	return 0;
}

在这里插入图片描述

案例2 归约,求和运算,对1+2+3+…

归约:多个输入但是最后只有一个输出

**
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
//#include <device_functions.h>

using namespace std;
__global__ void global_reduce_kernel(float* d_out, float* d_in)
{
    int myId = threadIdx.x + blockDim.x * blockIdx.x;  //有多个block.在所有线程里面的位置,全局变量的位置
    int tid = threadIdx.x; //在当前block里面的位置

    // do reduction in global mem
    //	//循环的条件是 s 大于 0,每次循环迭代结束后,s 的值右移一位(相当于将其除以 2),以便在下一次迭代中处理更小的数据块
//	//这里是实现每一个block里面线程的对半累加
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
        if (tid < s)
        {
            d_in[myId] += d_in[myId + s];
        }
        __syncthreads();        // make sure all adds at one stage are done! //每个线程块每次循环结束之后才进行下一次运行
    }

    //最后要的是线程块的所有计算的求和结果
    // only thread 0 writes result for this block back to global mem
    if (tid == 0)   //等于0说明当前block的所有线程对折运算结束
    {
        d_out[blockIdx.x] = d_in[myId];
    }
}

使用共享内存实现归约
__global__ void shmem_reduce_kernel(float* d_out, const float* d_in)
{
    // sdata is allocated in the kernel call: 3rd arg to <<<b, t, shmem>>>
    extern __shared__ float sdata[];  //给每一个block定义他们各自的内存;

    int myId = threadIdx.x + blockDim.x * blockIdx.x;   //有多个block.在所有线程的位置
    int tid = threadIdx.x;

    // load shared mem from global mem
    //	//然后将全局内存的数据复制给共享内存, 之后就不需要再从全局内存读数据,
//	//共享数据读取的速度快于全局内存
    sdata[tid] = d_in[myId];
    __syncthreads();            // make sure entire block is loaded!

    // do reduction in shared mem
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
        if (tid < s)
        {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();        // make sure all adds at one stage are done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
        d_out[blockIdx.x] = sdata[0];
    }
}

void reduce(float* d_out, float* d_intermediate, float* d_in,
    int size, bool usesSharedMemory)
{
    // assumes that size is not greater than maxThreadsPerBlock^2
    // and that size is a multiple of maxThreadsPerBlock
    const int maxThreadsPerBlock = 1024;
    int threads = maxThreadsPerBlock;
    int blocks = size / maxThreadsPerBlock;
    if (usesSharedMemory)
    {
        shmem_reduce_kernel << <blocks, threads, threads * sizeof(float) >> >
            (d_intermediate, d_in);
    }
    else
    {
        global_reduce_kernel << <blocks, threads >> >
            (d_intermediate, d_in);
    }
    // now we're down to one block left, so reduce it
    //	//这里的blocks设为1使因为,假设有1024个线程块,每个线程块有1024个线程,之前在每个线程块里面完成了各自的对半加,最后每个线程块将得到
//	//一个数据,总的得到1024个数据,刚好放置在一个block里面完成最后的对半累加
    threads = blocks; // launch one thread for each block in prev step
    blocks = 1;
    if (usesSharedMemory)
    {
        shmem_reduce_kernel << <blocks, threads, threads * sizeof(float) >> >
            (d_out, d_intermediate);
    }
    else
    {
        global_reduce_kernel << <blocks, threads >> >
            (d_out, d_intermediate);
    }
}

int main(int argc, char** argv)
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0) {
        fprintf(stderr, "error: no devices supporting CUDA.\n");
        exit(EXIT_FAILURE);
    }
    int dev = 0;
    cudaSetDevice(dev);

    cudaDeviceProp devProps;
    if (cudaGetDeviceProperties(&devProps, dev) == 0)
    {
        printf("Using device %d:\n", dev);
        printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
            devProps.name, (int)devProps.totalGlobalMem,
            (int)devProps.major, (int)devProps.minor,
            (int)devProps.clockRate);
    }

    const int ARRAY_SIZE = 1 << 10;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

    // generate the input array on the host
    float h_in[ARRAY_SIZE];
    float sum = 0.0f;
    for (int i = 0; i < ARRAY_SIZE; i++) {
        // generate random float in [-1.0f, 1.0f]
        h_in[i] = -1.0f + (float)rand() / ((float)RAND_MAX / 2.0f);
        printf("%f ", h_in[i]);
        sum += h_in[i];
    }

    // declare GPU memory pointers
    float* d_in, * d_intermediate, * d_out;

    //	//在GPU上创建存储待处理数据和处理结果的数组
    // allocate GPU memory
    cudaMalloc((void**)&d_in, ARRAY_BYTES);
    cudaMalloc((void**)&d_intermediate, ARRAY_BYTES); // overallocated
    cudaMalloc((void**)&d_out, sizeof(float));

    //	//将CPU输入数据输入到GPU
    // transfer the input array to the GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);

    int whichKernel = 0;
    if (argc == 2) {
        whichKernel = atoi(argv[1]);
    }

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    // launch the kernel
    switch (whichKernel) {
    case 0:
        printf("Running global reduce\n");
        cudaEventRecord(start, 0);
        for (int i = 0; i < 100; i++)
        {
            reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, false);
        }
        cudaEventRecord(stop, 0);
        break;
    case 1:
        printf("Running reduce with shared mem\n");
        cudaEventRecord(start, 0);
        for (int i = 0; i < 100; i++)
        {
            reduce(d_out, d_intermediate, d_in, ARRAY_SIZE, true);
        }
        cudaEventRecord(stop, 0);
        break;
    default:
        fprintf(stderr, "error: ran no kernel\n");
        exit(EXIT_FAILURE);
    }

    cudaEventSynchronize(stop);
    float elapsedTime;
    cudaEventElapsedTime(&elapsedTime, start, stop);
    elapsedTime /= 100.0f;      // 100 trials

    // copy back the sum from GPU
    float h_out;
    cudaMemcpy(&h_out, d_out, sizeof(float), cudaMemcpyDeviceToHost);

    printf("average time elapsed: %f\n", elapsedTime);

    // free GPU memory allocation
    cudaFree(d_in);
    cudaFree(d_intermediate);
    cudaFree(d_out);

    return 0;
}
**

案例3 扫描并行,也是实现求和运算,但是使用场景不同,使用场景是每输入一个数,完成当数和之前所有数据的累加 ??

算法思想:从n = 0,1,2,…依次以0,2^n 然后为间隔做累加,如下图,第一次累加是以0为间隔,相邻两个数做累加,
0前面没有可以和他直接相加的数,所以直接搬下去,第二次累加是以1为间隔的两个数做累加,0,1,前面没有和累加的数,所以,0,1,直接搬下去
在这里插入图片描述


#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
//#include <device_functions.h>

using namespace std;
__global__ void global_scan_kernel(float* d_out, float* d_in)
{
    int idx = threadIdx.x; //在当前block里面的位置

    float out = 0.00f;
    d_out[idx] = d_in[idx];
    __syncthreads();  //这里同步是希望上面的赋值操作全部完成之后,才能开始下面的操作,这是扫描计算的前提,所有数据都要准备好

    for (int interval = 1; interval < sizeof(d_in); interval *= 2)
    {
        if (idx - interval >= 0)
        {
            out = d_out[idx] + d_out[idx - interval];
        }
        __syncthreads();   //当前所有数据计算完毕,才能进入下一轮的计算,也就是要求每一轮的数据都是准备完毕的

        if (idx - interval >= 0)
        {
            d_out[idx] = out;
            out = 0.00f;
        }
    }
}

__global__ void shmem_scan_kernel(float* d_out, const float* d_in)
{
    // sdata is allocated in the kernel call: 3rd arg to <<<b, t, shmem>>>
    extern __shared__ float sdata[];  //给每一个block定义他们各自的内存;

    int idx = threadIdx.x; //在当前block里面的位置
    sdata[idx] = d_in[idx];
    float out = 0.00f;

    __syncthreads();            // make sure entire block is loaded!
    for (int interval = 1; interval < sizeof(d_in); interval *= 2)
    {
        if (idx - interval >= 0)
        {
            out = sdata[idx] + sdata[idx - interval];
        }

        __syncthreads();   //当前所有数据计算完毕,才能进入下一轮的计算,也就是要求每一轮的数据都是准备完毕的

        if (idx - interval >= 0)
        {
        //这里为什么能直接把out赋值到idx索引,因为在此之前加了同步语句,
       // 也就是多个线程同时运行,同时执行out = sdata[idx] + sdata[idx - interval];语句,
        //然后计算完毕就同时存在多个out,所以才能把各自的out赋值到idx,至于为什么要存到idx索引,
        //是因为算法需要它存到对应的位置
            sdata[idx] = out;
            out = 0.00f;
            //?????????????????????????  为什么打印不了中间过程
            //printf("*************\n");
            //printf("%f  ", sdata[idx]);
            //printf("*************\n");  
        }
    }
    d_out[idx] = sdata[idx];
}


int main(int argc, char** argv)
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0) {
        fprintf(stderr, "error: no devices supporting CUDA.\n");
        exit(EXIT_FAILURE);
    }
    int dev = 0;
    cudaSetDevice(dev);

    cudaDeviceProp devProps;
    if (cudaGetDeviceProperties(&devProps, dev) == 0)
    {
        printf("Using device %d:\n", dev);
        printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
            devProps.name, (int)devProps.totalGlobalMem,
            (int)devProps.major, (int)devProps.minor,
            (int)devProps.clockRate);
    }

    const int ARRAY_SIZE = 8;  //数字1 << 左移,1 << i = pow(2,i)
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

    // generate the input array on the host
    float h_in[ARRAY_SIZE];
    printf("-------------input-----------\n ");
    for (int i = 0; i < ARRAY_SIZE; i++) {
        // generate random float in [-1.0f, 1.0f]
        h_in[i] = i;
        printf("%f  ", h_in[i]);
    }

    // declare GPU memory pointers
    float* d_in, * d_out;

    //在GPU上创建存储待处理数据和处理结果的数组
    // allocate GPU memory
    cudaMalloc((void**)&d_in, ARRAY_BYTES);
    cudaMalloc((void**)&d_out, ARRAY_BYTES);

    //	//将CPU输入数据输入到GPU
    // transfer the input array to the GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);


    cudaEvent_t start, stop;
    cudaEventCreate(&start);  //用于创建事件,用于测量GPU操作的时间
    cudaEventCreate(&stop);
 

    int whichKernel = 1;
    if(!whichKernel)
    //global_scan_kernel << <blocks, threads >> > (d_in);
       global_scan_kernel << <1, ARRAY_SIZE >> > (d_out,d_in);
    else
       shmem_scan_kernel << <1, ARRAY_SIZE,ARRAY_SIZE*sizeof(float) >> > (d_out, d_in);
   
    // copy back the sum from GPU
    float h_out[ARRAY_SIZE];
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);


    cudaEventSynchronize(stop);
    float elapsedTime;
    cudaEventElapsedTime(&elapsedTime, start, stop);
    elapsedTime /= 100.0f;      // 100 trials


    printf("\n");
    if (whichKernel == 0)
    {
        printf("global_scan_kernel\n");
    }
    else
    {
        printf("share_scan_kernel\n");
    }
    printf("-------------out-----------\n "); 
    for (int i = 0; i < ARRAY_SIZE; i++) {
        // generate random float in [-1.0f, 1.0f]
        printf("%f ", h_out[i]);
    }
    printf("-------------out-----------\n ");
    printf("\n");
    printf("average time elapsed: %f\n", elapsedTime);

    // free GPU memory allocation
    cudaFree(d_in);
    cudaFree(d_out);

    return 0;
}





在这里插入图片描述
在这里插入图片描述

案例4 直方图的并行计算

(1)错误1:对数据同时操作读和写 ??

下面这种写法,对于第一个bin来说,假设输入数据是1:66655,数据1和17都会同时读bin里面的初始计数,然后进行累加,正常来说应该是两个数据同时读了bin1的0,然后同时写一个1进去,最后结果总是1,但是可能机器会出问题,导致最多也就写一个2,先后加了1.
在这里插入图片描述

![在这里插入图片描述](https://img-blog.csdnimg.cn/direct/1e4f145d9a4c407fa8828aa467786fc0.png pic_center =700x)

(2)对于(1)的改进,用原子相加,将读写并行做成串口,数据1读写完成之后,17才能进行读写。这样处理的并行化程度是取决于bin的数量的(2)的并行主要在累加的地方,写入的时候是串行,如果bin很多,平摊下来的时间就短,并行化程度就高了。

AtomicAdd就是原子操作,原子操作就相当于各个线程排队操作。 bin的数量就相当于队列的数量,队列多了,肯定就快了,队列多,等待是并行的。
在这里插入图片描述

(3)局部直方图,读写本身就是串行的,

在这里插入图片描述
在这里插入图片描述


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

//假设是8个block,8个线程,也就是每个block要处理16个数据
__global__ void local_histo(float* d_out, float* d_in, const int BIN_COUNYT, const int Threads_count) {
    int idx = threadIdx.x + blockDim.x*blockIdx.x;  //有多个block.在所有线程里面的位置
    int tid = threadIdx.x; //在当前block里面的位置
    //总共128个数,分成3个bin,但是先要分成8组同时用8个线程来对数据进行处理,也就是说每个线程处理16个数据,处理的结果最后各自形成各自的3个bin,
    //d_out里面,相连的三个数据认为是一组不重复且不连续的16个数的统计直方图,依次存储
    for (int interpre = 1; interpre < sizeof(d_in)/ Threads_count; interpre +=8)
    {
        int binID = d_in[idx] % BIN_COUNYT;
        d_out[tid] ++;
    }
}

int main(int argc, char** argv) {
    //**********************************
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0) {
        fprintf(stderr, "error: no devices supporting CUDA.\n");
        exit(EXIT_FAILURE);
    }
    int dev = 0;
    cudaSetDevice(dev);

    cudaDeviceProp devProps;
    if (cudaGetDeviceProperties(&devProps, dev) == 0)
    {
        printf("Using device %d:\n", dev);
        printf("%s; global mem: %dB; compute v%d.%d; clock: %d kHz\n",
            devProps.name, (int)devProps.totalGlobalMem,
            (int)devProps.major, (int)devProps.minor,
            (int)devProps.clockRate);
    }
    //**********************************

    const int ARRAY_SIZE = 128;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

    // generate the input array on the host
    float h_in[ARRAY_SIZE];
    for (int i = 0; i < ARRAY_SIZE; i++) {
        h_in[i] = float(i);
    }
    float h_out[ARRAY_SIZE];

    // declare GPU memory pointers
    float* d_in;
    float* d_out;

    // allocate GPU memory
    cudaMalloc((void**)&d_in, ARRAY_BYTES);
    cudaMalloc((void**)&d_out, ARRAY_BYTES);

    // transfer the array to GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);

    // launch the kernel
    int USED_threadNum = 8;
    int USED_BLOCK = 8;
    /*local_histo << <1, ARRAY_SIZE >> > (d_out, d_in);*/
    local_histo << < USED_BLOCK, threadNum >> > (d_out, d_in, 3, USED_threadNum)

    // copy back the result array to the GPU
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);

    // print out the resulting array
    for (int i = 0; i < ARRAY_SIZE; i++) {
        printf("%f", h_out[i]);
        printf(((i % 4) != 3) ? "\t" : "\n");
    }

    // free GPU memory allocation
    cudaFree(d_in);
    cudaFree(d_out);

    return 0;


}

案例5 压缩与分配

在这里插入图片描述

稀疏型:52个线程同时运行,进行判读,不满足条件的就终止运算了,也就是说52个线程的运行时间是不相同的,有些线程中途会被闲置
密集型:先52个线程一起判断是不是方块;然后再用13个线程完成卡片的计算
在这里插入图片描述

案例6 分段扫描

对每个数据段里面的数据,实现各自的扫描
在这里插入图片描述

分段扫描的应用案例:稀疏矩阵的处理
在这里插入图片描述

案例7 排序

在这里插入图片描述
归并排序
在这里插入图片描述在这里插入图片描述

点赞(0) 打赏

评论列表 共有 0 条评论

暂无评论

微信公众账号

微信扫一扫加关注

发表
评论
返回
顶部