cuda编程知识普及

本帖经过多方整理,大多来自各路书籍《GPGPU编程技术》《cuda高性能》

1 gridblock都可以用三元向量来表示:

grid的数组元素是block

  block的数组元素是grid

但是1.x计算能力的核心,grid的第三元必须为1.block的X和Y索引最大尺寸为512

2 通过__launch_bounds__(maxBlockSize,minBlocksPerMp)来限制每个block中最大的线程数,及每个多处理器上最少被激活的block数

3 SM streaming multiprocessor 多流处理器

SP scalar processor cores 标量处理核心

一个Block中的所有线程在一个多处理器上面并发执行。当这个Block的所有线程执行完后,再激活其他等待的Block.一个多处理器上也可以执行多个block。但是一个block却不能拆分为多个处理器上面执行

对于同一个Block里面的线程:

    1 同一个Block里的线程可以被同步

    2 可以共同访问多处理器里的共享存储器

到2.x为止,多处理器 执行任务时,以32个并行线程为单位,称为一个wrap。

当以个block到来的时候,会被分成线程号连续的多个wrap,然后多处理器上的SIMT控制器以wrap为单位控制调度线程。所以block中的线程数要是以32的整数倍来设计,就不会出现空闲的SP。组织WARP的时候,从线程号最小的开始

4 各个存储器存储位置及作用 

5 寄存器放在SP中,如果溢出,会被放在设备处理器上面,发生严重滞后,影响性能。

1.0   4KB 
2.0   16kb

 共享存储器位于SM中,大约两个时钟周期读写4B,静态分配 __shared__ int shared[16];

1.0   16KB 
2.0   48kb

6 共享存储器,是以4个字节为单位的16个存储器组

  bank冲突:半个warp中的多线程访问的数组元素处于同一个bank时,访问串行化,发生冲突

  避免冲突:最多的数据类型是int、float等占用4个字节的类型

7线程设计

float shared=data[base+tid];
base访问的起始元素下标 tid线程号

如果要是char类型,每个元素占1个字节,就会冲突

float shared = data[base+4*tid];

8 共享存储器广播访问:半个warp线程都访问一个数据

9 补白策略

    shared[tid]=global[tid];
 
    int number = shared[tid*16];
    int nRow = tid/16;
    int nColumn = tid%16;
    shared[nColumn*17+nRow] = global[tid];
 
    int number = shared[17*tid];

10 一次性访问全局存储器:数据的起始地址应为每个线程访问数据大小的16倍的整数倍

11 主机锁页存储器cudaHostMalloc()分配。

不参与操作系统分页管理的存储空间,访问锁页文件不会耗费主机内存分页管理方面的开销。不会被操作系统放到硬盘的页面文件中,因此比访问普通的主机存储器更快。

12 计算能力2.x的GPU上面,每个SM有独立的一级缓存,有唯一的二级缓存

13 异步并发

主机上的计算、

设备上的计算、

主机到设备上的传输、

设备到主机上的传输共同执行

14 设备存储器 类型是DRAM,动态随机存储器。使用它最高效的方式就是顺序读取。为了保证顺序:

__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
    const int tid = threadIdx.x;
    int nSum = 0;
    int i;
    clock_t clock_tStart;
    if(tid == 0) clock_tStart = clock();
 
    for(i = tid;i<DATA_SIZE;i+=THREAD_NUM){
        nSum += pnNumber[i]*pnNumber[i];
}
 
    pnResult[tid] = nSum;
    if(tid == 0)
        *pclock_tTime = clock()-clock_tStart;
}

每个block 在1.x的计算能力的GPU下,最多只有512的线程数

__global__ static void sumof(int *pnNumber,int* pnResult,clock_t* pclock_tTime){
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    int nSum = 0;
    int i;
    clock_t clock_tStart;
    if(tid == 0) pclock_tTime[bid] = clock();
 
    for(i = bid*THREAD_NUM+tid;i<DATA_SIZE;i+=BLOCK_NUM*THREAD_NUM){
        nSum += pnNumber[i]*pnNumber[i];
}
 
    pnResult[bid*THREAD_NUM+tid] = nSum;
 
    if(tid == 0)
        *pclock_tTime[bid+BLOCK_NUM] = clock();
}

