CUDA学习2-编程部分

CUDA编程

函数声明

host:主机端,通常指CPU

device:设备端,通常指GPU(数据可并行)

kernel:数据并行处理函数,在主机端调用kernel可以在设备端创建大量轻量级线程

host 和 device拥有各自的存储器

CUDA编程包括主机端和设备端两部分代码

执行位置 调用位置
_device_float DeviceFunc() device device
_global_void KernelFunc() device host
_host_float HostFunc() host host

_global_定义一个kernel函数

  • 入口函数,CPU上调用,GPU上执行

  • 必须返回void

_device_and_host_可以同时使用

CUDA核函数(kernels)

在N个不同的CUDA线程上并行执行

//定义kernel_global_ void VecAdd(float* A, float* B, float* C){int i = threadIdx.x;    C[i] = A[i]   B[i];}int main(){//...    //在N个不同的CUDA线程上并行执行    VecAdd<<<1, N>>>(A, B, C);}

线程层次(Thread Hierarchies)

块索引:blockIdx

维度:blockDim

  • 一维、二维或三维

//单线程块//定义kernel_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){int i = threadIdx.x;    int j = threadIdx.y;    C[i][j] = A[i][j]   B[i][j];}int main(){//...    //在N*N*1个不同的CUDA线程上并行执行    int numBlocks = 1;    dim3 threadsPerBlock(N, N);    VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);}
//多线程块//定义kernel_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){int i = threadIdx.x * blockDim.x   threadIdx.x;    int j = threadIdx.y * blockDim.y   threadIdx.y;    if(i<N && j<N)    C[i][j] = A[i][j]   B[i][j];}int main(){//...    //并行执行    dim3 threadsPerBlock(16, 16);    dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y);    VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);}/*N = 32i = [0,1] * 16   [0,15]*/

CUDA内存传输

主机端可以从设备端往返传输数据

Global memory 全局存储器

Constant memory 常量存储器

cudaMalloc():在设备端分配global memory

cudaFree():释放存储空间

float *Md;//指向设备端上的一个存储空间int size = Width * Width * sizeof(float);cudaMalloc((void**)&Md, size);//...cudaFree(Md);

cudaMemcpy():内存传输

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);//参数:目的地址 源地址 大小 传输方向
  • host to host

  • host to device

  • device to host

  • device to device

例子:矩阵相乘

//CPU实现void MatrixMulOnHost(float* M, float* N, float* P, int width){for(int i=0; i<width;   i)        for(int j=0; j<width;   j)        {float sum = 0;            for(int k=0; k<width;   k)            {float a = M[i * width   k];                float b = N[k * width   j];                sum  = a*b;            }            p[i * width   j] = sum;        }}
//cuda算法框架(3布)int main(void){//1.管理整个内存,为数据分配空间,将数据拷贝到GPU上        //2.在GPU上并行处理计算        //3.将结果拷贝回CPU}
//GPU实现void MatrixMulOnDevice(float* M, float* N, float* P, int Width){int size = Width * Width * sizeof(float);    //1.管理整个内存,为数据分配空间,将数据拷贝到GPU上    //分配输入    cudaMalloc(Md, size);    cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);    cudaMalloc(Nd, size);    cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);    cudaMalloc(Pd, size);        //2.在GPU上并行处理计算    _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)    {//访问一个matrix,采用二维block        int tx = threadIdx.x;        int ty = threadIdx.y;                //每个kernel线程计算一个输出        float Pvalue = 0;                //计算        for(int k=0; k<Width;   k)        {float Mdelement = Md[ty*Md.width   k];            float Ndelement = Nd[k*Nd.width   tx];            Pvalue  = Mdelement   Ndelement;        }                Pd[ty*Width   tx] = Pvalue;    }        //3.将结果拷贝回CPU    //1个block含width*width个线程    dim3 dimBlock(WIDTH, WIDTH);    dim3 dimGrid(1, 1);    MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd);    //传送数据    cudaMemcpy(Pd, P, size, cudaMemcpyDeviceToHost);    //释放    cudaFree(Md);    cudaFree(Nd);    cudaFree(Pd);}

