cuda基础
一:cuda编程模型
1:主机与设备
主机---CPU 设备/处理器---GPU
CUDA编程模型如下:
GPU多层存储空间结构如图:
2:Kernel函数的定义与调用
A:运行在GPU上,必须通过__global__函数类型限定符定义且只能在主机端代码中调用;
B:在调用时必须声明内核函数的执行参数----<<<>>>。
C:先为内核函数中用到的变量分配好足够空间再调用kernel函数
D:每个线程都有自己对应的id----由设备端的寄存器提供的内建变量保存,且是只读的。
3:线程结构
1)线程标识
dim3类型(基于uint3定义的矢量类型----由三个unsigned int组成的结构体)的内建变量threadIdx和blockIdx。
2)一维block
线程threadID----threadIdx.x.
3)二维block---(Dx,Dy)
线程threadID----threadIdx.x+threadIdx.y*Dx;
4)三维block---(Dx,Dy,Dz)
线程threadID----threadIdx.x+threadIdx.y*Dx+threadIdx.z*Dx*Dy;
4:硬件映射
1)计算单元
SM---流多处理器 SP---流处理器
A:一个SM包含8个SP,共用一块共享存储器
2)warp
线程束在采用Tesla架构的gpu中:一个线程束由32个线程组成,且其线程只和threadID有关
A:warp才是真正的执行单位
3)执行模型
SIMT---单指令多线程 SIMD---单指令多数据
4)deviceQuery实例
#include <stalib.h>
#include <stdio.h>
#include<string.h>
#include <cutil.h> int main()
{
int deviceCount;
CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));
if( == deviceCount)
{
printf("no deice\n");
}
int dev;
for(dev = ;dev <deviceCount;dev++)
{
cudaDeviceProp deviceProp;
CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp,dev));
print();
}
}
5)cuda程序编写流程
A:主机端
启动CUDA,使用多卡时需加上设备号,或使用cudaSetDevice()设置
为输入数据分配空间
初始化输入数据
为GPU分配显存,用于存放输入数据
将内存中的输入数据拷贝到显存
为GPU分配显存,用于存放输出数据
调用device端的kernel进行计算,将结果写到显存中对应区域
为CPU分配内存,用于存放GPU传回来的输出数据
使用CPU对数据进行其他处理
释放内存和显存空间
退出CUDA
B:设备端
从显存读数据到GPU片内 对数据进行处理 将处理后的数据写回显存
(1)在显存全局内存分配线性空间--cudaMalloc()/cudaFree()
(2)拷贝存储器中的数据 --cudaMemcpy()
拷贝操作类型:cudaMemcpyDeiceToHost cudaMemcpyHostToDevice cudaMemcpyDeviceToDevice
(3)网格定义
<<<Dg,Db,Ns,S>>>
Dg----grid纬度与尺寸 Db---block维度与尺寸 Ns--可分配动态共享内存大小 s--stream_t类型的可选参数
(4)设备端内建变量
gridDim blockIdx blockDim threadIdx warpSize
6)内核实例
A:与shared memory有关
__global__ void
testKernel(float* g_idata,float* g_odata)
{
//分配共享内存 将全局内存的数据写入共享内存 进行计算,将结果写入共享内存 将结果写回全局内存
extern __shared__ float sdata[];//动态分配共享内存空间--__device__ __global__函数中
//动态分配大小是执行参数中的第三个参数。当静态分配时必须指明大小 const unsigned int bid = blockIdx.x;
const unsigned int tid_in_block = threadIdx.x;
const unsigned int tid_in_grid = blockIdx.x*blockDim.x+threadIdx.x;
sdata[tid_in_block] = g_idata[tid_in_grid];
__syncthreads(); sdata[tid_in_block] *= (float)bid; __syncthreads(); g_odata[tid_in_grid] = sdata[tid_in_block];
}
最新文章
- AngularJS之ng-class(十一)
- jquery实现表格的搜索功能
- Linux命令(ntp)
- java笔记--反射机制之基础总结与详解
- 隐藏gvim中的工具栏和菜单栏
- Git CMD - push: Update remote refs along with associated objects
- Socket 理解
- 两台linux机器文件传输之scp
- Eclipse插件基础篇一
- Linux 下 Error: Could not find or load main class Hello
- Quartz定时调度
- Linux基础-最基础
- Python3+Selenium2完整的自动化测试实现之旅(一):自动化测试环境搭建
- Python虚拟环境的安装和配置-virtualenv与windows下多个python版本共存
- mysql 中表和数据库名称不要使用 &#39;-&#39; 命名
- wamp升级php5.3.10到5.4.31版本
- jmeter正则表达式提取器多模块相互调用
- javascript 实例 静态 公共 私有
- 汇编 循环位移指令 ROL, 循环位移指令 ROR
- photoshop cc 2018破解补丁(pscc2018注册机) 附使用方法