OpenHero 开勇's profile开勇 OpenHeroPhotosBlogListsMore  |  Tools Help |
开勇 OpenHerohttp://blog.csdn.net/openhero |
|||||||||||
|
October 21 《GPU高性能运算之CUDA》我们的书总算出来了
http://www.dearbook.com.cn/book/255585 http://www.china-pub.com/48582&ref=ps http://product.dangdang.com/product.aspx?product_id=20705176 http://product.dangdang.com/product.aspx?product_id=20705187 历时应该有一年多时间，从最开始组织吆喝一群朋友帮忙写书，到中途一些事情的变故，再到后来张舒和艳利同学的不懈努力，终于经历一年多的时间把这本书完成了。 张舒和艳利在书的编写中作出的不懈努力值得我好好学习。张钰博同学已经去了UC 戴维斯 读博士，非常感谢晓炜同学那天能收留我在他们宿舍过夜，也是在那天晚上和张钰博同学聊到了CPU加速，聊到了GPU加速，聊到了CUDA，也才有了最终的这本书。很感谢张钰博同学能带我进入GPU的世界，也是沾了浙大的光…… 这一年的光影仿佛历历在目，百感交集，有好多话想讲，却不知道从哪里开始。 生活，学习，工作…… 继往开来，步步为营，踏踏实实，是对自己的督促，也是给朋友们的勉励。 CUDA让我结识了很多朋友，从全国的巡讲： 广州 （2009年3月28日，星期六） 华南理工大学 A4204/学术报告厅（圆满结束） 到中科院的GPU研讨会： 再到Nvidia的GTC全球研讨会； 查看自己Email收发的CUDA相关的问题，超过几千封~回答过不止上千个的问题…… 很感谢CSDN能给我提供这样的平台，也很感谢Nvidia的，邓老师，魏鸣，谢强，还有很多帮助过我的朋友，由衷的感谢…… 中科院的葛蔚老师，张云泉老师，王龙博士，万宁博士，张先轶同学，陈实同学，刘伟峰...还有很多中科交流的时候认识的老师，同学…… 清华的邓仰东老师…… 还要感谢浪潮公司能给我和很多公司企业交流的机会,感谢王总，刘总对我的支持，还有team里面的兄弟的支持……； 在GTC遇到的联想的中国和俄罗斯大区的VP，魏总，也是我北理工的师兄，对我的鼓励，还有史博士，真的好感激…… 吉星吉达的佟小龙…… 全国各地好多的老师，同学，朋友…… 还有国外的朋友…… 还有找我写作业的国外的同学……让我知道了，不只是有国内的同学有找别人帮忙写作业的情况，其实全世界都有~~只是多少问题…… 还有台湾国立中正大学电机所得 陈永维 同学，到香港来旅游的时候，还送我一盒月饼，真的好感谢…… 还有参加GTC的时候送书给我的Steve Worley，虽然我知道他看不懂中文…… 还有这一年来一直陪在我身边的朋友…… 由于CUDA，我认识了好多好多的朋友，非常感谢，是你们的鼓励和支持，才让我感觉到了温暖……谢谢！ September 19 向量乘以矩阵(vector_matrix_multiplication)/********************************************************************
/************************************************************************/ bool InitCUDA(void){return true;} #else cudaGetDeviceCount(&count); for(i = 0; i < count; i++) { printf("CUDA initialized.\n"); #endif A = A + threadIdx.x; float subC = 0.0; for (int i = 0; i < wA; i+=64) #pragma unroll C[0] = subC; //__global__ static void vector_matrix_mult_kernel_32T(float* A, long wA, float* B, long wB, float* C) #define RUN_TEST #define WA 32 if(!InitCUDA()) { srand(2009); long wA = 64* WA; long size_A = wA; float *hA = (float*)malloc(sizeof(float) * size_A); for (int i = 0; i < size_A; i++) for (int i = 0; i < size_B; i++) float *dA = 0; CUDA_SAFE_CALL( cudaMalloc((void**) &dA, sizeof(float) * size_A)); CUDA_SAFE_CALL( cudaMemcpy(dA, hA, sizeof(float)*size_A, cudaMemcpyHostToDevice)); unsigned int timer = 0; dim3 threads = 64; CUT_CHECK_ERROR("Kernel execution failed\n"); CUDA_SAFE_CALL( cudaMemcpy(hC, dC, sizeof(float) * size_C, cudaMemcpyDeviceToHost)); for (int i = 0; i < wB; i++) CUT_SAFE_CALL( cutStopTimer( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); #ifdef RUN_TEST CUTBoolean res = cutCompareL2fe(testhC, hC, size_C, 1e-6f); CUDA_SAFE_CALL( cudaFree(dA)); return 0; CUDA 2009编程 大赛结果http://cuda.csdn.net/Contest/pro/nvidia_results.html 历时5个月的NVIDIA CUDA 在8月落下帷幕，大赛举办期间。有300多个学校700多名学生报名。 最终结果 特别贡献奖 奖 别 姓 名 学 校 参赛作品 特别贡献奖 彭江锋 施少怀 漆舒汉 杨植群 华南理工大学/软件学院 基于GPU的多模式网页精确匹配系统 自选赛奖项 奖 别 姓 名 学 校 参赛作品 一等奖 陈实富 中国科学院/深圳先进技术研究院 Glirt:基于CUDA的多模态三维医学图像配准 二等奖 陈伯君 北京大学/工学院 基于CUDA的DSMC高性能并行计算 二等奖 喻勤 电子科技大学 /电子工程学院 基于CUDA的高性能SAR成像模拟 三等奖 胡晓菡 上海交通大学 /生命科学技术学院 CUDA平台下的复杂疾病全基因组基因相互作用计算 三等奖 陆扬 中国科学技术大学 /工程科学学院 使用CUDA模拟顶点、几何着色器-原型演示 命题赛奖项 奖 别 姓 名 学 校 参赛作品 一等奖 翟艳堂 中国科学院 /计算机网络信息中心 稀疏大矩阵与矢量的乘积运算 二等奖 包南森 上海大学/计算机学院 粒子邻居搜索 二等奖 尚书杰 谢迪 浙江大学 /计算机学院 稀疏大矩阵与矢量的乘积运算 三等奖 褚艳利 邵帅 年华 西安电子科技大学 /计算机学院 基于CSR表示的稀疏矩阵-矢量乘法的CUDA实现 三等奖 张广勇 内蒙古大学 /计算机学院 基于GPU 平台的稀疏矩阵与矢量乘积 优秀作品奖 姓 名 学 校 参赛作品 朱小松 大连理工大学 /土木水利学院 基于CUDA的有限差分法求解简单热传递问题 冯涛 沈乐 清华大学 /工程物理系 CUDA 平台下三维反应堆中子通量分布和有效增 尹康学 长安大学 /信息学院 稀疏大矩阵与矢量相乘基于CUDA的一种可行算法 朱德东 复旦大学 /软件学院 基于CUDA的多分辨率图像融合算法 杨植群 华南理工大学 /计算机科学与工程学院 粒子邻居搜索 石丹 范正娟 赵安元 四川大学 /软件学院 CUDA平台下的超声图像斑点噪声抑制 陈晓熹 厦门大学 /信息科学与技术学院 CuParcone 李斌 段元泽 云南大学 /资源环境与地球科学学院 即时计算的动态分形屏保 优秀作品奖得主将会获得由NVIDIA 中国公司颁发的荣誉证书 恭喜以上获奖选手，请保持联系方式畅通，近期会与您联系相关事宜。 July 14 10. CUDA cosnstant使用（一）------GPU的革命10. CUDA cosnstant使用（一）------GPU的革命 序言：最近的事情无比的多，差点就找不到回家的路了，都快忘记出发的起点的时候，冷静下来，侧夜未眠，事情再多，都要一件一件的做好，做不好的，就不接，我想别人也可以接受的，我的个人的能力也是有限的，尽人事，听天命。有的时候，更多的时候是在听天命，而不是在尽人事。不过有时候累得不能再累的时候，真的好想什么都不做，就想回家好好的躺上几天，什么都不用思考，安安静静的就好。回过神来，一件事情，一件事情的做，做好一件，在做另一件，给自己信心，只有自己能把一件一件的事情都步步的做好的时候，就是给自己最大的信心。信心不是别人给的，是自己给自己的。正文： 书接上回《9.CUDA shared mem使用》讲了shared memory的使用，最近有几个朋友都在问我cosntant的使用的问题，这次首先先讲一下cosntant的使用，下一章节才讲一下cosntant的使用中性能的体现；下面是一个简单的代码:/******************************************************************** * cosntant_test.cu * This is a example of the CUDA program. * author: zhao.kaiyong(at)gmail.com *********************************************************************/ #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #include <cutil.h> /************************************************************************/ /* Init CUDA */ /************************************************************************/ #if __DEVICE_EMULATION__ bool InitCUDA(void){return true;} #else bool InitCUDA(void) { int count = 0; int i = 0; cudaGetDeviceCount(&count); if(count == 0) { fprintf(stderr, "There is no device.\n"); return false; } for(i = 0; i < count; i++) { cudaDeviceProp prop; if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) { if(prop.major >= 1) { break; } } } if(i == count) { fprintf(stderr, "There is no device supporting CUDA.\n"); return false; } cudaSetDevice(i); printf("CUDA initialized.\n"); return true; } #endif /************************************************************************/ /* Example */ /************************************************************************/ __constant__ char p_HelloCUDA[11];// = "Hello CUDA!"; __constant__ int t_HelloCUDA[11]={0,1,2,3,4,5,6,7,8,9,10}; __constant__ int num = 11; __global__ static void HelloCUDA(char* result) { int i = 0; for(i = 0; i < num; i++) { result[i] = p_HelloCUDA[i]+t_HelloCUDA[i]; } } /************************************************************************/ /* HelloCUDA */ /************************************************************************/ int main(int argc, char* argv[]) { if(!InitCUDA()) { return 0; } char helloCUDA[] = "Hdjik CUDA!"; char *device_result = 0; char host_result[12] ={0}; CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(char) * 11)); CUDA_SAFE_CALL( cudaMemcpyToSymbol(p_HelloCUDA, helloCUDA, sizeof(char)*11) ); unsigned int timer = 0; CUT_SAFE_CALL( cutCreateTimer( &timer)); CUT_SAFE_CALL( cutStartTimer( timer)); HelloCUDA<<<1, 1, 0>>>(device_result); CUT_CHECK_ERROR("Kernel execution failed\n"); CUDA_SAFE_CALL( cudaThreadSynchronize() ); CUT_SAFE_CALL( cutStopTimer( timer)); printf("Processing time: %f (ms)\n", cutGetTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); CUDA_SAFE_CALL( cudaMemcpy(&host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost)); printf("%s\n", host_result); CUDA_SAFE_CALL( cudaFree(device_result)); CUT_EXIT(argc, argv); return 0; } 这里写了两种使用cosntant的方法： 1. 一种方法是直接在定义的时候初始化constant： __constant__ int t_HelloCUDA[11]={0,1,2,3,4,5,6,7,8,9,10}; __constant__ int num = 11; 在kernel里面直接使用就可以了； 2. 第二种方法是定义一个cosntant数组，然后使用函数初始化它； __constant__ char p_HelloCUDA[11];// = "Hello CUDA!"; CUDA_SAFE_CALL( cudaMemcpyToSymbol(p_HelloCUDA, helloCUDA, sizeof(char)*11) ); 前面一句是定义constant，后面一句是初始化这个常量，在kernel里面使用的时候，就按照定义的方式使用就可以了； 当然也有朋友想定义自己的结构体，当然是可以的，只是在初始化的时候，copy相应的结构体就可以了。这个只是一个初步的使用方法，希望对家有用。 下一次会对cosntant的使用中体现出来的性能优势做一个简单的分析。 June 13 绿坝 图像过滤比对算法有感最近大家都在说绿坝，其实核心算法部分，就是图像识别和，图像比对，现在通用的图形库有 OpenCV/paintlib/CImg/FreeImage/CxImage/SILLY/Corona ，这个是当前个人感觉比较强大的四大图形库，但是这些都是架构在CPU上的，当然，也有一些，现在正在加入gpu的支持，加入多核的支持，当然看过一些对绿坝的分析，当然也下载了绿坝，看了一下，没有时间用IDA反汇编去看里面的图形比对的算法，是否是使用的是OpenCV的库，或者伪造一下OpenCV的库，在里面设定一些断点，（ps：OpenCV是开源的代码）当然可以查看， 不过不用管那么多，至少可以肯定的是在绿坝里面用到了OpenCV的dll，还有Cximage的库；感叹一下，真的不希望这个成为一个闹剧；毕竟是政 府的工程；一个朋友在密歇根大学的分析基础上一晚上写了一个漏洞的测试工具，不卡、不弹、不闪，，，，这个是难得的网页溢出漏洞~- - 哎= =记得这个朋友当初测试vista 0day的时候= =vista 那个时候还没出来多久= =囧啊~~下面都是转载~~~认识绿坝的朋友的，也可以给他们善意的给一些建议吧~~希望能做得更好，更完善= =真正的安全，放心，可靠，真正的能让“父母”放心的软件。
一下内容从一个朋友blog上转载：http://hi.baidu.com/hex1337/blog/item/13cc5bd014fff2da572c840d.html 《跨上草尼马,挥刀斩河蟹-绿坝V3.17超长URL溢出漏洞exploit by seer》 28度的冰 19:29:31 哪里有绿坝下载啊。我正需要这样一款软件。 28度的冰 19:29:39 平时上网总是有暴力 28度的冰 19:29:45 色情网站蹦出来 28度的冰 19:29:48 可讨厌了 也云 19:29:54 2 28度的冰 19:29:58 正好用绿坝过滤一下 seer 19:30:01 2 28度的冰 19:30:29 这么好的软件就算收费也是值得的 28度的冰 19:30:49 况且还免费预装 seer 19:30:48 我们要相信ZF是在保护孩子们 28度的冰 19:30:51 太好了 seer 19:31:12 比如，它只拦截IE 这就从小教会了孩子们使用Firefox ============用群里的某段聊天来开篇============== 为了监视用户的浏览网页行为,绿坝注入了IE进程,在处理超长的URL时,就会在IE的进程空间中触发一个典型的栈溢出.具体原理可以看看密歇根大学的牛人们给出的分析http://www.cse.umich.edu/~jhalderm/pub/gd/, 不错的学校... Exploit中使用了.net部署shellcode的方式,(BlackHat2008上公布的那招),这是一种比HeapSpray更稳定的方式,而且不卡IE,还能绕过Vista下的DEP和ASLR. 懒得写生成器了,我给出的利用程序shellcode是弹出XP sp2的计算器,所以更接近一个PoC,要换shellcode的话需要手动改下.net中的shellcode,然后用VS重新编译exploit.dll,这也算是防止被菜鸟们滥用的一个小小的门槛吧... Tips1:由于带有.net控件,这个exploit必须在IIS上发布(需要服务器装有.net framework,IIS不需要特别的设置),直接点开是不行的 Tips2:具体的编译方式放在附件中了,复制粘贴吧,只要你有VS开发环境... exploit下载:http://www.namipan.com/d/40c1b70f9ed40b2baf3e0fbe826254a324c2823755090000 Enjoy Hacking~;-) ps:转载 某 群 某 某 某 的 随便之作：）也是一个朋友，不方便透露消息了~ Enjoy Hacking~~- == =社会需要和谐： 绿坝的核心的人脸比对技术到底是不是使用opencv，并不一定能肯定，但是可以肯定的是使用了OpenCV的开源库，OpenCV是BSD的授权，但是绿坝的软件中好像一点都没提到BSD授权的事情，也可以申请专利？不知道是以什么申请专利的，这里面的细节也就不用太多的探讨，心知肚明就行。不过要是真正想做这样的过滤软件，做一个好的过滤软件，还是希望静下心花时间做真正自己国产的~= =和谐，和谐，低调，低调~~ 只是自己也一直想做一些开源的图像处理的库，设想一下，如果自己写的开源的东西，被别人拿来申请专利，拿来做商品……心里却是会有点不舒服，~~不过毕竟这样的人也是少数吧~更多的还是学习研究用~~：）只有Open的才是能经得起考验的~~ 下面两个连接是对绿坝的详细分析： 这个是密歇根大学对绿坝漏洞的详细分析：http://www.cse.umich.edu/~jhalderm/pub/gd 也希望做绿坝的人也可以好好看看，做这样的安全产品也要真正的安全：）善意的提醒·~我可不像树敌~= = 这个是对花季护航的 分析：https://docs.google.com/Doc?docid=dczkbptk_0ffc2hvc9&hl=en 这个是对绿坝的软件分析：https://docs.google.com/view?id=afk7vnz54wt_12f8jzj9gw 《绿坝-花季护航软件技术分析（不断更新中）》也只是希望政 府 如果真正希望做好的软件，好的产品，也得做出一个样子来：）善意的提醒。！ ps：不知道这篇文章，会不会被和谐掉~~吼吼~~ PS：顺便推广一下GPU，哈哈，，，绿坝的核心算法就是图片比对，这块要是用GPU来做，不用openCV开源库，性能会提升差不多1到两个量级，让大家尽情的yy一下吧~程序必须使用CUDA开发，然后政府要求，必须使用NVidia的显卡~Nvidia的卡支持CUDA~~~然后用stream开发的，必须使用ATI的卡~- = 每一个省用不同的卡~= =安装不同的软件= =炯炯有神了 那intel咋办？咱用Ct或者OpenCL，Intel Nvidia AMD都可以使用= = 好吧，给绿坝的开发团队，真诚而善意的提一些建议：使用GPU开发吧，性能会快很多，自己写比对算法，还不用OpenCV的库，不用考虑BSD这样的协议问题= = 现在已经有朋友可以利用URL的漏洞开发了一个远程注入的程序，希望也善意的提醒一下用户和开发团队，经快的修补漏洞，或者出来辟谣一下……毕竟这是 一个 政 府 项目，也希望能引起足够的重视，谢谢~~ |
|
|||||||||
|
|
