|
|
September 19 /********************************************************************
created: 2009/09/19
created: 19:9:2009 12:00
filename: vector_matrix_multiplication.cu
file base: vector_matrix_multiplication
file ext: CUDA
author: zhao.kaiyong(at)gmail.com
purpose: vector matrix multiplication
copyright: everyone can use this code, please specify source.
任意使用，请注明出处;
http://www.hpctech.com
http://openhero.net
*********************************************************************/
#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 */
/************************************************************************/
__global__ static void vector_matrix_mult_kernel(float* A, long wA, float* B, long wB, float* C)
{
__shared__ float subA[64]; A = A + threadIdx.x;
B = B + blockIdx.x * 64 + threadIdx.x;
C = C + blockIdx.x * 64 + threadIdx.x; float subC = 0.0; for (int i = 0; i < wA; i+=64)
{
subA[threadIdx.x] = A[i];
__syncthreads(); #pragma unroll
for (int j = 0; j < 64; j++, B += wB)
{
subC += subA[j] * B[0];
}
__syncthreads();
} 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
#define WB 64
/************************************************************************/
/* HelloCUDA */
/************************************************************************/
int main(int argc, char* argv[])
{ if(!InitCUDA()) {
return 0;
} srand(2009); long wA = 64* WA;
long wB = 64* WB; long size_A = wA;
long size_B = wA*wB;
long size_C = wB; float *hA = (float*)malloc(sizeof(float) * size_A);
float *hB = (float*)malloc(sizeof(float) * size_B);
float *hC = (float*)malloc(sizeof(float) * size_C);
float *testhC = (float*)malloc(sizeof(float) * size_C); for (int i = 0; i < size_A; i++)
{
hA[i] = (float)rand()/(float)RAND_MAX;
} for (int i = 0; i < size_B; i++)
{
hB[i] = (float)rand()/(float)RAND_MAX;
} float *dA = 0;
float *dB = 0;
float *dC = 0; CUDA_SAFE_CALL( cudaMalloc((void**) &dA, sizeof(float) * size_A));
CUDA_SAFE_CALL( cudaMalloc((void**) &dB, sizeof(float) * size_B));
CUDA_SAFE_CALL( cudaMalloc((void**) &dC, sizeof(float) * size_C)); CUDA_SAFE_CALL( cudaMemcpy(dA, hA, sizeof(float)*size_A, cudaMemcpyHostToDevice));
CUDA_SAFE_CALL( cudaMemcpy(dB, hB, sizeof(float)*size_B, cudaMemcpyHostToDevice)); unsigned int timer = 0;
CUT_SAFE_CALL( cutCreateTimer( &timer));
CUT_SAFE_CALL( cutStartTimer( timer)); dim3 threads = 64;
dim3 blocks = wB/64;
vector_matrix_mult_kernel<<<blocks, threads>>>(dA, wA, dB, wB, dC); CUT_CHECK_ERROR("Kernel execution failed\n"); CUDA_SAFE_CALL( cudaMemcpy(hC, dC, sizeof(float) * size_C, cudaMemcpyDeviceToHost));
CUT_SAFE_CALL( cutStopTimer( timer));
printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
CUT_SAFE_CALL( cutResetTimer( timer)); for (int i = 0; i < wB; i++)
{
float subC = 0.0;
for (int j = 0; j <wA; j++)
{
subC += hA[j] * hB[j*wB + i];
}
testhC[i] = subC;
} CUT_SAFE_CALL( cutStopTimer( timer));
printf("Processing time: %f (ms)\n", cutGetTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); #ifdef RUN_TEST CUTBoolean res = cutCompareL2fe(testhC, hC, size_C, 1e-6f);
printf("Test %s \n", (1 == res) ? "PASSED":"FAILED");
#endif CUDA_SAFE_CALL( cudaFree(dA));
CUDA_SAFE_CALL( cudaFree(dB));
CUDA_SAFE_CALL( cudaFree(dC));
free(hA);
free(hB);
free(hC);
CUT_EXIT(argc, argv); return 0;
} 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的革命 序言：最近的事情无比的多，差点就找不到回家的路了，都快忘记出发的起点的时候，冷静下来，侧夜未眠，事情再多，都要一件一件的做好，做不好的，就不接，我想别人也可以接受的，我的个人的能力也是有限的，尽人事，听天命。有的时候，更多的时候是在听天命，而不是在尽人事。不过有时候累得不能再累的时候，真的好想什么都不做，就想回家好好的躺上几天，什么都不用思考，安安静静的就好。回过神来，一件事情，一件事情的做，做好一件，在做另一件，给自己信心，只有自己能把一件一件的事情都步步的做好的时候，就是给自己最大的信心。信心不是别人给的，是自己给自己的。 正文： 书接上回《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的漏洞开发了一个远程注入的程序，希望也善意的提醒一下用户和开发团队，经快的修补漏洞，或者出来辟谣一下……毕竟这是 一个 政 府 项目，也希望能引起足够的重视，谢谢~~ June 01 2009-05-27参加了中科院的会议，本来准备了很多东西，不过只有15分钟的时间，即使用了最快的语速，好像速记员都不打字了……在会上学习了好多的技术，也了解了国内的GPU方向的比较有代表性的一些工作，不过还是有一点小小的感触，做底层研究的好少，对gpu性能本身挖掘的就更少了，好多的都是在上面跑应用，但是这样的应用能达到gpu本身的性能的多少啦？不过现在确实需要一些应用来驱动这个事情的发展，GPU只能看成是多核发展的一个代表而已，更应该看到的是整个多核芯片的发展，这个才是大方向，Nvidia，AMD，Intel各家都在向这个方向发展，硬件的发展将在这几年里面发生革命性的变换，但是这个时候更应该抓住的是并行算法的研究和硬件发展的趋势的研究，只有这两个抓住了才能适应最新的平台，如果只是跑应用，累的只是开发人员。 这次会议学习到很多，听到很多老教授的发言，很有感触，下面附上会议上的简单的报告~名字有点大~~~我也只是早一些接触GPU、多核开发而已，还需要更多的人的共同努力，真的好想中国在这块不落后别人，加油！年轻的一代共勉吧~~ GPU Computing Technology Overview.pdf April 08 前言：很久没写CUDA相关的文章了，其实也不是忙，只是零碎的事情比较多，不能抽出完整的时间写一些东西，在http://blog.csdn.net/openhero 上写本来想写一些列CUDA编程的文章，不过现在看来，很多朋友还是只是处在开发的初级阶段，一些基本的编程环节还需要讲解一下，其实像lib编程，dll编程，都不是CUDA的内容，这个只是windows，linux系统编程的内容，只要有时间，好好静下心来做几个项目就可以明白很多，而且网上就可以找到很多资料。 正文：CUDA的DLL开发其实和一般的C/C++的DLL开发是一个原理，当然，DLL的开发就有几种方式，这里就讲最容易理解的，也最直接的方式，然后把代码放出来。大家自己可以琢磨一下其它的方式。 1. 创建DLL 用我的Wizard 可以创建一个DLL项目工程： 
选择DLL项目，然后可以看到下面的项目文件结构： 
其中sampe_cu.h文件为头文件，stdafx.h和stdafx.cpp是windows的加载的相关文件，这里可以不用理会，DLL_Test.cpp是其中的一个导出函数文件，也是声明了DkkMain的入口函数，这里是在DLL创建，加载，卸载的时候需要处理的一些内容，这里如果不熟悉DLL也不用管它。 针对CUDA，这里我们就来看sample.cu文件： /************************************************************************/ /* HelloCUDA */ /************************************************************************/ int TestHelloCUDA(void) { if(!InitCUDA()) { return 0; } char *device_result = 0; char host_result[12] ={0}; CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(char) * 11)); unsigned int timer = 0; CUT_SAFE_CALL( cutCreateTimer( &timer)); CUT_SAFE_CALL( cutStartTimer( timer)); HelloCUDA<<<1, 1, 0>>>(device_result, 11); 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; } 这个就是sample.cu文件里面的实现的dll的接口函数，当然还有一个def文件，这个文件是用来做动态导出用的，看看def文件： LIBRARY "DLL_Test" EXPORTS TestHelloCUDA 这个是sample.def文件的内容，这个文件是为了能动态导出DLL里面的函数接口而用的。编译器通过这个文件把声明的TestHelloCUDA，接口导出。 编译，就可以在相应的目录里面得到一个xxxx.dll文件. 2. 调用DLL 这里就是测试DLL文件，当然我在同一个工程里面创建一个调用的测试工程： 
这里添加一个测试工程，当然你也可以再另外的地方创建一个测试工程，或者加载到自己的工程里面； 选择一个win32工作作为测试例子； 
这个例子就是一个简单的测试项目，都选默认得就可以； 
这里默认的就可以了，其他的自己去查找意思J不能偷懒到什么都告诉你…… 
这个时候就可以看到两个工程，一个是dll工程，一个是测试工程，然后在测试工程里面，我们采用动态加载DLL的方法来调用DLL； 
上面就是调用DLL的代码，这里为什么加__T()自己去找unicode的调用方法J 里面的__T()里面的dll的文件名就是我们要加载的DLL，当然这里是路径，入下图，我们的dll和exe在同一个文件夹里面，如果是不同的文件夹，那就用相对路径来找就okJ 
然后运行例子，就可以得到下图结果: 
这里的DLL就ok了~~~ 代码可以再这里下载： http://d.download.csdn.net/down/1188868/OpenHero http://groups.google.com/group/cudagroup/web/DLLTest.rar?hl=zh-CN March 19 从07年开始做CUDA到现在，都一年多了，从08年的4月份开始做CSDN的CUDA总版主到现在也差不多快一年，期间一直坚持做CUDA的学习和研究。其实CUDA不是我的主专业，gpu也不算是我的主专业，也只是业余爱好。但是业余爱好坚持下来做好，或许也能成为专业吧。不过CUDA对我的吸引到不是gpu的提速，而是给我提供了一个学习并行编程，并行计算的平台，因为我知道接下来的几年高薪能计算的方向应该是向单机HPC的方向发展，再加上集群化，分布式并行计算才是真正的发展方向，CUDA也只是我们学习并行计算的很好的切入点。 从08年初就开始和Nvidia的marketing的Ming和Perry有联系。可以看到CUDA在国内的发展，从开始的几个人，慢慢的增加到这么大的队伍，好几个QQ群都已经装不下，真的是很欣慰吧。只是觉得自己做的事情还不够，当初答应写书的事情，现在还没写完，本来想写一本CUDA的入门的教材，但是后来修改过了三版，反复琢磨，觉得要是写一本简单的入门的书，对于做CUDA的朋友来说，就是太不负责任了，要是CUDA不明白其中的软硬件平台，不明白优化，那就真的做不好CUDA编程，所以从去年5月份写，一直到现在，都在修改。也因为一直都太忙，自己也还是学生，一周有两门课，还有两个TA要带，还得写自己的paper，干boss的活，所以这件事情希望能在09年上半年搞定，也希望能给大陆学习CUDA的朋友一些一些帮助。 其实在08年刚开始做CUDA板块的总版的时候，就在和Gemin商量，做一个全国的性质的讲座，和一个全国性质的比赛，从那个时候开始，或许就在策划现在的巡讲和比赛吧。 前段时间一直在收集CUDA编程的题目，不过最终还是考虑到通用性的情况，通过评委的讨论，限定的题目就不一定要太多，就选了两个，其他的题目都是可以自己发挥的，能作出重大贡献的还可以有2w的奖励，这个是全国宣讲和比赛的网站入口 http://cuda.csdn.net/contest/pro/nvidia_home.html 这个是全国巡讲的网站 http://cudaevent.csdn.net/ 这个是CUDA全国比赛的网站 http://cuda.csdn.net/Contest/pro/index.aspx 其实对于全国巡讲，对于我而言，我自己真的是很希望能到更多的学校做推广，当然，不一定是只是CUDA编程，更多的应该看准的并行编程，CUDA是我们做并行编程的很好的切入点。 真的很希望中国的并行计算能作出一些让世人都瞩目的成绩来，我在其中能做一个穿针引线的工作已经很满足了~~ 知人者智 自知者明
胜人者有力 自胜者强
知足者富 强行者有志
不失其所者久 死而不亡者寿 ----《道德经》三十三章
ps：最近和boss又发了两篇文章，一篇Journal和一篇conference 有兴趣的可以到我的学校主页上看看 http://www.comp.hkbu.edu.hk/~kyzhao/ March 16 第一场景： 在宇宙飞船上，都是外星人……囧 3D实景的动画游戏，可以两人玩的对打的游戏……游戏现实不是平面的，而是空间立体显示的3d场景，很像日本动漫里面的太空堡垒里面的3D游戏……when will it be come true！很期待这样的游戏……动画场景，动画人物都可以在空间中显示出来……很帅…… 场景二： 飞船降落到地球上，外星人说自己是来之远古的地球人……囧…… 外星人就成地球人了…… 然后大家说吃早餐……吃包子…… 场景三： 理工良乡校区，张导带着他们外语学院的n多mm参加什么活动，也要说先吃早餐……结果他说只能吃方便面……囧，囧…… 场景四： 然后继续走，一群人来到一个寺庙里面，拜佛…… 听说长老还在打坐，需要我们在外面等…… 场景五： 长老的大徒弟从什么地方跳出来，直接就撞钟了，我还劝他别吵闹到长老了…… 场景六： 长老结果从寺院的大钟里面出来，还领着一个初中的同学……汗- -！ 场景七： 长老说我们几个是西游记里面的人物，说他自己是唐僧……我是悟空……还有八戒……还有沙僧……o(╯□╰)o 场景八： 我问长老，我要是悟空还能结婚吗？……超级汗- -！问他是不是经历了很多磨炼，把唐僧送到了，就成佛就不能娶老婆了……- -！o(╯□╰)o斗战胜佛能取老婆吗？ 场景九： 八戒和谁在理论什么，说什么辈分，好像唐僧转世过后的辈分比他低，还得叫他师傅，他有点不爽…… ………… 诡异的梦，我还在想着……突然醒了……- - ………………………………………华丽的分割线………………………………………………………………… 早上6点左右就醒了，再也睡不着……感觉很诡异的梦……经历了很多，难道就需要成佛，斩断所有的尘世？- -囧 最近发现自己有点找不到方向了……感觉自己有的时候变得有点不像自己那么干练…… 有很多担心的事情，担心什么？害怕什么啦？ 一个人吃饱全家人不饿的心态没有了，多了牵挂多了思念…… 有点放不开了？ = =好像自己从来没这样的小心过，小心什么啦？…… 或许就是因为这样，才那么能引起天蝎座的我痴迷吧…… …………………………………………………………分割线………………………………………… 出来的时候看到天空有晨晖……很美…… 透过白色的云，能看到香港蓝色的天空，很美，香港好像也很久没有晴天过了……前几天总是阴沉沉的，有点闷，或许就像前几天的感冒一样，不管是身体还是神经，都有点闷…… 在想是不是老外在见面的时候喜欢谈天气一类的，其实也是在说心情啦？ 今天天气不错…… 快到学校的时候，等红绿灯……对面看到一位大妈，咋眼一看，外套低胸快到肚脐了- -暗自感叹……哎，我真的out time 了吗？- -nnd~~我放不开？- -好奇心……打着胆子，提了又提胆子……在过红绿灯的时候，瞟了一下~- -- - --囧，人家穿的是肉色的毛衣……- - - ---------------------------分割线--------------------------------- 到office就想写点什么，结果看到gimy同学在英国时间的凌晨几点还在贴图片，她的男人女人……结果翻过来，再翻过来看，都没发现咱露脸的机会，哈哈~~gimy来一句，你太虚幻了，哈哈~~ 确实哈·~…… -----------------------------分割线----------------------------------- 早上洗漱完就开始找自己的运动装和运动鞋……nike的长裤，nike的短裤，还有adidas的不长不短的裤子，都从箱底翻出来……咱要开始跑步了~~为了恢复8块腹肌的目标……我要跑步了……= =！现在就剩下一块腹肌了- -…… 要是你们不信，我贴一张当年运动的我的照片…… 
顺便贴一个最近听的好玩的歌： March 13 导师招博士： http://www.comp.hkbu.edu.hk/~chxw/ 并行+分布，GPU+multicore 所以要求动手能力和解决问题的能力比较强 喜欢编程 MPI, CUDA，OPENMP，算是bonus 奖学金跟你们一样，是HK$13200/月 我会带CUDA和高性能编程 简历发到zhao.kaiyong(at)gmail.com March 10 3月2号，浪潮发布了CPU+GPU的Tesla架构的国产服务器，是对国内的用户的一个很好的信息。http://server.zol.com.cn/124/1241702.html 浪潮倚天服务器采用的是CPU+GPU异构架构，这样可以让cpu和gpu都能够完成自己最擅长的工作。gpu做为并行处理的强力工具，cpu做为串行和任务的分配的工作，这样可以使得系统在传统的cpu架构上有量级的提升。 单机高性能计算是一个大势所趋，将来的更高性能的芯片将会出现，浪潮算是走在了国内高性能计算的前列，希望更多的国产企业能在高性能计算中贡献一份自己的力量。GPGPU的计算在国外已经是从01，02，03就开始的工作，即使是CUDA也是在07年就开始在国外有很多研究机构，很多的国际会议也在08年看到大量的gpu计算的文章。自己写CUDA的blog，也只是希望能尽自己的一份微薄之力，把最新的国外的一些技术介绍给大家，希望真正的能为中国的计算机做点力所能及的事情，虽然我知道自己的能力有限，但是只是希望能写这样的文章，引起更多的朋友的共鸣。 高性能计算会在接下来的3-5年之内有很大的市场，但是纵观国内的软件，不论是做渲染还是做流体计算，很大一部分能用到实际生产的软件都是国外的，做流体的软件，几乎都是用得国外的，很多做军工的单位用的计算软件都是国外引进的，只是感觉有些揪心。看到国际上的很多会议都能有很多国内的朋友参加，很多高性能计算的会议也有国内的研究机构参加，为什么就不能作出产品级别的软件啦？研究所，或者研究单位都可以作出一些申请国家研究经费的项目，但是这些项目过后，这些软件怎么样了？就不能真正产业化，让更多的国人受益吗？当然这里或许愤青一下，也对国内的很多行业还不能太深入的了解，也不能做过多的评论，毕竟每个行业都有自己的苦衷。 高性能计算不只是硬件方面的竞争，还有软件，软件的竞争不是代码的累积，还需要更深层次算法的创新。高性能计算行业虽然是一个老行业，但是对于眼前的新芯片架构的出现，新计算机体系结构的出现，高性能计算又是一个新兴产业。中国改革开发30年，发展到现在积累了不少经验，也吃过不少亏，希望在这次计算机革命的时候，能少走弯路，在新兴行业起来的时候不是盲目的跟从，不是盲目的硬件的比拼和软件的累加，更需要的是对高性能计算的人才的培训和高性能计算标准的指定和行业规范的指定。现在做事情需要的抱着双赢或者三赢，多赢的态度去做，恶性竞争损害的只是自己的名声和自己的利益。 对于开发人员来说，积累技术经验比盲目的做软件开发重要。一般的技术的爆发是需要5-10年的技术积累。就像这次的gpu的革命，是早在02年左右就有的架构设想，做事不能太浮躁，积累经验，坚持下来。（《道德经》57章：“以正治国，以奇用兵，以无事取天下。吾何以知其然哉？以此：
天下多忌讳，而民弥贫；人多利器，国家滋昏；人多技巧，奇物滋起；法令滋彰，盗贼多有。
故圣人云：“我无为，而民自化；我好静，而民自正；我无事，而民自富；我无欲，而民自朴。” ）以不变应万变。这里的不变，无为，更应该理解为对事物的更深层次的理解，参悟事物发展的规律，找到其中不变的定律，变就是不变，不变就是变。积累自己的经验，扩展自己的眼界，稳重求精，精中求进，循序渐进。俞老师在《大家》节目里面的一句很经典的话：底蕴的厚度决定生命的高度。这里的底蕴是需要生活的磨炼和用心去感悟，去学习。最近看近现代文学的书，读到王蒙的《组织部来了个年轻人》，我不知道大家的感受是什么，我的感受是，看看作者写那篇文章的时候是1956年，那个时候对社会的描写和现在又有多大的区别啦？现在各种样式的人物在那个时候照样有，不是现在的社会才有的那样的人物。社会在变，有其不变的规律，我这样80后的人或许读了几天所谓的现代科学，就以为自己是最科学的，最现代的，就对社会抱不平？现在的浮躁的社会风气造就的只是愤青。作为年轻的一代，我们已经不是很年轻，不是能用年轻的理由来为自己的将来找一个理由。中国改革开发30年，还年轻吗？30而立！ 这次金融危机或许是对国人的一次机遇，也是一次考验，在这样的情况下，国家指定的发展目标一样是8%的增长速度！这个目标是外国想都不敢想的，当然在国际社会经济实力决定说话权的时候，我们有这个能力发好自己的言吗？各位工业界或者学术界的朋友，准备好了吗？有能力指定新的标准，新的规范没？ 寥寥几言，或许是对高性能计算的关注，也是对浪潮公司发布国产单机高性能计算服务器有感。 事主而不盡力，則有刑。事父母而不盡力，則不親。受業問學而不加務，則不成。故朝不勉力務進，夕無見功，故曰：「朝忘其事，夕失其功。」----《管子》 共勉！ 附上浪潮发布会上的沙画视频：http://v.youku.com/v_show/id_XNzU5NTYwNDg=.html 视频展现了服务器发展的过程，背景音乐很优美，从算盘到高性能计算机的发布，是一个一个的过程，其中的生物，航天，从小到大，可以看成是对未知领域的探索，也可以作为新服务器可以做的范围，可以为基因，为航天，提供一个高性能的廉价的开发平台！ February 16 【CUDA比赛】题目征集 准备举行全国的CUDA比赛，希望大家都踊跃参加； 现在是征集题目，有什么好的想法，有没解决的问题，都可以作为题目，看看大家都会用什么方法来解决； 有已经解决的问题吗？看看是否别人能用更优化的方法解决？ 欢迎投递任何题目； 可以和我联系： zhao.kaiyong(at)gmail.com February 14 9. CUDA shared memory使用------GPU的革命 序言：明年就毕业了，下半年就要为以后的生活做打算。这半年，或许就是一个抉择的时候，又是到了一个要做选择的时候。或许是自己的危机意识比较强，一直都觉得自己做得不够好，还需要积累和学习。或许是知足常乐吧，从小山沟，能到香港，一步一步，自己都比较满足，只是心中一直抱着一个理想，坚持做一件事情，坚持想做点事情，踏踏实实，曾经失败过，曾经迷茫过，才学会了坚持，学会了坚毅，才体会了淡定和从容。人生路上，一路走，一路看，一路学，抱着感恩的心，帮助别人，就是帮助自己，未来的路才会更宽…… 正文：书接上文《8. CUDA 内存使用 global 二------GPU的革命》 讲了global内存访问的时候，需要对齐的问题，只有在对齐的情况下才能保证global内存的高效访问。这一章节准备写一下shared memory的访问的问题，首先是讲一下shared的memory的两种使用方法，然后讲解一下shared memory的bank conflict的问题，这个是shared memory访问能否高效的问题所在； Shared memory的常规使用： 1. 使用固定大小的数组： /************************************************************************/ /* Example */ /************************************************************************/ __global__ void shared_memory_1(float* result, int num, float* table_1) { __shared__ float sh_data[THREAD_SIZE]; int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx %BANK_CONFLICT]; } result[idx] = ret; } 这里的sh_data就是固定大小的数组； 2. 使用动态分配的数组： extern __shared__ char array[]; __global__ void shared_memory_1(float* result, int num, float* table_1, int shared_size) { float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址，就可以动态分配空间 float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小； int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx %BANK_CONFLICT]; } result[idx] = ret; } 这里是动态分配的空间，extern __shared__ char array[];指定了shared的第一个变量的地址，这里其实是指向shared memory空间地址；后面的动态分配float* sh_data = (float*)array;让sh_data指向array其实就是指向shared memory上的第一个地址； 后面的float* sh_data2 = (float*)&sh_data[shared_size];这里的sh_data2是指向的第一个sh_data的shared_size的地址，就是sh_data就是有了shared_size的动态分配的空间； 入下图： 
3. 下面是讲解bank conflict 我们知道有每一个half-warp是16个thread，然后shared memory有16个bank，怎么分配这16个thread，分别到各自的bank去取shared memory，如果大家都到同一个bank取款，就会排队，这就造成了bank conflict，上面的代码可以用来验证一下bank conflict对代码性能造成的影响： /************************************************************************/ /* Example */ /************************************************************************/ __global__ void shared_memory_1(float* result, int num, float* table_1) { __shared__ float sh_data[THREAD_SIZE]; int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx %BANK_CONFLICT]; } result[idx] = ret; } // 1,2,3,4,5,6,7.....16 #define BANK_CONFLICT 16 这里的BANK_CONFLICT 定义为从1到16的大小，可以自己修改，来看看bank conflict对性能的影响；当BANK_CONFLICT为2的时候，就会通用有8个thread同时访问同一个bank，因为idx%2的取值只有2个0和1，所以16个都会访问bank0和bank1，以此类推，就可以测试整个的性能； 下面为示意图： 
当然我们还可以利用16bank conflict，大家都访问同一个bank的同一个数据的时候，就可以形成一个broadcast，那样就会把数据同时广播给16个thread，这样就可以合理利用shared memory的broadcast的机会。 下面贴出代码，最好自己测试一下； /******************************************************************** * shared_memory_test.cu * This is a example of the CUDA program. * Author: zhao.kaiyong(at)gmail.com * http://blog.csdn.net/openhero * http://www.comp.hkbu.edu.hk/~kyzhao/ *********************************************************************/ #include <stdio.h> #include <stdlib.h> #include <cutil.h> #include <cutil_inline.h> // 1,2,3,4,5,6,7.....16 #define BANK_CONFLICT 16 #define THREAD_SIZE 16 /************************************************************************/ /* static */ /************************************************************************/ __global__ void shared_memory_static(float* result, int num, float* table_1) { __shared__ float sh_data[THREAD_SIZE]; int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx%BANK_CONFLICT]; } result[idx] = ret; } /************************************************************************/ /* dynamic */ /************************************************************************/ extern __shared__ char array[]; __global__ void shared_memory_dynamic(float* result, int num, float* table_1, int shared_size) { float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址，就可以动态分配空间 float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小； int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx%BANK_CONFLICT]; } result[idx] = ret; } /************************************************************************/ /* Bank conflict */ /************************************************************************/ __global__ void shared_memory_bankconflict(float* result, int num, float* table_1) { __shared__ float sh_data[THREAD_SIZE]; int idx = threadIdx.x; float ret = 0.0f; sh_data[idx] = table_1[idx]; for (int i = 0; i < num; i++) { ret += sh_data[idx % BANK_CONFLICT]; } result[idx] = ret; } /************************************************************************/ /* HelloCUDA */ /************************************************************************/ int main(int argc, char* argv[]) { if ( cutCheckCmdLineFlag(argc, (const char**) argv, "device")) { cutilDeviceInit(argc, argv); }else { int id = cutGetMaxGflopsDeviceId(); cudaSetDevice(id); } float *device_result = NULL; float host_result[THREAD_SIZE] ={0}; CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(float) * THREAD_SIZE)); float *device_table_1 = NULL; float host_table1[THREAD_SIZE] = {0}; for (int i = 0; i < THREAD_SIZE; i++ ) { host_table1[i] = rand()%RAND_MAX; } CUDA_SAFE_CALL( cudaMalloc((void**) &device_table_1, sizeof(float) * THREAD_SIZE)); CUDA_SAFE_CALL( cudaMemcpy(device_table_1, host_table1, sizeof(float) * THREAD_SIZE, cudaMemcpyHostToDevice)); unsigned int timer = 0; CUT_SAFE_CALL( cutCreateTimer( &timer)); CUT_SAFE_CALL( cutStartTimer( timer)); shared_memory_static<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1); //shared_memory_dynamic<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1, 16); //shared_memory_bankconflict<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1); CUT_CHECK_ERROR("Kernel execution failed\n"); CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost)); CUT_SAFE_CALL( cutStopTimer( timer)); printf("Processing time: %f (ms)\n", cutGetTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); for (int i = 0; i < THREAD_SIZE; i++) { printf("%f ", host_result[i]); } CUDA_SAFE_CALL( cudaFree(device_result)); CUDA_SAFE_CALL( cudaFree(device_table_1)); cutilExit(argc, argv); } 这里只是一个简单的demo，大家可以测试一下。下一章节会将一些shared memory的更多的特性，更深入的讲解shared memory的一些隐藏的性质； 再在接下来的章节会讲一些constant和texture的使用； 写的内容一直都是文字比较多，代码比较少，其实学习的过程更重要的思想，实践的代码，最好是自己写，唯一可以学习的是思想，学习更重要的也是思想的交流，知识的传播，最好的是思想的传播，代码，方法，都是只是一些工具而已。但是工具的熟练层度，就得靠自己下来多练习。 February 13 - 在论坛里面讨论到一个问题，__global__函数里面传递的参数，到底是怎么传输到每一个thread的，然后做了以下的一些分析；
这个是问题讨论帖子：http://topic.csdn.net/u/20090210/22/2d9ac353-9606-4fa3-9dee-9d41d7fb2b40.html - C/C++ code
__global__ static void HelloCUDA(char* result, int num)
{
__shared__ int i;
i = 0;
char p_HelloCUDA[] = "Hello CUDA!";
for(i = 0; i < num; i++) {
result[i] = p_HelloCUDA[i];
}
}
-
- PTX code
.const .align 1 .b8 __constant432[12] = {0x48,0x65,0x6c,0x6c,0x6f,0x20,0x43,0x55,0x44,0x41,0x21,0x0};
.entry _Z9HelloCUDAPci
{
.reg .u16 %rh<3>;
.reg .u32 %r<16>;
.reg .pred %p<4>;
.param .u32 __cudaparm__Z9HelloCUDAPci_result;
.param .s32 __cudaparm__Z9HelloCUDAPci_num;
.local .align 4 .b8 __cuda___cuda_p_HelloCUDA_168[12];
.shared .s32 i;
.loc 14 15 0
$LBB1__Z9HelloCUDAPci:
mov.u32 %r1, __constant432; //
mov.u32 %r2, __cuda___cuda_p_HelloCUDA_168; //
ld.const.u32 %r3, [%r1+0]; // id:17 not_variable+0x0
st.local.u32 [%r2+0], %r3; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
ld.const.u32 %r4, [%r1+4]; // id:17 not_variable+0x0
st.local.u32 [%r2+4], %r4; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
ld.const.u32 %r5, [%r1+8]; // id:17 not_variable+0x0
st.local.u32 [%r2+8], %r5; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
.loc 14 20 0
mov.s32 %r6, 0; //
ld.param.s32 %r7, [__cudaparm__Z9HelloCUDAPci_num]; // id:16 __cudaparm__Z9HelloCUDAPci_num+0x0
mov.u32 %r8, 0; //
setp.le.s32 %p1, %r7, %r8; //
@%p1 bra $Lt_0_9; //
mov.s32 %r9, %r7; //
mov.u32 %r10, __cuda___cuda_p_HelloCUDA_168; //
mov.u32 %r11, __cuda___cuda_p_HelloCUDA_168; //
add.u32 %r12, %r7, %r11; //
ld.param.u32 %r13, [__cudaparm__Z9HelloCUDAPci_result]; // id:19 __cudaparm__Z9HelloCUDAPci_result+0x0
mov.s32 %r14, %r9; //
$Lt_0_7:
//<loop> Loop body line 20, nesting depth: 1, estimated iterations: unknown
.loc 14 21 0
ld.local.s8 %rh1, [%r10+0]; // id:20 __cuda___cuda_p_HelloCUDA_168+0x0
st.global.s8 [%r13+0], %rh1; // id:21
add.u32 %r13, %r13, 1; //
add.u32 %r10, %r10, 1; //
setp.ne.s32 %p2, %r10, %r12; //
@%p2 bra $Lt_0_7; //
st.shared.s32 [i], %r7; // id:22 i+0x0
bra.uni $Lt_0_5; //
$Lt_0_9:
st.shared.s32 [i], %r6; // id:22 i+0x0
$Lt_0_5:
.loc 14 23 0
exit; //
$LDWend__Z9HelloCUDAPci:
} // _Z9HelloCUDAPci
-
- Cubin code
architecture {sm_10}
abiversion {1}
modname {cubin}
consts {
name = __constant432
segname = const
segnum = 0
offset = 0
bytes = 12
mem {
0x6c6c6548 0x5543206f 0x00214144
}
}
code {
name = _Z9HelloCUDAPci
lmem = 12
smem = 28 // 我们注意这里的smem 的数量
reg = 3
bar = 0
bincode {
0x10000001 0x2400c780 0xd0000001 0x60c00780
0x10000201 0x2400c780 0xd0000801 0x60c00780
0x10000401 0x2400c780 0x307ccbfd 0x6c20c7c8
0xd0001001 0x60c00780 0x10014003 0x00000280
0x1000f801 0x0403c780 0x1000c805 0x0423c780
0x00000005 0xc0000780 0xd4000009 0x40200780
0x20018001 0x00000003 0xd00e0209 0xa0200780
0x3000cbfd 0x6c2147c8 0x20018205 0x00000003
0x1000a003 0x00000280 0x1000ca01 0x0423c780
0x00000c01 0xe4200780 0x30000003 0x00000780
0x00000c01 0xe43f0781
}
}
这个是一段加了shared memory的也有constant 的ptx代码~还有cubin
从cubin来看，确实有可能是通过global memory进来的，不过一定是分发到各自的.param变量里面去的，因为在每一个thread 里面都是可以修改传进来的参数的；
从这点来看：大体应该是 参数( global memory--> (constant memory or shared memory) 然后broadcast to each thread--> 每一个thread都各自的register里面有了参数
大体应该是这样一个过程； February 05 1 ．天行健，君子以自强不息。 —— 《周易》
译：作为君子，应该有坚强的意志，永不止息的奋斗精神，努力加强自我修养，完成并发展自己的学业或事业，能这样做才体现了天的意志，不辜负宇宙给予君子的职责和才能。
2 ．勿以恶小而为之，勿以善小而不为。 —— 《三国志》
译：对任何一件事，不要因为它是很小的、不显眼的坏事就去做；相反，对于一些微小的。却有益于别人的好事，不要因为它意义不大就不去做它。
3 ．见善如不及，见不善如探汤。 —— 《论语》
译：见到好的人，生怕来不及向他学习，见到好的事，生怕迟了就做不了。看到了恶人、坏事，就像是接触到热得发烫的水一样，要立刻离开，避得远远的。
4 ．躬自厚而薄责于人，则远怨矣。 —— 《论语》
译：干活抢重的，有过失主动承担主要责任是“ 躬自厚” ，对别人多谅解多宽容，是“ 薄责于人” ，这样的话，就不会互相怨恨。
5 ．君子成人之美，不成人之恶。小人反是。 —— 《论语》
译：君子总是从善良的或有利于他人的愿望出发，全心全意促使别人实现良好的意愿和正当的要求，不会用冷酷的眼光看世界。或是唯恐天下不乱，不会在别人有失败、错误或痛苦时推波助澜。小人却相反，总是“ 成人之恶，不成人之美” 。
6 ．见贤思齐焉，见不贤而内自省也。 —— 《论语》
译：见到有人在某一方面有超过自己的长处和优点，就虚心请教，认真学习，想办法赶上他，和他达到同一水准；见有人存在某种缺点或不足，就要冷静反省，看自己是不是也有他那样的缺点或不足。
7 ．己所不欲，勿施于人。 —— 《论语》
译：自己不想要的（痛苦、灾难、祸事…… ），就不要把它强加到别人身上去。
8 ．当仁，不让于师。 —— 《论语》
译：遇到应该做的好事，不能犹豫不决，即使老师在一旁，也应该抢著去做。后发展为成语“ 当仁不让” 。
9 ．君子欲讷于言而敏于行。 —— 《论语》
译：君子不会夸夸其谈，做起事来却敏捷灵巧。
10 ．二人同心，其利断金；同心之言，其臭如兰。 —— 《周易》
译：同心协办的人，他们的力量足以把坚硬的金属弄断；同心同德的人发表一致的意见，说服力强，人们就像嗅到芬芳的兰花香味，容易接受。
11 ．君子藏器于身，待时而动。 —— 《周易》
译：君子就算有卓越的才能超群的技艺，也不会到处炫耀、卖弄。而是在必要的时刻把才能或技艺施展出来。
12 ．满招损，谦受益。 —— 《尚书》
译：自满于已获得的成绩，将会招来损失和灾害；谦逊并时时感到了自己的不足，就能因此而得益。
13 ．人不知而不愠，不亦君子乎？ —— 《论语》
译：如果我有了某些成就，别人并不理解，可我决不会感到气愤、委屈。这不也是一种君子风度的表现吗？
14 ．言必信 ，行必果。 —— 《论语》
译：说了的话，一定要守信用；确定了要干的事，就一定要坚决果敢地干下去。
15 ．毋意，毋必，毋固，毋我。 —— 《论语》
译：讲事实，不凭空猜测；遇事不专断，不任性，可行则行；行事要灵活，不死板；凡事不以“ 我” 为中心，不自以为是，与周围的人群策群力，共同完成任务。
16 ．三有我师焉，择其善者而从之，其不善者而改之。—— 《论语》
译：三个人在一起，其中必有某人在某方面是值得我学习的，那他就可当我的老师。我选取他的优点来学习。
17 ．君子求诸己，小人求诸人。 —— 《论语》
译：君子总是责备自己，从自身找缺点，找问题。小人常常把目光射向别人，找别人的缺点和不足。
18 ．君子坦荡荡，小人长戚戚。 —— 《论语》
译：君子心胸开朗，思想上坦率洁净，外貌动作也显得十分舒畅安定。小人心里欲念太多，心理负担很重，就常忧虑、担心，外貌、动作也显得忐忑不安，常是坐不定，站不稳的样子。
19 ．不怨天，不尤人。 —— 《论语》
译：遇到挫折与失败，绝不从客观上去找借口，绝不把责任推向别人，后来发展为成语“ 怨天尤人” 。
20 ．不迁怒，不贰过。 —— 《论语》
译：犯了错误，不要迁怒别人，并且不要再犯第二次。
21 ．小不忍，则乱大谋。 —— 《论语》
译：不该干的事，即使很想去干，但坚持不干，叫“ 忍” 。对小事不忍，没忍性，就会影响大局，坏了大事。
22 ．小人之过也必文。 —— 《论语》
译：小人对自己的过错必定加以掩饰。
23 ．过而不改，是谓过矣。 —— 《论语》
译：有了过错而不改正，这就是真的过错了。
24 ．君子务本，本立而道生。 —— 《论语》
译：君子致力于根本，确立了根本，“ 道” 也就自然产生。
25 ．君子耻其言而过其行。 —— 《论语》
译：君子认为说得多做得少是可耻的。
26 ．三级笮小? —— 《论语》
译：每做一件事情必须要经过反复的考虑后才去做。
27 ．多行不义必自毙。 —— 《左传》
译：坏事做得太多，终将自取灭亡。
28 ．人谁无过，过而能改，善莫大焉。 —— 《左传》
译：人都有可能犯错误，犯了猎误，只要改正了仍是最好的人。
29 ．不以一眚掩大德。 —— 《左传》
译：评价一个人时，不能因为一点过失就抹杀他的功劳。
30 ．人一能之，己百之；人十能之，己千之。 —— 《中庸》
译：人家一次就学通的，我如果花上百次的功夫，一定能学通。人家十次能掌握的，我要是学一千次，也肯定会掌握的。
31 ．知耻近乎勇。 —— 《中庸》
译：知道什么是可耻的行为，那就是勇敢的好表现。
32 ．以五十步笑百步。 —— 《孟子》
译：以为自己的错误比别人的小，缺点比别人少而沾沾自喜。
33 ．君子莫大乎与人为善。 —— 《孟子》
译：君子最大的长处就是用高尚、仁义的心去对待别人。
34 ．人皆可以为尧舜。 —— 《孟子》
译：只要肯努力去做，人人都可以成为尧舜那样的大圣人。
35 ．千丈之堤，以蝼蚁之穴溃；百尺之室，以突隙之烟焚。 —— 《韩非子》
译：千里大堤，因为有蝼蚁在打洞，可能会因此而塌掉决堤；百尺高楼，可能因为烟囱的缝隙冒出火星引起火灾而焚毁。
36 ．言之者无罪，闻之者足以戒。 —— 《诗序》
译：提出批评意见的人，是没有罪过的。听到别人的批评意见要仔细反省自己，有错就改正，无错就当作是别人给自己的劝告。
37 ．良药苦于口而利于病，忠言逆于耳而利于行。 —— 《孔子家语》
译：好的药物味苦但对治病有利；忠言劝诫的话听起来不顺耳却对人的行为有利。
38 ．良言一句三冬暖，恶语伤人六月寒。 —— 明代谚语
译：一句良善有益的话，能让听者即使在三冬严寒中也倍感温暖；相反，尖酸刻薄的恶毒语言，伤害别人的感情和自尊心，即使在六月大暑天，也会让人觉得寒冷。
39 ．千经万典，孝悌为先。 —— 《增广贤文》
译：千万种经典讲的道理，孝顺父母，友爱兄弟是最应该先做到的。
40 ．善恶随人作，祸福自己招， —— 《增广贤文》
译：好事坏事都是自己做的，灾祸 幸福 也全是由自己的言行招来的。
41. 学而不思罔，思而不学则殆。 —— 《论语》
译：只学习却不思考就不会感到迷茫，只空想却不学习就会疲倦而没有收获。
42. 知之为知之，不知为不知，是知也。 —— 《论语》
译：知道就是知道，不知道应当说 不知道，不弄虚作假，这才是明智的行为。
43. 业精于勤，荒于嬉；行成于思，毁于随。 —— 韩愈
译：事业或学业的 成功 在于奋发努力，勤勉进取。太贪玩，放松要求便会一事无成；做人行事，必须谨慎思考，考虑周详才会有所成就。任性、马虎、随便只会导致失败。
44. 读书有三到：谓心到，眼到，口到。 —— 朱熹
译：用心思考，用眼仔细看，有口多读，三方面都做得到位才是真正的读书。
45. 学而不厌，诲人不倦。 —— 《论语》
译：努力学习却不感到满足，教导别人不感到厌倦。
46. 不积跬步，无以至千里，不积小流，无以成江海。 —— 《荀子》
译：不把半步、一步积累起来，就不能走到千里远的地方，不把细流汇聚起来，就不能形成江河大海。
47. 欲穷千里目，更上一层楼。 —— 王之涣
译：想看到更远更广阔的景物，你就要再上一层楼。想学到更多更深的知识，你就要比原来更努力。
48. 强中自有强中手，莫向人前满自夸。 —— 《警世通言》
译：尽管你是一个强者，可是一定还有比你更强的人，所以不要在别人面前骄傲自满，自己夸耀自己。
49. 玉不琢，不成器；人不学，不知道。 —— 《礼记. 学记》
译：玉石不经过雕琢，不能成为有用的玉器；人不经过学习，就不懂得事理。
50. 黑发不知勤学早，白首方悔读书迟。 —— 《劝学》
译：年轻的时候不知道抓紧 时间 勤奋学习，到老了想读书却为时已晚。
51. 知不足者好学，耻下问者自满。 —— 林逋《省心录》
译：知道自己的不足并努力学习就是聪明的人，不好问又骄傲自满的人是可耻的。
52. 学不可以已。 —— 《荀子》
译：学习是不可以停止的。
53. 学而时习之，不亦悦乎？ —— 《论语》
译：学过的知识，在适当的时候去复习它，使自己对知识又有了新的认识，这不是令人感到 快乐 的事吗？
54. 温故而知新，可以为师矣。 —— 《论语》
译：学了新的知识又常常温习已学过的知识，不断地学习，温习，学问和修养一定会很快得到提高，这样的人就可以成为老师了。
55. 读书破万卷，下笔如有神。 —— 杜甫
译：读书读得多，写起文章来就会笔下生花，像有神助一样。
56. 少壮不努力，老大徒伤悲。 —— 《汉乐府. 长歌行》
译：年轻时不努力学习，年老了只能后悔、歎息。
57. 读书百遍而义自见。 —— 《三国志》
译：读书必须反复多次地读，这样才能明白书中所讲的意思。
58. 学而不化，非学也。 —— 杨万里
译：学习知识但不能灵活运用，不能称为学习。
59. 好学而不贰。 —— 《左传》
译：爱好学习但不三心二意。
60. 学如不及，犹恐失之。
译：学习知识时生怕追不上，追上了又害怕再失去。 61. 人而不学，其犹正墙面而立。 —— 《尚书》
译：人如果不学习，就像面对墙壁站著，什么东西也看不见。
62 、知而好问，然后能才。 —— 《荀子》
译：聪明的人还一定得勤学好问才能成才。
63 、学之广在于不倦，不倦在于固志。 —— 葛洪
译：学问的渊博在于学习时不知道厌倦，而学习不知厌倦在于有坚定的目标。
64. 学而不知道，与不学同；知而不能行，与不知同。 —— 黄睎
译：学习知识不能从中明白一些道理，这和不学习没什么区别；学到了道理却不能运用，这仍等于没有学到道理。
65. 博观而约取，厚积而薄发。 —— 苏轼
译：广泛阅读，多了解古今中外的人和事，把其中好的部分牢牢记住；积累了大量的知识材料，到需要用时便可以很自如恰当地选择运用。
66. 差之毫厘，缪以千里。 —— 陆九渊
译：做任何事情，开始一定要认真地做好，如果做差了一丝一毫，结果会发现相差很远。
67. 盛年不重来，一日难再晨。 —— 陶渊明
译：美好的青春年华过去了就不会再来，一天不可能有两个早晨，要珍惜时光啊！）
68 、言之无文，行而不远。 —— 《左传》
译：文章讲的内容或题材都很好，可表达不适当，那么欣赏的人就不会很多，难以传播千古。
69. 人之为学，不可自小，又不可自大。 —— 顾炎武
译：学习时不要在渊博浩翰的知识面前感到自卑，也不能因为学到一点点知识而骄傲自满。
70. 好学近乎知，力行近乎仁，知耻近乎勇。 —— 《中庸》
译：勤奋好学就接近智，做任何事情只要努力就接近仁，懂得了是非善恶就是勇的一种表现。
71. 书到用时方恨少，事非经过不知难。 —— 陈廷焯
译：知识总是在运用时才让人感到太不够了，许多事情如果不亲身经历过就不知道它有多难。
72 、笨鸟先飞早入林，笨人勤学早成材。 —— 《省世格言》
译：飞得慢的鸟儿提早起飞就会比别的鸟儿早飞入树林，不够聪明的人只要勤奋努力，就可以比别人早成材。
73. 书山有路勤为径，学海无涯苦作舟。 —— 《增广贤文》
译：勤奋是登上知识高峰的一条捷径，不怕吃苦才能在知识的海洋里自由遨游。
74. 学如逆水行舟，不进则退。 —— 《增广贤文》
译：学习要不断进取，不断努力，就像逆水行驶的小船，不努力向前，就只能向后退。
75. 吾生也有涯而，知也无涯。 —— 《庄子》
译：我的生命是有限的，而人类的知识是无限的。
76. 天下兴亡，匹夫有责。 —— 顾炎武
译：国家的兴旺、衰败，每一个人都负有很大的责任。
77. 生于忧患，死于安乐。 —— 孟子
译：逆境能使人的意志得到磨炼，使人更坚强。相反，时常满足于享受，会使人不求上进而逐渐落后。
78. 位卑未敢忘国。 —— 陆游《病起书怀》
译：虽然自己地位低微，但是从没忘掉忧国忧民的责任。
79. 人生自古谁无死，留取丹心照汉青。 —— 文天祥《过零丁洋》
译：自古以来，谁都难免会死的，那就把一片爱国的赤胆忠心留在史册上吧！
80. 先天下之忧而忧，后天下乐而乐。 —— 范仲淹《岳阳楼记》
译：为国家分忧时，比别人先，比别人急；享受幸福，快乐时，却让别人先，自己居后。
81. 小来思报国，不是爱封侯。 —— 唐. 岑参《关人赴安西》
译：从小就想著报效祖国，而不是想著要封侯当官。
82. 有益国家之事虽死弗避。 —— 明. 吕坤《呻吟语. 卷上》
译：对国家有利的事情要勇敢地去做，就算有死亡的危险也不躲避。
83. 一寸山河一寸金。 —— 金. 左企弓语
译：祖国的每一寸山河比一寸黄金还要宝贵，是绝不能让给外人的。
84. 欲安其家，必先安于国。 —— 武则天
译：如果想建立个人幸福的小家，必须先让国定安定，繁荣起来。
85. 捐躯赴国难，视死忽如归。 —— 三国. 曹植《白马篇》
译：在国家有危难的时候要敢于挺身而出，把死当作回家一样。
86. 风声、雨声、读书声，声声入耳；家事、国事、天下事，事事关心。 —— 明. 顾宪成
译：风声、雨声、琅琅读书声，都进入我们的耳朵，所以，作为一个读书人，家事、国事，天下的事情，各种事情都应该关心，不能只是死读书。
87. 生当作人杰，死亦为鬼雄。 —— 李清照《夏日绝句》
译：活著的时候要做英雄，死后也要当英雄。
88. 利于国者爱之，害于国者恶之。 —— 《晏子春秋》
译：对于国家有利的事就要热心地去做，对国家有害的事就要憎恶它，远离它。
89. 读书本意在元元。 —— 宋. 陆）
译：读书的目的应该是掌握了知识后为社会和大众服务，而不是为了自己的升官发财。
90. 时穷节乃现，一一垂丹青。 —— 宋. 文天祥
译：历史上许多忠臣义士，在国家有难时，他们的节操就显现出来，一个个名垂史册。
91. 哀哀父母，生我劬劳。 —— 《诗经》
译：想起父母，做子女的是多么为他们感到心痛啊！他们生我育我，花费了多少辛勤的劳动啊！
92. 报国之心，死而后已。—— 宋. 苏轼
译：报效祖国的志向到死都不会变。
93. 忧国忘家，捐躯济难，忠臣之志也。 —— 三国. 曹植《求自诚表》
译：忧虑国家大事忘记小家庭，为拯救国家危难而捐躯献身，这都是忠臣的志向。
94. 大丈夫处世，当扫除天下，安事一室乎？ —— 汉. 陈蕃语
译：有志气的人活在世上，应当敢于跟各种不利于国家的行为作斗争，哪能只满足于处理好自己小家的小事呢？
95. 君子之交淡如水，小人之交甘若醴。 —— 《庄子》
译：君子之间的交往，像水一样的平淡、纯净，这样的友谊才会持久；往小人之间的交像甜酒一样的又浓又稠，但不会长久。
96. 老吾老，以及人之老；幼吾幼，以及人之幼。 —— 《孟子》
译：尊敬、爱戴别人的长辈，要像尊敬、爱戴自己长辈一样；爱护别人的儿女，也要像爱护自己的儿女一样。
97. 见侮而不斗，辱也。 —— 《公孙龙子》
译：当正义遭到侮辱、欺凌却不挺身而出，是一种耻辱的表现。
98. 天下皆知取之为取，而莫知与之为取。 —— 《后汉书》
译：人们都认为只有获取别人的东西才是收获，却不知道给予别人也是一种收获。
99. 人固有一死，死或重于泰山，或轻于鸿毛。 —— 汉. 司马迁
译：人终究免不了一死，但死的价值不同，为了人民正义的事业而死就比泰山还重，而那些自私自利，损人利已的人之死就比鸿毛还轻。
100. 羊有跪乳之恩，鸦有反哺之义。 —— 《增广贤文》
译：羊羔有跪下接受母乳的感恩举动，小乌鸦有衔食喂母鸦的情义，做子女的更要懂得孝顺父母。 January 13 2.1的release版本在前几天发布了，比2.1beta稳定一些，减少了一两个bug，没有太多更新，guide上有一个前景图，这个倒是有点感觉； 
下面是从网站上转载过来的，安装之前都最好自己好好读一遍；最近在回答很多人的问题，发现很多学习的时候都太急于求成，很多基础的东西都不太弄清楚的时候，就像很快出成绩，我还是对这些朋友善意的提醒，好好把基础弄扎实，磨刀不误砍柴工，没有谁可以一步登天，学习是需要积累的过程，慢慢积累这个很重要； 还有一些问题是粗心造成的，如果是这样，希望大家都能仔细的先自己分析一下问题； 遇到问题的时候，首先需要分析问题，定位问题，才能找到问题的根源，才能找到问题的解决方法，不能只在哪里看错误，为什么错了，一直不去找，优秀的测试人员比开发人员更重要。 我们需要debug，而不是需要bug maker。占用别人时间的时候，请别人回答问题的时候，首先得考虑一下别人的时间也是很宝贵的：）每个人的时间都差不多是平等的。 这里也没必要苦口婆心的说什么劝言。 CUDA的学习需要有C或者C++语言的基础，需要对并行算法有深入的了解，要想在这个基础上有更深入的研究，需要对device的架构，芯片的architecture有深入的了解，一点一点的积累； 我只是说一下我学习CUDA的过程，从CUDA开始，到现在，看的文章不下200篇，QQ群里面有的文章，几乎都看过，需要积累的过程，只有大量的投入，才会有产出。 要写高性能的程序，device + architecture + algorithm 这三者必须都要很深入的了解，device 是硬件的发展，architecture是芯片的架构，这三者都是相互促进的，对于编程语言本身，这个都是应该掌握的基本功，C/C++这样的都是基本功，我在学习C/C++的时候，用的方法也很简单，找了一本数据结构的书，把上面的每个数据结构都自己实现一遍。 大学毕业的时候，统计过大学写过的代码，大概有8w多行，离李开复先生提出来的10w行的代码还有很远的距离，我不是计算机专业的，想想，如果是计算机专业的人，应该可以达到10w行的。 看过少林的一个纪录片，少林释德建大师在三皇寨的生活，其中他讲到他对徒弟的教育的时候，说道十多年要求他徒弟练习的都是一些基本招式，都是一些基本功，没有教他们很多很花的招式。 编程也是一样的，基础太重要了…… http://forums.nvidia.com/index.php?showtopic=84440 NVIDIA CUDA FAQ version 2.1
This document is intended to answer frequently asked questions seen on the CUDA forums. General questions - What is NVIDIA CUDA?
NVIDIA® CUDA™ is a general purpose parallel computing architecture that leverages the parallel compute engine in NVIDIA graphics processing units (GPUs) to solve many complex computational problems in a fraction of the time required on a CPU.
With over 100 million CUDA-enabled GPUs sold to date, thousands of software developers are already using the free CUDA software development tools to solve problems in a variety of professional and home applications – from video and audio processing and physics simulations, to oil and gas exploration, product design, medical imaging, and scientific research.
CUDA allows developers to program applications in high level languages like C and will support industry standard APIs such as Microsoft® DirectX® and OpenCL, to seamlessly integrate into the development environments of today and tomorrow. - What's new in CUDA 2.1?
- TESLA devices are now supported on Windows Vista
- Support for using a GPU that is not driving a display on Vista (was already supported on Windows XP, OSX and Linux)
- VisualStudio 2008 support for Windows XP and Windows Vista
- Just-in-time (JIT) compilation, for applications that dynamically generate CUDA kernels
- New interoperability APIs for Direct3D 9 and Direct3D 10 accelerate communication with DirectX applications
- OpenGL interoperability improvements
- Debugger support using gdb for CUDA on RHEL 5.x (separate download)
- Linux Profiler 1.1 Beta (separate download)
- Support for recent releases of Linux including Fedora9, OpenSUSE 11 and Ubuntu 8.04
- Numerous bug fixes
- What is NVIDIA Tesla™?
With the world’s first teraflop many-core processor, NVIDIA® Tesla™ computing solutions enable the necessary transition to energy efficient parallel computing power. With 240 cores per processor and a standard C compiler that simplifies application development, Tesla scales to solve the world’s most important computing challenges—more quickly and accurately. - What kind of performance increase can I expect using GPU Computing over CPU-only code?
This depends on how well the problem maps onto the architecture. For data parallel applications, we have seen speedups anywhere from 10x to 200x.
See the CUDA Zone web page for many examples of applications accelerated using CUDA:
http://www.nvidia.com/cuda - What operating systems does CUDA support?
CUDA supports Windows XP, Windows Vista, Linux and Mac OS (including 32-bit and 64-bit versions).
CUDA 2.1 supports the following Linux distributions: - Red Hat Enterprise Linux 4.3
- Red Hat Enterprise Linux 4.4
- Red Hat Enterprise Linux 4.5
- Red Hat Enterprise Linux 4.6
- Red Hat Enterprise Linux 4.7*
- Red Hat Enterprise Linux 5.0
- Red Hat Enterprise Linux 5.1
- Red Hat Enterprise Linux 5.2*
- SUSE Linux Enterprise Desktop 10.0 SP1
- SUSE Linux 10.3
- SUSE Linux 11.0*
- Fedora 8
- Fedora 9*
- Ubuntu 7.10
- Ubuntu 8.04*
* new support with CUDA v2.1
Users have reported that CUDA works on other Linux distributions such as Gentoo, but these are not offically supported. - Can I run CUDA under DOS?
CUDA will not work in full-screen DOS mode since the display driver is not loaded. - What versions of Microsoft Visual Studio does CUDA support?
CUDA 2.1 supports Microsoft Visual Studio 2005 (v.8) and 2008 (v.9). Visual Studio 2003 (v.7) is no longer supported. - What GPUs does CUDA run on?
GPU Computing is a standard feature in all NVIDIA's 8-Series and later GPUs. We recommend that the GPU have at least 256MB of graphics memory. System configurations with less than the recommended memory size may not have enough memory to properly support all CUDA programs.
For a full list of supported products see:
http://www.nvidia.com/object/cuda_learn_products.html - What is the "compute capability"?
The compute capability indicates the version of the compute hardware included in the GPU.
Compute capability 1.0 corresponds to the original G80 architecture.
Compute capability 1.1 (introduced in later G8x parts) adds support for atomic operations on global memory. See "What are atomic operations?" in the programming section below.
Compute capability 1.2 (introduced in the GT200 architecture) adds the following new features: - Support for atomic functions operating in shared memory and atomic functions operating on 64-bit words in global memory
- Support for warp vote functions
- The number of registers per multiprocessor is 16384
- The maximum number of active warps per multiprocessor is 32
- The maximum number of active threads per multiprocessor is 1024
Compute capability 1.3 adds support for double precision floating point numbers.
See the CUDA Programming Guide for a full list of GPUs and their compute capabilities. - What are the technical specifications of the NVIDIA Tesla C1060 Processor ?
The Tesla C1060 consists of 30 multiprocessors, each of which is comprised of 8 scalar processor cores, for a total of 240 processors. There is 16KB of shared memory per multiprocessor.
Each processor has a floating point unit which is capable of performing a scalar multiply-add and a multiply, and a "superfunc" operation (such as rsqrt or sin/cos) per clock cycle.
The processors are clocked at 1.296 GHz. The peak computation rate accessible from CUDA is therefore around 933 GFLOPS (240 * 3 * 1.296). If you include the graphics functionality that is accessible from CUDA (such as texture interpolation), the FLOPs rate is much higher.
Each multiprocessor includes a single double precision multiply-add unit, so the double precision floating point performance is around 78 GFLOPS (30 * 2 * 1.296).
The card includes 4 GB of device memory. The maximum observed bandwidth between system and device memory is about 6GB/second with a PCI-Express 2.0 compatible motherboard.
Other GPUs in the Tesla series have the same basic architecture, but vary in the number of multiprocessors, clock speed, memory bus width and amount of memory.
See the programming guide for more details. - Does CUDA support multiple graphics cards in one system?
Yes. Applications can distribute work across multiple GPUs. This is not done automatically, however, so the application has complete control.
The Tesla S1070 Computing System supports 4 GPUs in an external enclosure:
http://www.nvidia.com/object/tesla_computing_solutions.html
There are also motherboards available that support up to 4 PCI-Express cards.
It is also possible to buy pre-built Tesla "Personal SuperComputer" systems off the shelf:
http://www.nvidia.com/object/personal_supercomputing.html
See the "multiGPU" example in the CUDA SDK for an example of programming multiple GPUs. - Where can I find a good introduction to parallel programming?
The course "ECE 498: Programming Massively Parallel Processors" at the University of Illinois, co-taught by Dr. David Kirk and Dr. Wen-mei Hwu is a good introduction:
http://courses.ece.uiuc.edu/ece498/al/Syllabus.html
The book "An Introduction to Parallel Computing: Design and Analysis of Algorithms" by Grama, Karypis et al isn't bad either (despite the reviews):
http://www.amazon.com/Introduction-Paralle...2044533-7984021
Although not directly applicable to CUDA, the book "Parallel Programming in C with MPI and OpenMP" by M.J.Quinn is also worth reading:
http://www.amazon.com/Parallel-Programming...3109&sr=1-1
For a full list of educational material see:
http://www.nvidia.com/object/cuda_education.html - Where can I find more information on NVIDIA GPU architecture?
- J. Nickolls et al. "Scalable Programming with CUDA" ACM Queue, vol. 6 no. 2 Mar./Apr. 2008 pp 40-53
http://www.acmqueue.org/modules.php?name=C...age&pid=532 - E. Lindholm et al. "NVIDIA Tesla: A Unified Graphics and Computing Architecture," IEEE Micro, vol. 28 no. 2, Mar.Apr. 2008, pp 39-55
- How does CUDA structure computation?
CUDA broadly follows the data-parallel model of computation. Typically each thread executes the same operation on different elements of the data in parallel.
The data is split up into a 1D or 2D grid of blocks. Each block can be 1D, 2D or 3D in shape, and can consist of up to 512 threads on current hardware. Threads within a thread block can coooperate via the shared memory.
Thread blocks are executed as smaller groups of threads known as "warps". The warp size is 32 threads on current hardware. - Can I run CUDA remotely?
Under Linux it is possible to run CUDA programs via remote login. We currently recommend running with an X-server.
CUDA does not work with Windows Remote Desktop, although it does work with VNC. - How do I pronounce CUDA?
koo-duh. Programming questions - I think I've found a bug in CUDA, how do I report it?
First of all, please verify that your code works correctly in emulation mode and is not caused by faulty hardware.
If you are reasonably sure it is a bug, you can either post a message on the forums or send a private message to one of the forum moderators.
If you are a researcher or commercial application developer, please sign up as a registered developer to get additional support and the ability to file your own bugs directly:
http://developer.nvidia.com/page/registere...er_program.html
Your bug report should include a simple, self-contained piece of code that demonstrates the bug, along with a description of the bug and the expected behaviour.
Please include the following information with your bug report: - Machine configuration (CPU, Motherboard, memory etc.)
- Operating system
- CUDA Toolkit version
- Display driver version
- For Linux users, please attach an nvidia-bug-report.log, which is generated by running "nvidia-bug-report.sh".
- What are the advantages of CUDA vs. graphics-based GPGPU?
CUDA is designed from the ground-up for efficient general purpose computation on GPUs. Developers can compile C for CUDA to avoid the tedious work of remapping their algorithms to graphics concepts.
CUDA exposes several hardware features that are not available via graphics APIs. The most significant of these is shared memory, which is a small (currently 16KB per multiprocessor) area of on-chip memory which can be accessed in parallel by blocks of threads. This allows caching of frequently used data and can provide large speedups over using textures to access data. Combined with a thread synchronization primitive, this allows cooperative parallel processing of on-chip data, greatly reducing the expensive off-chip bandwidth requirements of many parallel algorithms. This benefits a number of common applications such as linear algebra, Fast Fourier Transforms, and image processing filters.
Whereas fragment programs in graphics APIs are limited to outputting 32 floats (RGBA * 8 render targets) at a pre-specified location, CUDA supports scattered writes - i.e. an unlimited number of stores to any address. This enables many new algorithms that were not possible using graphics APIS to perform efficiently using CUDA.
Graphics APIs force developers to store data in textures, which requires packing long arrays into 2D textures. This is cumbersome and imposes extra addressing math. CUDA can perform loads from any address.
CUDA also offers highly optimized data transfers to and from the GPU. - Does CUDA support C++?
Full C++ is supported for host code. However, only the C subset of C++ is fully supported for device code. Some features of C++ (such as templates, operator overloading) work in CUDA device code, but are not officially supported. - Can the CPU and GPU run in parallel?
Kernel invocation in CUDA is asynchronous, so the driver will return control to the application as soon as it has launched the kernel.
The "cudaThreadSynchronize()" API call should be used when measuring performance to ensure that all device operations have completed before stopping the timer.
CUDA functions that perform memory copies and that control graphics interoperability are synchronous, and implicitly wait for all kernels to complete. - Can I transfer data and run a kernel in parallel (for streaming applications)?
Yes, CUDA 2.1 supports overlapping GPU computation and data transfers using streams. See the programming guide for more details. - Is it possible to DMA directly into GPU memory from another PCI-E device?
Not currently, but we are investigating ways to enable this. - Is it possible to write the results from a kernel directly to texture (for multi-pass algorithms)
Not currently, but you can copy from global memory back to the array (texture). Device to device memory copies are fast. - Can I write directly to the framebuffer?
No. In OpenGL you have to write to a mapped pixel buffer object (PBO), and then render from this. The copies are in video memory and fast, however. See the "postProcessGL" sample in the SDK for more details. - Can I read directly from textures created in OpenGL/Direct3D?
You cannot read directly from OpenGL textures in CUDA. You can copy the texture data to a pixel buffer object (PBO) and then map this buffer object for reading in CUDA.
In Direct3D it is possible to map D3D resources and read them in CUDA. This may involve an internal copy from the texture format to linear format. - Does graphics interoperability work with multiple GPUs?
Yes, although for best performance the GPU doing the rendering should be the same as the one doing the computation. - Does CUDA support graphics interoperability with Direct3D 10?
Yes, this is supported in CUDA 2.1. - How do I get the best performance when transferring data to and from OpenGL pixel buffer objects (PBOs)?
For optimal performance when copying data to and from PBOs, you should make sure that the format of the source data is compatible with the format of the destination. This will ensure that the driver doesn't have to do any format conversion on the CPU and can do a direct copy in video memory. When copying 8-bit color data from the framebuffer using glReadPixels we recommend using the GL_BGRA format and ensuring that the framebuffer has an alpha channel (e.g. glutInitDisplayMode(GLUT_RGBA_ | GLUT_ALPHA) if you're using GLUT). - What texture features does CUDA support?
CUDA supports 1D, 2D and 3D textures, which can be accessed with normalized (0..1) or integer coordinates. Textures can also be bound to linear memory and accessed with the "tex1Dfetch" function.
Cube maps, texture arrays, compressed textures and mip-maps are not currently supported.
The hardware only supports 1, 2 and 4-component textures, not 3-component textures.
The hardware supports linear interpolation of texture data, but you must use "cudaReadModeNormalizedFloat" as the "ReadMode". - What are the maximum texture sizes supported?
1D textures are limited to 8K elements. 1D "buffer" textures bound to linear memory are limited to 2^27 elements.
The maximum size for a 2D texture is 64K by 32K pixels, assuming sufficient device memory is available.
The maximum size of a 3D texture is 2048 x 2048 x 2048. - What is the size of the texture cache?
The texture cache has an effective 16KB working set size per multiprocessor and is optimized for 2D locality. Fetches from linear memory bound as texture also benefit from the texture cache. - Can I index into an array of textures?
No. Current hardware can't index into an array of texture references (samplers). If all your textures are the same size, an alternative might be to load them into the slices of a 3D texture. You can also use a switch statement to select between different texture lookups. - Are graphics operations such as z-buffering and alpha blending supported in CUDA?
No. Access to video memory in CUDA is done via the load/store mechanism, and doesn't go through the normal graphics raster operations like blending. We don't have any plans to expose blending or any other raster ops in CUDA. - What are the peak transfer rates between the CPU and GPU?
The performance of memory transfers depends on many factors, including the size of the transfer and type of system motherboard used.
We recommend NVIDIA nForce motherboards for best transfer performance. On PCI-Express 2.0 systems we have measured up to 6.0 GB/sec transfer rates.
You can measure the bandwidth on your system using the bandwidthTest sample from the SDK.
Transfers from page-locked memory are faster because the GPU can DMA directly from this memory. However allocating too much page-locked memory can significantly affect the overall performance of the system, so allocate it with care. - What is the precision of mathematical operations in CUDA?
All compute-capable NVIDIA GPUs support 32-bit integer and single precision floating point arithmetic. They follow the IEEE-754 standard for single-precision binary floating-point arithmetic, with some minor differences - notably that denormalized numbers are not supported.
Later GPUs (for example, the Tesla C1060) also. include double precision floating point. See the programming guide for more details. - Why are the results of my GPU computation slightly different from the CPU results?
There are many possible reasons. Floating point computations are not guaranteed to give identical results across any set of processor architectures. The order of operations will often be different when implementing algorithms in a data parallel way on the GPU.
This is a very good reference on floating point arithmetic:
What Every Computer Scientist Should Know About Floating-Point Arithmetic
The GPU also has several deviations from the IEEE-754 standard for binary floating point arithmetic. These are documented in the CUDA Programming Guide, section A.2. - Does CUDA support double precision arithmetic?
Yes. GPUs with compute capability 1.3 (those based on the GT200 architecture, such as the Tesla C1060 and later) support double precision floating point in hardware. - How do I get double precision floating point to work in my kernel?
You need to add the switch "-arch sm_13" to your nvcc command line, otherwise doubles will be silently demoted to floats. See the "Mandelbrot" sample in the CUDA SDK for an example of how to switch between different kernels based on the compute capability of the GPU.
You should also be careful to suffix all floating point literals with "f" (for example, "1.0f") otherwise they will be interpreted as doubles by the compiler. - Can I read double precision floats from texture?
The hardware doesn't support double precision float as a texture format, but it is possible to use int2 and cast it to double as long as you don't need interpolation:
texture<int2,1> my_texture;
static __inline__ __device__ double fetch_double(texture<int2, 1> t, int i)
{
int2 v = tex1Dfetch(t,i);
return __hiloint2double(v.y, v.x);
} - Does CUDA support long integers?
Current GPUs are essentially 32-bit machines, but they do support 64 bit integers (long longs). Operations on these types compile to multiple instruction sequences. - When should I use the __mul24 and __umul24 functions?
G8x hardware supports integer multiply with only 24-bit precision natively (add, subtract and logical operations are supported with 32 bit precision natively). 32-bit integer multiplies compile to multiple instruction sequences and take around 16 clock cycles.
You can use the __mul24 and __umul24 built-in functions to perform fast multiplies with 24-bit precision.
Be aware that future hardware may switch to 32-bit native integers, it which case __mul24 and __umul24 may actually be slower. For this reason we recommend using a macro so that the implementation can be switched easily. - Does CUDA support 16-bit (half) floats?
All floating point computation is performed with 32 or 64 bits.
The driver API supports textures that contain 16-bit floats through the CU_AD_FORMAT_HALF array format. The values are automatically promoted to 32-bit during the texture read.
16-bit float textures are planned for a future release of CUDART.
Other support for 16-bit floats, such as enabling kernels to convert between 16- and 32-bit floats (to read/write float16 while processing float32), is also planned for a future release. - Where can I find documentation on the PTX assembly language?
This is included in the CUDA Toolkit documentation ("ptx_isa_1.3.pdf"). - How can I see the PTX code generated by my program?
Add "-keep" to the nvcc command line (or custom build setup in Visual Studio) to keep the intermediate compilation files. Then look at the ".ptx" file. The ".cubin" file also includes useful information including the actual number of hardware registers used by the kernel. - How can I find out how many registers / how much shared/constant memory my kernel is using?
Add the option "--ptxas-options=-v" to the nvcc command line. When compiling, this information will be output to the console. - Is it possible to see PTX assembly interleaved with C code?
Yes! Add the option "--opencc-options -LIST:source=on" to the nvcc command line. - What is CUTIL?
CUTIL is a simple utility library designed used in the CUDA SDK samples.
It provides functions for: - parsing command line arguments
- read and writing binary files and PPM format images
- comparing arrays of data (typically used for comparing GPU results with CPU)
- timers
- macros for checking error codes
Note that CUTIL is not part of the CUDA Toolkit and is not supported by NVIDIA. It exists only for the convenience of writing concise and platform-independent example code. - Does CUDA support operations on vector types?
CUDA defines vector types such as float4, but doesn't include any operators on them by default. However, you can define your own operators using standard C++. The CUDA SDK includes a header "cutil_math.h" that defines some common operations on the vector types.
Note that since the GPU hardware uses a scalar architecture there is no inherent performance advantage to using vector types for calculation. - Does CUDA support swizzling?
CUDA does not support swizzling (e.g. "vector.wzyx", as used in the Cg/HLSL shading languages), but you can access the individual components of vector types. - Is it possible to run multiple CUDA applications and graphics applications at the same time?
CUDA is a client of the GPU in the same way as the OpenGL and Direct3D drivers are - it shares the GPU via time slicing. It is possible to run multiple graphics and CUDA applications at the same time, although currently CUDA only switches at the boundaries between kernel executions.
The cost of context switching between CUDA and graphics APIs is roughly the same as switching graphics contexts. This isn't something you'd want to do more than a few times each frame, but is certainly fast enough to make it practical for use in real time graphics applications like games. - Can CUDA survive a mode switch?
If the display resolution is increased while a CUDA application is running, the CUDA application is not guaranteed to survive the mode switch. The driver may have to reclaim some of the memory owned by CUDA for the display. - Is it possible to execute multiple kernels at the same time?
No, CUDA only executes a single kernel at a time on each GPU in the system. In some cases it is possible to have a single kernel perform multiple tasks by branching in the kernel based on the thread or block id. - What is the maximum length of a CUDA kernel?
The maximum kernel size is 2 million PTX instructions. - How can I debug my CUDA code?
Use device emulation for debugging (breakpoints, stepping, printfs), but make sure to keep checking that your program runs correctly on the device too as you add more code. This is to catch issues, before they become buried into too much code, that are hard or impossible to catch in device emulation mode. The two most frequent ones are: - Dereferencing device pointers in host code or host pointers in device code
- Missing __syncthreads().
See Section “Debugging using the Device Emulation Mode" from the programming guide for more details. - How can I optimize my CUDA code?
Here are some basic tips: - Make sure global memory reads and writes are coalesced where possible (see programming guide section 6.1.2.1).
- If your memory reads are hard to coalesce, try using 1D texture fetches (tex1Dfetch) instead.
- Make as much use of shared memory as possible (it is much faster than global memory).
- Avoid large-scale bank conflicts in shared memory.
- Use types like float4 to load 128 bits in a single load.
- Avoid divergent branches within a warp where possible.
We recommend using the CUDA Visual Profiler to profile your application. - How do I choose the optimal number of threads per block?
For maximum utilization of the GPU you should carefully balance the the number of threads per thread block, the amount of shared memory per block, and the number of registers used by the kernel.
You can use the CUDA occupancy calculator tool to compute the multiprocessor occupancy of a GPU by a given CUDA kernel:
http://www.nvidia.com/object/cuda_programming_tools.html - What is the maximum kernel execution time?
On Windows, individual GPU program launches have a maximum run time of around 5 seconds. Exceeding this time limit usually will cause a launch failure reported through the CUDA driver or the CUDA runtime, but in some cases can hang the entire machine, requiring a hard reset.
This is caused by the Windows "watchdog" timer that causes programs using the primary graphics adapter to time out if they run longer than the maximum allowed time.
For this reason it is recommended that CUDA is run on a GPU that is NOT attached to a display and does not have the Windows desktop extended onto it. In this case, the system must contain at least one NVIDIA GPU that serves as the primary graphics adapter. - Why do I get the error message: "The application failed to initialize properly"?
This problem is associated with improper permissions on the DLLs (shared libraries) that are linked with your CUDA executables. All DLLs must be executable. The most likely problem is that you unzipped the CUDA distribution with cygwin's "unzip", which sets all permissions to non-executable. Make sure all DLLs are set to executable, particularly those in the CUDA_BIN directory, by running the cygwin command "chmod +x *.dll" in the CUDA_BIN directory. Alternatively, right-click on each DLL in the CUDA_BIN directory, select Properties, then the Security tab, and make sure "read & execute" is set. For more information see:
http://www.cygwin.com/ml/cygwin/2002-12/msg00686.html - When compiling with Microsoft Visual Studio 2005 why do I get the error "C:\Program Files\Microsoft Visual Studio 8\VC\INCLUDE\xutility(870): internal error: assertion failed at: "D:/Bld/rel/gpgpu/toolkit/r2.0/compiler/edg/EDG_3.9/src/lower_name.c"
You need to install Visual Studio 2005 Service Pack 1:
http://www.microsoft.com/downloads/details...;displaylang=en - I get the CUDA error "invalid argument" when executing my kernel
You might be exceeding the maximum size of the arguments for the kernel. Parameters to __global__ functions are currently passed to the device via shared memory and the total size is limited to 256 bytes.
As a workaround you can pass arguments using constant memory. - What are atomic operations?
Atomic operations allow multiple threads to perform concurrent read-modify-write operations in memory without conflicts. The hardware serializes accesses to the same address so that the behaviour is always deterministic. The functions are atomic in the sense that they are guaranteed to be performed without interruption from other threads. Atomic operations must be associative (i.e. order independent).
Atomic operations are useful for sorting, reduction operations and building data structures in parallel.
Devices with compute capability 1.1 support atomic operations on 32-bit integers in global memory. This includes logical operations (and, or, xor), increment and decrement, min and max, exchange and compare and swap (CAS). To compile code using atomics you must add the option "-arch sm_11" to the nvcc command line.
Compute capability 1.2 devices also support atomic operations in shared memory.
Atomic operations values are not currently supported for floating point values.
There is no radiation risk from atomic operations. - Does CUDA support function pointers?
No, current hardware does not support function pointers (jumping to a computed address). You can use a switch to select between different functions. If you don't need to switch between functions at runtime, you can sometimes use C++ templates or macros to compile different versions of your kernels that can be switched between in the host code. - How do I compute the sum of an array of numbers on the GPU?
This is known as a parallel reduction operation. See the "reduction" sample in the CUDA SDK for more details. - How do I output a variable amount of data from each thread?
This can be achieved using a parallel prefix sum (also known as "scan") operation. The CUDA Data Parallel Primitives library (CUDPP) includes highly optimized scan functions:
http://www.gpgpu.org/developer/cudpp/
The "marchingCubes" sample in the CUDA SDK demonstrates the use of scan for variable output per thread. - How do I sort an array on the GPU?
The "particles" sample in the CUDA SDK includes a fast parallel radix sort.
To sort an array of values within a block, you can use a parallel bitonic sort. See the "bitonic" sample in the SDK.
The CUDPP library also includes sort functions. - What do I need to distribute my CUDA application?
Applications that use the driver API only need the CUDA driver library ("nvcuda.dll" under Windows), which is included as part of the standard NVIDIA driver install.
Applications that use the runtime API also require the runtime library ("cudart.dll" under Windows), which is included in the CUDA Toolkit. It is permissible to distribute this library with your application under the terms of the End User License Agreement included with the CUDA Toolkit. CUFFT - What is CUFFT?
CUFFT is a Fast Fourier Transform (FFT) library for CUDA. See the CUFFT documentation for more information. - What types of transforms does CUFFT support?
The current release supports complex to complex (C2C), real to complex (R2C) and complex to real (C2R). - What is the maximum transform size?
For 1D transforms, the maximum transform size is 16M elements in the 1.0 release. CUBLAS - What is CUBLAS?
CUBLAS is an implementation of BLAS (Basic Linear Algebra Subprograms) on top of the CUDA driver. It allows access to the computational resources of NVIDIA GPUs. The library is self contained at the API level, that is, no direct interaction with the CUDA driver is necessary.
See the documentation for more details. January 10 如果要做研究，数学是基础，算法也是基础 其实算法是一个思维的过程，数学是思维的工具 往往都是发现对新的事物有很敏锐的目光，但是往往发现实际做的时候，要想达到看得到的，得看垫在自己脚下的基础有多高…… 什么语言，只是工具而已，C c++？ 还是python 还是java 还是haskell……or…… just tool 如果要做研究，用算法的思维考虑分析问题，用数学的逻辑去组织问题，用数学的公式来呈现过程，编程语言语言只是实现验证的工具； 对新鲜事物的铭感度决定了创新的探索思维，对数字的敏锐，能决定是否能从新的事物里面总结出规律，才能真正创新，而不是落后在别人的背后 晓炜说得挺有道理，boss对数字的铭感，或许源于他ACM这样的系统的培训，boss也是清华科班计算机系出身，港科大计算机boss~对计算机本身的铭感度就很高，再加上对数字超敏锐… 要做研究，真的还是得好好的培养基础知识的，积累，积累，学习的过程 或许理科能让人在数学的敏锐，或许能做工程，但是要需要创新，还需要文科，太理科的思想，会让人的思想比较在某些方面比较禁锢…… 这里的文科应该更广的人文，更广的人文文化 就像武术，少林的禅宗讲究的武术与禅，达到禅武合一；武当的武术讲究的武术与道，达到天人合一…… 思考，实践，学习，思考，再实践…… 人生是旅程，是经历，是感受，是体验…… view in the sky， mind fly January 09 今天听了《the 7 habits of Highly effective people 》 里面讲了人生的一些成功的过程，成功的条件，怎么才能成功，养成什么样的习惯，才能取得成功 想想现在很多的 e文式的思想，什么方法，到处都在讲学，到处都在宣扬，还有n多的人拿高消费去学习，殊不知，自己老祖宗的东西，早就有这些精华了 记得师兄说过，华中科大的老校长，会让他的研究生都读 《道德经》太有道理了，人生也不过如此，一本道德经真实思想的根源，人生的悟道，成功，也只是人生的一个小小的过程而已。 《武经七书》里面的兵法战略，用到商场上，用到人生上，那一条不管用？ 《易经》可以说就是当时诸子百家的top 1 conference ranking是 最高的顶尖的会议…… 国学很博大精深，还有那么多人去批斗糟粕，那点时间去研究精华吧~~够用一辈子，下辈子，子孙都够用了……
|
Tools 