主要性能问题:访存

来源:https://www.icode9.com/content-1-820301.html

(0)

相关推荐

  • CUDA-GDB安装+环境配置

    在GPU上开发大规模并行应用程序时,需要一个调试器,GDB调试器能够处理系统中每个GPU上同时运行的数千个线程.CUDA-GDB提供了无缝的调试体验,可以同时调试应用程序的CPU和GPU部分. 就像G ...

  • GPU编程(二): GPU架构了解一下!

    目录 前言 GPU架构 GPU处理单元 概念GPU GPU线程与SM GPU线程 SM 加法 统一内存 乘法 最后 前言 在实际CUDA编程之前, 先来了解下GPU的结构. 和CPU相比显得粗暴又强大 ...

  • C 使用类调用 CUDA 核函数

    (给CPP开发者加星标,提升C/C++技能) 来源:小小一步 导读:CUDA是用于GPU编程的框架,在深度学习高速发展的今天,应用十分普遍.然而虽然CUDA对C语言有很好的支持,但是在C++的类函数中 ...

  • Py之cupy:cupy的简介、安装、使用方法之详细攻略

    Py之cupy:cupy的简介.安装.使用方法之详细攻略 cupy的简介 CuPy: NumPy-like API accelerated with CUDA.CuPy是NumPy兼容多维数组在CUD ...

  • 收藏,7个学习Python编程的最佳开源库!

    来源丨网络 1.learn-python3 这个存储库一共有19本Jupyter笔记本.它涵盖了字符串和条件之类的基础知识,然后讨论了面向对象编程,以及如何处理异常和一些Python标准库的特性等.每 ...

  • ABAP初学者如何系统地学习ABAP编程?

    有很多 ABAP 的从业人员,在各种社交媒体平台上,向我询问过这个问题. 先说说我自己的实际情况.我是2007年硕士毕业加入 SAP 成都研究院的,本科和研究生做的项目,一直是用 C/C++ 开发.进 ...

  • 学习Shell编程要掌握哪些基础?Linux运维

    学习Shell编程需要掌握哪些基础知识?学习Shell编程并不是只有学好这些基础知识,而是只有掌握这些基础知识才能更深入理解Shell编程.shell是Linux运维人员都会遇到的知识点,通过Shel ...

  • 零基础到底是否可以学习电脑编程?答案扎心了!

    我们偶尔推送过Python方面的推广,大白并不否认那是推广,不过也有朋友留言问零基础是否可以学习,还有些朋友留言说是骗人的,这里大白给大家解答一下. 首先那肯定不是骗人的,留言者所谓的骗人应该指的交了 ...

  • 学习SCL编程,先要搞懂这9种程序控制指令(有示例)!

    SCL作为一种编程语言,可以实现LAD/FBD所有的功能,大多数的指令与LAD/FBD都是相同的,只是在编辑器中的外形不同.只有一些指令使用是不太一样的甚至LAD/FBD没有的,这里只介绍这些不同的. ...

  • 学习PLC编程需要哪些器材,什么型号和配置?

    学习PLC编程需要哪些器材,什么型号和配置?

  • 这篇心得,值得所有学习mastercam编程的人一看!

    数控编程教学 数控基础编程教学,数控编程初学者的必备学习平台. 17篇原创内容 公众号 导读 学习masterCAM和学习其它的软件一样,首先要有学习的恒心,另外要多思考和总结.下面一篇过来人的心得希 ...

  • 运维为什么要学开发?学习linux编程难度有多大

    互联网技术的飞速发展,推动了Linux运维的发展前行.Linux技术越来越受到人们的重视.近年来,运维自动化也成为一个比较热门的趋势.运维一定要学开发吗?运维为什么要学开发? 运维为什么要学开发? 近 ...

  • 学习中文编程

    学习中文编程

  • 学习PLC编程的经验之谈

    鄙人是工业自动化专业毕业的,我刚开始学习PLC时,也是一头雾水, 后来在来了技成之后,在专家老师的学习指导下,才有了更加明确的学习方向.在此,我想粗略的总结一下自己的学习之路,供大家参考. 1.学习P ...