15 用缩减树避免bank冲突:

  bank冲突指的是,一个warp内的线程同时访问一个bank列,导致串行读取数据

    noffset = THREAD_NUM/2;
    while(noffset > 0){
        if(tid < offset)
            nshared[tid] += nshared[tid+noffset];
    }
    noffset >>= 1;
 
    __syncthreads();

16 CPU有强大的分支预测、程序堆栈、循环优化等针对控制采取的复杂逻辑。

    GPU相对简单,适合处理顺序的,单一的,少循环,少跳转的语句。

17  #progma unroll 5下面的程序循环5次

18 cuda中的同步

1》__syncthreads()同步

  同一个warp内的线程总是被一同激活且一同被分配任务,因此不需要同步。因此最好把需要同步的线程放在同一个warp内,这样就减少了__syncthreads()的指令

2》__threadfence() __threadfence_block()同步

  前者针对grid的所有线程,后者针对block内的所有线程。告知线程,全局存储器或共享存储器已经被改变

3》cudaThreadSynchronize() 主机与设备间的同步

  在主机程序里同步线程。该函数以上的设备线程完成后,控制权才交给cpu

4》volatile关键字

  使用这个关键字定义数组,设备会知道这个数组随时都会改变,就会自动重新读取数组(但是不能保证线程间读取的数据一致)

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

我来说两句

0 条评论
登录 后参与评论

相关文章

  • cuda by example

    int offset= x+y*dim x 线程块内的线程索引 y 线程块索引 dim 线程块的维度 tid = threadIdx.x+blockIdx.x*...

    用户1154259
  • 最大子段和

    最大子段和:给出一个数组,计算其中连续的最大的子段和 运行代码,及运行思想: /** * 动态规划:计算最大子段和 * 算法描述: * 数组a 有n个元素...

    用户1154259
  • RequireJS实例分析

      随着JS越来越庞大,已经不仅仅是以前复制粘贴做特效的时代了,JS越来越偏向于业务逻辑与应用。恰逢Node的流行,JS在web开发中占有越来越重要的地位。由...

    用户1154259
  • apache开启重写模式

    现在的好多的框架都使用有路由机制,但是如果在apache下,没有开启重写模式,服务器不会读取路由

    lin_zone
  • 如何在Ubuntu 14.04上使用Fail2Ban保护Apache服务器

    在操作Web服务器时,必须实施安全措施来保护您的站点和用户。使用防火墙策略保护您的网站和应用程序并使用密码身份验证限制对某些区域的访问是保护系统安全的一个很好的...

    丰一川
  • 小程序添加激励视频广告演示代码详解

    这里官方没有给任何dome,所以很多对js不熟悉的朋友很迷惑,但是仔细看看官方文档,其实写的很详细了,下面我直接展示我得dome

    许坏
  • Java基础:一、每个对象都提供服务(3)

    在开发或理解一个程序设计时,最好的方法之一就是将对象想象为“服务的提供者”,程序本身将向用户提供服务,它将通过调用其他对象提供的服务来实现这一目的。

    桑鱼
  • phpmyadmin新姿势getshell

    在一个有WAF、并且mysql中的Into outfile禁用的情况下,我该如何getshell? 首先环境如下: OS:Windows 2003 ...

    奶糖味的代言
  • ASP.NET MVC使用Bootstrap系列(1)——开始使用Bootstrap

    作为一名Web开发者而言,如果不借助任何前端框架,从零开始使用HTML和CSS来构建友好的页面是非常困难的。特别是对于Windows Form的开发者而言,更...

    用户1161731
  • 性能工具之linux三剑客awk、grep、sed详解

    linux 有很多工具可以做文本处理,例如:sort, cut, split, join, paste, comm, uniq, column, rev, ta...

    高楼Zee

扫码关注云+社区

领取腾讯云代金券