以亭
以亭_个人页
以亭
文章
56
问答
7
视频
0
个人介绍
CSDN卜居是也
擅长的技术
-
发表了文章 2017-11-08
CUDA编程模型【中科院课件】
中科院高性能计算系统第二课，介绍了CUDA执行模型：核函数，线程，线程块，网格的层次结构。
-
发表了文章 2017-11-08
CUDA存储器【中科院课件】
中科院高性能计算系统第三课，CUDA存储器模型，介绍了寄存器、共享内存、私有内存和全局内存、常量内存、纹理内存的概念和特点。
-
发表了文章 2017-08-07
Jetson TX2 初体验
0. 简介 Jetson TX2【1】是基于 NVIDIA Pascal™ 架构的 AI 单模块超级计算机，性能强大（1 TFLOPS），外形小巧，节能高效（7.5W），非常适合机器人、无人机、智能摄像机和便携医疗设备等智能终端设备。Jatson TX2 与 TX1 相比，内存和 eMMC 提高了一倍，CUDA 架构升级为 Pascal，每瓦性能提高一倍，支持 Jetson TX1 模块的所有功能，支持更大、更深、更复杂的深度神经网络。 TX2 内部结构如下： 1. 开箱 过程细节不展开，板卡上电后来张照片： 2. 刷机 TX2 出厂时，已经自带了 Ubuntu 16.04 系统，可以直接启动。但一般我们会选择刷机，目的是更新到最新的 JetPack L4T，并自动安装最新的驱动、CUDA Toolkit、cuDNN、TensorRT。 刷机注意以下几点： Host 需要安装 Ubuntu 14.04，至少预留 15 GB 硬盘空间，不要用 root 用户运行 JetPack-${VERSION}.run，我用的是 JetPack-L4T-3.1-linux-x64.run TX2 需要进入 Recovery Mode，参考随卡自带的说明书步骤 刷机时间大概需要 1~2 小时，会格式化 eMMC，主要备份数据 3. 运行视频目标检测 Demo 刷机成功后，重启 TX2，连接键盘鼠标显示器，就可以跑 Demo 了。 nvidia@tegra-ubuntu:~/tegra_multimedia_api/samples/backend$ ./backend 1 ../../data/Video/sample_outdoor_car_1080p_10fps.h264 H264 --trt-deployfile ../../data/Model/GoogleNet_one_class/GoogleNet_modified_oneClass_halfHD.prototxt --trt-modelfile ../../data/Model/GoogleNet_one_class/GoogleNet_modified_oneClass_halfHD.caffemodel --trt-forcefp32 0 --trt-proc-interval 1 -fps 10 视频截图如下： 4. 运行 TensorRT Benchmark TensorRT 【3】是 Nvidia GPU 上的深度学习 inference 优化库，可以将训练好的模型通过优化器生成 inference 引擎 将 TX2 设置为 MAXP （最高性能）模式，运行 TensorRT 加速的 GoogLeNet、VGG16 得到处理性能如下： 5. TX2 不支持的 feature 不支持 int8 待发现 参考文献 【1】嵌入式系统开发者套件和模块 | NVIDIA Jetson | NVIDIA【2】Download and Install JetPack L4T【3】TensorRT 附录 deviceQuery nvidia@tegra-ubuntu:~/work/TensorRT/tmp/usr/src/tensorrt$ cd /usr/local/cuda/samples/1_Utilities/deviceQuery nvidia@tegra-ubuntu:/usr/local/cuda/samples/1_Utilities/deviceQuery$ ls deviceQuery deviceQuery.cpp deviceQuery.o Makefile NsightEclipse.xml readme.txt nvidia@tegra-ubuntu:/usr/local/cuda/samples/1_Utilities/deviceQuery$ ./deviceQuery ./deviceQuery Starting... CUDA Device Query (Runtime API) version (CUDART static linking) Detected 1 CUDA Capable device(s) Device 0: "NVIDIA Tegra X2" CUDA Driver Version / Runtime Version 8.0 / 8.0 CUDA Capability Major/Minor version number: 6.2 Total amount of global memory: 7851 MBytes (8232062976 bytes) ( 2) Multiprocessors, (128) CUDA Cores/MP: 256 CUDA Cores GPU Max Clock rate: 1301 MHz (1.30 GHz) Memory Clock rate: 1600 Mhz Memory Bus Width: 128-bit L2 Cache Size: 524288 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 32768 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 1 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: Yes Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 0 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) > deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 8.0, CUDA Runtime Version = 8.0, NumDevs = 1, Device0 = NVIDIA Tegra X2 Result = PASS 内存带宽测试 nvidia@tegra-ubuntu:/usr/local/cuda/samples/1_Utilities/bandwidthTest$ ./bandwidthTest [CUDA Bandwidth Test] - Starting... Running on... Device 0: NVIDIA Tegra X2 Quick Mode Host to Device Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(MB/s) 33554432 20215.8 Device to Host Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(MB/s) 33554432 20182.2 Device to Device Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(MB/s) 33554432 35742.8 Result = PASS NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled. GEMM 测试 nvidia@tegra-ubuntu:/usr/local/cuda/samples/7_CUDALibraries/batchCUBLAS$ ./batchCUBLAS -m1024 -n1024 -k1024 batchCUBLAS Starting... GPU Device 0: "NVIDIA Tegra X2" with compute capability 6.2 ==== Running single kernels ==== Testing sgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0xbf800000, -1) beta= (0x40000000, 2) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 0.00372291 sec GFLOPS=576.83 @@@@ sgemm test OK Testing dgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0x0000000000000000, 0) beta= (0x0000000000000000, 0) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 0.10940003 sec GFLOPS=19.6296 @@@@ dgemm test OK ==== Running N=10 without streams ==== Testing sgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0xbf800000, -1) beta= (0x00000000, 0) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 0.03462315 sec GFLOPS=620.245 @@@@ sgemm test OK Testing dgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0xbff0000000000000, -1) beta= (0x0000000000000000, 0) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 1.09212208 sec GFLOPS=19.6634 @@@@ dgemm test OK ==== Running N=10 with streams ==== Testing sgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0x40000000, 2) beta= (0x40000000, 2) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 0.03504515 sec GFLOPS=612.776 @@@@ sgemm test OK Testing dgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0xbff0000000000000, -1) beta= (0x0000000000000000, 0) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 1.09177494 sec GFLOPS=19.6697 @@@@ dgemm test OK ==== Running N=10 batched ==== Testing sgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0x3f800000, 1) beta= (0xbf800000, -1) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 0.03766394 sec GFLOPS=570.17 @@@@ sgemm test OK Testing dgemm #### args: ta=0 tb=0 m=1024 n=1024 k=1024 alpha = (0xbff0000000000000, -1) beta= (0x4000000000000000, 2) #### args: lda=1024 ldb=1024 ldc=1024 ^^^^ elapsed = 1.09389901 sec GFLOPS=19.6315 @@@@ dgemm test OK Test Summary 0 error(s)
-
发表了文章 2016-12-23
如何让你的机器学习玩超级玛丽
【1】 http://www.ehrenbrav.com/2016/08/teaching-your-computer-to-play-super-mario-bros-a-fork-of-the-google-deepmind-atari-machine-learning-project/【2】 https://youtu.be/C-BY3JhXTiE
-
发表了文章 2016-07-14
从Caffe开始深度学习实战，徒手hack代码乐趣更多
阿里云工程师卜居接受CSDN记者专访，介绍了他在深度学习领域的实践经验和成长心得，以及完成Caffe著作背后的原因和故事。 从Geoffrey Hinton 2006年的论文算起，这一波深度学习（DL）浪潮才10年，而如果只算国内，深度学习的流行也不过5年，尽管如此，深度学习现在已经扎根中国互联网，成为BAT、京东、360、今日头条等公司的基础技术和战略技术，与之伴随的，则是深度学习技术人员的快速成长，例如，阿里云工程师卜居（赵永科）（博客：http://blog.csdn.net/kkk584520 ），2014年才开始接触深度学习实战，如今已在深度学习及计算优化方面方面有很独到的见解。卜居在最近写了一本浓缩其深度学习实战经验的书——《深度学习—21天实战Caffe》，该书获得了机器学习前辈的肯定。日前，卜居接受CSDN记者专访，介绍了他在深度学习领域的实践经验和成长心得，以及完成Caffe著作背后的原因和故事。 卜居认为，深度学习具备强大的表达能力和灵活多变的模型结构，并在各种硬件加速方案的支撑下不断成熟，而Caffe具有高效的C++/CUDA实现、Matlab/Python接口、独特的网络描述方式、清晰的代码框架等优势，徒手hack代码的乐趣更多，同时Caffe框代码于稳定，掌握了阅读技巧可以事半功倍，因而可以作为初学者学习的第一个深度学习框架，由此逐步深入了解使用C++/CUDA代码实现深度学习的计算过程。 谈到新书《深度学习—21天实战Caffe》，卜居表示，这是一本透过源码解读深度学习的书，也是一本注重“实战”的书。读者可以从本书中学习Caffe设计模式、编程技巧，以及深度学习最新的进展和生产环境批量部署等内容。而书中的一些思考题需要深入实践和思考之后才能得到答案，这可以让读者养成独立思考的习惯，从而更加从容地面对实际问题。 此外，对于不同的硬件加速方案，卜居认为，深度学习本身在不断演进，没有哪个计算架构能够一劳永逸，得到某方面优势会丧失另一部分特性，最终起决定作用的仍然是应用需求，例如批量离线处理更适合利CPU/GPU集群的规模优势，而在线应用、移动端应用更适合利用FPGA/ASIC的低功耗、低延迟特性。 以下为采访实录：卜居与深度学习CSDN：请介绍您自己，您是如何与深度学习结缘的，它的哪些特质吸引您？ 卜居：大家好，我博客笔名是卜居（出自《楚辞》），在阿里花名是以亭（《镜花缘》探花唐敖之字），真名赵永科，2014年毕业于中国科学院大学，目前就职于阿里云计算有限公司高性能计算团队，从事CPU/GPU/FPGA上的计算优化工作。 我最早接触“神经网络”、“深度学习”还是在中科院微波成像实验室从事现代信号处理工作期间，通过阅读Simon Heykin的大部头著作《Adaptive Filter Theory》《Neural Networks : A Comprehensive Foundation》，初次了解其设计理念，但当时仅仅停留在一些抽象理论上，一直希望有机会将它们应用到实际项目中。 直到2014年我到（当时我们团队的名称）阿里巴巴集团核心系统部专用计算组实习，看到周围同事在用Caffe做深度学习算法优化，而且已经有很多内部项目（即后来的拍立淘、OCR等）在使用该框架。当时首先被名字吸引：“快速特征植入的卷积结构”（Convolutional Architecture for Fast Feature Embedding），谐音“咖啡”。当时的Caffe还只是雏形，我看了一遍代码后深深被其设计所吸引，高效的C++/CUDA实现、Matlab/Python接口、独特的网络描述方式、清晰的代码框架……当时对深度学习的认识一下子从教科书上的抽象概念落地到实实在在的代码实现，甚至对其稍作修改就可以应用于线上生产环境，直接面向广大用户！也是从那个时候，我开始关注每年的ILSVRC比赛，看到计算机视觉领域逐渐接受深度学习方法，并随后在语音识别、自然语言处理取得巨大成功，意识到这次深度学习高潮不是泡沫，而是逐渐迈向成熟的标志。 现在，深度学习强大的表达能力、灵活多变的模型结构、在CPU/GPU/FPGA/ASIC上加速方案等，都是吸引我的闪光点。 CSDN：阿里巴巴的很多业务都使用了深度学习技术，能否介绍和您的工作有直接关系的技术环节包括哪些？ 卜居：在阿里云我们组主要是计算平台支持和计算效率提升两个环节。 计算平台支持方面，我们的HPC支持当前大多数主流计算框架（Caffe/Torch/Theano/MxNet/TensorFlow/……），利用IB+双万兆网络实现多机互联，提供多机多卡并行训练系统，最近使用专门为深度学习定制的Tesla M40 GPU加速器将单节点计算能力提升到16 TFLOPS。 计算效率提升方面，我们自研了CUDA Native汇编器（ https://bbs.aliyun.com/read/277465.html?spm=5176.bbsl355.0.0.dEWHEi ），进一步挖掘硬件计算能力，实现更高计算效率。我的同事 @念鸿 @耘路 在今年 GTC 上做了分享，得到与会同仁的广泛关注。 深度学习及实践经验CSDN：目前深度学习的论文非常多，您会关注哪类论文？神经网络最近的哪些进展让您印象深刻？ 卜居：我比较感兴趣的还是计算优化方面，包括计算体系结构优化和计算软件优化两类。 最近的ResNet实现了深达1000层的模型结构，这不禁让人思考：“深度学习”还可以多深？除了加深之外，网络结构也逐渐表现出了生物特征，例如出现了分形结构（ https://arxiv.org/abs/1605.07648 ），这些让人愈发感受到自然之美、生命之美。 CSDN：ICML、CVPR上来自Deepmind等大公司的论文比较出彩，这其中不乏大公司投入比较大要求少的因素，比如LeCun说目前Facebook对FAIR还没有盈利要求，能否介绍您的团队如何要求DL研究的投入？ 卜居：我不敢妄谈公司战略，单从个人角度谈下。 DL不仅是我在阿里云高性能计算团队工作时需要了解的内容，同时也是我个人爱好，对于论文和前沿的技术一般是用业余时间学习，工作时间更多是面向用户需求。在阿里云有得天独厚的研究条件，我们有大规模GPU集群可以折腾，周围许多同事自己攒机体验最新的硬件（GPU、FPGA），研发氛围浓厚，讨论问题热烈，都是自发行为。 总体来看，大公司在数据规模、基础设施方面有较大优势，初创团队一般在业务算法、模型调参方面更有优势。CSDN：论文虽多，深度学习理论体系却不似SVM那么优美，在很多实际应用场景中的效果也还不尽人意，从开发人员的角度，您如何看待当前深度学习技术的缺陷？成功应用深度学习有没有一些“黄金法则”？ 卜居：深度学习不需要专门特征工程师，降低了模型设计门槛，但对于调超参“黑科技”方面，经验仍然很重要，不然会经常出现使用了深度学习技术效果反而不如传统方法的尴尬局面。 模型参数远大于数据量时，相当于求解一个欠定方程，存在多解的可能性大，容易产生过拟合问题。模型参数远小于数据量时，相当于求解超定方程，可能无解，或者有解但准确率很低，这属于欠拟合问题。模型参数与数据量匹配时，相当于求解恰定方程，既能避免过拟合，又能兼顾准确率，但模型参数量和数据量怎样才能做到匹配，是一个工程问题。所以，如果你选择用某个模型处理数据，那么应该考虑这个因素，越大的模型越难训练，因为需要与之匹配的数据量、一系列避免过拟合的方法才能训练得到一个较为理想的模型。幸运的是，我们可以将大模型首先在较大的数据集（如ImageNet）上预训练，得到模型，再对特定数据集（如人脸数据）进行精调（fine-tuning），即可得到较为理想的结果。 深度学习目前的应用场景只是冰山一角，对于更多传统领域的应用，一方面需要引起该领域工程师对深度学习技术的关注，另一方面要重视数据，作为待开发的矿产资源。 CSDN：您曾经做过CNN硬件加速的分享，重点谈了阿里云用到的FPGA方案，但目前还有很多ASIC的尝试，包括Google的TPU、中科院的寒武纪等，同时考虑到神经网络的应用越来越广，比如LSTM在语音、机器翻译中的应用，您认为未来DL平台应当是什么样的架构，主流的加速方案是什么，应当具备哪些特性？ 卜居：阿里云HPC团队是一个既注重技术创新、又贴近用户实际业务的团队，我们在平台选择方面，会首先根据用户需求进行优化和定制，而一些通用的模块则逐渐沉淀成为公共服务。 DL本身在不断演进，没有哪个计算架构能让DL一劳永逸，得到某方面优势会丧失另一部分特性，最终起决定作用的仍然是应用需求，例如批量离线处理更适合利CPU/GPU集群的规模优势，而在线应用、移动端应用更适合利用FPGA/ASIC的低功耗、低延迟特性。 CSDN：阿里云已经在用Docker来提供DL服务，您认为DL任务中使用Docker有哪些坑需要填好？ 卜居：Docker 的优点在于一次开发，处处部署，简化了开发测试环境与生产环境的迁移。目前已经有越来越多用户选择Docker，越来越多开源DL框架开始支持Docker部署，进一步降低了这些框架使用难度。Docker本身也在不断发展中，存在一些坑也是正常的，对于个人用户而言Docker基本坑都已经填好，是时候学习和掌握这门技术了。 关于“21天实战Caffe” CSDN：深度学习开源工具众多，您的新书写的是Caffe，而您最近的博客在写TensorFlow，能否介绍您尝试过哪些DL框架？如何看待它们的优缺点？ 卜居：在阿里云HPC平台上我们支持当前大多数主流计算框架（Caffe/Torch/Theano/MxNet/TensorFlow/……），我个人都使用过这些工具，还通读过其中几个框架的源码。 Caffe可以作为初学者学习的第一个深度学习框架，因为代码组织较规范，徒手hack代码的乐趣更多。但相比之下所需依赖库比较多，适合在固定计算平台（如服务器、带GPU的台式机）上运行。 TensorFlow提供了更灵活的设计，不仅可以部署在服务器端，对资源受限的硬件支持也更好，例如内置了定点化运算库，适合做移动、嵌入式设备上的DL应用。 CSDN：从KDnuggets的统计来看，用得最多的似乎是Pylearn2而不是Caffe，能否进一步介绍您的新书选择写Caffe的原因？ 卜居：诚然，从上述统计结果看，Caffe受欢迎程度可能不及Pylearn2、Theano。但我们在企业一线生产环境统计的结果来看，直接使用Pylearn2、Theano接生产任务的应用几乎是凤毛麟角。性能是一个关键的因素。 新书选择Caffe，是希望读者不仅知其然，更要知其所以然，知道怎样使用 C++/CUDA代码实现深度学习的计算过程。 CSDN：Caffe的架构还在发展，社区文档也在不断完善，您为什么会在现在写Caffe的书？这本书适合的阅读对象是谁？读者需要一些知识基础吗？ 卜居：从诞生到现在已经过去2年时间，Caffe经历过很多变化，框架代码其实变动不大，掌握了阅读技巧可以事半功倍，写这本书是希望更多读者可以加入hack Caffe、改进Caffe的行列中。 建议读者具备高等数学、线性代数、概率论等知识水平。编程方面，建议读者至少具备C++、Python、MATLAB基础。对于本科生，建议大二以上阅读本书。研究生、博士生理论基础达标，可以从本书中学习Caffe设计模式、编程技巧方面内容，充实自己的科研生活。企业一线工程师，可以通过本书了解深度学习最新的进展和生产环境批量部署等内容。 CSDN：您认为这本书的最大的亮点是什么？ 卜居：这是第一本透过源码解读深度学习的书，也是一本注重“实战”的书。 CSDN：如果开发者能完全吸收书中的营养，21天之后的DL水平是怎么样的？ 卜居：本书不是一部让读者从0开始，21天之后精通 Caffe 的“武林秘籍”，更像是游戏攻略，可以让读者更快地找到捷径，避免在原地徘徊。每章后面的思考题，有些是需要深入实践和思考之后才能得到答案的。相信通过这些锻炼，读者可以养成独立思考的习惯，面对实际问题时更加胸有成竹。 CSDN：开发人员在读完这本书之后，如何继续提升DL技能？还有一些DL书籍可以推荐？ 卜居：DL修行道阻且长，要有扎实的数学基础、精湛的编程技术，还要有广阔的视野，不断吸取其他领域的营养。推荐继续阅读其他ML/DL大部头著作（例如Youshua Bengio的《Deep Learning》值得深入阅读），以及该领域的经典论文。 笔耕收获CSDN：能否介绍您写这本书的最大挑战和收获？ 卜居：克服惰性。写书后，我养成了早起的习惯，每天精神了许多。 CSDN：从您2008年发表第一篇CSDN博客文章，到现在已经8年，能否介绍您坚持写博客的动力和收获、心得？博客创作对您的技术工作和完成著作有多大的帮助？ 卜居：2008年我还是在大一、大二刚刚学习编程的阶段，每期《程序员》、《电脑报》、《黑客 X 档案》都要细细阅读。偶然的机会看到CSDN上面的一些博客写得很好，于是也注册了一个号，尝试自己写文章。在此之前我写文章都是发表在QQ空间、人人网，发一些编程相关的文章不太“合群”，于是逐渐淡出了这些圈子，全面转向CSDN。 早期的文章不成体系，很少坚持。真正开始写系列文章是从研究生开始，结合自己学习CUDA的心得体会，整理了十篇左右，反响不错。后来又整理了关于Zynq开发的系列文章，访问量暴增，进一步增加了写博客的动力和信心。再后来毕业后参加工作，接触 DL，学习Caffe、TensorFlow，这时的博客更像笔记，比较随性。 写博客最大的收获，是查资料时可以先从自己博客中找，如果Cache命中，直接返回；否则再找其他人博客，学习后收藏。博客相比互联网上零散知识点而言更有凝聚性，可以集中解决一类问题，习惯从博客中找答案可以触类旁通。 此外，写博客也会认识更多同行者，可以在交流之间取长补短，相互促进，并打开视野，这对工作和学习有很大帮助。 原文链接：http://geek.csdn.net/news/detail/88598?utm_source=tuicool&utm_medium=referral 阿里云 HPC 产品页：https://www.aliyun.com/product/hpc Nvidia Kepler GPU上的性能极致优化（念鸿）：http://www.atatech.org/articles/58029 为什么要写一个GPU的汇编器（长仁）：http://www.atatech.org/articles/58106 【阿里集团卜居深度解析】卷积神经网络的硬件加速：https://zhuanlan.zhihu.com/p/21430333?refer=dlclasshttp://mp.weixin.qq.com/s?__biz=MzI1NTE4NTUwOQ==&mid=2650324711&idx=1&sn=8d833edfa7d856f0fc8db1d82ba0ad11&scene=0#wechat_redirect
-
发表了文章 2016-03-28
Windows 上 Caffe开发环境搭建
Caffe是目前非常流行的深度学习框架，使用C++/CUDA编写，使用方便，性能优异，适合线上环境部署。 原生Caffe是在Linux下编译部署的。对于初学者而言，大量依赖包需要花非常大代价才能编译成功，让人望而生畏。况且在Linux下阅读代码是一件非常头疼或蛋疼的事情。 为此，研究在Windows下搭建Caffe开发环境是利国利民的好事。微软不负众望在Github上开源了Microsoft版Caffe，链接：Microsoft Caffe 只有代码还不够，需要准备一台机器。 如果手头没有机器而且也不想攒机，可以考虑买一台阿里云HPC（购买链接：阿里云HPC），配备双 Tesla K40m 或 双 Tesla M40，专为深度学习定制机型。 购买后默认操作系统是Cent OS 7，通过提交工单可以重置为Windows 系统。为了完成本文实验，推荐重置为 Windows Server 2012 R2 64bit版。 编译环境：Visual Studio 2013 Ultimate版， 获取地址：http://download.microsoft.com/download/9/3/E/93EA27FF-DB02-4822-8771-DCA0238957E9/vs2013.5_ult_chs.iso?type=ISO CUDA Toolkit版本：7.5，获取地址：https://developer.nvidia.com/cuda-downloads CUDNN版本：4.0， 获取地址：https://developer.nvidia.com/cudnn 依次安装GPU驱动、Visual Studio 2013、CUDA Toolkit。将Microsoft/caffe代码下载到本地磁盘，本文路径为C:\Users\Administrator\Desktop\caffe-master。 将CUDNN解压到C:\Users\Administrator\Desktop\cuda。准备工作完毕。 进入C:\Users\Administrator\Desktop\caffe-master\windows目录，将文件CommonSettings.props.example 重命名为CommonSettings.props，修改其内容如下：
-
发表了文章 2015-11-26
Mac下安装Caffe
最近将工作笔记本从X1换成了MBA，真是不习惯，连复制粘贴都不那么顺手了。 好在命令行和Linux差不多，于是想搞个Caffe在Mac上跑跑，哪怕CPU版本的也行。 参考：http://caffe.berkeleyvision.org/install_osx.html 首先安装homebrew工具，相当于Mac下的yum或apt $ ruby -e "$(curl -fsSL https://raw.githubusercontent.com/Homebrew/install/master/install)" 等待片刻，安装成功，然后安装Caffe依赖： $ brew install -vd snappy leveldb gflags glog szip lmdb $ brew tap homebrew/science $ brew install hdf5 opencv $ brew install protobuf boost 下载Caffe源码： $ git clone https://github.com/bvlc/caffe.git $ cd caffe/ $ mv Makefile.config.example Makefile.config 修改Makefile.config，打开CPU_ONLY选项，保存。 $ make -j 编译成功。 运行一个例子，我们需要下载几个额外文件。 $ cd data/ilsvrc12/ $ vi get_ilsvrc_aux.sh 将这行： wget http://dl.caffe.berkeleyvision.org/caffe_ilsvrc12.tar.gz 改为： curl http://dl.caffe.berkeleyvision.org/caffe_ilsvrc12.tar.gz > caffe_ilsvrc12.tar.gz 然后保存，执行。下载数据到当前目录。 下载模型： $ cd ../../models/bvlc_reference_caffenet/ $ curl http://dl.caffe.berkeleyvision.org/bvlc_reference_caffenet.caffemodel > bvlc_reference_caffenet.caffemodel 回到根目录： $ cd ../.. 执行： $./build/examples/cpp_classification/classification.bin models/bvlc_reference_caffenet/deploy.prototxt models/bvlc_reference_caffenet/bvlc_reference_caffenet.caffemodel data/ilsvrc12/imagenet_mean.binaryproto data/ilsvrc12/synset_words.txt examples/images/cat.jpg ---------- Prediction for examples/images/cat.jpg ---------- 0.3132 - "n02123045 tabby, tabby cat" 0.2380 - "n02123159 tiger cat" 0.1235 - "n02124075 Egyptian cat" 0.1005 - "n02119022 red fox, Vulpes vulpes" 0.0716 - "n02127052 lynx, catamount"
-
发表了文章 2015-11-25
Zynq UltraScale+ MPSoC新鲜出炉
手头必备的TRM UG1085（http://www.xilinx.com/support/documentation/user_guides/ug1085-zynq-ultrascale-trm.pdf） 寄存器手册（http://www.xilinx.com/support/documentation/registers/ug1087/ug1087-zynq-ultrascale-registers.html） 抢先看几张框图 更多请参考： References Zynq UltraScale+ MPSoC Product Overview (DS891) Zynq UltraScale+ MPSoC Data Sheet: DC and AC Switching Characteristics (DS925) Zynq UltraScale+ MPSoC Register Reference (UG1087) Zynq UltraScale+ MPSoC Software Developer’s Guide (UG1137) Zynq UltraScale+ MPSoC Packaging and Pinout User Guide (UG1075) UltraScale Architecture SelectIO Resources User Guide (UG571) UltraScale Architecture Clocking Resources User Guide (UG572) UltraScale Architecture Memory Resources User Guide (UG573) UltraScale Architecture Configurable Logic Block User Guide (UG574) UltraScale Architecture GTH Transceivers User Guide (UG576) UltraScale Architecture GTY Transceivers User Guide (UG578) 12. UltraScale Architecture System Monitor User Guide (UG580) 13. UltraScale Architecture DSP Slice User Guide (UG579) 14. UltraScale Architecture PCB Design User Guide (UG583) ARM References ARM® CoreSight® SoC-400 Technical Reference Manual, r3p1ARM document numberDDI 0480F. ARM CoreSight SoC-400 User Guide, r3p1ARM document number DUI 0563F. ARM CoreSight SoC-400 System Design Guide, r3p1ARM document number DGI 0018E. ARM CoreSight SoC-400 Implementation Guide, r3p1ARM document number DII 0267F. ARM CoreSight SoC-400 Integration Manual, r3p1ARM document number DIT 0037E. ARM CoreSight STM-500 System Trace Macrocell Technical Reference Manual, r0p0ARMdocument number DDI 0528A. CoreSight Trace Memory Controller Technical Reference Manual, r0p1ARM documentnumber DDI 0461B. ARM Cortex®-A53 MPCore Processor Technical Reference Manual, r0p2ARM documentnumber DDI 0502D. CoreSight ETM-R5 Technical Reference Manual, r0p0ARM document number DDI 0469A. ARM CoreSight Architecture Specification, v2.0ARM document number IHI 0029D. ARM System Memory Management Unit Architecture Specification, SMMU Architectureversion 2.0, ARM IHI 0062C (ID091613).
-
发表了文章 2015-10-28
互联网是一巨大的FPGA
互联网和FPGA的几个相似点： 1. 网络互联，相当于FPGA内的走线； 2. 存储服务，相当于Flash或其他非易失存储器； 3. 数据库，相当于LUT； 4. 缓存服务器，相当于FPGA内部寄存器； 5. 网站逻辑状态机，与FPGA内部RTL实现的状态机并无二致。 6. 互联网服务可以抽象为IP核（黑盒）； 7. 请求，即输入； 8. 响应，即输出； 9. 对网站测试，很像写testbench； 10. 对网站故障进行诊断，需要借助服务器log，而对RTL调试，需要借助仿真时序图，或用chipscope、signal tap抓取IP的输入/输出波形log； 二者从本质上仅是宏观与微观，规模与个例，星系和原子的区别。 也许有天可以用js进行FPGA布局布线，也许有天可以用SystemVerilog设计网站架构。
-
发表了文章 2015-08-26
使用专用硬件加速深度卷积神经网络
英文论文链接：http://research.microsoft.com/apps/pubs/default.aspx?id=240715 翻译：卜居 转载请注明出处：http://blog.csdn.net/kkk584520/article/details/47711755 【摘要】 最近在多层卷积神经网络的突破导致了识别任务（如大量图片分类和自动语音识别）准确率的大幅提升【1】。这些多层神经网络很大，很复杂，需要大量计算资源来训练和评估【2】。然而这些需求发生在目前这样一个尴尬的时刻，商业处理器性能增长十分缓慢。 专用硬件形式有GPGPU，FPGA和ASIC，提供了通往处理能力和高能效的坦途。微软通过使用FPGA增强型服务器（类似已经集成到微软数据中心的硬件【3】）加速深度卷积神经网络驾驭了专用硬件。 开始只是在中等规模FPGA上实现了单节点CNN加速器，展示了优于遗忘FPGA设计和高端GPGPU的性能，降低了功耗。未来会通过低延迟通信互联多个FPGA，进一步可能训练和评估模型。 【背景】 目前深度卷积神经网络一般组织为交替的卷积层、最大池化层后面加一系列稠密的全连接层，如图1中经典的拓扑图所示。 每个3D立方体表示一层的输入，送入下一层会变换为新3D立方体。在例子中有5个卷积层，3个最大池化层，以及3个全连接层。 本文中我们主要讨论3D卷积问题，尽管其他操作如pooling和全连接层也会涉及。图2展示了3D卷积的基本模型。一个N x N x D的3D输入立方体被H个k x k x D维的卷积核卷积，输出间隔为S，每个3D核用类似滑动窗口的形式移动（移动偏移量由参数S定义）。每次移动，3D卷积核的每个权值同映射到3D输入立方体的输入值进行乘加，卷积后，可选地进行pooling操作（由参数p和s定义）用来下采样卷积输出，选择窗口的最大值或平均值 【在数据中心加速深度卷积神经网络】 2014年微软宣布了Catapult项目，成功展示了用FPGA在数据中心使Bing Ranking加速了近2倍【3】。利用这个基础，我们团队在微软研究院开发了高吞吐CNN FPGA加速器，在很低的服务器功耗下获得了优异性能。图3给出了用于高效计算卷积层前向传播的CNN FPGA加速器高层次概览。 关键特征如下： （1）软件可配置的引擎，支持多层运行时配置（无需硬件重编程）； （2）高效数据缓冲体制和片上分发网络，将片外访存降至最低； （3）处理单元（PE）构成的空间分布阵列，可很容易扩展到上千个单元； 在正常操作下，CNN加速器可以获取输入图像，连续处理多个卷积层。在初始阶段，输入图像像素从本地DRAM流入片上，存储到多个bank输入缓冲区。之后，数据流入多个PE阵列，实现3D卷积步骤中的独立点乘操作。顶层控制器完成序列化、寻址、分发数据到每个PE阵列。最终，累加结果发送到特定片上网络，将计算输出循环送入输入缓冲区用于下一轮的计算（虽然图3中没有显示，但存在额外的逻辑处理pooling和ReLU操作）。 图3中高亮的加速器位于双插槽Xeon服务器，装备一个Catapult FPGA卡，包括中等规模Stratix D5 FPGA和8GB DDR3-1333[3]。每个FPGA卡通过PCIe 3x8支持高达8GB/s带宽，同时本地DRAM支持21.3GB/s带宽。更多硬件描述可参考Catapult论文【3】。 表1显示了使用著名模型（如基于cuda-convnet【4】的CIFAR-10，基于Alex的ImageNet-1K【1】等）进行图像分类（只有前向传播过程）的吞吐情况。我们进一步评估了最大的和最具挑战性的模型：ImageNet-22K，在微软ADAM项目中训练的DCNN【2】。 总体上看，我们目前在装备中等规模Stratix V D5 FPGA的Catapult服务器上获得了最高处理吞吐。相比最近发表的FPGA实现【5】和运行在高端GPGPU上的Caffe+cuDNN【6】。注意GPGPU解决方案需要高达235W功耗【7】，使得部署到功耗限制的数据中心不太现实。相反，FPGA解决方案功耗不超过25W，占整个服务器功耗的不到10%。我们的实现获得了近3倍加速，相对最近发表的利用Virtex 7 485T FPGA的CNN加速器【5】。 （卜居注：【5】只说了处理能力达到61.62GFLOPS，并没有说处理图片速度，这篇文章怎么算的？） 我们参数化的CNN加速器可灵活扩展到更新和更快的FPGA。我们团队正在将设计移植到Altera的最新Arria 10 FPGA上，提供了对浮点处理支持，可以在高能效情况下达到1TFLOPS处理能力【8】。表1列出了我们的预期性能。 总而言之，本文描述了使用FPGA加速深度卷积神经网络的研究。前期结果是很有希望的，显示出专用硬件具有低功耗、高性能的优势。未来我们希望在Arria 10和Stratix 10上设计更高性能的加速器，并考虑多片互联并行训练和评估。
-
发表了文章 2015-08-07
一种Cache Line向量处理机概念
目前多核CPU都是保持各自的上下文，相互通信则依赖加锁/解锁以及共享内存实现。这种架构对于运行多任务、多线程的操作系统来说比较适合，而对于高性能计算则有点浪费。 其实，可以将多个计算核心组织为向量方式，每次访存都从Cache Line中抓取一整行数据，然后分头处理，类似SIMD方式，处理结果也是以Cache Line方式写回。这样的好处是不存在交叉访问，可以实现最高效率的计算/吞吐，如果待处理数据都是严格按照一个个Cache Line顺序排好，那么理论上可以实现最高效率的处理，主频为F，核心数为N，那么处理性能为N*F OPS。 当然，目前内存带宽可能远远跟不上CPU处理能力，所以这只是一个概念，对于具体实现还有很多工程性的问题。
-
发表了文章 2015-04-16
Altera OpenCL SDK与Xilinx SDAccel的使用对比
出于工作需求，申请了这两家的高级语言综合工具，对典型算法进行了实现和评估（数据暂时保密）。 简要谈谈使用体验。 1. Altera OpenCL SDK 首先需要安装Quartus（13.1版本以上）和配套的SoC EDS，分别申请两个license，一个用于OpenCL SDK，一个用于SoCEDS，缺一不可。 然后需要有实现平台，我用的是DE1-SoC开发板。该平台提供了Open CL BSP，用给出的examples改改就能实现自己的算法。 SoCEDS用来编译Host代码，这里是ARM； Open CLSDK用来编译device代码，生成.aocx二进制文件，运行时需要借助ARM对FPGA进行配置。 开发周期较长的是编译Open CL工程，在我的笔记本（CPU 酷睿i5-4300，8GB内存）上大约需要40min。 优化Open CL可以通过编译选项、编译指导语句#pragma进行。 Open CL BSP自带的硬件工程可以用Quartus打开查看，但里面的Open CL相关逻辑是加密的，无法修改。 Open CL SDK不带图形界面，只能在命令行下运行，自动调用quartus_map, quartus_fit, quartus_sta等工具。 2. Xilinx SDAccel 前面文章介绍过Xilinx Vivado和Vivado HLS工具。按照我的猜想，SDAccel只是一层包装，里面内容还是HLS。果然如此。 软件安装比较省事，只申请SDAccel license就能使用所有Xilinx软件功能。这点比Altera做得人性化。 用过HLS工具的童鞋都知道，这玩意开发起来飞快，但其实只完成了10%进度，麻烦事在后头呢，系统集成会花掉剩下90%的时间。 SDAccel就是这样一个用于开发完整项目的工具。 利用Open CL编写的代码会先转化为HLS工程，其优化策略也都一一映射为HLS的优化，所以前面学过HLS的童鞋可以很容易转到SDAccel上来。 最重要的步骤都有工具自动完成。只需build_system, package_system，等上几十分钟就能得到一个可直接运行的安装包。 SDAccel不需要你真的有一块开发板，它可以直接用CPU仿真，便于调试。 SDAccel既能使用图形界面开发（类似HLS），又能在命令行执行。上传一张谍照。 小结 使用OpenCL的一大优势就是，只需将原来CPU、GPU上的C、C++、CUDA代码做些许改动就能运行在FPGA上。 两种工具都带来了一定开发上的便利，但真想用好这些工具，仍然需要对计算架构、算法本身有深入的理解。
-
发表了文章 2015-02-11
cuda-convnet2与caffe对比
出于对性能和多GPU训练CNN的考虑，这段时间一直在研究cuda-convnet2。 搜了下，网上居然一篇像样的研究cuda-convnet2 代码的文章都找不到，看来假期有的忙了。 Caffe作者贾扬清也在一些场合表达了对Convnet2作者Alex的仰慕之情，可见两个CNN实现的差距。 Caffe比较符合大众的口味，而convnet2符合GPU发烧友的追求。 convnet2代码风格不如Caffe那样有条理。 Caffe本质是单线程的，或者是CPU思维方式。convnet2是多线程的，属于GPU思维方式。 Caffe过度依赖库函数（glob, gflags, leveldb, lmdb, mkl/blas……），而convnet2几乎全都自己搞定。 Caffe参数设置更自由，而convnet2出于性能考虑，参数设置约束较多。 Caffe接近软件，而convnet2接近硬件。 Caffe适合懒人，convnet2适合极客。
-
发表了文章 2014-12-30
Fedora + Eclipse + CUDA6.0
为了运行Caffe，在Fedora 17系统中安装了CUDA6.0，某天突然发现开始菜单中有NSight Eclipse Edition，于是好奇地打开看看和Visual Studio有什么区别。 打开时过场动画如下： 弹出对话框如下，我们选择默认工程位置。 进入Eclipse后，选择File -> New -> CUDA C/C++ Project，弹出对话框如下： 工程名输入first，工程类型选择Executable/Import CUDA Sample，工具链选择CUDA Toolkit 6.0，下一步。 这里选择gencode选项，由于已经检测到GeForce610M的计算能力为2.1，我们只选PTX 2.0和GPU 2.1两项，点下一步。 这里保持默认，下一步。 仍然默认，完成。 这里打开的工程是convolutionFFT2D。在左侧Project Explorer中first文件夹上右键，选择Build Project，状态如下： 编译无误，我们可以仍然在first文件夹上右键，选择Run As... Local Application，在Console中得到输出结果如下： [/root/cuda-workspace/first/Debug/first] - Starting... GPU Device 0: "GeForce 610M" with compute capability 2.1 Testing built-in R2C / C2R FFT-based convolution ...allocating memory ...generating random input data ...creating R2C & C2R FFT plans for 2048 x 2048 ...uploading to GPU and padding convolution kernel and input data ...transforming convolution kernel ...running GPU FFT convolution: 130.005202 MPix/s (30.768000 ms) ...reading back GPU convolution results ...running reference CPU convolution ...comparing the results: rel L2 = 8.130692E-08 (max delta = 5.342852E-07) L2norm Error OK ...shutting down Testing custom R2C / C2R FFT-based convolution ...allocating memory ...generating random input data ...creating C2C FFT plan for 2048 x 1024 ...uploading to GPU and padding convolution kernel and input data ...transforming convolution kernel ...running GPU FFT convolution: 40.667351 MPix/s (98.359001 ms) ...reading back GPU FFT results ...running reference CPU convolution ...comparing the results: rel L2 = 8.405842E-08 (max delta = 5.613083E-07) L2norm Error OK ...shutting down Testing updated custom R2C / C2R FFT-based convolution ...allocating memory ...generating random input data ...creating C2C FFT plan for 2048 x 1024 ...uploading to GPU and padding convolution kernel and input data ...transforming convolution kernel ...running GPU FFT convolution: 46.281297 MPix/s (86.428001 ms) ...reading back GPU FFT results ...running reference CPU convolution ...comparing the results: rel L2 = 8.385063E-08 (max delta = 5.613083E-07) L2norm Error OK ...shutting down Test Summary: 0 errors Test passed 可见运行成功。本文工程可在这里下载。
-
发表了文章 2014-12-06
NESASM教程——第十三天——背景
【简介】 你应该知道什么是背景，它就是静态图片，主角在其间穿梭。背景不难做，但缺少用于NES的背景制作工具确实是个头疼的问题。 【制作背景】 在Tile Layer Pro中打开你的"our.bkg"，绘制一些瓷砖，然后在纸上绘制一个用瓷砖编号排列的地图。注意瓷砖编号从0开始。假设你的Tile Layer Pro每行有16个瓷砖，那么它们的编号排列如下： $0 $1 $2 $3 $4 $5 $6 $7 $8 $9 $A $B $C $D $E $F $10 $11 $12 $13 $14... 如果1号瓷砖是一个笑脸的左上四分之一，2号瓷砖是右上，$11=17是左下，$12=18是右下四分之一，你在纸上或文本文件中写出来应该就像这样： 1 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 11 12 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 注意：NES屏幕是32X32个瓷砖。 下面就是把这个地图放到一个二进制文件中，以便载入PPU内存，就像你做调色板一样。 我已经写好一个VB小程序用来将一个类似上面的地图文件（不包括‘...’）转换为二进制文件。你不必写全不32行，除非你确实有需要。 从这里下载，运行它，在第一个框中输入地图文件（文本），第二个框输入你想创建的文件名。如果地图文件包括很多地图数据，程序会创建多于1个文件。 【命名表】 命名表这个东西让我实在是不理解，也是我不想写NES汇编的最大原因。技术手册中常常出现。 简单说来，”名“就是我们上面的瓷砖编号，”表“就是地图。所以命名表就是我们在PPU中想放地图的位置。这个地址是$2000，但屏幕上第一行和最后一行常常是不显示的，所以实际开始地址是$2020。 【构建代码】 假设你已经有一个源文件，包含了以下基本内容： （1）段1 有$FFFA那些东西（复位向量， VBlank向量等） （2）段2 包含两个文件our.spr和our.bkg，顺序正确； （3）段0 开始于$8000，有代码设置PPU，并载入调色板； （3.5）二进制调色板文件包含语句在死循环后面； 在调色板载入循环后面加入这些代码： ;;--- 代码开始 ---;; lda #$20 sta $2006 ; 写起始地址$2020 sta $2006 ldx #$00 loadNames: lda ourMap, X ; 从地址(ourMap + X)处取一个字节，加载到A inx sta $2007 cpx #64 ; 我们上面一共写了64个字节； bne loadNames ; 循环 在.incbin 调色板文件语句后面加入这些代码： ourMap: .incbin "our.map" ; 假设our.map是生成的二进制map文件 好了，汇编，运行，让背景显示吧！ 【一些其他事情】 如果map2bin.exe产生了多个文件，你需要为每个文件准备一个循环和一个.incbin语句，同时循环次数变为$FF次。 map2bin，由于用VB写的，使得每个文件大了2字节，因为VB每次写完文件都加一个新行。 还有一种表叫做属性表，我们后面讲它。
-
发表了文章 2014-12-06
NESASM教程——第十二天——VBlank和中断
【VBlank？】 如果你不知道什么是VBlank， 那么假装我给了你一条鳟鱼； 如果你知道什么是VBlank，那么假装我给了你一条鳟鱼； 如果你不确定一条鳟鱼和NES有什么关系，那么假装我给了你一条鳟鱼； 既然我们房间里有一股想象中的鱼臭味，那么让我们学习更好的方式来与VBlank同步。 【我们怎么用它？】 首先我们有个标签，比如VBlank_Routine，无哦一我们把它放在段1的地址$FFFA处，就像这样： .bank 1 .org $FFFA .dw VBlank_Routine ; VBlank中断向量 .dw Start ; 复位向量 .dw 0 ; 当BRK指令执行时的中断向量，改天再谈 .bank 0 .org $0000 ;保留一个变量VBlankOrNo VBlankOrNo .db 0 .org $8000 ;code VBlank_Routine: ;VBlank子程序开始 inc VBlankOrNo ; 使变量VBlankOrNo增1 rti ; RTI 意思是中断返回 (Interrupt RETurn or ReTurn from Interrupt) Start: ;主程序开始 ; ;waitforvblank: 旧代码 ; lda $2002 旧代码 ; bpl waitforvblank 旧代码 ; 新代码如下 WaitForVBlank: lda VBlankOrNo ; A = VBlankOrNO cmp #1 ; if A == 1 说明发生了VBlank bne WaitForVBlank ; 没有发生VBlank，循环 dec VBlankOrNo ; 发生了VBlank，使变量VBlankOrNo减一，变为0 ; 代码如旧 希望你懂了总体思想。 哦对了，别忘了将PPU控制寄存器中的“Generate VBlank Bit"设为1（我记得是bit7，可以复习第三天内容）。 修改后的代码比之前旧的循环更稳定。 晚安~
-
发表了文章 2014-12-06
NESASM教程——第十一天——JSR指令
【JSR？】 JSR代表Jump to SubRoutine，跳转到子程序，作用就是，嗯，跳转到子程序。JSR代表Jump to SubRoutine，跳转到子程序，作用就是，嗯，跳转到子程序。 【我们怎么用它？】 我们用它时只需要给出我们想跳转到的标签即可。例如：我们用它时只需要给出我们想跳转到的标签即可。例如： jsr OurSub ; 主程序 OurSub: ; 子程序 ; 返回指令 【怎样返回？】 当然使用RTS（ReTurn from Subroutine）指令了！所以完整版代码应该为： jsr OurSub ; 主程序 OurSub: ;子程序 rts ; 返回到主程序 【更多】 记住一件事情，如果你正好想把一部分代码分离为单独的模块，那么上面描述的就够了。 【今日回顾】 子程序不得不说，否则中断你更不懂了。你不妨试试把我们前面讲过的几个功能模块分为若干子程序，代码会清晰很多~ 编程愉快！
-
发表了文章 2014-12-06
NESASM教程——第十天——主角DMA
【主角DMA？】 是的，主角DMA。你还记得我们怎么利用$2003和$2004寄存器写入SPR-RAM（OAM）的吗？实际上真实系统中这种做法是不可靠的。我们应该利用内存作为OAM，然后向一个寄存器写入值，所有内容都自动拷贝到真实OAM。如果你不懂，那我们做一遍就懂了。 【什么内存？】 就像我上面说的，我们需要使用“变量”内存来复制一份OAM。本教程使用$0300用于复制的OAM。注意NES上只有64个主角（占用64*4B=256B，0x100）。所以基本上你自己定义的变量尽可能放在$0000-$0200内。从$0300-$0400的布局应该同OAM完全一致。 另外记住，$300是一个内存地址，而不像内存变量。我们必须在每次读写后将地址加一，而不是持续读写同一个地址的内容。 我们用偶数在百位用于我们OAM拷贝。我会告诉你为什么的。 好了，希望你理解了上述理论，下面看汇编代码。 ;;--- 代码开始 ---;; .inesmap 0 ; .inesprg 1 ; .ineschr 1 ; .inesmir 1 ; .bank 1 .org $FFFA .dw 0 .dw Start .dw 0 .bank 0 ; 代码段 .org $0000 ; ; 普通变量在这里定义 .org $0300 ; OAM镜像从这里开始 Sprite1_Y: .db 0 ; 1号主角的纵坐标 Sprite1_T: .db 0 ; 1号主角的瓷砖编号 Sprite1_S: .db 0 ; 1号主角的特殊待遇 Sprite1_X: .db 0 ; 1号主角的横坐标 Sprite2_Y: .db 0 ; 2号主角的纵坐标 Sprite2_T: .db 0 ; 不用说了吧？ Sprite2_S: .db 0 ; Sprite2_X: .db 0 ; ; 依此类推去吧。。。 .org $8000 ; 代码开始 Start: ; ; 先卖个关子 ; 后面会给出详细代码 infin: jmp infin ; 死循环 ;;--- 代码结束 ---;; 如果你不懂，email我告诉我你到底哪不懂。 【DMA寄存器】 DMA寄存器地址是$4014，我们需要写3进去。为什么是3？因为我们内存OAM在$300处。你向$4014写入n，那么就从$0n00处拷贝内容到真正OAM。也就是说，如果从$0400开始那就写4，如果从$0500开始那就写5，明白？ 下面看怎么搬运我们的OAM到真实OAM中： lda #$3 ; 也可以写成 #3, 因为显然3的十进制和十六进制表示是相同的 sta $4014 ; 一旦写入, 拷贝就开始 就是这样，比我们老办法可靠多了，而且更简单！ 【怎样按照上述方法修改第九天的代码？】 我们需要做几件事。 首先，拷贝 .org $0300和后面那堆东西到我们旧的变量区。我们再也不用那些变量了，因为我们的主角X，Y都使用OAM拷贝。 其次，使用查找替换功能吧所有X_Pos和Y_Pos改为Sprite1_X和Sprite1_Y。 再次，找到写$2003的代码块，替换为： lda #$3 ; sta $4014 ; 就是它了！我们用这种方法也节省了几个字节代码空间。 【今日回顾】 希望你喜欢主角DMA，我已经将它尽可能口语化了。我们今天学到了更好的写OAM数据的方法，明天还要看下更好的方法来捕获VBlank。接下来是。。。。。中断！ 希望你的代码没有bug。
-
发表了文章 2014-12-02
Caffe代码导读（0）：路线图
【Caffe是什么？】 Caffe是一个深度学习框架，以代码整洁、可读性强、运行速度快著称。代码地址为：https://github.com/BVLC/caffe 【博客目的】 从接触Caffe、编译运行、阅读代码、修改代码一路走来，学习到不少内容，包括深度学习理论，卷积神经网络算法实现，数学库MKL，计算机视觉库OpenCV，C++模板类使用，CUDA程序编写…… 本博客目的是为初学者清除代码阅读中的障碍，结合官网文档、融入个人理解、注重动手实践。 【如何开始】 在开始阅读Caffe代码之前，应该做好下面几件事： （1）下载Caffe源码； （2）配置开发环境（安装CUDA、OpenCV、boost、leveldb、lmdb、Python等，安装步骤参考http://tutorial.caffe.berkeleyvision.org/installation.html）； （3）编译； （4）运行例子（如MNIST、CIFAR10、ImageNet等）； 【路线图】 （1）Caffe源码阅读路线图应该是从CAFFE_ROOT/src/caffe/proto/caffe.proto开始，了解各类数据结构，主要是内存对象和序列化磁盘文件的一一对应关系，知道如何从磁盘Load一个对象到内存，以及如何将内存对象Save到磁盘，中间的过程实现都是由Protobuf自动完成的。 （2）第二步就是看头文件，不用急于去看cpp文件，先理解整个框架。Caffe中类数目众多，但脉络十分清晰。在Testing时，最外层的类是Caffe::Net，包含了多个Caffe::Layer对象，而Layer对象派生出神经网络多种不同层的类（DataLayer, ConvolutionLayer, InnerProductionLayer, AccurancyLayer等），每层会有相应的输入输出（Blob对象）以及层的参数（可选，Blob对象）；Blob中包括了SyncedMemory对象，统一了CPU和GPU存储器。自顶向下去看这些类，结合理论知识很容易掌握使用方法。 （3）第三步就是有针对性地去看cpp和cu文件了。一般而言，Caffe框架不需要修改，只需要增加新的层实现即可。例如你想自己实现卷积层，只需从ConvolutionLayer派生一个新类MyConvolutionLayer，然后将几个虚函数改成自己的实现即可。所以这一阶段关注点在算法上，而不是源码本身。 （4）第四步就很自由了，可以编写各类工具，集成到Caffe内部。在CAFFE_ROOT/tools/下面有很多实用工具，可以根据需要修改。例如从训练好的模型中抽取参数进行可视化可以用Python结合matplot实现。 （5）接下来，如果想更深层次学习，最好是自己重新写一遍Caffe（时间充裕的情况）。跳出现有的框架，重新构建自己的框架，通过对比就能学到更多内容。
-
发表了文章 2014-11-16
NESASM教程——第九天——运动的主角
【啥？】 今天不是教程，更像是前天留的作业的解决方案，即用手柄控制主角运动。我相信我们实际上已经掌握了所有必要的技能。下面给出完整代码。 【代码】 ;;--- 代码开始 ---;; .inesprg 1 .inesmap 0 .inesmir 1 .ineschr 1 .bank 1 .org $FFFA .dw 0 ; .dw Start ; 复位向量 .dw 0 ; .bank 0 .org $0000 X_Pos .db 20 ; 主角X坐标 Y_Pos .db 20 ; 主角Y坐标 .org $8000 ; Start: lda #%00001000 ; sta $2000 ; lda #%00011110 ; 典型的PPU设置代码 sta $2001 ; ldx #$00 ; 清零X ;; 开始载入调色板 lda #$3F ; sta $2006 ; lda #$00 ; sta $2006 loadpal: ; lda tilepal, x ; sta $2007 ; inx ; cpx #32 ; bne loadpal ; ;; 载入调色板结束 infinite: ; 死循环 waitblank: lda $2002 ; bpl waitblank ; 等待VBlank lda #$00 ; 开始写SPR-RAM sta $2003 ; lda #$00 ; sta $2003 ; lda Y_Pos ; sta $2004 ; 纵坐标 lda #$00 ; sta $2004 ; tile编号0 lda #$00 ; sta $2004 ; lda X_Pos ; sta $2004 ; 横坐标 ; 务必注意顺序 lda #$01 ; sta $4016 ; lda #$00 ; sta $4016 ; 设置/选通手柄 lda $4016 ; A lda $4016 ; B lda $4016 ; Select lda $4016 ; Start lda $4016 ; UP and #1 ; bne UPKEYdown ; ;如果按下“UP”，跳转到UPKEYdown lda $4016 ; DOWN and #1 ; bne DOWNKEYdown lda $4016 ; LEFT and #1 ; bne LEFTKEYdown lda $4016 ; RIGHT and #1 ; bne RIGHTKEYdown jmp NOTHINGdown ; 没有按键按下，跳到NOTHINGdown UPKEYdown: lda Y_Pos ; 将Y坐标载入A sbc #1 ; 减1 sta Y_Pos ; 保存 jmp NOTHINGdown ; 处理结束 DOWNKEYdown: lda Y_Pos adc #1 ; Y坐标加1 sta Y_Pos jmp NOTHINGdown ; LEFTKEYdown: lda X_Pos sbc #1 sta X_Pos jmp NOTHINGdown ;左键按下，X坐标减1 RIGHTKEYdown: lda X_Pos adc #1 sta X_Pos ; 右键按下，X坐标加1 NOTHINGdown: jmp infinite tilepal: .incbin "our.pal" ; .bank 2 .org $0000 .incbin "our.bkg" .incbin "our.spr" ;;--- 代码结束 ---;; 希望你还保存着第五天的our.pal, our.bkg, our.spr文件。你应该理解所有这些必要的内容。 注意：在一些模拟器上你可能得按下好几次按钮，主角才出现。我不知道为啥。 【今日回顾】 感谢那些发给我运动主角代码的人，至少我知道有人读了我的教程。
-
发表了文章 2014-11-16
NESASM教程——第七天——使用内存
【从哪里获得？】 我们做任何其他事情之前，最好展示怎么使用内存用于我们自己的数据。这个数据可能是任何东西，例如你需要存储的数字或主角的X/Y坐标值。我们将使用自由内存：位于CPU $0000。 【怎样使用？】 代码段0用ORG伪指令定位到$0000，为一些内存贴上标签，然后ORG定位到$8000开始我们的代码。 创建一个带标签的内存（也就是创建一个变量），可使用汇编操作.db，就像这样： .bank 0 .org $0000 label1: .db 0 byte1: .db $A5 ; .db后面的值不需要加# .org $8000 Start: ; 正常写代码 由于label1和byte1实际上都是标签，它们不需要缩进。 【载入变量值】 lda label1 ; label1处变量的值赋给A ldx label1 ; label1处变量的值赋给X ldy label1 ; 你懂得 lda byte1 ; 你懂得 【保存到变量】 sta label1 ; 将A的值保存至label1处的变量 stx label1 ; 将X的值保存至label1处的变量 sty label1 ; 你懂得 stx byte1 ; 你懂得 【一些重要的现实】 我想提的主要内容是，我们除了字节，还可以定义字，但由于NES的CPU只能工作在8bit，所以字不能一次读取或写入。 我还没试，你可以这样： inc label1; 增1 dec label1; 减1 【今日回顾】 我不知道明天讲什么，因为我需要完成GBA教程的第13天。我希望你有愉快的一天。为什么你不试试用手柄控制主角移动呢？？？
-
发表了文章 2014-11-16
NESASM教程——第五天——主角出场
【主角】 主角和背景比起来，你肯定觉得背景会更简单。错！主角才简单得要命。 制造主角只需要向SPR-RAM写入一点点内容（主角的x，y坐标，tile编号，等等）。 等等，我觉得咱们得先绘制主角的图片才行，走起~：） 【Tile Layer Pro】 存储主角像素块的表格称为Pattern Table，大小4KB，位于PPU $0000~$0FFF或$1000~1FFF（由PPU 控制寄存器设置，见第三天内容）。 绘制像素块的工具可以用Tile Layer Pro，我们从一个空白的像素块开始。 运行TLP.exe，打开MT.spr。点击View菜单，设置Format为NES。 点击大的网格窗口的左上角第一个小方块，看到右边有个Pallete Editor窗口，这就对了，我们只能用4种颜色。现在在Tile Editor窗口中画一个小笑脸或者其他什么东东。保存文件为our.spr。很好，我们拥有了一个主角！你要是实在做不来，那就下载我做的这个。 之后，重新下载MT.spr，重命名为our.bkg。顾名思义，这个是用于背景的。今天我们不讨论背景的制作，但为了让今天程序运行，我们必须有这个文件。 将our.spr和our.bkg放到同一个文件夹下，你的源代码也放在这里。 【段2变得更重要了】 前一天，我们告诉PPU，我们将使用段2的$0000作为背景图片数据区，$1000作为主角图片数据区。这样一来段2就拥有8KB，即$2000字节。4KB就是4096B或十六进制下的$1000B。代码中我们这样实现： .bank 2 .org $0000 .incbin "our.bkg" .incbin "our.spr" .incbin包含一个二进制文件。 【写SPR-RAM】 SPR-RAM也称“主角信息区”，写入通过两个寄存器：$2003和$2004。类似前面载入调色板时写$2006和$2007，我们同样需要先写两次$2003，告诉PPU地址，然后将数据写入$2004。 lda #$00 sta $2003 sta $2003 这样设置SPR-RAM地址为$0000。 写$2004的内容包括4字节，具体如下： Y —— 屏幕上的垂直位置 N —— 8x8 tile编号，我们包含的二进制.spr文件里面有256个不同tile，编号0~255 C —— 一些特殊信息，我还没细看 X —— 屏幕上的水平位置 PPU可以放置64个主角信息。如果我们想让主角在坐标（20，50）出场，tile编号为0，C不知是啥也写0，程序如下： lda #50 ; sta $2004 ; 设置纵坐标为50 lda #00 ; sta $2004 ; tile编号为0 sta $2004 ; C = 0. lda #20 sta $2004 ; 设置横坐标为20 咣当！主角出现（希望如此）！ 【VBlank】 理论上我们应该在场消隐期间写入SPR-RAM。这段时间称为VBlank。只给代码，不解释。我们重点是主角。 waitblank: lda $2002 ; bpl waitblank ; 不用担心，我们马上就把代码组合到一起。 【第一个完整代码】 新建一个our.asm文件，用记事本或其他编辑器打开，拷贝以下代码： ;;--- 程序开始 ---;; ; INES 文件头 .inesprg 1 ; 1 段代码 .ineschr 1 ; 1 段数据 .inesmir 1 ; 总是1 .inesmap 0 ; 使用mapper 0 .bank 1 ; .org $FFFA ; .dw 0 ; dw 表示定义字，NMI中断向量 = 0（禁止） .dw Start ; 复位向量，代码从Start标签处开始 .dw 0 ; VBlank 中断向量，0表示禁止 ; .bank 0 ; bank 0 - 代码段 .org $8000 ; 代码从 $8000开始 Start: ;还记得大明湖畔的复位向量吗？ lda #%00001000 ; 就像那天说的一样设置PPU sta $2000 ; lda #%00011110 ; sta $2001 ; ldx #$00 ; 准备载入调色板 lda #$3F ; sta $2006 ; lda #$00 ; sta $2006 loadpal: ; lda tilepal, x ; sta $2007 ; inx ; cpx #32 ; bne loadpal ; waitblank: ; 等待 VBlank lda $2002 ; 读取 $2002的值 bpl waitblank ; 如果bit7 == 0，那就接着waitblank lda #$00 ; 写入SPR-RAM sta $2003 ; lda #$00 ; sta $2003 ; lda #50 ; sta $2004 ; lda #$00 ; sta $2004 ; lda #$00 ; sta $2004 ; lda #20 ; sta $2004 ; ; 注意顺序呀！ infin: jmp infin ; 死循环 tilepal: .incbin "our.pal" ; 包含调色板，贴标签 .bank 2 ; 数据段 .org $0000 ; .incbin "our.bkg" ; 空白背景数据 .incbin "our.spr" ; 我们绘制的主角数据 ; 注意顺序呀！ ;;--- 代码结束 ---;; 上面代码如果有不懂的，再看看前面几天的教程。 【汇编】 把所有原材料准备好（our.pal, our.bkg, our.spr, nesasm.exe, our.asm），放到同一个文件夹，打开DOS，CD到这个文件夹，输入： nesasm our.asm 命令结束后，在当前文件夹下产生our.nes文件。用模拟器运行这个nes文件试试。 【今日回顾】 哇，这么多！别担心，看上去很多，一旦你懂了NES，这些东西都会乖乖听你的。明天我们也许会讲背景。睡个好觉~
-
发表了文章 2014-11-16
NESASM教程——第三天——NES编程相关
【PPU】 NES上发生的一切都需要通过图形方式输出，我们需要对PPU（图像处理单元）编程。其实很简单，就是向特定内存地址写一些数值，然后PPU的设置就被修改为你设置的值。NES所有编程都使用所谓的内存映射寄存器（MMR）实现。如果你给GBA编过程序，肯定知道MMR。但对于Intel芯片知识来说，MMR是个外星人。 【二进制注意事项】 为了绝对清晰，我们将按以下顺序给出内存寄存器的比特（0x63）： bit： 7 6 5 4 3 2 1 0 val：0 1 1 0 0 0 1 1 所以，第7比特位于最左边，第0比特位于最右边。 【设置PPU】 注意首先要做的一件事是设置PPU。我们只需向存储器$2000和$2001写一些值，这两个地址对应两个PPU控制寄存器。为了解释$2000和$2001，下面从YOSHi的文档里面摘抄的寄存器描述： 2000h - PPU 控制寄存器 1 (只写) Bit7 发生VBlank时执行NMI (0=Disabled, 1=Enabled) Bit6 PPU 主从模式选择 (0=主模式, 1=从模式) (NES中没用) Bit5 主角尺寸 (0=8x8, 1=8x16) Bit4 背景Pattern Table 首地址 (0=VRAM 0000h, 1=VRAM 1000h) Bit3 主角Pattern Table 首地址 (0=VRAM 0000h, 1=VRAM 1000h) Bit2 端口0x2007 VRAM 地址增量 (0=自动增1, 1=自动增32) Bit1-0 Name Table卷轴首地址 (0-3=VRAM 2000h,2400h,2800h,2C00h) (也就是说，Bit0=水平卷轴滚动256, Bit1=垂直卷轴滚动240) 2001h - PPU 控制寄存器 2 (只写) Bit7-5 颜色增强 (0=正常, 1-7=增强) Bit4 主角可见 (0=不显示, 1=显示) Bit3 背景可见 (0=不显示, 1=显示) Bit2 主角切除 (0=切除左边8个像素列, 1=不切除) Bit1 背景切除 (0=切除左边8个像素列, 1=不切除) Bit0 灰度模式 (0=彩色模式, 1=灰度模式) 为了设置PPU，我们进行两次写内存操作，代码为： lda #%00001000 sta $2000 lda #%00011110 sta $2001 我们写入$2000的值表示告诉PPU以下内容： NMI禁止了； 主角尺寸8x8； 背景Pattern Table起始地址$0000； 主角Pattern Table起始地址$1000； 地址自动增1； NameTable起始地址$2000； 写入$2001的值告诉PPU： 不要影响调色板； 显示主角； 显示背景； 显示所有位置的主角； 显示最左边8像素； 彩色模式； 你自己应该能从二进制码读懂上述具体含义。 【为了获得输出我们还需要做什么？】 如果你说的是背景，我们还要做： 制作一个背景； a. 绘制一些tile（用Tile Layer Pro）； b. 在段2首先包含上面的tile文件； c. 创建并载入一个调色板； d. 在Name Table中设置显示的tile号码 这样就显示了背景！ 如果是主角，那么顺序应该是这样： 制作一个主角： a. 在Tile Layer Pro中绘制主角； b. 段2中包含该文件作为第二个； c. 创建并载入调色板； d. 在主角数据区（类似GBA中的OAM）设置主角属性； 主角出场！ 我们明天学习怎样载入调色板。调色板很重要，一半是背景调色，一半是主角调色。
-
发表了文章 2014-11-15
NESASM教程——第二天——源文件的结构
【发生了什么？】 嗯，你是个NES小白程序员，今天我们将学习一个NESASM代码文件的结构。不幸的是，我们的汇编器对缩进非常讲究，一行的开始位置只能放标签，放其他所有内容都必须加一个【TAB】键缩进。尽管这样听上去很糟糕，但这样确实提高了可读性，让你更容易理解代码，尤其是大程序。 【关于段（Bank）】 不不不，Bank不会帮你拿着你的钱，它们帮你拿着你的程序和数据。我们将经常用三个段： 段 0 —— 放我们的代码，起始于$8000 段 1 —— 中断向量表，很重要，起始于$FFFA 段 2 —— 我们将主角和背景数据点阵信息放在这里，起始于$0000 我不确定一共有多少段，但显然至少3个。我们将用.bank指令来移动段，用.org指令来告诉汇编器在那个段我们的起始地址是什么。 【INES文件头】 INES文件头放在每个ROM文件的开头，告诉模拟器一些信息，它们是： .inesprg —— 告诉模拟器有多少个代码段 .ineschr —— 告诉模拟器有多少图片数据段 .inesmir —— 告诉模拟器……我忘了是什么，但总是1 .inesmap —— 我们总是用Mapper 0 我们常用设置为： .inesprg 1 ; 一个代码段 .ineschr 1 ; 一个数据段 .inesmap 0 ; 使用mapper 0 .inesmir 1 ; 总是1 这四行将放在（几乎所有）代码文件的最前面。 【段0和.org】 我们使用段0来放代码，起始地址为$8000。代码将这么写： .bank 0 ; 段 0. .org $8000 ; 去 $8000. ; 真正代码从这里开始 就这么多。注意分号（；）后面的内容为注释，汇编器忽略一行分号后面的所有内容。 【段1和三个中断向量】 不用长篇大论，直接来点美味小吃代码如何？ .bank 1 ; 切到段 1 .org $FFFA ; 从 $FFFA开始 .dw 0 ; NMI中断向量 .dw Start ; 复位向量，复位时从这个地址开始运行代码,我们给出了Start标签的地址，该标签最终放在段0中 .dw 0 ; VBlank 中断向量，目前我们不需要 段1内容就这么多，简单！ 【段2和图片数据】 段2，我们将从地址$0000开始，里面包含我们的图片数据，可用于背景和主角显示。代码如下： .bank 2 ; 切到段 2 .org $0000 ; 从 $0000开始 .incbin "our.bkg" ; 包含二进制文件，内容为我们的背景图片数据 .incbin "our.spr" ; 包含二进制文件，内容为我们的主角图片数据 【今天内容复习】 今天真的就这么多。我觉得有点慢了，因为NES编程相对GBA有点痛苦。明天我们会弄明白更多内容。明儿见！
-
发表了文章 2014-11-13
Caffe代码导读（3）：LevelDB例程
Caffe自带例子Cifar10中使用leveldb存储输入数据，为此我们研究一下怎样使用它。安装步骤可以参考http://blog.csdn.net/kangqing2003/article/details/6658345 Leveldb库提供了一种持续的键值对存储方式。键和值可以为任意字节数组。键存储顺序可由用户定义的比较函数决定。 打开一个数据库 Leveldb数据库有个与文件系统目录相对应的名字。数据库的所有内容都保存在这个目录中。下面例子展示了怎样打开一个数据库，必要时创建它： #include <assert> #include "leveldb/db.h" leveldb::DB* db; leveldb::Options options; options.create_if_missing = true; leveldb::Status status = leveldb::DB::Open(options,"/tmp/testdb", &db); assert(status.ok()); 如果你想在数据库已经存在情况下报错，只需要在leveldb::DB::Open调用前增加以下代码 options.error_if_exists = true; 状态 你可能注意到了上面的leveldb::Status类型。Leveldb中大多数可能遇到错误的函数返回该类型的值。你可以检查返回值是否为ok，必要时可打印相应的错误信息： leveldb::Status s = ...; if(!s.ok()) cerr << s.ToString() << endl; 关闭数据库 当你操作完一个数据库，只需delete掉数据库对象。例子： ...open the db as described above ... ... dosomething with db ... deletedb; 读和写 数据库提供Put，Delete和Get方法来修改/检索数据库。例如，下面代码将key1键下的值value移动到key2键下： std::string value; leveldb::Status s = db->Get(leveldb::ReadOptions(), key1,&value); if(s.ok()) s = db->Put(leveldb::WriteOptions(), key2, value); if(s.ok()) s = db->Delete(leveldb::WriteOptions(), key1); 原子更新 注意到如果进程在key2 Put操作后、key1 delete操作前终止，那么相同的值value可能留存在多个键下。这类问题可以使用WriteBatch类避免，该类可以原子地应用一系列更新： #include "leveldb/write_batch.h" ... std::string value; leveldb::Status s = db->Get(leveldb::ReadOptions(), key1,&value); if(s.ok()) { leveldb::WriteBatch batch; batch.Delete(key1); batch.Put(key2, value); s =db->Write(leveldb::WriteOptions(), &batch); } WriteBatch持有一系列针对数据库的编辑操作，这些操作将在一个batch内顺序执行。注意到我们在Put前调用Delete，这样如果key1恰好等于key2时，最终我们不会错误地丢掉整个value。 除了原子操作的优点，WriteBatch也可以用于加速批量更新操作，只需要将大量独立的改动操作放到同一个batch中。 同步写 默认情况下，每次写到leveldb都是异步的：进程一旦将写操作推送给操作系统就返回。操作系统内存到非易失存储的传输将异步发生。在某次写入中可将标志位sync使能，这样会使写操作直到数据写入非易失存储后才返回。（在采用了Posix的系统中，写操作返回前调用fsync(), fdatasync(),msync(…,MS_SYNC)）。 leveldb::WriteOptions write_options; write_options.sync = true; db->Put(write_options, ...);
-
发表了文章 2014-11-12
Caffe代码导读（2）：LMDB简介
闪电般的内存映射型数据库管理（LMDB） 简介 LMDB是基于二叉树的数据库管理库，建模基于伯克利数据库的应用程序接口，但做了大幅精简。整个数据库都是内存映射型的，所有数据获取返回数据都是直接从映射的内存中返回，所以获取数据时没有malloc或memcpy发生。因此该数据库仍是非常简单的，因为它不需要自己的页面缓存层，并且非常高效、省内存。它在语义上完全符合ACID（原子性、一致性、隔离性、持久性）。当内存映射为只读时，数据库完整性不会被应用程序的迷失指针写破坏。 该库也是线程可见的，支持来自多进程/线程的并发读/写访问。数据页使用写时复制策略，故没有活动数据页被覆盖写入。这也提供了保护机制，经历系统崩溃后不需要特殊恢复过程。写入过程为完全串行的；一次只有一个写会话是活动的，这保证了写入者不可能死锁。数据库结构是多个版本，所以读出者运行时不加锁。写入这不会阻塞读出者，读出者也不会阻塞写入者。 不像其他熟知的数据库机制（使用写前会话日志或数据仅追加写），LMDB操作时不需要保持会话。前面两种都需要周期性地检查或者压缩他们的日志或数据库文件，否则会无限增长。LMDB记录数据库内的空页面，在新的写入操作时重用他们，所以正常使用时数据库尺寸不会无限增加。 内存映射可以用作只读映射或读写映射。默认为只读映射，这提供了对破坏完全的免疫力。使用读写模式提供了更高的写性能，但增加了被恶意写入破坏数据库的可能性。当然如果你的应用代码是已知无bug的，那么这不是个严重的问题。
-
发表了文章 2014-11-12
NESASM教程——第六天——跳转和按键处理
【读取按键】 我们假设你要读取的是一个普通的方形NES手柄，而不是一些其他乱七八糟的东西。为了知道一个按键是否按下，你一次一键地读取0x4016（手柄1）或0x4017（手柄2）。如果按键按下，右边的bit0将被置位（1）。你用1同它相“与”，若结果不为0则跳转。在你读取任何内容之前，你需要先复位手柄（选通）。 【选通/复位】 为了选通/复位手柄，我们写入一个1然后一个0到0x4016（手柄1）或0x4017（手柄2）。可以这么搞： lda #$01 sta $4016 lda #$00 sta $4016 对，就这样！之后可以真正读取手柄信息。 【读取顺序】 在对0x4016或0x4017进行的的每个读取操作，你获得了不同按键的状态，这些被读出按键的顺序为： 读取序号 | 对应按键 1. A 2. B 3. SELECT 4. START 5. UP 6. DOWN 7. LEFT 8. RIGHT 那么在我们搞些事情之前先学习一些条件和跳转指令怎么样？ 【条件跳转】 我们在大多数指令中设置了一些特定“条件”，他们是： EQ——EQual-Zero， 等于0 NE——Not Equal，不等于0 LT——Less Than，小于0 GT——Greater Than， 大于0 PL——Plus，正数 MI——Minus，负数 CC——Carry Clear，进位标志为0 CS——Carry Set，进位标志为1 VC——oVerflow Clear，溢出标志为0 VS——oVerflow Set， 溢出标志为1 这些条件会被大多数指令设置，甚至包括数据载入。所有这些条件都有一个分支指令。在条件码前加上“B”就构成了条件跳转指令，例如： ;假设有个标签叫做Loopto beq Loopto;跳到Loopto如果上个比较结果为相等或者上条指令结果为0 bpl Loopto;跳到Loopto如果上次结果bit7为0 bmi Loopto;你懂的 我希望你已经发现规律了。我相信只有NE和EQ条件能被一个非比较指令影响。 以防万一你不知道，偷偷告诉你，一个label（标签）就是一个名称后面跟上冒号，例如 Loopto: aslkdfj: Hello_a: 我希望你懂了。记住：标签很好玩。 【读取按键的一个小栗子】 说够了。上代码： lda #$01 ; | sta $4016 ; \ lda #$00 ; - 设置手柄用来读 sta $4016 ; _/ lda $4016 ; 读取按键A and #1 ; 判断是否按下 bne WasDown ; 按下就走 ; 我不确定为什么是个BNE，但就是这样，所以就用吧！ lda $4016 ; 读取按键B lda $4016 ; 读取按键SELECT lda $4016 ; 读取按键START and #1 ; 看看是否按下了？ bne StartDown ; 按下了，走起~ lda $4016 ; UP lda $4016 ; DOWN lda $4016 ; LEFT lda $4016 ; RIGHT jmp NothingDown ; 滚远点 StartDown: ; 开始了开始了 WasDown: ; AAAAAAAAAAAAAAAA NothingDown: ; 啥都不干 希望你不认为这太难。 个人认为GBA是更容易编程的。所以对我而言，NES程序需要更多学习来搞明白。 【复习今天内容】 首先我想提一点，你可能知道我的GBA 汇编系列教程了，或许你也知道我的x86(DOS）汇编教程？在三门汇编语言中穿梭有时让人抓狂，所以如果你看到一些奇怪的或者错误，让我知道就行。谢谢。 按键检测不那么激情澎湃，但我可以告诉你后面更精彩。。。
-
发表了文章 2014-11-12
Caffe代码导读（1）：Protobuf例子
Protobuf是一种可以实现内存与外存交换的协议接口。这是由谷歌开发的开源工具，目前研究Caffe源码时用到。 一个软件项目 = 数据结构 + 算法 + 参数，对于数据结构和算法我们都已经有较多研究，但不同开发者对参数管理却各有千秋。有人喜欢TXT格式化的参数文件，有人喜欢BIN简单高效，也有人喜欢图形化界面的直观。不一致的参数管理带来很多问题，例如一个项目组内不同成员必须约定一套统一的参数方案，或者称为通信协议，这样便于模块集成。而Protobuf工具就完美解决了这个问题，关键部分代码自动生成，节省了大量的开发、调试时间。 首先下载protobuf，地址（打不开？……不解释） 这里用Linux版本2.5.0 解压： tar zxvf protobuf-2.5.0.tar.gz 切到主目录： cd protobuf-2.5.0 编译： ./configure make sudo make install 添加环境变量： export PKG_CONFIG_PATH=$(pwd) 编译examples: cd examples/ make cpp 这里我们只编译C++代码。 编译完成，生成了以下可执行文件： add_person_cpp list_people_cpp 这是个通讯录的例子。我们首先运行add_person_cpp： ./add_person_cpp zyk zyk: File not found. Creating a new file. Enter person ID number: 123 Enter name: zhaoyongke Enter email address (blank for none): zhaoyongke@yeah.net Enter a phone number (or leave blank to finish): 188188188 Is this a mobile, home, or work phone?（回车） Unknown phone type. Using default. Enter a phone number (or leave blank to finish):（回车） 然后运行list_people_cpp： ./list_people_cpp zyk Person ID: 123 Name: zhaoyongke E-mail address: zhaoyongke@yeah.net Home phone #: 188188188 可见我们生成了新的通讯录zyk，里面保存了相应的信息。 例子运行结束了，我们看下代码是如何生成的。 protobuf使用前，先编写proto文件，这是描述我们需要配置参数的数据结构。这个例子里面的proto如下： // See README.txt for information and build instructions. package tutorial; option java_package = "com.example.tutorial"; option java_outer_classname = "AddressBookProtos"; message Person { required string name = 1; required int32 id = 2; // Unique ID number for this person. optional string email = 3; enum PhoneType { MOBILE = 0; HOME = 1; WORK = 2; } message PhoneNumber { required string number = 1; optional PhoneType type = 2 [default = HOME]; } repeated PhoneNumber phone = 4; } // Our address book file is just one of these. message AddressBook { repeated Person person = 1; } 前几行是定义包的，可以忽略。 message Person{...}定义了一个需要传输的参数结构体，可见包括这么几个单元：name（string类型）、id（int32类型）、email（string类型）、phone（PhoneNumber类型，嵌套在Person内的类）。前面标记为“required”是必须有值的，而“optional“则为可选项，”repeated“表示后面单元为相同类型的一组向量。 有了如上定义，我们可以用protobuf工具生成接口代码，命令如下： protoc --cpp_out=. addressbook.proto 运行后生成了两个文件：addressbook.pb.cc 和addressbook.pb.h，代码比较长就不贴了。我们的应用程序可以通过自动生成的接口实现参数的序列化/反序列化，代码如下： //add_person.c #include <iostream> #include <fstream> #include <string> #include "addressbook.pb.h" using namespace std; // This function fills in a Person message based on user input. void PromptForAddress(tutorial::Person* person) { cout << "Enter person ID number: "; int id; cin >> id; person->set_id(id); cin.ignore(256, '\n'); cout << "Enter name: "; getline(cin, *person->mutable_name()); cout << "Enter email address (blank for none): "; string email; getline(cin, email); if (!email.empty()) { person->set_email(email); } while (true) { cout << "Enter a phone number (or leave blank to finish): "; string number; getline(cin, number); if (number.empty()) { break; } tutorial::Person::PhoneNumber* phone_number = person->add_phone(); phone_number->set_number(number); cout << "Is this a mobile, home, or work phone? "; string type; getline(cin, type); if (type == "mobile") { phone_number->set_type(tutorial::Person::MOBILE); } else if (type == "home") { phone_number->set_type(tutorial::Person::HOME); } else if (type == "work") { phone_number->set_type(tutorial::Person::WORK); } else { cout << "Unknown phone type. Using default." << endl; } } } // Main function: Reads the entire address book from a file, // adds one person based on user input, then writes it back out to the same // file. int main(int argc, char* argv[]) { // Verify that the version of the library that we linked against is // compatible with the version of the headers we compiled against. GOOGLE_PROTOBUF_VERIFY_VERSION; if (argc != 2) { cerr << "Usage: " << argv[0] << " ADDRESS_BOOK_FILE" << endl; return -1; } tutorial::AddressBook address_book; { // Read the existing address book. fstream input(argv[1], ios::in | ios::binary); if (!input) { cout << argv[1] << ": File not found. Creating a new file." << endl; } else if (!address_book.ParseFromIstream(&input)) { cerr << "Failed to parse address book." << endl; return -1; } } // Add an address. PromptForAddress(address_book.add_person()); { // Write the new address book back to disk. fstream output(argv[1], ios::out | ios::trunc | ios::binary); if (!address_book.SerializeToOstream(&output)) { cerr << "Failed to write address book." << endl; return -1; } } // Optional: Delete all global objects allocated by libprotobuf. google::protobuf::ShutdownProtobufLibrary(); return 0; } 可见只需要调用addressbook.pb.h中声明的tutorial::AddressBook类、Person类中的接口（add_person(), add_phone(), set_number(), set_email()等）就能操作相应的参数，最后将内存中的参数序列化为文件只需要执行SerializeToOstream()。相应的读取参数文件的操作为ParseFromIstream()。这里贴出例子中的第二个程序如下： // list_people.c #include <iostream> #include <fstream> #include <string> #include "addressbook.pb.h" using namespace std; // Iterates though all people in the AddressBook and prints info about them. void ListPeople(const tutorial::AddressBook& address_book) { for (int i = 0; i < address_book.person_size(); i++) { const tutorial::Person& person = address_book.person(i); cout << "Person ID: " << person.id() << endl; cout << " Name: " << person.name() << endl; if (person.has_email()) { cout << " E-mail address: " << person.email() << endl; } for (int j = 0; j < person.phone_size(); j++) { const tutorial::Person::PhoneNumber& phone_number = person.phone(j); switch (phone_number.type()) { case tutorial::Person::MOBILE: cout << " Mobile phone #: "; break; case tutorial::Person::HOME: cout << " Home phone #: "; break; case tutorial::Person::WORK: cout << " Work phone #: "; break; } cout << phone_number.number() << endl; } } } // Main function: Reads the entire address book from a file and prints all // the information inside. int main(int argc, char* argv[]) { // Verify that the version of the library that we linked against is // compatible with the version of the headers we compiled against. GOOGLE_PROTOBUF_VERIFY_VERSION; if (argc != 2) { cerr << "Usage: " << argv[0] << " ADDRESS_BOOK_FILE" << endl; return -1; } tutorial::AddressBook address_book; { // Read the existing address book. fstream input(argv[1], ios::in | ios::binary); if (!address_book.ParseFromIstream(&input)) { cerr << "Failed to parse address book." << endl; return -1; } } ListPeople(address_book); // Optional: Delete all global objects allocated by libprotobuf. google::protobuf::ShutdownProtobufLibrary(); return 0; } 相信做完这个实验，你将不再对Caffe代码中的参数初始化、参数保存操作感到陌生，一切都很自然。 除了上述简单功能，Protobuf还可以用来传递不同语言(C/C++与Java、Python)之间的参数，省去了自己手动维护数据结构的繁琐工作。也可以支持客户端/服务器模式，在主机/从机之间传递参数。
-
发表了文章 2014-11-10
NES文件利用MATLAB可视化
NES是Nintendo Entertainment System的缩写，记录了NES小游戏的所有代码和数据。像超级玛丽、忍者龙剑传、热血格斗、007等游戏都有精彩纷呈的背景图片和形象生动的人物造型，我们是否能提取出这些素材，经过加工，用于其他UI设计呢？今天我们用MATLAB研究下具体内容。 NES文件结构分为3大部分：文件头、CPU代码区、PPU数据区。通过文件头可以获得代码区、数据区的大小。这里只需要数据区。 数据区组织为4KB大小的数据块，填充在PPU地址区间0x0000~0x1000（或者0x1000~0x2000），称为Pattern Table，其实就是一个个8x8的点阵数据，用户程序写入Name Table、SPR-RAM中相应的索引就能实现在屏幕指定位置上描绘相应的点阵信息，组合多个点阵可以实现复杂背景和人物造型的显示。这里只将原始点阵描出。代码如下： clear; clc; close all; path = './rom/'; [nes_file_name, path, filterindex] = uigetfile('*.nes', 'Pick an NES file',path);%选择NES文件 fid = fopen(fullfile(path,nes_file_name),'rb'); rawdata = fread(fid,'uint8'); fclose(fid); len_header = 16;<span style="white-space:pre"> </span>%文件头，16Bytes len_pgm = 1024*16; %16KB len_ppu = 1024*8; %8KB header = rawdata(1:len_header); if and(and(and((header(1) == uint8('N')),header(2) == uint8('E')),header(3) == uint8('S')),header(4) == 26) fprintf('NES File Loaded...\n'); else fprintf('Not an NES File!\n'); return; end num_pgm = header(5); num_ppu = header(6); pgm = rawdata(len_header + (1:len_pgm)); for ppu_idx = 1:num_ppu ppu = rawdata(len_header + len_pgm * num_pgm + (ppu_idx - 1) * len_ppu + (1:len_ppu)); ppu1 = ppu(0 + (1:len_ppu/2)); % 4KB ppu2 = ppu(len_ppu/2 + (1:len_ppu/2)); % 4KB ppu1 = reshape(ppu1,16,[]); ppu2 = reshape(ppu2,16,[]); showimag1 = zeros(128,128); showimag2 = zeros(128,128); and_matrix = ones(8,1)*2.^(7:-1:0); zz1 = zeros(8,8); zz2 = zeros(8,8); for k = 1:size(ppu1,2) x = mod(k-1,16); y = floor((k-1)/16); rdata1l = logical(bitand(ppu1(1:8,k)*ones(1,8),and_matrix)); rdata1h = logical(bitand(ppu1(9:16,k)*ones(1,8),and_matrix)); rdata2l = logical(bitand(ppu2(1:8,k)*ones(1,8),and_matrix)); rdata2h = logical(bitand(ppu2(9:16,k)*ones(1,8),and_matrix)); zz1 = rdata1l + rdata1h * 2; zz2 = rdata2l + rdata2h * 2; showimag1((y)*8 + (1:8), (x)*8 + (1:8)) = zz1 * 64; showimag2((y)*8 + (1:8), (x)*8 + (1:8)) = zz2 * 64; end figure;subplot(121);imshow(uint8(showimag1));title('Pattern Table #1'); subplot(122);imshow(uint8(showimag2));title('Pattern Table #2'); end 下面欣赏几幅游戏中提取出来的Pattern Table
-
发表了文章 2014-11-10
一位工程师的FPGA项目开发经验总结
转自：http://bbs.ednchina.com/BLOG_ARTICLE_3018248.HTM?click_from=8800024401,6106445608,2014-11-10,EDNCOL,NEWSLETTER 1. 要和人配合 以我们做硬件的工程师为例，测试的时候一般都需要软件的配合，一个对硬件来说无比复杂的工作，可能在软件工程师看来就是几行简单的代码。所以要和人配合，多听听别人的意见，这样必然可以产生新的 know-how，从而加快测试和开发的速度，退一步讲，至少没有坏处。 2. 测试还是要别人来做 开发者看待自己的产品有如看待自己，大多是没有勇气去发现缺点的。一是源自自尊心，二是为了避免额外的工作。所以就算有问题，如果不严重就藏着掖着。但是这对项目来说是不行的，所以测试，verification，一定要旁人来做。 3. 多点时间思考 出现问题后，不要急着修改。要思考推测可能的原因，想清楚后把这些可能的原因都用debug pin或者chipscope引出来。 4. 注意复用已有的debug pin 很多时候，在测试过程中产生了一大堆测试信号，但是时间一长就忘了复用。实际上，当一个问题产生的时候，通过反复观察已有的debug-pin或许足以发现问题根源，而无需再引出新的pin，并浪费时间去综合和PAR。 5. 仿真加时序足矣 数字电路在时钟同步的设计原则下，其功能通过simulation就可以验证。simulation的结果和PAR后产生的FPGA-image 完全等价。当然FPGA也要遵循同样的设计原则，即时钟同步。所以对于PAR的结果首先就要确保其时钟同步的特性。体现为寄存器之间的path必须在一个 时钟周期内完成。（当然有其他约束的例外。）同时要满足FPGA器件的setup和hold要求。一旦出现timing-error必须通过各种途径消除 error，因为error的存在，意味着时钟同步的大前提已经被破坏，这时，simulation取得的结果和FPGA是不等价的，继续测试也毫无意义 了。 6. 注意不可控的接口部 分 FPGA内部的寄存器之间的timing完全可以通过PAR报告来确认是否有问题。但是和外界的接口部分却充满了疑问。我们一般通过假定的 input-delay和output-delay来对接口部分进行约束。由于从一开始就施加的是假定的delay，所以即使没有timing- error，其结果也存在诸多疑问。以我正在进行的测试为例，模块内 部loopback测试完全正常，但是一过cable，传到对方FPGA，则马上产生很多误码。由于simulation没有问题，所以必然是我们的某个 假定出现了问题，尤其是时钟同步的假定会得不到满足。这时候，就要想尽一切办法，使接口也满足假定的条件，或者调整设计，将不理想的接口adapting 成理想的接口。 7. 向直接上司汇报情况，寻求各种可能的许可。 懒得向直接上司汇报情况时，万一出现进度或者结果不符，所有责任都需要本人承担。如果提前向上司汇报情况并取得 许可，则一切后果都在可控范围内。比如，工作繁忙时又被派给新的任务，则不能一味逆来顺受。应该向上司说明困难，并提前想好一个可行的解决方案供上司参考。 8. 外部接口是最大障碍 如前所述，FPGA内部如果timing没有问题的话，一般和仿真结果是一致的，问题是外部的接口，包括cable连线等，不在我们 确切控制的范围内，比如其延时特性在40MHz下仍然正常，但是在80MHz时可能出现不可预料的情况。所以应该尽量使用经过验证的“cable- frequency”组合。或者通过设备测量并确认外部接口的延时特性。这样可以进行有针对性的调整。我最近的教训就是花了整整一个月调整并测试内部的结 构，但是仍然失败。结果发现由于cable的问题，80MHz的信号（数据+使能+others）无法正常并行传输。如果换成40MHz的信号就通过了。 9. 综合PR后的结果要和代码等价 前面提到仿真加时序足矣，这里面的前提是PR的结果和原始代码要等价。为了确认这一点，就要把握syn和pr过程中的所有 warning以及error，warning的内容不是完全可以忽略的。要特别关注综合报表中的以下内容：unused ports, removal of redundant logic, latch inference，simulation mismatch等等。在报表中输入关键字查找即可。
-
发表了文章 2014-11-07
利用C++ Boost编写扩展Python模块
Python很强大，但已有的模块可能满足不了人民日益增长的物质文化需求，于是有时需要编写扩展模块进行完善。 可行的方案有很多：SWIG、Weave、ctypes、BOOST…… BOOST无疑是开发最快的一种方案。下面介绍下最简单的C++ helloworld程序如何变为Python的一个模块。 1. 安装Python、Boost 这里用Linux环境。Python和Boost都用源码安装，网址为： Python2.6：https://www.python.org BOOST1.57.0：http://sourceforge.net/projects/boost/?source=typ_redirect 2. 编写helloworld.cpp #define BOOST_PYTHON_SOURCE #include <boost/python.hpp> #include <iostream> using namespace std; using namespace boost::python; void hello_func() { cout<<"hello boost python"<<endl; } BOOST_PYTHON_MODULE(boostpy) { def("Hello", hello_func, "Function 's targets..."); } 3. 编译为动态库 命令行中执行： g++ -shared -o boostpy.so -fPIC -I/YourPythonIncludePath/ helloworld.cpp -lpython2.6 -lboost_python 生成了动态链接库boostpy.so 4. Python环境中调用Hello >>> import boostpy >>> boostpy.Hello() hello boost python >>>help(boostpy)Help on module boostpy: NAME boostpy FILE /...../boostpy.so FUNCTIONS Hello(...) Hello() -> None : Function 's targets... C++ signature : void Hello() (END) 总结：这里只做了个最简单的调用，没有参数传递的问题。后面会继续研究怎样在C++和Python之间共享数据。
-
发表了文章 2014-11-06
Python入门基础
1. Python支持C/C++的运算符有： = += -= *= /= %= Python特有的运算符： **（乘方） **=（乘方赋值） 不支持的运算符：前缀、后缀形式的++，-- 2. Python数值类型： int long float complex 3. 字符串 单引号、双引号里面的内容都是字符串。 第一个字符索引下标为0； 最后一个字符索引下标为-1； 字符串可以用“+”进行拼接； >>> str1 = 'hello' >>> str2 = 'world' >>> str1+str2 'helloworld' 字符串可以用“*”进行重复； >>> str1*5 'hellohellohellohellohello' >>> str1[0] 'h' >>> str1[-1] 'o' >>> str1[0:-1] 'hell' 4. 列表(list)和表列(tuple) 这是Python里自带容器。里面可以承载任意个数的Python对象。 列表用[]定义，其元素和长度可以变化。如 >>> alist = [1,2,3,4] >>> alist [1, 2, 3, 4] >>> alist[0:-1] [1, 2, 3] >>> alist[2:] [3, 4] >>> alist[:] [1, 2, 3, 4] >>> alist[:3] [1, 2, 3] >>> alist[1] = 5 >>> alist [1, 5, 3, 4] 表列用()定义，元素不能再次赋值。其他操作同列表。如 >>> atuple = ('robots',77,93,'try') >>> atuple ('robots', 77, 93, 'try') >>> atuple[0] 'robots' 5. 字典 字典就是Python里的hash table数据类型。用{}定义，内容包含key 和 value。key通常是数字或字符串。value可为任意基本数据类型。 >>> adict = {} >>> adict['host'] = 'earth' >>> adict['port'] = 80 >>> adict {'port': 80, 'host': 'earth'} >>> adict.keys() dict_keys(['port', 'host']) >>> adict['host'] 'earth' >>> adict.values() dict_values([80, 'earth']) 6. 缩进风格 Python代码段是用缩进标识的。 7. 文件操作 handle = open(file_name,access_mode = 'r') 几种模式： r:只读（默认） w:只写 a:append,添加内容到末尾 +：读写 b：二进制方式 8. 异常 try :try_block except someError:processing_block raise：明确引发一个异常。 9. 函数 def function_name([arguments]):'optional documentation string'function_suite ##注意：函数所有参数都是以引用方式传递的，因此参数在函数中的任何变化都会影响到原始对象。（貌似3.3有变化了） def adder2(x,y) :return (x+y) 参数允许默认值。 10. 类 class class_name[(base_classes_if_any)]:"optional documentation string"static_member_declarationsmethod_declarations 一个例子： class Cat :"Hello world, I am a cat, miao miao miao..."count = 0def __init__(self, str = "Xiaomei") :'Constructor'self.name = strcount += 1print("My name is : %s" % str)def show_info(self)'Show some class info'print("My class name is :%s" % self.__class__)print("My name is : %s" % self.name)print("I have %d companies" % self.count) __class__内建变量显示类名，这里是__main__.Cat __doc__内建变量显示类声明时的documentation string 11. 模块 import 模块名 模块名.函数() 模块名.变量()
-
发表了文章 2014-11-06
NumPy学习笔记（1）
1. 程序开头导入库 import numpy as np 2. 创建数组(ndarray) a = [1,2,3,4] aa = np.array(a)--------->生成一维数组 bb = np.array([a,a])----->生成二维数组 aa.shape bb.shape------------------>获得尺寸(2, 4) bb.shape = (1,8) ---------> Reshape为1X8数组，8可以用-1表示，这样自动计算尺寸 c = bb.reshape(2,4)------>用reshape函数获得新数组，注意这里c与bb共享存储空间。若想隔离，需要用copy()函数 3. 数组的类型 bb.dtype------------------------>获得bb的数值类型，可能为int32, int64, float64, complex128等 生成不同数值类型的数组： d = np.array([1,2,3,4], dtype=np.complex) 完整的类型列表： >>> set(np.typeDict.values()) set([<type 'numpy.bool_'>, <type 'numpy.int8'>, <type 'numpy.void'>, <type 'numpy.int32'>, <type 'numpy.uint32'>, <type 'numpy.float128'>, <type 'numpy.string_'>, <type 'numpy.uint8'>, <type 'numpy.float32'>, <type 'numpy.complex256'>, <type 'numpy.int64'>, <type 'numpy.uint64'>, <type 'numpy.complex64'>, <type 'numpy.unicode_'>, <type 'numpy.uint64'>, <type 'numpy.int16'>, <type 'numpy.uint16'>, <type 'numpy.float64'>, <type 'numpy.object_'>, <type 'numpy.complex128'>, <type 'numpy.int64'>]) 4. 几个专门创建数组的函数 np.arrange(start, end, step)---------------->生成等差数列，不包括end值 np.linspace(start, end, num)--------------------------->生成等差数列，默认包括end值 np.linspace(start, end, num, endpoint=False)-------->不包括end值 np.logspace(start, end, num, base, endpoint)------------------------->生成等比数列 np.zeros(size, dtype)， np.zeros_like(other) np.ones(size, dtype)，np.ones_like(other) np.random.rand(size, dtype)------------------------------------>生成随机数组 np.empty(size, dtype)--------------------------------------------->仅分配空间，不初始化，最快！ np.fromstring(str, dtype)-------------------------------------------->从字符串创建数组 np.frombuffer() np.fromfile() np.fromfunction(funcname, size) 5. 数组的切片和下标访问 一维数组的索引如下： 0,1,2，。。。-3，-2，-1，即负数会从后往前数 a[start:end:step]----------------------->取数组的一个切片，与原数组共享内存。不包括end值。 省略start，表示start = 0； 省略end，表示end = -1； 省略step，表示step = 1； 单独生成切片可以用slice函数 slice(start, end, step) = start:end:step slice(None, end, step) = :end:step slice(None, None, None) = : 另外也可以用s_对象生成slice >>> np.s_[::2,2:] (slice(None, None, 2), slice(2, None, None)) 二维数组取一次下标，得到一维数组（想想C中的数组） 多次下标可以用元组方式，如x[2][2]等价于x[2,2]，或者直观一点x[(2,2)] 整数索引列表查表 idx = [1, 4, 5, 7] y = x[idx]------------------->取出x中[1,4,5,7]位置处的值，放入y，不和原数组共享内存 整数数组也是一样的 布尔数组（np.array([True, False, True,...])）：只取数组中相应位置为True的元素； 布尔列表[True, False, True, ...]：将按照整数方式取， True = 1, False = 0.
-
发表了文章 2014-11-06
python 函数参数的传递(参数带星号的说明)
原创地址：http://www.cnblogs.com/smiler/archive/2010/08/02/1790132.html python中函数参数的传递是通过赋值来传递的。 函数参数的使用又有俩个方面值得注意： 1.函数参数是如何定义的 2.在调用函数的过程中参数是如何被解析 先看第一个问题，在python中函数参数的定义主要有四种方式：1.F(arg1,arg2,...)这 是最常见的定义方式，一个函数可以定义任意个参数，每个参数间用逗号分割，用这种方式定义的函数在调用的的时候也必须在函数名后的小括号里提供个数相等的 值（实际参数），而且顺序必须相同，也就是说在这种调用方式中，形参和实参的个数必须一致，而且必须一一对应，也就是说第一个形参对应这第一个实参。例 如： def a(x,y): print x,y 调用该函数，a(1,2)则x取1，y取2，形参与实参相对应，如果a(1)或者a(1,2,3)则会报错。 2.F(arg1,arg2=value2,...) 这种方式就是第一种的改进版，提供了默认值 def a(x,y=3): print x,y 调用该函数，a(1,2)同样还是x取1，y取2，但是如果a(1)，则不会报错了，这个时候x还是1，y则为默认的3。上面这俩种方式，还可以更换参数位置，比如a(y=8,x=3)用这种形式也是可以的。 3.F(*arg1) 上 面俩个方式是有多少个形参，就传进去多少个实参，但有时候会不确定有多少个参数，则此时第三种方式就比较有用，它以一个*加上形参名的方式来表示这个函数 的实参个数不定，可能为0个也可能为n个。注意一点是，不管有多少个，在函数内部都被存放在以形参名为标识符的tuple中。 >>> def a(*x): if len(x)==0: print 'None' else: print x >>> a(1) (1,) #存放在元组中 >>> a() None >>> a(1,2,3) (1, 2, 3) >>> a(m=1,y=2,z=3) Traceback (most recent call last): File "<pyshell#16>", line 1, in -toplevel- a(m=1,y=2,z=3) TypeError: a() got an unexpected keyword argument 'm' 4.F(**arg1) 形参名前加俩个*表示，参数在函数内部将被存放在以形式名为标识符的dictionary中，这时调用函数的方法则需要采用arg1=value1,arg2=value2这样的形式。 >>> def a(**x): if len(x)==0: print 'None' else: print x >>> a() None >>> a(x=1,y=2) {'y': 2, 'x': 1} #存放在字典中 >>> a(1,2) #这种调用则报错 Traceback (most recent call last): File "<pyshell#25>", line 1, in -toplevel- a(1,2) TypeError: a() takes exactly 0 arguments (2 given) 上面介绍了四种定义方式，接下来看函数参数在调用过程中是怎么被解析的,其实只要记住上面这四种方法优先级依次降低，先1，后2，再3，最后4，也就是先把方式1中的arg解析，然后解析方式2中的arg=value，再解析方式3，即是把多出来的arg这种形式的实参组成个tuple传进去，最后把剩下的key=value这种形式的实参组成一个dictionary传给带俩个星号的形参，也就方式4。 >>> def test(x,y=1,*a,**b): print x,y,a,b >>> test(1) 1 1 () {} >>> test(1,2) 1 2 () {} >>> test(1,2,3) 1 2 (3,) {} >>> test(1,2,3,4) 1 2 (3, 4) {} >>> test(x=1,y=2) 1 2 () {} >>> test(1,a=2) 1 1 () {'a': 2} >>> test(1,2,3,a=4) 1 2 (3,) {'a': 4} >>> test(1,2,3,y=4) Traceback (most recent call last): File "<pyshell#52>", line 1, in -toplevel- test(1,2,3,y=4) TypeError: test() got multiple values for keyword argument 'y'
-
发表了文章 2014-03-18
从研究生到程序员
研究生阶段主要任务是掌握学习方法，针对课题进行资源搜索，整合，处理，综合性很强，有些问题比较复杂，需要多方查证，用理论铺平道路，设计实验，选择合适的仿真参数，最后才到编码实现。研究生的主要评估指标就是论文，包括会议论文、期刊论文、专利、报告PPT等，对逻辑性要求较高，需要对自己对课题从背景到应用各个层次都有深入理解。 毕业后马上就要成为程序员。程序员的评估指标，首先就是编程语言掌握程度，研究生期间主要使用matlab，对于其他语言多少都有点陌生了，而在工作中，要求效率，要求准确性，要求性能，这是在校期间很少或较少关注的，应有所准备。说白了，程序员就要做一些实现层面的事，这些事在研究生在学校期间所占比重很低，因此思维要有所转变。如果能保持自己在工作之余积极看书，看感兴趣领域的最新论文，那么对自己的知识更新也是有积极意义的。 理解代码，首先理解架构。学习算法，首先了解历史背景。开始可能进度很慢，但后面速度会越来越高。
-
发表了文章 2014-01-27
ZED Board从入门到精通系列（七）——Vivado+SDK实现MP3播放
本文将给出通过Vivado IDE开发Zynq平台上PS裸机应用程序的流程。通过与本系列博客（三）对比，读者将看到Vivado开发更高效、快捷。 MP3我们都听过，现在我们可以用ZED-Board来听。板子上有音频芯片ADAU1761，可以实现录音、放音，但不具有MP3解码功能。Zynq 双核ARM9做MP3软件解码应该是可以实现的，但是博主本人有一颗VS1003，可以实现MP3硬件解码，软件将得以简化，对MP3解码原理感兴趣的可以深入研究如何利用CortexA9+ADAU1761实现MP3播放。电路图如下： 利用Zynq MIO实现VS1003控制，这样只和PS有关，PL完全可以丢弃。在本节基础上，读者可以尝试将SPI模块移到PL上实现，这样可以降低PS部分IO读写频率，提高CPU利用率。实物连接图如下： Zynq板子外接用排母，为了使用杜邦线，需要一个双公排针，可以用普通单排2.54mm排针压制而成 下面介绍软件开发流程。建立Vivado工程，命名为MP3Player，过程遵循上节Vivado建立工程步骤，略。 进入IDE后，点击左侧流程管理器中的IPI Integrator下的Create Block Design。 这个工具是2013.1版本后才出现的，将取代XPS完成系统集成。 在编辑区右键，选择Add IP...，名称保持默认design_1.bd 搜索框中输入zynq，双击第一个，添加IP到电路图中。 添加完成后，自动进行布线连接，点下图中圆圈区域 Run Block Automation。 等待完成，结果如下图所示。 可以看到，DDR和固定IO自动进行了连接。这是因为我们建立工程时选择了ZedBoard DVK，这样就能按照板子描述自动连接引脚到相应外设。 另外看到，默认状态下使能了M_AXI_GP0，可以将PL部分带AXI从接口的IP连接到PS进行控制。本节不需要，所以必须禁用，否则验证设计时会报错。双击方块，见下图 看到了熟悉又陌生的画面，有些像XPS中Zynq视图，但精简了很多。单击左侧“PS-PL Configuration"，界面如下： 将AXI GP0接口后的勾取消选择，确认，回到IPI。 验证设计，在空白处右键，点击Validate Design。无误，点确认即可。 在上图位置点Generate Block Design，确认。 在Sources窗口中找到design_1，右键选择生成顶层HDL包装。确认。 直接点左侧流程中的Generate Bitstream，一步到位。完成比特流大约需要5~8min。 完成后，先Open Implementated Design，再导出到SDK。 完成后，先Open Implementated Design，再导出到SDK。如果没有做这一步，上图中第二项会变成灰色。 后面就是SDK开发了，和本系列教程（三）中相同。建立Application工程，C工程，模板helloworld。将代码改为下面： #include <stdio.h> #include "platform.h" #define MIO_BASE 0xE000A000 #define DATA0 0x40 #define DATA0_RO 0x60 #define DIRM_0 0x204 #define OEN_0 0x208 void delay(unsigned int t) { unsigned int i,j; for(j=0;j<t;j++) { for(i=0;i<600;i++); } } /*---------------------------------------------------------------------------------------------------------*/ /* MAIN function */ /*---------------------------------------------------------------------------------------------------------*/ #define VS_XRESET_0 DrvGPIO_ClrBit(MIO_BASE + DATA0,12) #define VS_XRESET_1 DrvGPIO_SetBit(MIO_BASE + DATA0,12) #define VS_DREQ DrvGPIO_GetBit(MIO_BASE + DATA0_RO,11) #define VS_XDCS_0 DrvGPIO_ClrBit(MIO_BASE + DATA0,10) #define VS_XDCS_1 DrvGPIO_SetBit(MIO_BASE + DATA0,10) #define VS_XCS_0 DrvGPIO_ClrBit(MIO_BASE + DATA0,13) #define VS_XCS_1 DrvGPIO_SetBit(MIO_BASE + DATA0,13) #define SPI_MOSI_0 DrvGPIO_ClrBit(MIO_BASE + DATA0,0) #define SPI_MOSI_1 DrvGPIO_SetBit(MIO_BASE + DATA0,0) #define SPI_SCL_0 DrvGPIO_ClrBit(MIO_BASE + DATA0,9) #define SPI_SCL_1 DrvGPIO_SetBit(MIO_BASE + DATA0,9) void DrvGPIO_ClrBit(volatile unsigned int * p,int idx); void DrvGPIO_SetBit(volatile unsigned int * p,int idx); unsigned char DrvGPIO_GetBit(volatile unsigned int * p,int idx); void init_vs1003(void); void VS_Reset(void); //VS1003软复位及初始化 void VS_Write_Reg(unsigned char addr,unsigned char hdat,unsigned char ldat); //向VS1003的功能寄存器写入一个字 unsigned int VS_Read_Reg(unsigned char addr); //从VS1003的功能寄存器读取一个字 void VS_Send_Dat(unsigned char dat); //向VS1003发送音频数据 void VS_Flush_Buffer(void); //清空VS1003的数据缓冲区 void VS_sin_test(unsigned char x); //正弦测试 void LoadPatch(void); //为VS1003打补丁 void SPI_WriteByte(unsigned char x); #include "mp3.h" void print(char *str); int main() { init_platform(); print("Hello World\n\r"); unsigned int i; init_vs1003(); VS_Reset(); //VS1003复位初始化 VS_sin_test(200); //正弦测试，可以听到一声滴 VS_Flush_Buffer(); for(i = 0;i<sizeof(mp3_table);i++) { VS_Send_Dat(mp3_table[i]); } while(1) { DrvGPIO_ClrBit(MIO_BASE + DATA0,7); delay(40000); DrvGPIO_SetBit(MIO_BASE + DATA0,7); delay(40000); } return 0; } void DrvGPIO_ClrBit(volatile unsigned int * p,int idx) { (*p) &= ~(1<<idx); } void DrvGPIO_SetBit(volatile unsigned int * p,int idx) { (*p) |= (1<<idx); } unsigned char DrvGPIO_GetBit(volatile unsigned int * p,int idx) { return (((*p)&(1<<idx))>>idx); } void init_vs1003(void) { DrvGPIO_SetBit(MIO_BASE + OEN_0,7); DrvGPIO_SetBit(MIO_BASE + DIRM_0,7); DrvGPIO_SetBit(MIO_BASE + OEN_0,0); DrvGPIO_SetBit(MIO_BASE + DIRM_0,0); DrvGPIO_SetBit(MIO_BASE + OEN_0,9); DrvGPIO_SetBit(MIO_BASE + DIRM_0,9); DrvGPIO_SetBit(MIO_BASE + OEN_0,10); DrvGPIO_SetBit(MIO_BASE + DIRM_0,10); DrvGPIO_SetBit(MIO_BASE + OEN_0,12); DrvGPIO_SetBit(MIO_BASE + DIRM_0,12); DrvGPIO_SetBit(MIO_BASE + OEN_0,13); DrvGPIO_SetBit(MIO_BASE + DIRM_0,13); } void SPI_WriteByte(unsigned char x) { unsigned char i=0; for(i=0;i<8;i++) { if(x&0x80) { SPI_MOSI_1; } else { SPI_MOSI_0; } SPI_SCL_0; SPI_SCL_1; x<<=1; } } /****************************************************************** - 功能描述：向VS1003的功能寄存器中写入数据（一个字，即两个字节） - 隶属模块：VS1003B模块 - 函数属性：外部，用户可调用 - 参数说明：addr是功能寄存器的地址 hdat是要写入的高字节 ldat是要写入的低字节 - 返回说明：无返回 ******************************************************************/ void VS_Write_Reg(unsigned char addr,unsigned char hdat,unsigned char ldat) { while(!VS_DREQ); //VS1003的DREQ为高电平时才接收数据 VS_XCS_0; //打开片选，SCI有效，这样才能对功能寄存器进行读写 SPI_WriteByte(0x02); //写入操作码0x02 00000010 （功能寄存器写操作） SPI_WriteByte(addr); //写入寄存器地址 SPI_WriteByte(hdat); //写入高字节 SPI_WriteByte(ldat); //写入低字节 VS_XCS_1; //关闭片选，SCI无效 } /****************************************************************** - 功能描述：VS1003软复位及初始化（设置时钟频率及音量） - 隶属模块：VS1003B模块 - 函数属性：外部，用户可调用 - 参数说明：无 - 返回说明：无 ******************************************************************/ void VS_Reset(void) { VS_XRESET_1; delay(100); VS_XRESET_0; delay(100); VS_XRESET_1; //硬件复位，XRESET低电平有效 delay(100); VS_Write_Reg(0x00,0x08,0x04);//软件复位，向0号寄存器写入0x0804 SM_SDINEW为1 SM_RESET为1 VS_Write_Reg(0x03,0x98,0x00);//时钟设置，向3号寄存器写入0x9800 SC_MULT 为4 SC_ADD 为3 SC_FREQ为0 VS_Write_Reg(0x0b,0x00,0x00);//音量设置，左右声道均最大音量 VS_XDCS_0; //打开数据片选，注意此时XCS（片选）为高电平，SDI有效 SPI_WriteByte(0); //写入数据，这里写入4个0，是无关数据，用来启动数据传输 SPI_WriteByte(0); SPI_WriteByte(0); SPI_WriteByte(0); VS_XDCS_1; //关闭数据片选，SDI无效 } /****************************************************************** - 功能描述：向VS1003写入一个字节的音频数据（即用于播放的数据） 注：调用前先将VS_XDCS置为0，打开数据片选 - 隶属模块：VS1003B模块 - 函数属性：外部，用户可调用 - 参数说明：dat是要写入的字节 - 返回说明：无 ******************************************************************/ void VS_Send_Dat(unsigned char dat) { VS_XDCS_0; //打开SDI，此时可以向VS1003写入音频数据 while(!VS_DREQ); //VS1003的DREQ为高才能写入数据 SPI_WriteByte(dat);//通过SPI向VS1003写入一个字节的音频数据 VS_XDCS_1; //关闭SDI } /****************************************************************** - 功能描述：向VS1003写入2048个0，用于清空VS1003的数据缓冲区 注：在播放完一个完整的音频（如一首完整的MP3）后，调用 此函数，清空VS1003数据缓冲区，为下面的音频数据（如下 一首MP3）作准备。 - 隶属模块：VS1003B模块 - 函数属性：外部，用户可调用 - 参数说明：无 - 返回说明：无 ******************************************************************/ void VS_Flush_Buffer(void) { unsigned int i; VS_XDCS_0; //打开数据片选，即开启SDI传输 for(i=0;i<2048;i++) { VS_Send_Dat(0); } VS_XDCS_1; //关闭数据片选 } /****************************************************************** - 功能描述：正弦测试，这是测试VS1003芯片是否正常的有效手段！！ - 隶属模块：VS1003B模块 - 函数属性：外部，用户可调用 - 参数说明：x决定了正弦测试中产生的正弦波的频率，直接影响听到的 声音的频率 - 返回说明：无 ******************************************************************/ void VS_sin_test(unsigned char x) { VS_Write_Reg(0x00,0x08,0x20);//启动测试，向0号寄存器写入0x0820 SM_SDINEW为1 SM_TEST为1 while(!VS_DREQ); //等待DREQ变为高电平 VS_XDCS_0; //打开数据片选 SDI有效 SPI_WriteByte(0x53);//写入以下8个字节,进入正弦测试 SPI_WriteByte(0xef); SPI_WriteByte(0x6e); SPI_WriteByte(x); //参数x用来调整正弦测试中正弦波的频率 FsIdx (b7~b5):采样率表索引 S (b4~b0):正弦波的跃速 频率F=Fs X S / 128 SPI_WriteByte(0); //比如x=126 (0b 011 11110) FsIdx=011=3 Fs=22050Hz S=11110=30 F=22050Hz X 30 /128 =5168 Hz SPI_WriteByte(0); SPI_WriteByte(0); SPI_WriteByte(0); delay(6000); //这里延时一段时间，为了听到“正弦音” SPI_WriteByte(0x45);//写入以下8个字节，退出正弦测试 SPI_WriteByte(0x78); SPI_WriteByte(0x69); SPI_WriteByte(0x74); SPI_WriteByte(0); SPI_WriteByte(0); SPI_WriteByte(0); SPI_WriteByte(0); VS_XDCS_1; //关闭数据片选 ，SDI无效 } 音频文件需要转换为C头文件，可以用matlab实现： clear; clc; close all; f = fopen('222.mp3','rb'); a = fread(f,'uint8'); fclose(f); fb = fopen('D:\Tutor_My\MP3Player\MP3Player.sdk\SDK\SDK_Export\mp3\src\mp3.h','w'); fprintf(fb,'const unsigned char mp3_table[] = {\r\n'); fprintf(fb,'0x%02x,\r\n',a(1:end)); fprintf(fb,'\r\n};'); fclose(fb); 下载比特流，运行。通过耳机可以听到你转换的mp3。 完成上述工程，只需要10min，操作完全由Vivado+SDK完成，操作十分简单集中。
-
发表了文章 2013-10-23
关于图像识别，我们是否可以用一种编程语言来描述？
今天在整理照片时，发现有张图片上拍了一个表格，手动将内容抄写下来不太省力，于是想如果有一个软件可以抓取图像中的表格文字该多好。。。 从头设计一套软件完成这个工作量有点大，想到已经有了一些手写体识别的应用，只要在该技术基础上再增加一层接口，其实就能完成上述任务。 该接口要灵活，能够适应各种应用的文字识别，比如汉字手写体，英文手写体，宋体，以及上面的表格内容识别。因为一幅图像中的内容是千差万别的，我们可以借助一种图像描述原语（Image Description Primitives, IDP）来完成这些内容的区别。 假设有如下IDP: Def "Table" p1; //表示该图像中有一个表格，声明为对象p1 Set Origin(x0,y0); //表示表格的左上角像素坐标 Set End(x1,y1); //表示表格的右下角像素坐标 p1.cols = [5-10]; //表示该表格的列数可能在5~10列之间 p1.rows = [100-200]; //行数在100~200行之间 Table tb = p1.GetContent(); //创建表格对象，从p1对象中抽取内容，以字符串形式存入tb中 //打印输出字符串 for(int i = 0;i<tb.rows;i++) { for(int j = 0;j<tb.cols;j++) { printf("%s\t",tb[i][j]); } printf("\r\n"); } End p1; //销毁对象 除了表格对象，我们还可创建纯文本对象，手写体对象，英文手写体对象等，识别原语都封装到p1.GetContent()中。 只是一时想到的，不知有没有价值。本人只是对文字识别感兴趣，没有相关经验，一家之言，欢迎讨论。
-
发表了文章 2013-08-17
ZED-Board从入门到精通系列例程——全局定时器
本文系ZED-Board从入门到精通（三）：从传统ARM开发到PS开发的转变之后增加的PS例程。由于原文较长，在原帖后面添加例程会使阅读不便，于是单独开一帖。 实际项目中几乎离不开时间的测量。定时器是硬件系统运行状态的忠实记录者，它不受CPU直接干预，自己独立运行，可以完成计时、定时、中断、实时时钟等功能。 ARM Cortex-A9内部有一个64bit全局定时器，特性包括： 64bit，增计数； 内存映射至私有内存空间； 只有复位后，在安全模式下才能访问； 可被所有Cortex-A9核访问，每个核有私有比较器； 时钟源为PERIPHCLK； 定时器的精度是由其时钟源决定的，而时钟源来自ARM系统时钟。我们先来看一下硬件系统时钟分配情况， 系统PS_CLK为板上的晶振输入，频率为33.3333MHz PS-CLK进入芯片后，又做如下分配（摘自Zynq-7000-TRM）： 可见经过了3个PLL，最终生成的系统时钟有cpu_6x4x，cpu_3x2x，cpu_2x，cpu_1x。具体的系统时钟频率值我们可以查看XPS中的时钟选项，这里不再详述，只要知道全局定时器的输入时钟为cpu_3x2x，它的频率为CPU时钟的一半（333.333MHz），定时精度为3ns，又由于其具有64bit范围，最大定时值可达3e34s。 操作定时器需要访问其对应寄存器，我们看一下TRM中的描述： 这里只给出了基地址，具体寄存器的分布需要查看ARM文档cortex_a9_mpcore_r4p1_trm： 其中前两个为定时器的计数值存放寄存器，两个32bit凑成一个64bit实现连续增计数。 第三个寄存器为控制寄存器，位定义如下： 我们需要关注的是最低位（b0)，即定时器使能位，该位为0时，定时器停止，这时可以读写计数值；而该位为1时，定时器运行，不能写入计数值（只能读出）。 其它的寄存器我们暂时不用，不加解释。需要的话可以自己翻一翻手册。 相比基于操作系统的软件计时器，我们采用硬件计时器具有非常高的精度，可以精确到ns级别！对于非常窄的脉冲，我们照样可以通过计时器完成其脉宽测量。程序中有时需要精确延时（例如红外通信，DS18b20单总线读写），我们先写一个精确延时1s的函数，然后把它用在我们第一个流水灯实验中。本节例程仍基于第一个例程进行，硬件部分不需要改动，只需要改软件，打开helloworld.c，将内容改为： /* * Copyright (c) 2009 Xilinx, Inc. All rights reserved. * * Xilinx, Inc. * XILINX IS PROVIDING THIS DESIGN, CODE, OR INFORMATION "AS IS" AS A * COURTESY TO YOU. BY PROVIDING THIS DESIGN, CODE, OR INFORMATION AS * ONE POSSIBLE IMPLEMENTATION OF THIS FEATURE, APPLICATION OR * STANDARD, XILINX IS MAKING NO REPRESENTATION THAT THIS IMPLEMENTATION * IS FREE FROM ANY CLAIMS OF INFRINGEMENT, AND YOU ARE RESPONSIBLE * FOR OBTAINING ANY RIGHTS YOU MAY REQUIRE FOR YOUR IMPLEMENTATION. * XILINX EXPRESSLY DISCLAIMS ANY WARRANTY WHATSOEVER WITH RESPECT TO * THE ADEQUACY OF THE IMPLEMENTATION, INCLUDING BUT NOT LIMITED TO * ANY WARRANTIES OR REPRESENTATIONS THAT THIS IMPLEMENTATION IS FREE * FROM CLAIMS OF INFRINGEMENT, IMPLIED WARRANTIES OF MERCHANTABILITY * AND FITNESS FOR A PARTICULAR PURPOSE. * */ /* * helloworld.c: simple test application */ #include <stdio.h> #include "platform.h" #define MIO_BASE 0xE000A000 //MIO基地址 #define DATA1_RO 0x64 #define DATA2 0x48 #define DATA2_RO 0x68 #define DIRM_2 0x284 #define OEN_2 0x288 #define GTC_BASE 0xF8F00200 //Global Timer基地址 #define GTC_CTRL 0x08 //控制寄存器偏移量 #define GTC_DATL 0x00 //数据寄存器（低32bit） #define GTC_DATH 0x04 //数据寄存器（高32bit） #define CLK_3x2x 333333333 //定时器输入时钟频率 void print(char *str); void delay_1s(int t) //t无实际意义 { int i = CLK_3x2x,j; *((volatile int*)(GTC_BASE+GTC_CTRL)) = 0x00; //清零定时器使能位，定时器停止 *((volatile int*)(GTC_BASE+GTC_DATL)) = 0x00000000; //写入计数值（低32bit） *((volatile int*)(GTC_BASE+GTC_DATH)) = 0x00000000; //写入计数值（高32bit） *((volatile int*)(GTC_BASE+GTC_CTRL)) = 0x01; //开启定时器 do { j=*((volatile int*)(GTC_BASE+GTC_DATL)); } while(j<i); //判断是否计时够1s？ } void print(char *str); int main() { int i; init_platform(); *((volatile int*)(MIO_BASE+OEN_2)) = 0xff; *((volatile int*)(MIO_BASE+DIRM_2)) = 0xff; print("Hello world!\r\nThe Leds are flowing...\r\n"); while(1) { for(i = 0;i < 8; i++) { *((volatile int*)(MIO_BASE+DATA2)) = 0x01<<i; delay_1s(1000); } } cleanup_platform(); return 0; } 上面例子中，将原来的delay_1s改成了利用64bit全局定时器实现的精确定时（虽然这样做有点浪费，呵呵）。 运行结果仍为流水灯，灯移一位的时间应该是标准的1s。 我们可以通过简单的编程，实现对程序性能的监测，例如在运行算法程序之前，先开启计时器，等算法程序结束，再停止计时，读取计时器的计数值从而计算算法运行时间，这样可以评估算法性能。这个功能有点像Matlab里面的tic，toc，为了方便程序编写，我们也如此定义函数： #define GTC_BASE 0xF8F00200 #define GTC_CTRL 0x08 #define GTC_DATL 0x00 #define GTC_DATH 0x04 #define CLK_3x2x 333333333 void tic(void) { *((volatile int*)(GTC_BASE+GTC_CTRL)) = 0x00; *((volatile int*)(GTC_BASE+GTC_DATL)) = 0x00000000; *((volatile int*)(GTC_BASE+GTC_DATH)) = 0x00000000; //清零定时器的计数值 *((volatile int*)(GTC_BASE+GTC_CTRL)) = 0x01; } double toc(void) { *((volatile int*)(GTC_BASE+GTC_CTRL)) = 0x00; long long j=*((volatile int*)(GTC_BASE+GTC_DATH)); double elapsed_time = j<<32; j=*((volatile int*)(GTC_BASE+GTC_DATL)); //读取64bit定时器值，转换为double elapsed_time+=j; elapsed_time/=CLK_3x2x; elapsed_time*=1000; printf("Elapsed time is %f ms.\r\n",elapsed_time); return elapsed_time; } 调用时非常简单： tic(); my_algorithm(); toc(); 运行时，程序输出和matlab完全一致。这里使用硬件计时，精度可以达到ns级别，具有普通软件计时无法比拟的特性，对于非常窄的脉冲，我们照样可以用上面的方法测量其脉宽。 通过本节定时器的例子，相信童鞋们对PS开发有种驾轻就熟的感觉。没错，真正基于Zynq的PS开发流程就是如此，首先查阅文档，知道硬件寄存器定义，然后按照说明进行底层软件编写，并为上层程序提供较为简洁和直观的接口。掌握了这个技巧，后面进行PS与PL协同开发时，只要根据PL相应内存映射地址和寄存器定义，就可以完成PS端控制软件的设计，从而为后面进一步编写基于操作系统的驱动程序打下坚实的基础。 大家可以读完本文后，进一步利用官方文档，熟悉一下PS的其他外设操作。
-
发表了文章 2013-07-28
ZED-Board从入门到精通（零）：写在前面
去年10月份有幸报名参加了OpenHW2012开源硬件与嵌入式设计大赛，今年2月底顺利得到一块ZED-Board，从此步入了ZYNQ All-Programmable的世界。晒晒板子如下： ZYNQ系列SoC在单颗芯片上集成了ARM Cortex A9双核与FPGA，不仅开发软件可裁剪，而且硬件设备也可定制、自主开发，具有更大的灵活性。FPGA部分采用Xilinx Artix-7架构，被称为Programmable Logic，简称PL。ARM部分则称为Processing System，简称PS。 下个月马上就要进行决赛了。参加完这个比赛之后，感觉应该趁热打铁把一些设计流程和经验记录下来，暑假期间刚好有空将这些内容进行系统整理，所以想再开一个ZYNQ设计专栏，将设计道路铺平，利于后人。由于精力有限，可能不能面面俱到，只是将设计过程中用到的工具、方法记录下来，一些没有用到的内容需要留有空白，让后人来填补了。 学习ZYNQ，需要具备的基础有：FPGA设计基础（最好有Xilinx FPGA开发经验）；ARM设计基础（裸机+嵌入式系统）；Linux基础；计算机体系结构基础等。如果某一方面基础不够，可以跟着本教程走一遍，然后有针对性地学习。 计划内容如下： 1. ZYNQ结构简介 2. AXI简介 3. 从传统ARM开发到PS开发的转变 4. 从传统FPGA开发到PL开发的转变 5. 软硬件协同设计 6. 基于嵌入式Linux开发流程 7. 基于Xilinux开发流程 8. 基于Linaro开发流程 9. 基于Android开发流程 10. 基于WinCE开发流程（待定）
-
发表了文章 2013-07-23
CUDA从入门到精通（六）：块并行
同一版本的代码用了这么多次，有点过意不去，于是这次我要做较大的改动，大家要擦亮眼睛，拭目以待。 块并行相当于操作系统中多进程的情况，上节说到，CUDA有线程组（线程块）的概念，将一组线程组织到一起，共同分配一部分资源，然后内部调度执行。线程块与线程块之间，毫无瓜葛。这有利于做更粗粒度的并行。我们将上一节的代码改为块并行版本如下： #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size); __global__ void addKernel(int *c, const int *a, const int *b) { int i = blockIdx.x; c[i] = a[i] + b[i]; } int main() { const int arraySize = 5; const int a[arraySize] = { 1, 2, 3, 4, 5 }; const int b[arraySize] = { 10, 20, 30, 40, 50 }; int c[arraySize] = { 0 }; // Add vectors in parallel. cudaError_t cudaStatus; int num = 0; cudaDeviceProp prop; cudaStatus = cudaGetDeviceCount(&num); for(int i = 0;i<num;i++) { cudaGetDeviceProperties(&prop,i); } cudaStatus = addWithCuda(c, a, b, arraySize); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addWithCuda failed!"); return 1; } printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]); // cudaThreadExit must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaThreadExit(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadExit failed!"); return 1; } return 0; } // Helper function for using CUDA to add vectors in parallel. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size) { int *dev_a = 0; int *dev_b = 0; int *dev_c = 0; cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // Allocate GPU buffers for three vectors (two input, one output) . cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. addKernel<<<size,1 >>>(dev_c, dev_a, dev_b); // cudaThreadSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaThreadSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } Error: cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); return cudaStatus; } 和上一节相比，只有这两行有改变，<<<>>>里第一个参数改成了size，第二个改成了1，表示我们分配size个线程块，每个线程块仅包含1个线程，总共还是有5个线程。这5个线程相互独立，执行核函数得到相应的结果，与上一节不同的是，每个线程获取id的方式变为int i = blockIdx.x；这是线程块ID。 于是有童鞋提问了，线程并行和块并行的区别在哪里？ 线程并行是细粒度并行，调度效率高；块并行是粗粒度并行，每次调度都要重新分配资源，有时资源只有一份，那么所有线程块都只能排成一队，串行执行。 那是不是我们所有时候都应该用线程并行，尽可能不用块并行？ 当然不是，我们的任务有时可以采用分治法，将一个大问题分解为几个小规模问题，将这些小规模问题分别用一个线程块实现，线程块内可以采用细粒度的线程并行，而块之间为粗粒度并行，这样可以充分利用硬件资源，降低线程并行的计算复杂度。适当分解，降低规模，在一些矩阵乘法、向量内积计算应用中可以得到充分的展示。 实际应用中，常常是二者的结合。线程块、线程组织图如下所示。 多个线程块组织成了一个Grid，称为线程格（经历了从一位线程，二维线程块到三维线程格的过程，立体感很强啊）。 好了，下一节我们介绍流并行，是更高层次的并行。
-
发表了文章 2013-07-23
CUDA从入门到精通（五）：线程并行
多线程我们应该都不陌生，在操作系统中，进程是资源分配的基本单元，而线程是CPU时间调度的基本单元（这里假设只有1个CPU）。 将线程的概念引申到CUDA程序设计中，我们可以认为线程就是执行CUDA程序的最小单元，前面我们建立的工程代码中，有个核函数概念不知各位童鞋还记得没有，在GPU上每个线程都会运行一次该核函数。 但GPU上的线程调度方式与CPU有很大不同。CPU上会有优先级分配，从高到低，同样优先级的可以采用时间片轮转法实现线程调度。GPU上线程没有优先级概念，所有线程机会均等，线程状态只有等待资源和执行两种状态，如果资源未就绪，那么就等待；一旦就绪，立即执行。当GPU资源很充裕时，所有线程都是并发执行的，这样加速效果很接近理论加速比；而GPU资源少于总线程个数时，有一部分线程就会等待前面执行的线程释放资源，从而变为串行化执行。 代码还是用上一节的吧，改动很少，再贴一遍： #include "cuda_runtime.h" //CUDA运行时API #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size); __global__ void addKernel(int *c, const int *a, const int *b) { int i = threadIdx.x; c[i] = a[i] + b[i]; } int main() { const int arraySize = 5; const int a[arraySize] = { 1, 2, 3, 4, 5 }; const int b[arraySize] = { 10, 20, 30, 40, 50 }; int c[arraySize] = { 0 }; // Add vectors in parallel. cudaError_t cudaStatus; int num = 0; cudaDeviceProp prop; cudaStatus = cudaGetDeviceCount(&num); for(int i = 0;i<num;i++) { cudaGetDeviceProperties(&prop,i); } cudaStatus = addWithCuda(c, a, b, arraySize); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addWithCuda failed!"); return 1; } printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]); // cudaThreadExit must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaThreadExit(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadExit failed!"); return 1; } return 0; } // 重点理解这个函数 cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size) { int *dev_a = 0; //GPU设备端数据指针 int *dev_b = 0; int *dev_c = 0; cudaError_t cudaStatus; //状态指示 // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); //选择运行平台 if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // 分配GPU设备端内存 cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // 拷贝数据到GPU cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // 运行核函数 addKernel<<<1, size>>>(dev_c, dev_a, dev_b); // cudaThreadSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaThreadSynchronize(); //同步线程 if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost); //拷贝结果回主机 if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } Error: cudaFree(dev_c); //释放GPU设备端内存 cudaFree(dev_a); cudaFree(dev_b); return cudaStatus; } 红色部分即启动核函数的调用过程，这里看到调用方式和C不太一样。<<<>>>表示运行时配置符号，里面1表示只分配一个线程组（又称线程块、Block），size表示每个线程组有size个线程（Thread）。本程序中size根据前面传递参数个数应该为5，所以运行的时候，核函数在5个GPU线程单元上分别运行了一次，总共运行了5次。这5个线程是如何知道自己“身份”的？是靠threadIdx这个内置变量，它是个dim3类型变量，接受<<<>>>中第二个参数，它包含x,y,z 3维坐标，而我们传入的参数只有一维，所以只有x值是有效的。通过核函数中int i = threadIdx.x;这一句，每个线程可以获得自身的id号，从而找到自己的任务去执行。 下节我们介绍块并行。
-
发表了文章 2013-07-23
CUDA从入门到精通（四）：加深对设备的认识
前面三节已经对CUDA做了一个简单的介绍，这一节开始真正进入编程环节。 首先，初学者应该对自己使用的设备有较为扎实的理解和掌握，这样对后面学习并行程序优化很有帮助，了解硬件详细参数可以通过上节介绍的几本书和官方资料获得，但如果仍然觉得不够直观，那么我们可以自己动手获得这些内容。 以第二节例程为模板，我们稍加改动的部分代码如下： // Add vectors in parallel. cudaError_t cudaStatus; int num = 0; cudaDeviceProp prop; cudaStatus = cudaGetDeviceCount(&num); for(int i = 0;i<num;i++) { cudaGetDeviceProperties(&prop,i); } cudaStatus = addWithCuda(c, a, b, arraySize); 这个改动的目的是让我们的程序自动通过调用cuda API函数获得设备数目和属性，所谓“知己知彼，百战不殆”。 cudaError_t 是cuda错误类型，取值为整数。 cudaDeviceProp为设备属性结构体，其定义可以从cuda Toolkit安装目录中找到，我的路径为：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\driver_types.h，找到定义为： /** * CUDA device properties */ struct __device_builtin__ cudaDeviceProp { char name[256]; /**< ASCII string identifying device */ size_t totalGlobalMem; /**< Global memory available on device in bytes */ size_t sharedMemPerBlock; /**< Shared memory available per block in bytes */ int regsPerBlock; /**< 32-bit registers available per block */ int warpSize; /**< Warp size in threads */ size_t memPitch; /**< Maximum pitch in bytes allowed by memory copies */ int maxThreadsPerBlock; /**< Maximum number of threads per block */ int maxThreadsDim[3]; /**< Maximum size of each dimension of a block */ int maxGridSize[3]; /**< Maximum size of each dimension of a grid */ int clockRate; /**< Clock frequency in kilohertz */ size_t totalConstMem; /**< Constant memory available on device in bytes */ int major; /**< Major compute capability */ int minor; /**< Minor compute capability */ size_t textureAlignment; /**< Alignment requirement for textures */ size_t texturePitchAlignment; /**< Pitch alignment requirement for texture references bound to pitched memory */ int deviceOverlap; /**< Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. */ int multiProcessorCount; /**< Number of multiprocessors on device */ int kernelExecTimeoutEnabled; /**< Specified whether there is a run time limit on kernels */ int integrated; /**< Device is integrated as opposed to discrete */ int canMapHostMemory; /**< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer */ int computeMode; /**< Compute mode (See ::cudaComputeMode) */ int maxTexture1D; /**< Maximum 1D texture size */ int maxTexture1DMipmap; /**< Maximum 1D mipmapped texture size */ int maxTexture1DLinear; /**< Maximum size for 1D textures bound to linear memory */ int maxTexture2D[2]; /**< Maximum 2D texture dimensions */ int maxTexture2DMipmap[2]; /**< Maximum 2D mipmapped texture dimensions */ int maxTexture2DLinear[3]; /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */ int maxTexture2DGather[2]; /**< Maximum 2D texture dimensions if texture gather operations have to be performed */ int maxTexture3D[3]; /**< Maximum 3D texture dimensions */ int maxTextureCubemap; /**< Maximum Cubemap texture dimensions */ int maxTexture1DLayered[2]; /**< Maximum 1D layered texture dimensions */ int maxTexture2DLayered[3]; /**< Maximum 2D layered texture dimensions */ int maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */ int maxSurface1D; /**< Maximum 1D surface size */ int maxSurface2D[2]; /**< Maximum 2D surface dimensions */ int maxSurface3D[3]; /**< Maximum 3D surface dimensions */ int maxSurface1DLayered[2]; /**< Maximum 1D layered surface dimensions */ int maxSurface2DLayered[3]; /**< Maximum 2D layered surface dimensions */ int maxSurfaceCubemap; /**< Maximum Cubemap surface dimensions */ int maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */ size_t surfaceAlignment; /**< Alignment requirements for surfaces */ int concurrentKernels; /**< Device can possibly execute multiple kernels concurrently */ int ECCEnabled; /**< Device has ECC support enabled */ int pciBusID; /**< PCI bus ID of the device */ int pciDeviceID; /**< PCI device ID of the device */ int pciDomainID; /**< PCI domain ID of the device */ int tccDriver; /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */ int asyncEngineCount; /**< Number of asynchronous engines */ int unifiedAddressing; /**< Device shares a unified address space with the host */ int memoryClockRate; /**< Peak memory clock frequency in kilohertz */ int memoryBusWidth; /**< Global memory bus width in bits */ int l2CacheSize; /**< Size of L2 cache in bytes */ int maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */ }; 后面的注释已经说明了其字段代表意义，可能有些术语对于初学者理解起来还是有一定困难，没关系，我们现在只需要关注以下几个指标： name：就是设备名称； totalGlobalMem：就是显存大小； major,minor：CUDA设备版本号，有1.1, 1.2, 1.3, 2.0, 2.1等多个版本； clockRate：GPU时钟频率； multiProcessorCount：GPU大核数，一个大核（专业点称为流多处理器，SM，Stream-Multiprocessor）包含多个小核（流处理器，SP，Stream-Processor） 编译，运行，我们在VS2008工程的cudaGetDeviceProperties()函数处放一个断点，单步执行这一函数，然后用Watch窗口，切换到Auto页，展开+，在我的笔记本上得到如下结果： 可以看到，设备名为GeForce 610M，显存1GB，设备版本2.1（比较高端了，哈哈），时钟频率为950MHz（注意950000单位为kHz），大核数为1。在一些高性能GPU上（如Tesla，Kepler系列），大核数可能达到几十甚至上百，可以做更大规模的并行处理。 PS：今天看SDK代码时发现在help_cuda.h中有个函数实现从CUDA设备版本查询相应大核中小核的数目，觉得很有用，以后编程序可以借鉴，摘抄如下： // Beginning of GPU Architecture definitions inline int _ConvertSMVer2Cores(int major, int minor) { // Defines for GPU Architecture types (using the SM version to determine the # of cores per SM typedef struct { int SM; // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version int Cores; } sSMtoCores; sSMtoCores nGpuArchCoresPerSM[] = { { 0x10, 8 }, // Tesla Generation (SM 1.0) G80 class { 0x11, 8 }, // Tesla Generation (SM 1.1) G8x class { 0x12, 8 }, // Tesla Generation (SM 1.2) G9x class { 0x13, 8 }, // Tesla Generation (SM 1.3) GT200 class { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class { -1, -1 } }; int index = 0; while (nGpuArchCoresPerSM[index].SM != -1) { if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor)) { return nGpuArchCoresPerSM[index].Cores; } index++; } // If we don't find the values, we default use the previous one to run properly printf("MapSMtoCores for SM %d.%d is undefined. Default to use %d Cores/SM\n", major, minor, nGpuArchCoresPerSM[7].Cores); return nGpuArchCoresPerSM[7].Cores; } // end of GPU Architecture definitions 可见，设备版本2.1的一个大核有48个小核，而版本3.0以上的一个大核有192个小核！ 前文说到过，当我们用的电脑上有多个显卡支持CUDA时，怎么来区分在哪个上运行呢？这里我们看一下addWithCuda这个函数是怎么做的。 cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } 使用了cudaSetDevice(0)这个操作，0表示能搜索到的第一个设备号，如果有多个设备，则编号为0,1,2...。 再看我们本节添加的代码，有个函数cudaGetDeviceCount(&num)，这个函数用来获取设备总数，这样我们选择运行CUDA程序的设备号取值就是0,1,...num-1，于是可以一个个枚举设备，利用cudaGetDeviceProperties(&prop)获得其属性,然后利用一定排序、筛选算法，找到最符合我们应用的那个设备号opt，然后调用cudaSetDevice(opt)即可选择该设备。选择标准可以从处理能力、版本控制、名称等各个角度出发。后面讲述流并发过程时，还要用到这些API。 如果希望了解更多硬件内容可以结合http://www.geforce.cn/hardware获取。
-
发表了文章 2013-07-23
CUDA从入门到精通（三）：必备资料
刚入门CUDA，跑过几个官方提供的例程，看了看人家的代码，觉得并不难，但自己动手写代码时，总是不知道要先干什么，后干什么，也不知道从哪个知识点学起。这时就需要有一本能提供指导的书籍或者教程，一步步跟着做下去，直到真正掌握。 一般讲述CUDA的书，我认为不错的有下面这几本： 初学者可以先看美国人写的这本《GPU高性能编程CUDA实战》，可操作性很强，但不要期望能全看懂（Ps：里面有些概念其实我现在还是不怎么懂），但不影响你进一步学习。如果想更全面地学习CUDA，《GPGPU编程技术》比较客观详细地介绍了通用GPU编程的策略，看过这本书，可以对显卡有更深入的了解，揭开GPU的神秘面纱。后面《OpenGL编程指南》完全是为了体验图形交互带来的乐趣，可以有选择地看；《GPU高性能运算之CUDA》这本是师兄给的，适合快速查询（感觉是将官方编程手册翻译了一遍）一些关键技术和概念。 有了这些指导材料还不够，我们在做项目的时候，遇到的问题在这些书上肯定找不到，所以还需要有下面这些利器： 这里面有很多工具的使用手册，如CUDA_GDB，Nsight，CUDA_Profiler等，方便调试程序；还有一些有用的库，如CUFFT是专门用来做快速傅里叶变换的，CUBLAS是专用于线性代数（矩阵、向量计算）的，CUSPASE是专用于稀疏矩阵表示和计算的库。这些库的使用可以降低我们设计算法的难度，提高开发效率。另外还有些入门教程也是值得一读的，你会对NVCC编译器有更近距离的接触。 好了，前言就这么多，本博主计划按如下顺序来讲述CUDA： 1.了解设备 2.线程并行 3.块并行 4.流并行 5.线程通信 6.线程通信实例：规约 7.存储模型 8.常数内存 9.纹理内存 10.主机页锁定内存 11.图形互操作 12.优化准则 13.CUDA与MATLAB接口 14.CUDA与MFC接口
-
发表了文章 2013-07-22
CUDA从入门到精通（二）：第一个CUDA程序
书接上回，我们既然直接运行例程成功了，接下来就是了解如何实现例程中的每个环节。当然，我们先从简单的做起，一般编程语言都会找个helloworld例子，而我们的显卡是不会说话的，只能做一些简单的加减乘除运算。所以，CUDA程序的helloworld，我想应该最合适不过的就是向量加了。 打开VS2008，选择File->New->Project，弹出下面对话框，设置如下： 之后点OK，直接进入工程界面。 工程中，我们看到只有一个.cu文件，内容如下： #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size); __global__ void addKernel(int *c, const int *a, const int *b) { int i = threadIdx.x; c[i] = a[i] + b[i]; } int main() { const int arraySize = 5; const int a[arraySize] = { 1, 2, 3, 4, 5 }; const int b[arraySize] = { 10, 20, 30, 40, 50 }; int c[arraySize] = { 0 }; // Add vectors in parallel. cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addWithCuda failed!"); return 1; } printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n", c[0], c[1], c[2], c[3], c[4]); // cudaThreadExit must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaThreadExit(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadExit failed!"); return 1; } return 0; } // Helper function for using CUDA to add vectors in parallel. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size) { int *dev_a = 0; int *dev_b = 0; int *dev_c = 0; cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // Allocate GPU buffers for three vectors (two input, one output) . cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. addKernel<<<1, size>>>(dev_c, dev_a, dev_b); // cudaThreadSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaThreadSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } Error: cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); return cudaStatus; } 可以看出，CUDA程序和C程序并无区别，只是多了一些以"cuda"开头的一些库函数和一个特殊声明的函数： __global__ void addKernel(int *c, const int *a, const int *b) { int i = threadIdx.x; c[i] = a[i] + b[i]; } 这个函数就是在GPU上运行的函数，称之为核函数，英文名Kernel Function，注意要和操作系统内核函数区分开来。 我们直接按F7编译，可以得到如下输出： 1>------ Build started: Project: cuda_helloworld, Configuration: Debug Win32 ------ 1>Compiling with CUDA Build Rule... 1>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\nvcc.exe" -G -gencode=arch=compute_10,code=\"sm_10,compute_10\" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --machine 32 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin" -Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT " -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\include" -maxrregcount=0 --compile -o "Debug/kernel.cu.obj" kernel.cu 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu 1>tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu 1>tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu 1>tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu 1>kernel.cu 1>kernel.cu 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp 1>tmpxft_000000ec_00000000-24_kernel.compute_10.ii 1>Linking... 1>Embedding manifest... 1>Performing Post-Build Event... 1>copy "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart*.dll" "C:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\Debug" 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll 1>已复制 2 个文件。 1>Build log was saved at "file://c:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.htm" 1>cuda_helloworld - 0 error(s), 105 warning(s) ========== Build: 1 succeeded, 0 failed, 0 up-to-date, 0 skipped ========== 可见，编译.cu文件需要利用nvcc工具。该工具的详细使用见后面博客。 直接运行，可以得到结果图如下： 如果显示正确，那么我们的第一个程序宣告成功！
-
发表了文章 2013-07-22
CUDA从入门到精通（零）：写在前面
在老板的要求下，本博主从2012年上高性能计算课程开始接触CUDA编程，随后将该技术应用到了实际项目中，使处理程序加速超过1K，可见基于图形显示器的并行计算对于追求速度的应用来说无疑是一个理想的选择。还有不到一年毕业，怕是毕业后这些技术也就随毕业而去，准备这个暑假开辟一个CUDA专栏，从入门到精通，步步为营，顺便分享设计的一些经验教训，希望能给学习CUDA的童鞋提供一定指导。个人能力所及，错误难免，欢迎讨论。 PS：申请专栏好像需要先发原创帖超过15篇。。。算了，先写够再申请吧，到时候一并转过去。
-
发表了文章 2013-07-10
【OpenHW参赛手记】AXI-Stream接口开发详细流程
下面讲一个例子，来加深对上面介绍内容的理解。笔者使用的软件版本为ISE 14.2。 1.建立PlanAhead工程，一直到进入XPS，具体流程见官方文档CTT[1]。 2.在XPS中，添加一个AXI-DMA模块，配置界面如图1所示。 图1 AXI-DMA模块配置 其余参数默认。SG模块如果选上，那么后面软件控制会相对复杂一些。这里不选，采用Simple模式，实现较为简单的传输。 3.选菜单Hardware->Createor Import Peripheral…，设计自定义IP。名称起为my_stream_ip，自动版本为1.00a。遇到Bus Interface选择AXI4-Stream类型，一直点下一步到最后结束。该类型IP的生成过程比AXI4-Lite和AXI4都要简单。 4.添加一个my_stream_ip到系统中，连接见图2。 图2 AXI Stream IP硬件连接 由XPS自动生成的my_stream_ip实现了先接收8个32bit字，然后求和，再将结果发送回去（连续发送8次）。上图连接方式说明是AXI-DMA模块发送数据给my_stream_ip，然后my_stream_ip又将结果发回AXI-DMA。同时看到AXI-DMA和PS的数据流连接是通过HP0传输，而控制流通过GP0传输。 5.上面连接在不做任何改动的情况下有问题（主要是XPS的bug），需要一项项手动修改。首先是HP0的地址区间报错，可以先点Zynq标签，然后单击HP0绿线，在弹出的配置对话框中将HP0的地址区间改为我们ZED Board 上DDR2区间0x00000000~0x1FFFFFFF，像图3一样。 图3 修正bug1 在较高版本软件ISE14.5中，这个bug已经修复，不需要改。 第二个bug就是AXI-DMA和my_stream_ip的连线问题。本来都是Stream 接口，按理说是标准接口，不应该有差异。但事实就是这样，XPS界面掩饰的问题层出不穷。我们右击my_stream_ip，选择View MPD，将内容改为： BEGIN my_stream_ip ## Peripheral Options OPTION IPTYPE = PERIPHERAL OPTION IMP_NETLIST = TRUE OPTION HDL = VERILOG ## Bus Interfaces BUS_INTERFACE BUS=M_AXIS, BUS_STD=AXIS, BUS_TYPE=INITIATOR BUS_INTERFACE BUS=S_AXIS, BUS_STD=AXIS, BUS_TYPE=TARGET ## Parameters PARAMETER C_S_AXIS_PROTOCOL = GENERIC, DT = string, TYPE = NON_HDL, ASSIGNMENT= CONSTANT, BUS = S_AXIS PARAMETER C_S_AXIS_TDATA_WIDTH = 32, DT = integer, TYPE = NON_HDL, ASSIGNMENT =CONSTANT, BUS = S_AXIS PARAMETER C_M_AXIS_PROTOCOL = GENERIC, DT = string, TYPE = NON_HDL, ASSIGNMENT= CONSTANT, BUS = M_AXIS PARAMETER C_M_AXIS_TDATA_WIDTH = 32, DT = integer, TYPE = NON_HDL, ASSIGNMENT =CONSTANT, BUS = M_AXIS ## Peripheral ports PORT ACLK = "", DIR=I, SIGIS=CLK, BUS=M_AXIS:S_AXIS PORT ARESETN = ARESETN, DIR=I, INITIALVAL = VCC PORT S_AXIS_TREADY = TREADY, DIR=O, BUS=S_AXIS PORT S_AXIS_TDATA = TDATA, DIR=I, VEC=[31:0], BUS=S_AXIS PORT S_AXIS_TLAST = TLAST, DIR=I, BUS=S_AXIS PORT S_AXIS_TVALID = TVALID, DIR=I, BUS=S_AXIS PORT M_AXIS_TVALID = TVALID, DIR=O, BUS=M_AXIS PORT M_AXIS_TDATA = TDATA, DIR=O, VEC=[31:0], BUS=M_AXIS PORT M_AXIS_TLAST = TLAST, DIR=O, BUS=M_AXIS PORT M_AXIS_TREADY = TREADY, DIR=I, BUS=M_AXIS PORT M_AXIS_TKEEP = TKEEP, DIR=O, VEC=[3:0], BUS=M_AXIS END 这里存在两个问题：一个是ARESETN，在连接时AXI-DMA上没有合适的引脚与之相连，默认接地。这里显式声明接VCC。另一个问题是TKEEP信号，在我的博客文章《AXI-Stream调试日记（三）》里说过了，这里加上这个引脚，才能准确地将数据发回AXI-DMA。 保存MPD文件，关闭。再次右击my_stream_ip，选择Browse HDL Sources，打开my_stream_ip.v（或my_stream_ip.vhd），添加TKEEP信号并设置TLAST信号。 module my_stream_ip ( ACLK, ARESETN, S_AXIS_TREADY, S_AXIS_TDATA, S_AXIS_TLAST, S_AXIS_TVALID, M_AXIS_TVALID, M_AXIS_TDATA, M_AXIS_TLAST, M_AXIS_TREADY, M_AXIS_TKEEP ); input ACLK; input ARESETN; output S_AXIS_TREADY; input [31 :0] S_AXIS_TDATA; input S_AXIS_TLAST; input S_AXIS_TVALID; output M_AXIS_TVALID; output [31 :0] M_AXIS_TDATA; output M_AXIS_TLAST; input M_AXIS_TREADY; output [3:0] M_AXIS_TKEEP; localparamNUMBER_OF_INPUT_WORDS = 8; localparamNUMBER_OF_OUTPUT_WORDS = 8; localparam Idle =3'b100; localparam Read_Inputs = 3'b010; localparam Write_Outputs = 3'b001; reg [2:0] state; reg [31:0] sum; reg [NUMBER_OF_INPUT_WORDS -1:0] nr_of_reads; reg [NUMBER_OF_OUTPUT_WORDS - 1:0] nr_of_writes; assign S_AXIS_TREADY =(state == Read_Inputs); assign M_AXIS_TVALID = (state == Write_Outputs); assign M_AXIS_TDATA = sum; assign M_AXIS_TLAST = (nr_of_writes == 1); assign M_AXIS_TKEEP = 4'b1111; always @(posedge ACLK) begin // process The_SW_accelerator if(!ARESETN) // Synchronous reset (active low) begin state <= Idle; nr_of_reads <= 0; nr_of_writes <=0; sum <= 0; end else case (state) Idle: if (S_AXIS_TVALID== 1) begin state <= Read_Inputs; nr_of_reads <= NUMBER_OF_INPUT_WORDS - 1; sum <= 0; end Read_Inputs: if(S_AXIS_TVALID == 1) begin sum <= sum + S_AXIS_TDATA; if (nr_of_reads == 0) begin state <= Write_Outputs; nr_of_writes <= NUMBER_OF_OUTPUT_WORDS - 1; end else nr_of_reads <= nr_of_reads - 1; end Write_Outputs: if(M_AXIS_TREADY == 1) begin if (nr_of_writes == 0) state <= Idle; else nr_of_writes <= nr_of_writes - 1; end endcase end endmodule 到这里修正了已知的所有bug。VHDL代码见我的博客文章http://www.eeforum.com/附件，或通过邮件联系我获取。完成上述更改后，点XPS菜单Project->Rescan User Repositories，实现用户配置更新。 6.点Port标签，引脚连接。这里重点是将所有带CLK字样的都连接到PS7_FCLK_CLK0.如图4所示。 图4 PORT标签信号线连接 7.点击Addresses标签，看看AXI-DMA是否分配了控制端口地址 图5 地址分配 注意，如果你的axi_dma_0的地址和图中不一样，那么在后面软件编写时一定要修改成你的地址。 8.点Project->DesignRule Check；没错时，点Hardware->Generate Netlist，完成后关闭XPS。 9.在PlanAhead中完成综合、实现、生成Bit等步骤[12]。其实上一步已经完成了综合，所以这一步速度就会非常快。 10 导出SDK工程。建立Helloworld工程。将Helloworld.c里面的内容改为如下代码。 #include <stdio.h> #include <stdlib.h> #include "platform.h" #include "xil_cache.h" //必须包含该头文件，实现cache操作 #define sendram ((int *)0x10000000) //发送缓冲区 #define recvram ((int *)0x10001000) //接收缓冲区 #define sizeofbuffer 32 void print(char *str); #define WITH_SG 0 #define AXI_DMA_BASE 0x40400000 #define MM2S_DMACR 0 #define MM2S_DMASR 1 #if WITH_SG #define MM2S_CURDESC 2 #define MM2S_TAILDESC 4 #else #define MM2S_SA 6 #define MM2S_LENGTH 10 #endif #define S2MM_DMACR 12 #define S2MM_DMASR 13 #if WITH_SG #define S2MM_CURDESC14 #define S2MM_TAILDESC16 #else #define S2MM_DA 18 #define S2MM_LENGTH 22 #endif void debug_axi_dma_register(unsigned int *p) { printf("MM2S_DMACR = 0x%x\n",*(p+MM2S_DMACR)); printf("MM2S_DMASR = 0x%x\n",*(p+MM2S_DMASR)); #if WITH_SG printf("MM2S_CURDESC = 0x%x\n",*(p+MM2S_CURDESC)); printf("MM2S_TAILDESC = 0x%x\n",*(p+MM2S_TAILDESC)); #else printf("MM2S_SA = 0x%x\n",*(p+MM2S_SA)); printf("MM2S_LENGTH = 0x%x\n",*(p+MM2S_LENGTH)); #endif printf("S2MM_DMACR =0x%x\n",*(p+S2MM_DMACR)); printf("S2MM_DMACSR =0x%x\n",*(p+S2MM_DMASR)); #if WITH_SG printf("S2MM_CURDESC =0x%x\n",*(p+S2MM_CURDESC)); printf("S2MM_TAILDESC= 0x%x\n",*(p+S2MM_TAILDESC)); #else printf("S2MM_DA =0x%x\n",*(p+S2MM_DA)); printf("S2MM_LENGTH =0x%x\n",*(p+S2MM_LENGTH)); #endif } void init_axi_dma_simple(unsigned int * p) { *(p+MM2S_DMACR) = 0x04; //reset send axi dma while(*(p+MM2S_DMACR)&0x04); *(p+S2MM_DMACR) =0x04; //reset send axi dma while(*(p+S2MM_DMACR)&0x04); *(p+MM2S_DMACR)=1; while((*(p+MM2S_DMASR)&0x01)); *(p+S2MM_DMACR)=1; while((*(p+S2MM_DMASR)&0x01)); *(p+MM2S_SA) = (unsigned int )sendram; *(p+S2MM_DA) =(unsigned int )recvram; Xil_DCacheFlushRange((u32)sendram,sizeofbuffer); //将cache内容同步到DDR2 *(p+S2MM_LENGTH) =sizeofbuffer;//sizeof(recvram); *(p+MM2S_LENGTH) = sizeofbuffer;//sizeof(sendram); while(!(*(p+MM2S_DMASR)&0x1000)); //wait for send ok } void init_sendbuffer() { int i; for(i=0;i< sizeofbuffer/4;i++) { sendram[i]=i*2; } } void show_recvbuffer() { int i; printf("Recv contents are:\n"); for(i=0;i< sizeofbuffer/4;i++) { printf("%d\t",recvram[i]); } printf("\r\n"); } void show_sendbuffer() { int i; printf("Send contents are:\n"); for(i=0;i< sizeofbuffer/4;i++) { printf("%d\t",sendram[i]); } printf("\r\n"); } int main() { unsigned int status=0; int rxlen; init_platform(); init_sendbuffer(); init_axi_dma_simple((unsignedint *)AXI_DMA_BASE); printf("Hello World\n\rPlease input data:"); while(1) { scanf("%x",&status); printf("Got 0x%x\n",status); debug_axi_dma_register((unsigned int *)AXI_DMA_BASE); if(status==0) { break; } } show_sendbuffer(); Xil_DCacheInvalidateRange((u32)recvram,sizeofbuffer); //将DDR2内容同步到cache show_recvbuffer(); cleanup_platform(); return 0; } 保存，等待生成elf。然后连接板子，下载bit文件，Run App，打开串口终端，等待输出。由图6可见结果正确。 图6 程序输出 最终实现的my_stream_ip对外接口如下图所示。其中“M_AXIS”开头的信号线表示为AXI_Stream主机信号线，而“S_AXIS”开头的信号线表示为AXI_Stream从机信号线。自动生成的代码中没有M_AXIS_TKEEP信号，根据AXI4_Stream协议，这会导致该模块作为主机时发送的数据一直处于无效状态，影响数据传输。我们在my_stream_ip中添加了该信号，并使之有效，从而能够获得正确的处理数据。 图7 my_stream_ip对外接口 其中 Xil_DCacheFlushRange()和Xil_DCacheInvalidateRange()两个函数均在"xil_cache.h"中声明，用于将cache内容同步到DDR2或相反的操作。之前由于不了解cache，导致程序一直得不到正确的结果，总是怀疑硬件问题，后来通过forums.xilinx.com看到了相关的帖子才明白这一点，在此感谢论坛上国内外的技术大牛为社区提供的支持。
-
发表了文章 2013-07-10
【OpenHW参赛手记】AXI-Stream接口介绍
AXI4-Stream协议是一种用来连接需要交换数据的两个部件的标准接口，它可以用于连接一个产生数据的主机和一个接受数据的从机。当然它也可以用于连接多个主机和从机。该协议支持多种数据流使用相同共享总线集合，允许构建类似于路由、宽窄总线、窄宽总线等更为普遍的互联。AXI4-Stream接口的信号线定义如图1所示[1]。比较重要的信号线有： ACLK为时钟线，所有信号都在ACLK上升沿被采样； ARESETn为复位线，低电平有效； TVALID为主机数据同步线，为高表示主机准备好发送数据； TREADY为从机数据同步线，为高表示从机准备好接收数据；这两根线完成了主机与从机的握手信号，一旦二者都变高有效，数据传输开始。 TDATA为数据线，主机发送，从机接收。 TKEEP为主机数据有效指示，为高代表对应的字节为有效字节，否则表示发送的为空字节。 TLAST为主机最后一个字指示，下一clk数据将无效，TVALID将变低。 TID，TDEST，TUSER均为多机通信时的信号，这里不涉及，不予考虑。 看到这里，可能大家都还对Stream没有一个直观的认识。其实Stream并不陌生，在我们学c++编程时，一定会包含<iostream>，这样就可以完成控制终端对程序的输入输出了。如果还是不够直观，想象一下连续不断的水流，永远向着一个方向以固定的速度输送的接口。以我们看视频为例，视频文件本来是保存在硬盘里的，怎么播放呢，不能一下子把整个文件都显示到屏幕上，而是以一定的速度，连续不断地输出到屏幕上（每秒30~60帧），这个过程就是流Stream接口完成的。 Xilinx提供的流式IP核有很多用途，可以实现音频流、视频流、数据流到内存或者相反方向的传输。有人问了，内存是PS控制的，怎么才能把PS里DDR2的内容以Stream形式发出去呢（例如以固定速度送往DA，完成信号发生器的设计）？答案就是利用AXI总线做转换。ZYNQ的PS部分是ARM Cortex A9系列，支持AXI4，AXI-Lite总线。PL部分也有相应AXI总线接口，这样就能完成PS到PL的互联。仅仅这样还不够，需要PL部分实现流式转换，即AXI-Stream接口实现。Xilinx提供的从AXI到AXI-Stream转换的IP核有：AXI-DMA，AXI-Datamover，AXI-FIFO-MM2S以及AXI-VDMA等。这些IP核可以在XPS的IP Catalog窗口中看到。 AXI-DMA：实现从PS内存到PL高速传输高速通道AXI-HP到AXI-Stream的转换； AXI-FIFO-MM2S：实现从PS内存到PL通用传输通道AXI-GP到AXI-Stream的转换； AXI-Datamover：实现从PS内存到PL高速传输高速通道AXI-HP到AXI-Stream的转换，只不过这次是完全由PL控制的，PS是完全被动的； AXI-VDMA：实现从PS内存到PL高速传输高速通道AXI-HP到AXI-Stream的转换，只不过是专门针对视频、图像等二维数据的。 除了上面的还有一个AXI-CDMAIP核，这个是由PL完成的将数据从内存的一个位置搬移到另一个位置，无需CPU来插手。这个和我们这里用的Stream没有关系，所以不表。 这里要和大家说明白一点，就是AXI总线和接口的区别。总线是一种标准化接口，由数据线、地址线、控制线等构成，具有一定的强制性。接口是其物理实现，即在硬件上的分配。在ZYNQ中，支持AXI-Lite，AXI4和AXI-Stream三种总线，但PS与PL之间的接口却只支持前两种，AXI-Stream只能在PL中实现，不能直接和PS相连，必须通过AXI-Lite或AXI4转接。PS与PL之间的物理接口有9个，包括4个AXI-GP接口和4个AXI-HP接口、1个AXI-ACP接口，均为内存映射型AXI接口。 上面的IP是完成总线协议转换，如果需要做某些处理（如变换、迭代、训练……），则需要生成一个自定义Stream类型IP，与上面的Stream接口连接起来，实现数据输入输出。用户的功能在自定义Stream类型IP中实现。
-
发表了文章 2013-07-10
【OpenHW参赛手记】System Generator 与XPS连接的方法
由于项目需要用到DSP算法实现，考虑用System Generator辅助设计算法，但在参赛赠送的书里没有相关知识，需要自己动手摸索。 还好在Matlab Simulink 内包含库XilinxBlocksets里右键发现了帮助文档，进入之后了解了怎么用Xilinx器件来完成算法设计，并生成网表。最重要的器件应该是Gatewayin和gatewayout，将matlab内的模块与Xilinx模块隔离，实现数位精度变换。这两个元件相当于input和output，在生成实例的时候可以看到。 看了帮助文档后，做个实验，只用PL部分，不用PS部分，用ProjectNavigator新建一个工程，添加一个SysGen模块，里面什么都不用做，直接将Gatewayin连接到Gatewayout，位数设为8bit，无小数部分。然后新建顶层verilog模块，生成一个SysGen的实例（有模板可用），将输入通过UCF约束到电路板的SW，输出约束到LED，生成.bit，下载到PL上，很容易就成功实现了。 在此基础上，再次进入matlab，在Gatewayin和Gatewayout之间加入一个移位器模块，设为左移，重新生成.bit下载，也成功了。 单独的DSP模块显然功能很有限，必须要和双核ARM9连接起来才有价值。 新建planahead工程，参考懒兔子的自定义IP部分，生成mygpio模块，然后连接到AXI总线，这时不要连接到外部引脚，而是连接到顶层模块的内部信号线。 在planahead里面，新建SysGen模块，搭建自己的算法，生成模块，回到planahead。 之前的内部信号线连接到SysGen模块的输入。 还是像懒兔子一样导出软件工程，开发方法没有差别。下载.bit文件配置PL部分，然后运行程序。 这时就可以发现，你通过串口写到slave寄存器里的值，通过SysGen模块的处理，就能迅速反馈给你。这就说明，通过planahead顶层模块可以实现SysGen和XPS的互联，属于系统互联，这里涉及到两个系统，一个是DSP子系统，一个是PS子系统，二者都在planahead顶层模块被实例化，然后用信号线互联。 这是最近两天的摸索，算法实现总算有了点眉目。 另外看到在SysGen里面支持AXI4总线，而且可以导出为XPS工程，还有待进一步研究，找到更紧密的连接方式，这样可以进一步提高数据在DSP模块和XPS之间的传输效率。如果大家有什么好的建议，请不要吝啬及时告诉我一声~~~~ 时间真是快，还有很多值得去研究的内容，包括GAL均衡算法优化，多种算法(LMS，RLS，基于小波的。。。）性能对比，盲均衡技术实现（CMA，DF）等。。。除了算法实现外还有信号类型分析（DRM，DTV，GSM，4G），时不我与，不再多说，回去继续埋头学习。。。
-
发表了文章 2013-07-10
【OpenHW参赛手记】AXI初体验
去年报名参加了OpenHW2012开源硬件大赛，入围后收到了Xilinx赞助的ZED-Board。 自从今年3月拿到ZED板卡，就开始体验ARM+FPGA这个神奇的架构。 AXI，是PS与PL之间最佳的通信手段。 Xilinx XPS中用户自定义IP核可以拥有AXI-Lite，AXI4，AXI-Stream，PLB和FSL这些接口实现和PS通信。 其中AXI-Lite具有轻量级，结构简单的特点，适合小批量数据、简单控制场合。AXI4和AXI-Lite差不多，只是增加了一项功能就是突发传输，可以连续对一片地址进行一次性读写。上面两种均采用内存映射控制方式，即ARM将用户自定义IP编入某一地址进行访问，读写时就像在读写自己的片内RAM，编程也很方便，开发难度较低。代价就是资源占用过多，需要额外的读地址线、写地址线、读数据线、写数据线、写应答线这些信号线，而且传输速度受限（主要是因为采用AXI-GP物理接口，带宽很低）。 另外一种AXI接口就是AXI-Stream，这是一种连续传输的接口技术，速度能达到很高，而且不需要地址线（有点像FIFO，你就一直读或一直写就行）。这类IP不能通过上面的内存映射方式控制，必须有一个转换装置，例如AXI-Interconnector，AXI-DMA模块就能实现内存映射到流式接口的转换，但编程较复杂，调试起来没有内存映射方式直观，必须要通过芯片内部调试接口（Chipscope）来观察。 AXI-Stream适用的场合有很多：视频流处理；通信协议转换；数字信号处理；无线通信等。 其本质都是针对数值流构建的数据通路，从信源（例如ARM内存、DMA、无线接收前端等）到信宿（例如HDMI显示器、音频输出等）构建起连续的数据流。这种接口适合做实时信号处理。当然，实际处理中也有分块和不分块的情况，典型分块情况就是计算FFT。 后面两种貌似在ZED上面用处不大，都是Microblaze的接口。不过应该也有桥接IP。没有研究。 我们项目中属于典型的流式数据，从射频前端、ADC采集到信号传输到DDR2内存，组织为时分复用或并行通路来传输数值数据到自定义IP，并携带额外信息（当前帧均衡器的系数），利用控制流通道传输过去。处理结果仍传回DDR2中，交给主机显示或存储为文件。和DDR2的通信需要借助AXI-HP物理接口，PL部分为master，负责数据搬移。 通过以上论述，应该比较清楚的看到整个数据走向了。具体实施细节还需要进一步研究。
-
发表了文章 2013-07-07
数字地面电视中单频网介绍
为了实现更大区域覆盖，数字电视布站即沿袭了模拟电视的布站方式即多站多频网覆盖，这样可以避免不同站之间互相干扰。除了这种方式，还有一种数字电视专有的覆盖方式，即单频网覆盖。 单频网，英文名SFN（Single Frequency Network），是指多个发射站在同一时间、同一频率上发送同一节目内容，即“三个同步”。这样可以减少频谱资源占用，同时在某些覆盖交叠区能得到更好的接收效果。 单频网最大的问题是如何保证同步，尤其是时间同步。一般采用GPS进行同步，时间精度可达到us级别。其次问题是强多径干扰抑制，当接收机接收到不同时延的两个发射站信号时，相当于信号经过不同路径到达接收机，而且多径分量相对于直达分量的幅度相差可能不超过5dB，属于相当强的多径分量，相对时延可能超过100us，均衡根本不起作用，如果直接解调，多径分量对直达分量的干扰将导致信噪比极低，从而使接收失败。该问题可以采用OFDM技术解决。OFDM为正交频分复用技术，可将串行传输的码元调制到相互正交的不同子载波上同时传输，而每个子载波的持续时间变长，从而使码间串扰对其影响变小。通过发送端添加循环前缀，将接收机得到的信号进行子载波解调，可估计多径分量相对直达分量的时延，频域均衡技术可以把多径分量与直达分量合并，从而使接收信号不受多径影响，实现可靠通信。前提：循环前缀长度要覆盖多径可能分布的时延范围。 北京市数字地面电视采用单频网的频道号为32频道，频率666MHz，播放内容为6套标清（CCTV1，CCTV2，少儿，音乐，BTV1，CETV3）。
暂无更多信息
-
发表了文章
2017-11-08
CUDA编程模型【中科院课件】
-
发表了文章
2017-11-08
CUDA存储器【中科院课件】
-
发表了文章
2017-08-07
Jetson TX2 初体验
-
发表了文章
2016-12-23
如何让你的机器学习玩超级玛丽
-
发表了文章
2016-07-14
从Caffe开始深度学习实战，徒手hack代码乐趣更多
-
发表了文章
2016-03-28
Windows 上 Caffe开发环境搭建
-
发表了文章
2008-07-19
自己写的一个链表应用程序
-
发表了文章
2010-03-29
想考研
-
发表了文章
2010-07-16
课程设计——数字通信计算机仿真：2PSK调制仿真
-
发表了文章
2011-05-13
TMS320VC5509A上的无需CSL库的USB底层程序
-
发表了文章
2013-06-30
数字地面电视软件接收机设计
-
发表了文章
2013-07-01
中国数字地面广播电视DTTB（国标GB20600-2006）介绍
-
发表了文章
2013-07-07
数字地面电视中单频网介绍
-
发表了文章
2013-07-10
【OpenHW参赛手记】AXI初体验
-
发表了文章
2013-07-10
【OpenHW参赛手记】System Generator 与XPS连接的方法
-
发表了文章
2013-07-10
【OpenHW参赛手记】AXI-Stream接口介绍
-
发表了文章
2013-07-10
【OpenHW参赛手记】AXI-Stream接口开发详细流程
-
发表了文章
2013-07-22
CUDA从入门到精通（零）：写在前面
-
发表了文章
2013-07-22
CUDA从入门到精通（二）：第一个CUDA程序
-
发表了文章
2013-07-23
CUDA从入门到精通（三）：必备资料
滑动查看更多
-
回答了问题
2016-05-04
【求助】想在阿里云上购一台云服务器进行ae的渲染，不知可不可行
赞0 踩0 评论0 -
回答了问题
2016-04-03
-
回答了问题
2016-04-03
-
回答了问题
2016-04-02
-
回答了问题
2016-04-02
-
回答了问题
2016-04-02
-
回答了问题
2015-10-21
阿里云ECS是否支持GPU
Re阿里云ECS是否支持GPU2015年10月15日云栖大会上推出了GPU物理机，现在已经可以正式购买，链接为 https://www.aliyun.com/product/hpc赞0 踩0 评论0
滑动查看更多
暂无更多信息