(一)，安装程序

以前用vs制作过安装程序，现在把步骤写出来，有帮助的大家一定要顶哦

第一步:建立工程

1.打开vs，新建项目->其他项目类型->安装和部署(這個子项下面有安装项目和Web安装项目等,安装项目就是普通的桌面程序安装,Web安装就是安装网站,通常安装到IIS下,这里以普通桌面程序安装为例),新建安装项目，命名为SetupTest。

2.新建工程后，可以在“解决方案资源管理器”里面看到子项：文件系统编辑器，注册表编辑器，文件类型编辑器，用户界面编辑器，自定义操作编辑器，启动条件编辑器。下面将会举例来说明各个地方怎么使用。

3.最简单的一个安装程序：点击“文件类型编辑器”，找到“应用程序文件夹”，添加->文件，选择你要打包的程序文件（譬如我的程序有run.exe和lib.dll组成），确定。在项目“SetupTest”上点击生成，一个简单的安装包就做好了。生成后，点击安装，如果没有错误的话，你现在可以看到你的安装包的效果了。

4.如何自定义代码？在制作过程中，这些简单的功能不能满足你的要求，那么就自己写代码来完成吧。

解决方案->添加->新建项目->类库,命名为Library，在项目Library下面，删除Class1.cs，添加->新建项目->安装程序类（名称为InstallerTest.cs），打开InstallerTest.cs，你会发现这个class InstallerTest继承自Installer，里面就一个构造函数。好了，现在安装程序不能完成的功能，你可以在这个类写代码来自己实现了。

    举例：在构造函数里面添上this.BeforeInstall += new InstallEventHandler(InstallerTest_BeforeInstall)，实现方法InstallerTest_BeforeInstall(object sender, InstallEventArgs e)，就可以完成程序在安装之前要做的事情了，Installer主要要用到的事件不外乎这几个：BeforeInstall，AfterInstall，AfterRollback，AfterUninstall，BeforeRollback，BeforeUninstall，很明显的意义，就不用解释了吧

5.完成了4中的InstallerTest_BeforeInstall方法，是不是在安装之前就要执行InstallerTest_BeforeInstall中的代码呢，答案是不行。还要用到“自定义操作编辑器”。项目SetupTest->添加->项目输出->选择Library后确定，你会发现，在项目SetupTest下面多了一项“主输出来自Library(活动)”的字样。好，到正题了，来到“自定义操作”处，可以看到四个子项：安装，提交，回滚，卸载。在安装上面点击右键->添加自定义操作->选择“主输出来自Library(活动)”。这样，重新生成后，安装之前就会执行InstallerTest_BeforeInstall代码了。（当然，如果你在要在卸载前后做一些事情，就必须在“卸载”处也把自定义操作添加进去，否则即使监听了事件AfterUninstall，也不会执行的，完全之策是在安装，提交，回滚，卸载里面都把主输出来自Library(活动)添加进去就不怕了，嘿嘿）

6.如何得到用户所选安装路径？在实现代码的时候，用户所选安装路径大多数情况下都要用到，怎么得到他的值呢？在自定义操作->安装->“主输出来自Library(活动)”上面右键属性，在CustomActionData中填入/targetdir="[TARGETDIR]""， 然后在类InstallerTest中用这句话this.Context.Parameters["targetdir"] 就可以取得安装的目录了。

7.如何得到用户在安装的时候输入的更多的信息？用户在安装的时候，还有可能塡很多其他的信息，比如我以前做的一个安装文件中包含安装数据库，安装的时候用户就要填写数据库密码，用户名等信息。实现如下：用户界面->安装->启动->右键->添加对话框->选择"文本框(A)",可以看到文本框(A)有很多属性，如：Edit1Label，Edit1Property，Edit1Value，Edit1Visible等，Edit1Label就是這個输入框所要输入的内容的提示语，Edit1Value就是所输入的内容。

例如：将Edit1Label内容填上“请输入数据库用户名:”，将Edit1Property内容填上DBUSER,Edit1Visible选问true，另外在6中所提到的CustomActionData中，添加上 /dbuser="[DBUSER]",再在InstallerTest中用这句话this.Context.Parameters["dbuser"] 就可以或得用户在安装的时候所填写的用户名了（其他的单选框，复选框，button按钮框都差不多获取值的方法，不再累赘）.总之，添加了一些对话框后，把属性好，还要在CustomActionData添加一个/parm="[属性值]",才能在代码中获取改值。

8.其他:创建快捷方式,你可能在程序中要创建多个快捷方式，这里展示一种写代码的方式来创建快方式（各个选项意义就不说了，其中有些选项是可以不需要的，读者可自己测试选择使用）

添加引用：com名称为Windows Script Host Object Model

               using IWshRuntimeLibrary;

string dk = System.Environment.GetFolderPath(System.Environment.SpecialFolder.Desktop);//得到桌面文件夹

                IWshRuntimeLibrary.WshShell shell = new IWshRuntimeLibrary.WshShellClass();

                IWshRuntimeLibrary.IWshShortcut shortcut = (IWshRuntimeLibrary.IWshShortcut)shell.CreateShortcut(@dk + """快捷方式到数据中心SiteDataCenter.lnk");

shortcut = (IWshRuntimeLibrary.IWshShortcut)shell.CreateShortcut(@dk + """互联网公共上网服务场所信息安全管理系统.lnk");

                    shortcut.TargetPath = @"%HOMEDRIVE%/Program Files"Internet Explorer"IEXPLORE.EXE";

                    shortcut.Arguments = "http://localhost/web/Main.html";//参数

                    shortcut.Description = "火星网络安全有限责任公司";

                    shortcut.WorkingDirectory = pathWeb;//程序所在文件夹，在快捷方式图标点击右键可以看到此属性

                    shortcut.IconLocation = "notepad.exe, 0";//图标

                    shortcut.Hotkey = "CTRL+SHIFT+N";//热键

                    shortcut.WindowStyle = 1;

                    shortcut.Save();

9.如何在安装完成后自动运行某个程序呢（如run.exe）？

将这句代码写在AfterInstall里面：

                System.Diagnostics.ProcessStartInfo psiConfig = new System.Diagnostics.ProcessStartInfo(path + """run.exe");//path即是安装的目录

                System.Diagnostics.Process pConfig = System.Diagnostics.Process.Start(psiConfig);

(二)卸载（以下是转的，看客自己提取有用信息）

1.新建一C# Windows From项目，完成项目的工作

2.菜单-》文件-》添加项目-》新建项-》安装部署项目-》安装向导

选择上面那个项目的主输出。

在安装项目的属性栏中会显示出"XXX部署项目属性"，在第一个AddRemoveProgramsIcon属性后"浏览"到你的ICO图标，接着在下面Author输入作者名，也就是你的名字，其他Localization是语言属性，Manufacturer,ProductName的属性将是用户在安装时默认设置的文件夹路径一部分了，设置Title为你的应用程序标题名称。

3.右击"安装项目"上，选择文件系统视图

在应用程序文件夹，添加一个文件，如果是WIN98下的安装，请选择WIN98操作系统目录下的msiexec.exe文件（一般在c:"windows"system目录下）；如果是其他，比如WIN2K系列，请选择该操作系统目录下的msiexec.exe文件(一般在"winnt"system32目录下)。

4.在文件系统视图的"用户程序菜单"里，添加一个目录，命名为项目的名字，

并且，在这个目录里，添加一个快捷方式，指向项目的主输出，

再添加一个快捷方式，指向加入的文件Msiexec.exe，命名为："卸载"，

5.在解决方案里，选择安装项目，再打开属性视图，打开安装项目的属性面板，里面有一项：

ProductCode {5284694C-24C1-4FCF-A705-D4B9A0081274}

把后面的{5284694C-24C1-4FCF-A705-D4B9A0081274}，复制一下。

6.在"卸载"快捷方式的属性里，有一个叫Arguments的属性栏，在属性栏中首先填入"/x"后加一个空格，再"粘贴"，将ProductCode的属性值COPY至此。

7、再在VS.Net的顶部菜单栏中"生成"-"生成 XXXX"(XXXX为你的项目名称)

补充：合并模块项目的使用

若在某个大型的安装程序中，模块飞分为模块1，模块2等，每次都把模块N的文件加进来重新编译就很麻烦，这是就可用“合并模块项目(MergeModule)”。

1.新建一个合并模块项目，名字为MergeModule1，添加需要的文件，生成后，将看到MergeModule1.msm

2.在這個大型的安装程序中，添加“合并模块”，将MergeModule1.msm添加进去

3.生成，安装，你会发现MergeModule1的内容也被安装到了安装目录下

经史子集”是我国古代读书人对经典的分类法。

　　“经”是指古代社会中的政教、纲常伦理、道德规范的教条，主要是儒家的典籍，有儒学十三经：《周易》、《尚书》、《周礼》、《礼记》、《仪礼》、《诗经》、《春秋左传》、《春秋公羊传》、《春秋谷梁传》、《论语》、《孝经》、《尔雅》、《孟子》。

　　“史”是各种体裁历史著作，分为正史、编年、纪事本末、别史杂史、诏令奏议、传记、史钞、载记、时令、地理、职官、政书、目录、史评十五类。

　　“子”是诸子百家及释道宗教著作，分为儒家、兵家、法家、农家、医家、天文算法、术数、艺术、诸录、杂家、类书、小说家、释家、道家十四类。

　　“集”是收历代作家一人或多人的散文、骈文、诗、词、散曲等的集子和文学评论、戏曲等著作，分为楚辞、别集、诗文评、诗词五类。

　　注：以上分类依照　清《四库全书》。

-------------------------------------------------------

四书：亦称四子书，即《大学》、《中庸》、《论语》、《孟子》；五经：《诗》、《书》、《礼》、《易》、《春秋》。指儒家经典。

*对话框设计：在图形用户界面程序设计中，对话框是重要的信息显示和获取输入数据的用户界面对象。

1、公共对话框:

 公共对话框是利用windows资源的对话框，包括文件打开、文件保存、颜色设置、字体设置、打印设置等。

1) 文件打开对话框：用于打开文件

uigetfile

uigetfile(‘FilterSpec’)

uigetfile(‘FilterSpec’,’DialogTitle’)

uigetfile(‘FilterSpec’,’DialogTitle’,x,y)

[fname,pname]=uigetfile(…)

2) 文件保存对话框：用于保存文件

uiputfile

uiputfile(‘InitFile’)

uiputfile(‘InitFile’,’DialogTitle’)

uiputfile(‘InitFile’,’DialogTitle’,x,y)

[fname,pname]=uiputfile(…)

3) 颜色设置对话框：用于图形对象颜色的交互设置

c=uisetcolor(‘h_or_c,’DialogTitle’)

4) 字体设置对话框：用于字体属性的交互式设置

uisetfont

uisetfont(h)

uisetfont(S)

uisetfont(h,’DialogTitle’)

uisetfont(S,’DialogTitle’)

S=uisetfont(…)

5） 打印设置对话框：用于打印页面的交互式设置

dlg=pagesetupdlg(fig)

pagedlg

pagedlg(fig)

6) 打印预览对话框：用于对打印页面进行预览

printpreview

printpreview(f)

7) 打印对话框：

printdlg

printdlg(fig)

printdlg(‘-crossplatform’,fig)

printdlg(-‘setup’,fig)

2、MATLAB专用对话框

1)错误信息对话框：用于提示错误信息

errordlg  打开默认的错误信息对话框

errordlg(‘errorstring’) 打开显示’errorstring’信息的错误信息对话框

errordlg(‘errorstring’,’dlgname’) 打开显示’errorstring’信息的错误信息对话框,对话框的标题由‘dlgname’指定

erordlg(‘errorstring’,’dlgname’,’on’) 打开显示’errorstring’信息的错误信息对话框,对话框的标题由‘dlgname’指定. 如果对话框已存在，’on’参数将对话框显示在最前端。

h=errodlg(…)  返回对话框句柄

例：errordlg('输入错误,请重新输入','错误信息')

 2) 帮助对话框：用于帮助提示信息

 helpdlg    打开默认的帮助对话框

helpdlg(‘helpstring’)  打开显示’errorstring’信息的帮助对话框,

helpdlg(‘helpstring’,’dlgname’) 打开显示’errorstring’信息的帮助对话框, 对话框的标题由‘dlgname’指定

h=helpdlg(…)  返回对话框句柄

例：helpdlg('矩阵尺寸必须相等','在线帮助')

3) 输入对话框：用于输入信息

answer=inputdlg(prompt)  打开输入对话框，prompt为单元数组，用于定义输入数据窗口的个数和显示提示信息，answer为用于存储输入数据的单元数组。

answer=inputdlg(prompt,title)  与上者相同，title确定对话框的标题。

answer=inputdlg(prompt,title,lineNo)  参数lineNo可以是标量、列矢量或m×2阶矩阵，若为标量，表示每个输入窗口的行数均为lineNo；若为列矢量，则每个输入窗口的行数由列矢量lineNo的每个元素确定；若为矩阵，每个元素对应一个输入窗口，每行的第一列为输入窗口的行数，第二列为输入窗口的宽度。

answer=inputdlg(prompt,title,lineNo,defAns) 参数defans为一个单元数组，存储每个输入数据的默认值，元素个数必须与prompt 所定义的输入窗口数相同，所有元素必须是字符串。

answer=inputdlg(prompt,title,lineNo,defAns,Sesize) 参数resize决定输入对话框的大小能否被调整，可选值为on或off.

例：prompt={'Input Name','Input Age'};

title='Input Name and Age';

lines=[2 1]';

def={'John Smith','35'};

answer=inputdlg(prompt,title,lines,def);

4) 列表选择对话框：用于在多个选项中选择需要的值

[selection,ok]=listdlg(‘Liststring’,S,…) 输出参数selection为一个矢量，存储所选择的列表项的索引号，

输入参数为可选项’Liststring’（单元数组）, ’SelectionMode’（’single’或’multiple’,  ’ListSize’([wight,height]),  ’Name’ (对话框标题）等

5) 信息提示对话框：用于显示提示信息

msgbox(message)  打开信息提示对话框，显示message信息。

msgbox(message,title)  title确定对话框标题。

msgbox(message,title,’icon’) icon用于显示图标，可选图标包括：none(无图标)/error/help/warn/custom（用户定义）

msgbox(message,title,’custom’,icondata,iconcmap) 当使用用户定义图标时，icondata为定义图标的图像数据，iconcmap为图像的色彩图。

msgbox(…,’creatmode’)  选择模式creatmode,选项为：modal, non_modal, 和replace。

h=msgbox(…)   返回对话框句柄

6) 问题提示对话框：用于回答问题的多种选择

button=questdlg(‘qstring’) 打开问题提示对话框，有三个按钮，分别为：yes,no和cancel,’questdlg’确定提示信息。

button=questdlg(‘qstring’,’title’) title确定对话框标题。

button=questdlg(‘qstring’’title’,’default’) 当按回车键时，返回default值default 必须是yes,no或cancel 之一。

button=questdlg(‘qstring’,’title’,’str1’,’str2’,’default’) 打开问题提示对话框，有两个按钮，分别由str1和str2确定，’qstdlg’确定提示信息，default必须是str1或str2之一。

button=questdlg(‘qstring’, ’title’,’str1’,’str2’,’str3’,’default’) 打开问题提示对话框，有三个按钮，分别由str1， str2和str3确定，’qstdlg’确定提示信息，default必须是str1, str2或str3之一。

7) 进程条：以图形方式显示运算或处理的进程

h=waitbar(x,’title’)  显示以title为标题的进程条，x为进程条的比例长度，其值必须在0到1之间，h为返回的进程条对象的句柄。

waitbar(x,’title’,’creatcancelbtn’,’button_callback’) 在进程条上使用creatcancelbtn参数创建一个撤销按钮，在进程中按下撤销按钮将调用button_callback函数。

waitbar(…,property_name,property_value,…) 选择其它由prompt_name定义的参数，参数值由prompt_value指定。

例：

h=waitbar(0,'pleas wait...');

for i=1:10000

  waitbar(i/10000,h)

end

  close(h)

8）警告信息对话框：用于提示警告信息

 h=warndlg(‘warningstring’,’dlgname’) 打开警告信息对话框，显示warningstring信息，dlgname确定对话框标题，h为返回对话句柄。

建立一个简单的CDUA程序，需要设置的参数有很多，这样对于初学者来说增加了难度。在这里作者开发了一个针对VS2005的CUDA wizard 《CUDA VS2005 Wizard》，这样就可以很方便的在VS2005的环境中进行CUDA的开发。

下面是快速入门的流程：

1。安装 《CUDA VS2005 Wizard》就可以可以在VS的工程里面看到一个CUDAWin32App的工程向导。

点击这个就可以创建一个简单的CUDA工程，里面已经就有一个简单的example，可以进行CUDA开发了。

2.我这里创建的是一个测试Thread 的demo，所以生成的cu文件名字就是threaddemo1

可以看到这里已经由VS wizard向导自动生成了一些简单的hello cuda的程序。

3.我们再来看工程的环境配置，就可以发现Debug Release EmuDebug EmuRelease已经在系统设置里面了，已经由Wizard自动生成了工程的各个环节配置。

4.工程环境变量的查看和配置，右键点击工程，选择“属性” 或者 “properties" 就可以看到project的环境配置里面就多了一个CUDA选项，这里就可以对CUDA进行环境配置。

5.编译程序

6.运行 输出CUDA initialized. Hello CUDA。time！

7.调试，需要把cu文件关联到VC的project里面，不然调试的时候不能断点跟踪进去，所以必须在系统的VC++project 环境中加入*.cu文件，VS才能把你的程序和源文件关联起来，进行调试。
如果没找到关联cu，在调试的时候就会发现找不到源文件。

8.简单的快速开发就讲到这里了。我的字体这里显示高亮度，或许有些朋友会问，怎么才能让CU文件显示高亮度，可以查看我的另一篇翻译，《syntax highlighting when editing your .cu files in Visual Studio》

9.再下一节中，准备对CUDA的线程模型做一个简单的实用讲解;欢迎再来……hoho >_<!

ps：你的留言是对我的最大鼓励，Open Heart Breavly fly!

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. ");

        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. ");

        return false;

    }

    cudaSetDevice(i);

    printf("CUDA initialized. ");

    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 ");

    CUDA_SAFE_CALL( cudaThreadSynchronize() );

    CUT_SAFE_CALL( cutStopTimer( timer));

    printf("Processing time: %f (ms) ", cutGetTimerValue( timer));

    CUT_SAFE_CALL( cutDeleteTimer( timer));

    CUDA_SAFE_CALL( cudaMemcpy(&host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost));

    printf("%s ", 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的使用中体现出来的性能优势做一个简单的分析。

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 ");

    CUDA_SAFE_CALL( cudaThreadSynchronize() );

    CUT_SAFE_CALL( cutStopTimer( timer));

    printf("Processing time: %f (ms) ", cutGetTimerValue( timer));

    CUT_SAFE_CALL( cutDeleteTimer( timer));

    CUDA_SAFE_CALL( cudaMemcpy(&host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost));

    printf("%s ", 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了~~~

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 ");

    CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost));

    CUT_SAFE_CALL( cutStopTimer( timer));

    printf("Processing time: %f (ms) ", 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的使用；

写的内容一直都是文字比较多，代码比较少，其实学习的过程更重要的思想，实践的代码，最好是自己写，唯一可以学习的是思想，学习更重要的也是思想的交流，知识的传播，最好的是思想的传播，代码，方法，都是只是一些工具而已。但是工具的熟练层度，就得靠自己下来多练习。

8. CUDA 内存使用 global 二------GPU的革命

前面一章节已经写到了内存访问的问题，内存对齐的问题，不过在看到编程手册第五章的时候，还是会有很多朋友问到我关于CUDA的global内存访问的问题，怎么是访问的冲突，怎样才能更好的访问内存，达到更高的速度。下面先看几张图，这些图都是CUDA编程手册上的图，然后分别对这些图做解释，来理解硬件1.0，1.1 以及现在最新的硬件的访问内存的区别。

我们在这里再深入的讲解一下global内存对齐的问题，每次执行一条明命令的时候，都是会按照32个thread为一个warp，一起来执行，但是在执行的时候，又会按照硬件的条件（这里有两个限制条件，一个是内存访问的时钟和执行core的时钟不一样，第二个是为了细粒度的分支的问题）然后就会把16个thread组成的half-warp来一次访问global内存才能让访问内存的性能高一些，这个可以理解；

就像手册上说的那样，如果16个thread（half-warp）访问内存的时候，如果每一个thread访问32bits就是4个字节，那么就可以合并为一个64bytes的访问，手册上这点写得有点让人咋一看不太明白~4bytes（32bits）*16 = 64bytes，就是这么来的，如果每一个thread访问64bits（8个bytes），那么就可以合并为128bytes的访问；这里啦，由于合并访问的最大限制是128bytes，所以最大也按照128bytes一次访问来合并，如果超过，就得多次访问，或者如果没有按照这样的方式对齐访问，也会多次访问；下面是1.2device之前的访问的几个图，这里要把1.2device以前和以后的分开，是因为这里在对齐访问的方式的时候，有不同的策略；先看1.2device以前的能合并为一次访问的情况：

下图是编程手册上的图：

这里的每一个thread都是访问的对应的地址，是对齐的，所以可以合并为一个存储event；

下面这个图是没有对齐访问，就造成了non-coalesced访问的问题，下面可以看图说话：

左边的那个好理解，中间对应的thread访问的地址交叉了，thread3和thread4交叉访问了，在硬件1.2版本之前的都会造成Non-Coalesced访问；

详细的需要说明的是右边的为什么也造成了Non-Coalesced（非对齐）访问，这个是基础问题，大家理解的内存对齐是怎么样的？按照固定思路，或者教材上强调的都是中间过程的内存访问的对齐，但是内存是从offset 0x00000000位置开始的，就是偏移量0开始的，如果要真的满足内存对齐，严格的说起来就需要从内存的0地址开始算起，再加上我们知道的global内存的对齐方式有几种，4位bytes，8bytes，16bytes，这里说的对齐方式，注意区别关系；再来看看右边的那个图：thread0开始，从address128的位置开始向下便宜的位置是？132-128=4 偏移了4，16个threads整体访问的是16*4=64，是从132开始的，从0算起来，132-0 =132； 132/16 = 8…4,从整体上讲，从0偏移位置开始，偏移了4个位置，这里的就造成了访问的未对齐，这个是从整体角度上讲的，和左边的图比较一下，那个是按照局部对齐来说的，注意理解；

继续看图说话：

左边的图看看，算一下，局部的时候偏移了，从局部和整体来说，都会引起未对齐访问；

右边的图自己算一下，是不是超出了刚才我说的范围；所以造成了内存访问的未对齐情况；

前面我们看的图都是1.2版本前的硬件的情况下的内存访问情况，现在看看1.2版本以后的硬件；

这里解释一下，什么叫1.2版本的硬件，或许有些朋友也不太了解，g80架构的都是1.0或者1.1的硬件架构，现在的gtx200系列的都是1.3的架构，其实1.2的硬件架构，或许是Nvidia的一个内部的，没有推出产品，可能准备提供给低端的产品，但是我想没有推出低端的产品，直接就上1.3device了，市场需求吧~~如果下一步GTX的架构还是按照老路子，不改进的话，或许Intel的Lrb上来以后，对Nvidia的产品，就是一个很大的竞争了；

不说废话了，先看图：

1.2以后的硬件版本，弱化了threads之间交叉访问的时候，没对齐的情况，只要大家都在一次访问的64bytes的一个段里面，或者128bytes的一个段里面面，这样的段访问，那就可以不用多次访问，当然如果你16个threads分别跨过了16个段，那就得产生16个存储event~记住几个段的定义，这里说的段，就是我们常常理解的对齐的方式，全局的内存访问对齐方式，8个bits的是按照32bytes对齐，16bits的是按照64bytes对齐，32bits和64bits都是按照128bytes对齐；

在优化代码的时候，这个地方是一个值得注意的部分；

API函数里面有对齐访问的接口，会按照对齐的方式分配global内存给你，不过注意其中的一个offset值的使用，这个是为了解决全局情况下的对齐偏移的问题：）cudaMallocPitch，这个函数，注意使用~

(CUDA 编程7).CUDA内存访问（一）提高篇------按部就班

------GPU的革命

谈到内存的访问，其实也就是几个API函数的调用，感觉也没什么好讲的，知道cudaMalloc，知道cudaMemcpy，还有cudaFree，就应该可以对device上的内存进行分配，然后就可以在device上使用内存。但是就像我们看运动会一样，只看100米的那不到10秒的时间，也许你只会说，他跑得真快。又有多少人能清楚其过程中的磨炼啦？要真正理解device上内存的调度访问，让我们的程序达到更快的速度，就得对内存访问的情况做更深入的了解，明白其中的过程。

我记得大学刚开始上编程语言的课程的时候，都会讲到内存的分配和内存的释放，但是一般都是在书的最后面才会讲内存的对齐，内存的空间布局。就像C++一样，或许你学习了好几年，用了它好多年，但是你是否清楚class的内存布局？知道虚表是怎么一个内存访问过程？就像做网络，当做到一定地步的时候，才会发现数据从一段传递到另一段的时候数据发生了变化，内容不对了，定义的结构体传递到另一边的时候就错位了。当在使用SSE对数据处理加速的时候，是否为内存对齐的问题头疼？说了这么多，或许有人在烦了~~~讲正题吧，这些部分或许只有自己真正在实际运用中遇到问题，才会去考虑的，不过还是希望大家能在遇到问题之前，就能掌握解决问题的能力。不是在遇到美女的时候才开始刮胡子，平时的习惯就要养好。

前面这一段话或许只是生活的一些经历，其实可以跳过，当遇到的时候再回来看看：）我们还是按部就班的来讲解device上内存的访问吧。这么多年的发展，内存的价格是越来越便宜，但是有谁知道当你做编织内存的时候，很多人的眼睛都被弄瞎过啦，为她们当年为计算机作出的贡献bless一下。

还记得小时候经常看枪战片，看到小马哥端着枪一个劲的扫射，哪个帅啊~~不过后来稍微大一些以后，心理面一直有一个疑问，子弹用不完？一次能装多少子弹啊~就那么小的弹夹。8发子弹的左轮手枪在帅哥哥手里可以搞定十几个人~还不装弹- 囧！一般的自动手枪一般都是8发，14发，最多的驳壳枪（毛瑟枪）也才能装20发子弹。你要说人家是《第一滴血4》里面的史泰龙可以开着战车上面的M2HB12.7mm重机枪，子弹不用弹夹，一次可以装几千发的弹；对，人家是DMA直接内存访问，不是通过弹夹访问的，不需要处理器来中转数据。G80支持的内存访问的能力是一次访问4bytes，8byets或者16bytes，G80有三种弹夹，一种可以一次装4发子弹，一种是8发子弹，还有16发的。

Global Memory 在访问过程中没有Cache，就像以前的火枪一样，打一枪以后，装药，然后才能打下一枪。每次的访问时间是400-600个clock（core 跑的时钟）延迟。所以在CUDA编程中，其中一个瓶颈就是内存访问。利用SDK提供的bandwidthTest可以测试到host到device，device到host，device到device的传输性能。虽然PCIE有3.2G/s的理论值，但实际达不到这么多。Device to Device的传输能达到89G/s（GTX260）左右，理论值是90G/s（GTX260）也达到差不多了。这个地方每个人显卡不一样，主板不一样，设置环境不同，也不一定一样。

Device上active的一个warp有32个线程，但是实际是16个thread同时在运行，就是half warp。当half warp的16个threads访问内存的时候，最好让16个thread依次对着相对应的内存地址，这样的时候才能保证Coalesced Access。如下图http://www.isi.edu/~ddavis/GPU/Course/Slides/GPU+CUDA.pdf:

如果还没明白，看解释，每一个thread一次访问对应的内存地址，不互相交错，这样就不会卡壳了~。如果有交错了，就会出现下面的情况 Uncoalesced Access，

小时候玩过仿照左轮手枪的，自动的转动肯定会很快把子弹打完，没看到谁打完第1颗，再去打第3颗，再回来打第2颗子弹的~……这里也不多做解释了~

下面是一个例子，也是刚才那个连接上的：

把一个Uncoalesced float3 Code的代码优化为Coalesced的代码的过程：

Float3是12bytes，每一个thread就会读取3个float。还记得我们前面说过的每次换一个弹夹是4，8，或者16，但是明显float3不等于这些。还记得active thread是怎么工作的吧？如果不记得了最好参看前面的thread工作的章节。Warp工作的时候是16个threads一起运作的，所以同时就有16个threads访问内存，加起来就是16*3*4（16个threds，float3是3个float，每个float有4个byte）, 192个bytes。所以就造成了Uncoalesced Access。

下面怎么来解决这个问题，一个图：

看看这个图，然后再看一段代码：

自己先算一下，这里是怎么一个访问过程。

下面解释：我们设置block的thread为256，然后当大家都执行第一条存储命令s_data[threadIdx.x] = g_in[index]，还记得block内的thread的执行模型吗？SIMP，一条指令，同时256个threads都要执行完，才会执行下一个指令。

这里其实用到了shared memory来作为中转，避免了global内存访问的Uncoalesced。如果结构体不是size（4，8，16）还有内存强制对齐的方式来使得global访问的coalesced。例如用__align(x)来强制对齐内存，但是这里就会浪费一些空间，如果float3，的结构体，用__align(16)，对齐，就会有一个float的空白空间用来做对齐了~。

总结一下，就讲了一个global内存访问对齐的东东，如果访问不能连续，就采用两种方法来使得其对齐~其实看完前面的，理解到最后就这一句管用~

下一章节应该会讲到shared memory里面访问的bank conflict的部分，其实也很简单，画一个图就出来解决了~~

Ps：多画图，不明白的时候就画图出来：）

(CUDA 编程6)CUDA编程接口（二）------一十八般武器

------GPU的革命

4.     程序运行控制：像Stream，Event，Context, Module, Execution control这样的咱都把归类到运行管理的。这里也得分清楚有Runtime级别的，也有Driver级别的。

Stream：如果了解AGP时代的显卡，就知道在Device和Host之间进行数据交换的时候有作为中转数据的部分，叫着stream；发展到G8X以后就有了新的专门针对GPGPU统一设计的方便数据传输的Stream out的硬件设备。这个层的功能是将Vertex Shader和Pixel Shader（在G8X的设计中已经差不多不用把他们俩再分开，差不多统一起来）处理完成的数据输出给用户，由用户进行处理后再反馈给流水线继续处理。它可以直接读写本地显存。如果是内存对齐的，速度会更快，这个部分可以参考simpleStream的代码。

Event：和平时的大家在编程中使用的Event一个道理，都是起到通知的作用，创建一个事件，然后同步；Event在多线程编程中会经常遇到。

Context：上下文，什么叫叫上下文- -！（中文翻译叫上下文，实在的，俺觉得这个翻译不恰当，但是很多传统的书都这么翻译了，咱们也只能叫上下文）这里的Context和CPU里面使用的Context差不多，都是一个process里面需要用到的一些“资源”（系统资源：堆栈，内存……等等）。我倒觉得翻译为process包含的东东更为恰当。

Module：可以理解为linux里面的module意思；如果你不了解linux也不用担心，module就是一个个的专门针对Device的程序。知道很久很久以前的dos时代就有的.com文件吗？那个东东就像直接在CPU上可以调用的一个文件，加载到内存中就可以运行的。这里的module就是加载到Device上就可以运行。

Execution control：这个就是从驱动层面来控制线程怎么在Device上运行。

Stream和Event在Runtime API和Driver API中都有，函数接口也差不多。

Context, Module, Execution control就只有在Driver API层面才有。

上面的这些都没多少难点，其实API的讲解，一般都没多少难点，难点在于怎么能灵活的运用API，那就得通过多次的练习，多用API实践；其实有的时候也能发现API的不足，才会产生新的API。这里就不放具体的代码，代码可以参考编程手册上面。

5.     好了，这里就剩下OpenGL和Direct3D的接口函数了，他也有两个层次的API，有Runtime层次的，也有Driver层次的。Runtime和Driver层面都有调用的API，CUDA2.0在这里也做了一些优化，在内存交换方面。具体代码参考编程手册。

前面说过要讲解函数的定义，其实在以前的帖子里面正好已经翻译过一段，不过现在要讲解，也就提出来，专门再把那些东西翻出来：GPU NVIDIA_CUDA_Programming_Cuide_1.0 Chapter 4. Application Programming Interface

4.2.1 Function Type Qualifiers 函数类型

4.2.1.1 __device__

__device__ 规定的函数：

在device上执行。

只能在device上调用。

4.2.1.2 __global__

__global__ 定义了一个kernel函数：

在device上运行。

只能在host上调用

4.2.1.3 __host__

__host__定义的函数：

在host上运行，

只能在host上调用。

没有定义__host__，__device__或者__global__的函数等同于__host__函数，系统都会把函数编译成host函数。

另外，__host__定义可以和__device__定义一起使用，编译器会把这个函数编译为host和device通用的函数。

4.2.1.4 Restrictions (强调，限制)

__device__函数通常是inline，所以如果需要不inline 就需要加上uninline

__device__和__global__函数都不支持递归调用。

__device__和__global__函数都不能定义static变量在函数内部。

__device__和__global__函数不能使用可变参数。

__device__函数没有函数地址，也没有指向它的函数指针，但是__global__函数有。

__global__定义和__host__定义不能一起使用。

__global__函数必须是void返回类型。

任何调用__global__的函数都必须指明运行配置。（Section 4.2.3）也就是线程kernel的调用方法。

__global__函数是异步调用的。在运行结束前就会返回。

__global__函数的参数通常是通过shared memory调用到Device中，最多是265bytes。

(CUDA 编程5) CUDA编程接口（一）------一十八般武器

------GPU的革命

序言：所谓一十八般武器，不同的年代又有不同的说法，最早的汉武年间的：矛、镗、刀、戈、槊、鞭、锏、剑、锤、抓、戟、弓、钺、斧、牌、棍、枪、叉。到三国的：九长：刀、矛、戟、槊、镗、钺、棍、枪、叉；九短：斧、戈、牌、箭、鞭、剑、锏、锤、抓。再到明清的：弓、弩、枪、刀、剑、矛、盾、斧、钺、戟、黄、锏、挝、殳（棍）、叉、耙头、锦绳套索、白打（拳术）。《水浒传》里的：矛、锤、弓、弩、铳、鞭、锏、剑、链、挝、斧、钺、戈、戟、牌、棒、枪、扒。今天的武术届又有：刀、枪、剑、戟、斧、钺、钩、叉、鞭、锏、锤、抓、镗、棍、槊、棒、拐、流星。400多种古代冷兵器时代的武器，常用的也只有这么多种。也就像我们的API一样，API有无数多个，你自己都可以给自己造几个API出来，常用的，或者就那么多个。要打天下也不能扛着锄头，竹竿干吧。秦国之所以能统一六国，在武器上的统一，提供同一个的型号的武器装备（看秦的历史，就可以发现所有的兵器都是同一型号生产，弓弩上的器件可以互换，从兵马俑坑中找到的剑戟，箭头的尺寸误差很小，都可以互换），也是他能战胜其他六国的很好的基础。

正文：

子曰：工欲善其事，必先利其器。我们要把显卡作为通用并行处理器来做并行算法处理，就得知道CUDA给我提供了什么样的接口，就得了解CUDA作为通用高性能计算平台上的一十八般武器。（如果你想自己开发驱动，自己写开发库- -那我不得不佩服你很有时间，想必也不会有很多人想自己在去实现一个CUDA吧，呵呵，虽然实现一个也不是太难）。书接上回《 CUDA硬件实现分析（二）------规行矩步------GPU的革命》前面我们讲到了一些简单的CUDA的C语言扩展的规则，下面就具体来讲解CUDA给我听哦买提供了多少方便的API函数。在开发CUDA的时候，CDUA也给我们提供了一套完整的API函数。从一开始就在想，怎么把这些枯燥的API函数，或者这CUDA的一十八般武器说得清楚。如果按照中文的翻译的第四章那么讲解，或许晕的人更多，只知道这些是武器，而不知道什么武器是用来干嘛的。从序言看到，十八般武艺所列兵器大同小异，形式和内容却十分丰富。有长器械，短器械；软器械、双器械；有带钩的、带刺的、带尖的、带刀的；有明的、暗的；有攻的、防的；有打的、杀的击的；也有射的、挡的。我们来看CUDA的时候，一看到这么多的API函数，先来给他分一些类，然后才好徐徐道来。

一、            API总结：

1.     通用的一些接口，前一章节也有提高过：数学函数，时间函数，同步函数，原子操作；

2.     控制Device的函数；就是得到设备信息，管理设备信息的函数。设置那块显卡工作，得到那块显卡的性能。这里有分为driver级别的API和runtime级别的API；有人会问什么是driver级别和runtime级别请看图：

这个图我们在前一章已经看到过了。不会不清楚吧～- -！driver级别的API就是提供驱动级别的API，就像写驱动一样的感觉。Runtime级别的API就是封装了一些Driver级别的API，按照一些常规的方法封装了一些底层的API。其实这里就像我们平时生活中一样，最开始对汽车不熟悉的时候，买一辆车回来能开就ok了；能用熟悉Runtime级别的API就行了。慢慢的，感觉汽车自带的音箱不好，自己就开始买一些原始设备回来改装车；慢慢的感觉整车都有点不爽了，然后慢慢的发现想修改发动机，修改外形……就开始改装车了，这样的工作，就得从Driver级别开始做了，玩得更高级一些的就自己设计图纸，自己来用一些零部件来组装车了。这就是Driver API和Runtime API的关系。Runtime的在开始的用起来一般都比较方便，慢慢的发现如果高层（high-level）的Runtime API用起来不方便，就用底层（low-level）的Driver API来自己做改装的车……

3.     内存管理，host的内存，device的内存，global的内存，constant第，shared的，这里会分出来一章单独讲Texture（纹理）内存的使用，说实在的Texture也是内存～非要搞那么神秘，没办法，也只好拿出来单讲……PS：内存管理也分为Runtime级别的API接口，和Device级别的API接口。

4.     程序运行控制：像Stream，Event，Context, Module, Execution control这样的咱都把归类到运行管理的。这里也得分清楚有Runtime级别的，也有Driver级别的。

5.     好了，这里就剩下OpenGL和Direct3D的接口函数了，其实把，这也是为了方便做图来用的，主要是OpenGL和DX都已经成了图形显示方面的标准，so～显卡也得照顾这两个东东了，要不然显卡自己画……hoho要是真自己再来实现OpenGL或者DX，CUDA就真的会头大了，hoho～～还好就借用现有的图形显示的程序来做就行了。牛顿人家都说是站在“牛头人”（巨人）的肩膀上才能看得更远……咱也不要非自己费那么大的经去做一些无用功。想想吃不饱的时候没看到多少人减肥的，倒是现在吃得好了，减肥的人多了……长胖了去健身房减肥，始终感觉有点怪怪的……拿钱去做无用功，（有用功就是减肥）不过现在也有人在做实验，把多余的这些减肥的人的能量转化为电力……扯远了……提一下：他也有两个层次的API，有Runtime层次的，也有Driver层次的。

好了，差不多就这5个部分的API了。下面我们就来个个讲解CUDA的十八般武器。

二、            API讲解

记得小时候练过几天棍法，现在好像起式都记不起来啦，哈哈，只记得棍要齐眉，剑齐耳朵，不知道错没，hoho。大学的时候一直想给武术队的张锐学九节鞭，不过一直太忙，除了和他练习过几次散打对抗，看他刀+鞭，双刀，枪，剑，九节鞭练习过几次以外，到大学毕业，也都没机会学九节鞭，～现在也就留着那一根九节鞭在身边。看到它就能想起一些朋友吧：）大学毕业也就各奔东西，都不知道他现在在干嘛……

又走神了……想想前面五个部分，下面一部分一部分的讲解：

1.  一些通用的函数：

数学函数，前面章节已经讲了，提一下CUDA 2.0里面添加了一两个新的函数。

时间函数，clock()这个自己去查C的函数库；还有CUDA提供的几个时间函数，用起来也没什么难的，只是说一些，CUDA提供的几个时间函数在计算Device上的运行时间的时候，和CPU上的时间函数比起来还不是那么的完美的准确，所以在做时间数据的时候，最好多取几次求个平均。

同步函数前面章节讲到了， __syncthreads()函数，原子操作函数也就和以前的原子操作函数一个道理，也不用多加解释。不过其实我都觉得这__syncthreads()函数都应该归到程序运行控制部分。

2．Device管理

Runtime和Driver层面的API都提供了设备管理的函数，其实两个层面的API提供的功能都差不多，可以在API的说明中查到他的区别。这里需要指出来的是虽然Host主机上的多线程程序是应该可以同时访问同一块显卡（Device）的，毕竟显卡就是按照PCIE标准插在PCI插槽上的标准设备（也有AGP接口的），这就是一个多线程访问硬件的问题了。So～本来是应该可以的。但是由于CUDA的设计原因，这里的host上的当个多线程的线程每个线程都要执行CUDA kernel的时候，就必须执行在多个Device上面。保证每个线程访问的Device不是同一个。多个线程线程A不能分享线程B在Device上创建的资源。

Runtime API：cudaGetDeviceCount() 和 cudaGetDeviceProperties() 提供了遍历硬件设备，得到某个设备性能参数的功能。

int deviceCount;

cudaGetDeviceCount(&deviceCount);

int device;

for (device = 0; device < deviceCount; ++device) {

   cudaDeviceProp deviceProp;

   cudaGetDeviceProperties(&deviceProp, device);

}

cudaSetDevice() 设置某一块Device作为这个主机host上的某一个运行线程的设备:

cudaSetDevice(device);

这个函数必须要在使用 __global__ 的函数或者Runtime

的其他的API调用之前才能生效。 如果没有调用cudaSetDevice()，device 0 就会被设置为默认的设备，接下里的如果还有cudaSetDevice()函数也不会有效果。

Driver API：

cuDeviceGetCount()和cuDeviceGet()  看名字就知道干嘛的～（英语不好的这应该能看明白吧- -！不要被我这个那 国家四级都没过的人BS你哈～！～）

int deviceCount;

cuDeviceGetCount(&deviceCount);

int device;

for (int device = 0; device < deviceCount; ++device) {

   CUdevice cuDevice;

   cuDeviceGet(&cuDevice, device);

   int major, minor;

   cuDeviceComputeCapability(&major, &minor, cuDevice);

}

3．内存管理：

Device上的内存可以被分配成线性的，也可以分配为CUDA的数组形式的。CUDA的内存可以为1维，2维，还有3维（2.0版本）。内存的类型有unsigned8，16或者32位的int，16位（只有driver API可以做到）float，32位的float。这里分配的内存也只能通过kernel里面的函数通过处理纹理的方法来处理。这个地方也是GPU的历史原因了，以前都是处理图像的，所以这里叫纹理。……叫啥都是别人取得名字 - -！在计算机里面不就是内存嘛 - -！

Host的runtime的运行库也提供按照page-locked的内存管理的函数，page-locked的内存要比pageable方式快很多。好的东西往往比较少～page-locked就是很稀少的。如果通过减少分配pageable的内存来分配多的page-locked内存，系统需要的分页内存就少了，这也就会让系统的性能降低了。所以在处理这块的时候要合理。

Runtime API：

使用 cudaMalloc() 或者 cudaMallocPitch() 来分配线性内存，通过cudaFree()释放内存.

下面是分配一个大小为256 float数组的方法:

float* devPtr;

cudaMalloc((void**)&devPtr, 256 * sizeof(float));

在使用2D数组的时候最好用cudaMallocPitch()来分配，在guide的第五章在讲到内存之间的调度的时候，就会看到他的好处。下面是一个分配一个大小为width×height 2D float数组的例子:

// host code

float* devPtr;

int pitch;

cudaMallocPitch((void**)&devPtr, &pitch,

              width * sizeof(float), height);

myKernel<<<100, 512>>>(devPtr, pitch);

// device code

__global__ void myKernel(float* devPtr, int pitch)

{

   for (int r = 0; r < height; ++r) {

       float* row = (float*)((char*)devPtr + r * pitch);

       for (int c = 0; c < width; ++c) {

          float element = row[c];

       }

   }

}

CUDA 的数组方式，需要用 cudaMallocArray()和cudaFreeArray(). cudaMallocArray()又需要cudaCreateChannelDesc()来管理，这个其实可以在第guide的第五章里面可以看到，我们后面也会详细的介绍内存的调度和管理，和传统的GPU的内存方式不一样的地方.

分配 width×height 32位float的CUDA array例子:

cudaChannelFormatDesc channelDesc =

cudaCreateChannelDesc<float>();

cudaArray* cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);

cudaGetSymbolAddress用来在全局中定位一个数组的位置，然后 cudaGetSymbolSize()来回忆他分配的时候大小----如果有全局编程或者多线程编程经验的，或者用过几个函数同时处理一个数据的经验，都会为了把数据的独立性弄出来，不能让数据和函数耦合太大，一般都不会让函数直接牵扯上数据，只是在函数处理的时候重新定位数据----这地方有点绕～～

下面是一些例子，内存之间的拷贝，回想一下有几种内存～linear的有两个函数可以分配的，还有CUDA array的内存:

cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch,

                 width * sizeof(float), height,

                 cudaMemcpyDeviceToDevice);

The following code sample copies some host memory array to device memory:

float data[256];

int size = sizeof(data);

float* devPtr;

cudaMalloc((void**)&devPtr, size);

cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);

从host上面拷贝内存到device的constant上面:

__constant__ float constData[256];

float data[256];

cudaMemcpyToSymbol(constData, data, sizeof(data));

-#%#$^$^*^&系统int中断：(@$:突然发现已经是早上8点多了 - -！又一夜！@#￥%……！#……

Driver API：

cuMemAllocPitch()被推荐来作为2D的数组分配函数，会在内存对齐方面做一些check～然后保证在处理数据（像cuMemcpy2D()这样的拷贝函数)的时候达到最优的速度。下面就是一个float的width×height的2D的数组，并在循环里面处理的例子: ----（这个地方扩展提两个东东：一个是内存对齐，做程序优化的时候，或者处理网络问题的时候，都会遇到这些问题，内存对齐是一个普遍存在的问题，像SSE这样的编程的时候也是要求内存对齐的，看以后要是有机会以单独讲解内存对齐的问题：）第二个是CUDA的内存访问的问题，就是很多朋友都会在写kernel的时候，搞不明白里面的threadid，block id和传进来的内存的关系，这个地方必须要搞清楚的；内存和线程是CUDA编程必须搞明白的两个概念，不然到时候就会很混乱。其实看看前面的章节，应该能看明白的，不明白就用手画图～要是感觉还是有点不太清楚，也不要担心，接下来的章节会单独把一些难懂的问题更详细的讲解。）

// host code

CUdeviceptr devPtr;

int pitch;

cuMemAllocPitch(&devPtr, &pitch,

              width * sizeof(float), height, 4);

cuModuleGetFunction(&cuFunction, cuModule, “myKernel”);

cuFuncSetBlockShape(cuFunction, 512, 1, 1);

cuParamSeti(cuFunction, 0, devPtr);

cuParamSetSize(cuFunction, sizeof(devPtr));

cuLaunchGrid(cuFunction, 100, 1);

// device code

__global__ void myKernel(float* devPtr)

{

   for (int r = 0; r < height; ++r) {

       float* row = (float*)((char*)devPtr + r * pitch);

       for (int c = 0; c < width; ++c) {

          float element = row[c];

       }

   }

}

cuArrayCreate()和cuArrayDestroy()来创建和释放CUDA array类型的数组。下面是一个 width×height 32-bit float类型的CUDA array 分配的例子：

CUDA_ARRAY_DESCRIPTOR desc;

desc.Format = CU_AD_FORMAT_FLOAT;

desc.NumChannels = 1;

desc.Width = width;

desc.Height = height;

CUarray cuArray;

cuArrayCreate(&cuArray, &desc);

这个也是几个内存之间拷贝的例子:

CUDA_MEMCPY2D copyParam;

memset(&copyParam, 0, sizeof(copyParam));

copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY;

copyParam.dstArray = cuArray;

copyParam.srcMemoryType = CU_MEMORYTYPE_DEVICE;

copyParam.srcDevice = devPtr;

copyParam.srcPitch = pitch;

copyParam.WidthInBytes = width * sizeof(float);

copyParam.Height = height;

cuMemcpy2D(&copyParam);

拷贝host上面的内存到device上面：

float data[256];

int size = sizeof(data);

CUdeviceptr devPtr;

cuMemAlloc(&devPtr, size);

cuMemcpyHtoD(devPtr, data, size);

本来想一次把下面的4，5点也讲了～但是如果一下讲出来～lz帖子太长了～～这就不好了～～呵呵，实在的，看第四章的中文翻译的时候，就看了前面几个就不想看了～帖子不能老长老长的又不吸引人- -！hoho～ 所以后面的4，5就在下贴里面发了～～看了这么多也比较累的～好好的休闲一下～你还会发现出了4，5部分，还少了一个部分，那就是CUDA自己的函数，怎么定义，有device的，有global的，这个又有怎么区分啦～且听下回讲解。Hoho

ps：熬夜不好～熬夜很伤身……

有的时候，我们经常会用旧的东东来和新的东东比较，就像C和C++都不知道争论多少年了。其实很多时候我到觉得是没必要的争论，除非你是做C或者C++本身开发的。就像新的硬件不停的变化，以前的概念或者今天就不能用了，有的时候，我们关心架构在这个之上的程序开发就够了，没有太多的必要去问茴香豆的茴有几种写法。有的时候看到论坛里面的争论的时候，很多都不太清楚问题的也参加到争论里面，感觉就是明白事的没有几个，倒是起哄的不少 - -！人家说当局者迷，旁观者清，我看现在很多时候倒是当局者清，旁观者瞎起哄- -！呵呵，好像合乎了现在很多选秀节目的心理哈，起哄的人越多，人家的节目越红，哈哈～----扯远了……

Ps2：在学习API的时候，最好的方法就是多实验，多尝试API的性能，就像练习武术中的器械一样，经常用才会精通的，才能在这个基础上想出新的招式。就像有人在CSDN论坛问，怎么才能弄好ACM比赛，看到像刘汝佳写的《算法艺术与信息学竞赛》的书，都是一些方法，没有代码，就觉得很不解，就像问用什么代码来实现。我的回答就是：阅尽天下A×，心中自然无码。多做代码的练习，找一本数据结构的书，对照上面的代码都自己实现一遍，找一个代码练习的书，自己都重新写一遍。要不然就把像MSDN这样的介绍API的资料上的Demo能实现多少都去实现以下，呵呵。其实学习怎么编码都是一些基础工作。最重要的到最后都是算法的实现。其实到最后你会发现写程序就是数组的处理，就是字符串的处理……仅此而已～。所以不要小看了C语言那些书上的小程序例子，不要小看了输出（**）星星这样的例子，实际程序中很多时候都是处理这些星星～hoho～不要一上来就想去写什么游戏，写什么网络软件～先把字符串处理好了，就nb了，真的～你看ACM的题目，topcoder的题目，几乎都是字符串处理- -！哈哈。多实践，多去做一些demo例子。到真正用的时候就会觉得手到擒来，怎么用怎么顺手了。

还是那句话：约尽天下A×，心中自然无码。

（CUDA 编程4）CUDA硬件实现分析（二）------规行矩步

------GPU的革命

前面已经讲解了很多概念上的东西，其实CUDA的最重要的两个东西，就是线程和内存。只要掌握了这两个东西，CUDA的东西也就很简单了。它的编写语言是C扩展的，所以，就当C语言用就行了，只是主要它的特殊的几个标志就ok了。前面讲解了线程和内存的模型，大概，应该，似乎，可以在你的脑海里面有一个概念了吧。只要有这个概念，我的文章的目的就达到了。前面的《CUDA硬件实现分析（一）------安营扎寨-----GPU的革命》已经讲解了线程在CUDA的具体运行过程。下面我们一起来看看内存在CUDA的硬件实现中的一些规定。这也比较合理吧，大军安营扎寨了，就应该颁布规则制度，只有了解CUDA的规则制度，才能真正的把各个线程都管理好。才能在这个平台上让程序高效的运行。

这里我们先明确几个

一．            Threads，Warps, Blocks

1. 一个warp最多有32个threads。只有在总线程少于32的时候，才可能在一个warp里面少于32个线程。

2. 每一个block最多有16个warp。就是说一个一个block里面最多有512个thread。

3. 每一个Block在同一个SM上执行，也就是同一个block的warp都在同一个SM上运行。

4. G80有16个SM。

5. 所以最少16个blocks才能占全所有的SM。

6. 如果资源（看看前面讲解的线程都要从device哪里分什么资源）够线程分，一个SM上面可以跑多余一个block的线程。就是同时可以跑2个，3……个block的线程。

二．            访问速度

Register—HW 一个时间周期

Shared Memory------HW一个时钟周期

Local Memory --- DRAM，no cache， 慢

Global Memory --- DRAM， no cache， 慢

Constant Memory --- DRAM， cached， 1……10s……100s个周期，这个和cache的locality有关。

Texture Memory --- DRAM， cached， 1……10s……100s个周期，这个和cache的locality有关。

Instruction Memory（不可见）--- DRAM，cached

三．            CUDA程序架构

如图

四．            语言扩展

单从学习语言来说，我到觉得应该精学一种，然后其他的就触类旁通了。在学习新语言的时候，要想快速的入门，也有诀窍。1.变量的定义方式。2.函数的定义方式。3.逻辑控制方式（if，loop……）。只要把这3个东西弄明白了，管他啥新的语言，20分钟就可以入门……然后入门都可以了，那要慢慢的更入的研究，那就得看你对这门语言的了解了。其实万变不离其中。其实从计算机编程语言的角度出发，就是定义一些数据，然后对数据进行操作，so……学习语言就从这个角度入手，那就很简单了。像java或者C#等一些语言比C语言多的新的特性不外乎就是方便你开发而已。

所以我们这里再来说说CUDA的语言，不外乎就是扩展了C语言，为了方便在GPU显卡上运行，规定一个特定环境。就定义一些特定的变量，说明他们是在GPU上的。这里有说明内存，和函数，是在GPU上的……so，这样一来，CUDA就扩展了C语言的变量分配定义和函数定义。

上面这张图是来之Fall 2007 syllabus，上面已经说得很清楚各个变量定义的时候的位置和生存周期。其实就是在C语言的常规变量的时候，定义了变量的位置而已。

其中有一个约束限制，就是指针变量，在kernel里面的指针变量，只能指向从global上面分配的内存。

五．            内建的变量

所谓内建，就是CUDA自己在kernel里面定义的一些变量。就像我们以前计算线程id的时候，就利用了他的自己的变量，dim3 gridDim；dim3 blockDim；dim3 blockIdx；dim3 threadIdx；注意gridDim，他的gridDim.z在现在的CUDA1.1版本中没定义。

内建变量，[u]char[1..4], [u]short[1..4], [u]int[1..4], [u]long[1..4], float[1..4] 就是构建了有4个变量的struct；

             uint4 param; ---》等价位一个struct里面有4个int。

             int y = param.y;

dim3 就是unit3这样的struct。

六．            通用的数学函数

·          pow, sqrt, cbrt, hypot

·          exp, exp2, expm1

·          log, log2, log10, log1p

·          sin, cos, tan, asin, acos, atan, atan2

·          sinh, cosh, tanh, asinh, acosh, atanh

·          ceil, floor, trunc, round

但是这里要指出，有几个函数是不精确的，但是可以很快的运行：

–        __pow

–        __log, __log2, __log10

–        __exp

–        __sin, __cos, __tan

七．            在host部分的运行库（CUDA Runtime ）

1.       提供Device管理的api（有多个显卡的时候怎能来设置，这里我们还没讲当多显卡并行运行库）

2.       初始化调用的Runtime函数

3.       每一个host的线程（thread）只能调用一个device的函数在一个device上运行。就是同时不能有几个host（主机）线程在调用同一个device上的运行函数。

八．            内存管理函数

我们讲了那么多的内存，下面就来看看到底是什么样的函数：

–        cudaMalloc（），cudaFree（）；内存分配和释放（device上的）

–        内存copy：cudaMemcpy(), cudaMemcpy2D(), cudaMemcpyToSymbol(), cudaMemcpyFromSymbol()；

九．            线程同步函数

void __syncthreads();

同步同一个block里面的线程，让block里面的线程都允许到这一点的时候，就等待同一个block里面的其他线程，就像军训的时候，大家吃饭吃完了还不能一个个走，必须得一个桌子的人都吃完了才能走。还得列队一起走，呵呵，这就是同步。所以最好保证每一个kernel里面的处理都是很快的，这样才不会让其他thread等待太久，不然会挨骂的 - -！hoho

哥们来点实际的吧- - 肯定很多人都这么在吼了。嘿嘿，下面让我们看一段代码，example里面的transpose：

#define BLOCK_DIM 16

// This kernel is optimized to ensure all global reads and writes are coalesced,

// and to avoid bank conflicts in shared memory.  This kernel is up to 11x faster

// than the naive kernel below.  Note that the shared memory array is sized to

// (BLOCK_DIM+1)*BLOCK_DIM.  This pads each row of the 2D block in shared memory

// so that bank conflicts do not occur when threads address the array column-wise.

__global__ void transpose(float *odata, float *idata, int width, int height)

{

    __shared__ float block[BLOCK_DIM][BLOCK_DIM+1]; //(1)

    // read the matrix tile into shared memory

    unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;

    unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

    if((xIndex < width) && (yIndex < height)) // (2)

    {

        unsigned int index_in = yIndex * width + xIndex;

        block[threadIdx.y][threadIdx.x] = idata[index_in];

    }

    __syncthreads(); //(3)

    // write the transposed matrix tile to global memory

    xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;

    yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;

    if((xIndex < height) && (yIndex < width))

    {

        unsigned int index_out = yIndex * height + xIndex;

        odata[index_out] = block[threadIdx.x][threadIdx.y];

    }

}

是不是在代码面前都感觉是那么的亲切啊～嘿嘿，有时候看算法的书看久了以后，看到代码实现，那才叫亲切啊>_<!

下面对加红的3处代码分别讲解：

（1）      shared__ float block[BLOCK_DIM][BLOCK_DIM+1];

为啥要用shared啊？知道为啥不，快啊～shared在block里面，是大家公用的，不用再从global哪里调换。还记得军训的时候食堂吃饭的时候那个内存讲解吗？桌子上的菜，肯定比你从中间几个桶里面去拿菜要快啊～

（2）      if((xIndex < width) && (yIndex < height))

由于CUDA在生成线程的时候，是一块一块的生成的，所以有的时候，x和y有可能越界。数据都是均匀的分配到每一个block里面去的，很可能就会出现有的blcok里面分配到的数据不一定完整，要是想不明白的（先画图想想），等后面的章节，举例子说明。

（3）      __syncthreads();

同步，就是等每一个thread都运行到这里的时候，再接下来运行。这里也避免了数据访问的冲突。

（4）      还有一个得讲讲，为啥下面的还重新计算一下x和y，因为线程被挂起来了以后，重新进入warp的时候，分配的id就不一定是一样的了。这里就有点像铁打的营盘流水的兵，每一波人都可以分配到同一张桌子吃饭，但是不一定是上一波人了～。碗筷还是那些碗筷，位置还是那个位置，坐的人不一样了。

上面的都是C语言扩展库的东西，就是说，这些是扩展C语言的东东，只要是基于C语言的规则都可以运行的，不管是host还是device。

好了，今天要讲的部分应该差不多了，要是再写多一些，或者就有会像手册第4章那样，看了一半很多人就不想看了～- -！太长，lz 帖子太长，顶lz- -！我可不想看到这样的留言，hoho～

从小到大，我们听过好多规则，但是好多时候听进去的有多少？很多玩游戏的人，游戏的规则都不清楚- -ps，俺现在还不清楚StarCraft里面的某些兵的战斗力是多少 - -！就玩- -！小小的bs一下自己。只是为了消遣罢了·

我始终觉得自己不适合玩游戏，我也知道自己玩游戏不会完成专业玩家，so～还是老老实实做自己的工作，做一些能养活自己的事情，或许能让自己更充实一些。在游戏里面充实以后，或许在现实中就会空虚了- -！到时候就不是游戏人生，那就是人生被游戏了。呵呵。想起高中的数学老师，虽然在很多journal上发表了文章，也自己出版了数学的几本书了。但是和别人打麻将的时候总是输多赢少。问他为啥数学这么厉害，玩麻将不行～他说，我玩麻将的时候只是消遣，根本就没用脑袋想麻将这回事。

或许我们做事情的时候，就应该专注做自己的一些方面吧，只要坚持下来做了，做好了，做精了，那就够了。小时候就从长辈哪里听到，半灌水叮当响这样的词，或许我们更应该低调做人，高调做事吧。

（CUDA编程3）CUDA硬件实现分析（一）------安营扎寨

前面我们已经说到线程的一些概念，但是这些概念都是软环节的。我们常常会听到某某单位说他们的软硬件配置如何如何的好。软件再好，每个士兵都是可造之才，但是如果硬件条件跟不上，也没他们的勇武之地。就像国内做过一些跳槽原因的统计，很多人已经跳槽都是为了高工资，但实际统计结果表明，很多人都觉得在以前那个公司里面学不到东西，或者说得不到发挥自己长处的地方。那就得看公司有没有这样的机会让你发挥你的能力了。看到这里，扯远了～很多人都看得不耐烦了……书接上回《 CUDA 线程执行模型分析（二）大军未动粮草先行------GPU的革命》。已经讲到CUDA在线程模型是一个什么样子，经过几天的吸收，也应该在脑子里面有一些印象了。但是你会问一个问题，我们是不是可以开无数个线程来执行啦？或许招兵的人都想找到很多很多的人，但是你得也考虑一下你的粮食有多少，军营有多大。在这么我们得讨论一下现在的支持CUDA的nvidia的显卡现在的硬件情况。

空洞的讲解或许还是没怎么又说服力，下面以G80为例子。

1. G80里面有16个Multiprocessor.

2. 每个Multiprocessor都有一组（G80里面是8个）32位的Processor（每个Processor都是SIMD架构，什么叫SIMD架构：军训的时候，大家都到了食堂，不是像在学校里面，每个人那自己的碗筷就三三俩俩的去吃饭，那可要讲究纪律，啥叫纪律，一群人站在桌子面前，连长没发话，谁也不敢坐下来……连长一声令下：“坐下”。所有的人才按照敢坐下来，也是同时坐下来- -！要是谁没有同步坐下来，那就惨了- -！再来一次，一定要是同步坐下去的，都能听到声的，咵！恨不得把板凳给坐碎了------还是部队的东西结实，兄弟们怎么坐到军训结束只有坐坏屁股的，没有听说凳子坐坏的- -！所以啊，爱惜公物就是爱惜自己。记住了吧，这就是SIMD：Single Instruction Multiple Data 。 ）还有共用的Instruction Unit（这玩儿就不用翻译是啥了吧～看SIMD，自己理解去）。在G80里面有两个Instruction Unit模块。

3. 每一个时钟周期内，按照warp（这玩儿咋翻译啦？就理解为运行的时候，一个block里面一起运行的thread，例如block里面有512个thread，但是每次只有32个thread在运行，那么这32个thread就是一个运行的warp组- -! 还好不是rap- -！俺就真没法解释了）

4. 每一个warp里面包含的thread数量是有限的，现在的规定是32个。将来不知道会不会有变化？不知道，这个只有CUDA开发人员知道了。

我们还是按照我们的既定方式学习吧～看图说话- -！接下来又是一张图： 

刚才我们已经说到了硬件上的处理器的模型，现在我们来看看内存吧。前面的章节我们已经讲到了内存的register，local，shared，constant，texture，global，现在看到对应的硬件的位置了吧？哦，对了，有人会问，local和global跑那里去了，你看看，是不是多了一个Device Memory嘛。Local和Global 都是在Device Memory上的。只是在线程模型运行的时候，为了方便说明线程的模型，把它单独划分出来好说明。其实Constant和Texture都是在Device Memory里面的，但是人家是只读的Cache，所以比Global和local要快一些。不过Cache和shared Memory比起来，又显得慢一些了。

硬件的架构差不多就是这么多。如果还不是很清楚，我们再讲讲军营的生活。还是讲吃食堂吃饭吧（回想起大学的时候和几个好朋友一起成立了一个编程爱好者协会……结果后来变成了好吃协会- -！）。由于我们现在的连队人数众多，一个grid里面就可以有65535个thread，所以一下子让这么多人一起吃饭，是不可能的。我们有好多个grid怎么办？所以每次Device（就是显卡）只处理一个grid。但是每个grid里面又有那么多人。一来是咱们地方小，不能容纳那么多人一起吃饭，而来是也没必要再建立那么多食堂。大家可以轮流吃，早吃的早走，晚来的也晚不了多久。

这次是Grid1先到达食堂，然后他把士兵分成了几个block，（我们上次已经讲过的，block和grid的关系，应该还记得把）但是现在桌子还是少，为了大家不乱起来，每个block的人就指定到某一个multiprocessor（一个周期只能容纳32个人一起用餐）哪里。

这个block1的人就围着multiprocessor1这张桌子坐，但是这张桌子也不是太大，所以block里面的人又分成32个人一组（warp）来吃饭。那如果这个block没有32个人，就占不完这张桌子，所以一个multiprocessor又可以让别的block的thread一起来用餐。

因为一次就把所有的register和shared memory分配给了所有的block的所有的thread。所以怎么样合理的少用register和shared memory就可以同时让更多的thread得到处理。

讲了这么多，怎么来看自己的显卡那片军营如何拉？下面给出两种方法，一种是用一个小软件，GPU-Z的工具，这个工具可以得到很多显卡的参数，这个可以上网搜索到。

第二种方法是用CUDA的命令，这个只是针对支持CUDA的显卡的，下面给一段代码：

/********************************************************************

*  InitCUDA.cu

*  This is a init CUDA of the CUDA program.

*********************************************************************/

#include <stdio.h>

#include <stdlib.h>

#include <cuda_runtime.h>

/************************************************************************/

/* Init CUDA                                                            */

/************************************************************************/

bool InitCUDA(void)

{

   int count = 0;

   int i = 0;

   cudaGetDeviceCount(&count);

   if(count == 0) {

       fprintf(stderr, "There is no device. ");

       return false;

   }

   for(i = 0; i < count; i++) {

       cudaDeviceProp prop;

       if(cudaGetDeviceProperties(&prop, i) == cudaSuccess)

       {         

          printf("name:%s ",                       prop.name);

          printf("totalGlobalMem:%u ",             prop.totalGlobalMem);

          printf("sharedMemPerBlock:%u ",          prop.sharedMemPerBlock);

          printf("regsPerBlock:%d ",               prop.regsPerBlock);

          printf("warpSize:%d ",                   prop.warpSize);

          printf("memPitch:%u ",                   prop.memPitch);

          printf("maxThreadsPerBlock:%d ",         prop.maxThreadsPerBlock);

          printf("maxThreadsDim:x %d, y %d, z %d ",    prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);

          printf("maxGridSize:x %d, y %d, z %d ",    prop.maxGridSize[0],prop.maxGridSize[0],prop.maxGridSize[0]);

          printf("totalConstMem:%u ",              prop.totalConstMem);

          printf("major:%d ",                   prop.major);

          printf("minor:%d ",                   prop.minor);

          printf("clockRate:%d ",               prop.clockRate);

          printf("textureAlignment:%u ",           prop.textureAlignment);

          if(prop.major >= 1) {

              break;

          }

       }

   }

   if(i == count) {

       fprintf(stderr, "There is no device supporting CUDA 1.x. ");

       return false;

   }

   cudaSetDevice(i);

   printf("CUDA initialized. ");

   return true;

}

（CUDA编程2）CUDA 线程执行模型分析（二）大军未动粮草先行

------GPU的革命

书接上回--《CUDA 线程执行模型分析（一）招兵------ GPU的革命》，经过几天的征兵，将士也招了差不多五六千人，五六千人，就是五六千张嘴，得吃饭。这和我们的CUDA内存模型又有什么相识啦？或许有很多人都参加过军训，知道军队里面吃饭吗？到吃饭的时间，各处训练的队伍都拉到食堂前面。先到的就先进去吃，后到的就在外面拉歌，等待下一波，这里我们暂时还不讲解怎么安排每一对吃饭。我们先讲讲每一个Grid手下的兄弟们吃饭都需要那些家伙。下面来看看站长图：

这里我们有一个block的士兵（16个人）在吃饭，每个人有一双筷子（Registers），一个碗（local memory），大家公用一张桌子（Block 里面的shared memory）。记得军训那年，在河北宣化，八个人一张桌子，一起抢桌子上的吃的时候，那个景象。训练了一天，一个个都是如狼似虎的，嘴巴里面咬着，筷子里面夹着，碗里面盛着，研究还看着桌子上的那几根鱼骨头（鱼肉早就进肚子里面了- -！）那个时候就看谁的筷子夹得多，碗装得多了。不过再多，筷子就那么长，碗就那么大～咱们再来看看block的情况：

如图： ……不需要注释了吧……

餐厅里面肯定不止是一张桌子了，我已经不记得当时军训的时候，那个军营的食堂有多少个桌子，不过我们吃饭的时候可以装下几个连一起吃饭……想想那吃饭的阵势……

哦，对了，我还记得在食堂里面有几个大桶，那玩儿应该不能称之为桶了，太大了 - -！可以分为三类，一类是装菜的，一类是装饭的，一类是装水果的。这几个桶都是大家可以共用的，记得有好吃的菜，大家都最先吃完那个菜，然后再跑到中间这个几个桶的地方找这个菜，装完一份跑到桌子那边一起分享……很是怀念那段日子。 Long long ago了……

我能回想起来的那个时候的饭堂，如图（一个大的Grid吃饭图……）：

CUDA的内存模型也就这个样子了～

这些都是我们承诺给士兵的，每个人都要吃上饭。但是实际上我记得食堂再大，也不能一下子装下两个团的人吃饭，那个时候都是每个团里面的几个连队先吃饭，然后后到的就在外面拉歌，等里面的吃好了，后面的再进去吃……不是每个食堂都像北大的“万人坑”那样能装1w人- -实地考察，好像要是真装了一w人，大家也就不用坐着吃饭了～- -

对应手册上的第二章，运行模型就讲完了，接下来就要把队伍拉出来练习练习了……

对了，我们的SIMD还没讲啦～队伍人数那么多，现在就那么几杆枪～怎么练习，还没讲啦……接下来的章节更精彩……

PS: 题外话，或许是小的时候喜欢画图吧～现在还记得小学的美术老师给我们上的第一堂美术课，在黑板上画的正方形。在那个时候好像脑子里面就可以看穿物体一样，看到物体的这一面，就可以在大脑里面勾画出他的另一面～记得大学的时候参加定向越野，看到地图的时候，第一感觉，图形都好熟悉，然后再和实际的物体一对照，都是那么清晰，看到一栋建筑物，一对照地图，就可以知道他的后面是什么样，仿佛整个地图就在脑子里面构建起了地形图～呵呵很是奇妙……不知道大家有没有这样感觉的时候……

（CUDA 编程1） CUDA 线程执行模型分析（一）招兵

------ GPU的革命

 序：或许看到下面的内容的时候，你会觉得和传统的讲解线程，和一些讲解计算机的书的内容不是很相同。我倒觉得有关计算机，编程这些方面的内容，并不都是深奥难懂的，再深奥难懂的事情，其实本质上也是很简单的。一直以为计算机编程就像小时候搭建积木一样，只要知道游戏规则，怎么玩就看你自己了。或许是从小学那会，就喜欢在做数学题的时候用一些简便方法来解题，养成了一些习惯，喜欢 把复杂的问题都会尝试用最简单的额方法来解决，而不喜欢把简单的问题弄得很复杂。不再多说了，有的朋友已经看得不耐烦了……ps：再罗嗦一句，如果下面看不明白的，就当小说看了，要是觉得不像小说，那就当故事看，要是觉得故事不完整，写得太乱，那就当笑话看，在各位学习工作之余能博得大家一笑，也倍儿感荣幸……ps2：想好再说 ……突然想到了，确实是了一段时间再想到的，既然叫GPU革命，那就得招集队伍啊，下面我就开始招兵了。

正题：

要真正进入CUDA并行化开发，就必须先了解CUDA的运行模型，才能在这个基础上做并行程序的开发。

CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格（Grid）的概念在显卡硬件（GPU）上执行。每一个线程网格又可以包含多个线程块（block），每一个线程块中又可以包含多个线程（thread）。

在这里我们可以拿古时候的军队作为一个例子来理解这里的程序执行模型。每一个线程，就相当于我们的每一个士兵，在没有当兵之前，大家都不知道自己做什么。当要执行某一个大的军事任务的时候，大将军发布命令，大家来要把对面的敌人部队的n个敌人消灭了。然后把队伍分成M个部分，每一个部分完成自己的工作，有的是做侦查的工作，有的是做诱敌的工作，有的是做伏击的工作，有的是做后备的工作，有的是做后勤的工作……反正把一个大任务按照不同的类别，不同的流程不同，分别由M个部分来完成。

这里我们可以把大将军看着是Host，它把这次军事行动分解成一个一个的kernel：kernel_1，kernel_2……kernel_M，每一个kernel就交给每一个Grid（副将？千户？就看管的人多人少了，如果GPU硬件支持少一点，那就是千户；要是GPU硬件高级一些，管理的人多一些，那就副将？戚家军也不过四五千人，咱也不能太贪心，一下子就想统军百万，再说了，敢问世上韩信一样的将才又哪有那么多啦？）来完成。当要执行这些任务的时候，每一个Grid又把任务分成一部分一部分的，毕竟人太多，他一个官不过来，他只要管理几个团队中间的高级军官就可以了。Grid又把任务划分为一个个的Block（百户？），这里每一个Grid管理的Block也是有限的，(人就那么多……想管多少得看硬件的支持)。毕竟显卡上的GPU硬件还是很少，Thread（线程）相对于真正的军队来说人还是少了很多。所以到Block这个层的时候，就直接管理每一个Thread（士兵）。

由于古代通信不是很方便（从GPU的发展史来看，如果按照中国的历史，现在的GPU也就还处在战国时代吧……），所以每一个Block（百户）内部的Thread（士兵）才能方便的通信，按照既定的规则进行同步；而各个block之间就没那么方便了，大家不能互相通讯。不过同一个（千户）Grid管理的block之间是共享同一个任务分配的资源的。每一个Grid都可以从大将军那里分配到一些任务，和一些粮食，同一个Grid的block都可以分到这个Grid分配到的粮食。而每一个（千户）Grid本身的任务就不一样，所以Grid除了知道自己做的事情外，其他的Grid他都不会知道了。----这差不多就是一个运行模型。下面让我们来看看在GPU中东图例说明：

看到这张图，我们可以对应来讲解我们的Thread部队。一个大将军Host，分配了任务中的两个任务（Kernel1, Kernel2）给了千户（Grid1，Grid2）来完成。千户Grid1里面把自己的队伍分成了6个百户Block，然后每一个百户又把任务分配给了自己的士兵（Thread）来具体完成。这里得说明的是，由于千户拿到的任务Kernel是定了的，所以到每个士兵（Thread）也就那里就只会埋头做同样的事情（就像戚继光招的兵：在胡宗宪的幕僚郑若曾所著的《江南经略》中，有着这样一份详细的招生简章，如果不服气，大可以去对照一下： 凡选入军中之人，以下几等人不可用，在市井里混过的人不能用，喜欢花拳绣腿的人不能用，年纪过四十的人不能用，在政府机关干过的人不能用。 以上尚在其次，更神奇的要求还在下面： 喜欢吹牛、高谈阔论的人不能用，胆子小的人不能用，长得白的人不能用，为保证队伍的心理健康，性格偏激（偏见执拗）的人也不能用。……概括起来，戚继光要找的是这样一群人：四肢发达，头脑简单，为人老实，遵纪守法服从政府，敢打硬仗，敢冲锋不怕死，具备二愣子性格的肌肉男。----《明朝那些事儿》）。

为了方便统一管理，大家都去掉了自己的名字，按照Grid1，Block（x，y），Thread（x，y）这样的编号来称呼每一个Thread士兵。如果你要找到某一个Thread，你就跑到军营里面大叫：喂，Grid1手下的Block1管理的三排第二个Thread（1，2）出来。对于每个士兵自己来说，他要知道自己的位置，就得知道自己的长官都是谁。Thread（1，2）要知道自己再整个Grid手下算第几个兵（钢七年第……个兵），当Grid1叫到他的号了，他得马上回答：我在Block（1，1）的编号是：

unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;

   unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

如果要得到线性的编号，我们可以自己算一下：这个士兵是在Grid部队下的第xIndex行，yIndex列站着如果要得到线性的编号，我们可以自己算一下：这个士兵是在Grid部队下的第xIndex行，yIndex列站着（这里我们必须注意：blockDim.x （这里为5），blockDim.y（这里为3）不是block的坐标，这里是block的size，切记！）。如果从第一个士兵哪里算是线性编号0，那他的线性编号就是

unsigned int index = xIndex（6） + size_x * yIndex（5）; 这里的size_x就是一行一共有多少个士兵（Thread），例如上图，这里一行有3个block每一个block里面的每一行有5个Thread，所以size_x就应该为3×5=15，一个Grid的一行有15个士兵，那刚才叫道的那个人的线性编号就应该是……还要我算吗？如果不得81，自己再算一下……（编号是从0，开始的）计算过程：index = 6+5×15=81；

说了那么多，咱也不能光说不练假把式。下面给一个简单的Thread 测试的Demo。本来打算把整个代码都copy过来，但是考虑到又会被别人copy出去，这样copy来copy去，在编程中很容易出现错误，所以作者也不提倡在做编程中直接copy代码，这样很危险的，很多自己都不知道的bug就隐藏在copy的代码当中……so，截图……

作者这里做了一个简单的测试，测试了512个线程。这里只有一个grid，从这一个grid里面也只分了一个bock：

dim3   grid(size_x / BLOCK_DIM, 1);--》   dim3   grid(1, 1);一个Grid里面一个blcok

dim3   block(BLOCK_DIM, 1, 1);--》dim3 block(512, 1, 1);一个Block里面分配512个Thread；

这里的每一个任务kernel就是：

__global__ static void ThreadDemo1(unsigned int* ret)

{

   unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;

   unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

   if(xIndex < size_x && yIndex < size_y)

   {

       unsigned int index = xIndex + size_x * yIndex;

       ret[index] = xIndex;

       ret[index + size_x*size_y] = yIndex;

   }

}

计算自己的线性id然后把自己的坐标写入到线性id对应的数组里面。Ps：说明一下，这个记录id坐标的数组ret[],ret的前一半记录的是线程的x坐标，后一般是记录的y坐标。PS2：题外话，cu是C的扩展，这里的const定义的常量的用法在ANSIC C里面是行不通的，但是在C++中是可用的。

每个任务kernel都说好了，然后就是host下达命令

   ThreadDemo1<<<grid,block>>>(ret);

来运行程序。所有的士兵都开始工作了，把自己的坐标x，y写入到ret数组里面。

下面是得到的结果：

(0,0) (1,0) (2,0) (3,0) (4,0) (5,0) (6,0) (7,0) (8,0) (9,0) (10,0) (11,0) (12,0) (13,0) (14,0) (15,0) (16,0) (17,0) (18,0) (19,0) (20,0) (21,0) (22,0) (23,0) (24,0) (25,0) (26,0) (27,0) (28,0) (29,0) (30,0) (31,0) (32,0) (33,0) (34,0) (35,0) (36,0) (37,0) (38,0) (39,0) (40,0) (41,0) (42,0) (43,0) (44,0) (45,0) (46,0) (47,0) (48,0) (49,0) (50,0) (51,0) (52,0) (53,0) (54,0) (55,0) (56,0) (57,0) (58,0) (59,0) (60,0) (61,0) (62,0) (63,0) (64,0) (65,0) (66,0) (67,0) (68,0) (69,0) (70,0) (71,0) (72,0) (73,0) (74,0) (75,0) (76,0) (77,0) (78,0) (79,0) (80,0) (81,0) (82,0) (83,0) (84,0) (85,0) (86,0) (87,0) (88,0) (89,0) (90,0) (91,0) (92,0) (93,0) (94,0) (95,0) (96,0) (97,0) (98,0) (99,0) (100,0) (101,0) (102,0) (103,0) (104,0) (105,0) (106,0) (107,0) (108,0) (109,0) (110,0) (111,0) (112,0) (113,0) (114,0) (115,0) (116,0) (117,0) (118,0) (119,0) (120,0) (121,0) (122,0) (123,0) (124,0) (125,0) (126,0) (127,0) (128,0) (129,0) (130,0) (131,0) (132,0) (133,0) (134,0) (135,0) (136,0) (137,0) (138,0) (139,0) (140,0) (141,0) (142,0) (143,0) (144,0) (145,0) (146,0) (147,0) (148,0) (149,0) (150,0) (151,0) (152,0) (153,0) (154,0) (155,0) (156,0) (157,0) (158,0) (159,0) (160,0) (161,0) (162,0) (163,0) (164,0) (165,0) (166,0) (167,0) (168,0) (169,0) (170,0) (171,0) (172,0) (173,0) (174,0) (175,0) (176,0) (177,0) (178,0) (179,0) (180,0) (181,0) (182,0) (183,0) (184,0) (185,0) (186,0) (187,0) (188,0) (189,0) (190,0) (191,0) (192,0) (193,0) (194,0) (195,0) (196,0) (197,0) (198,0) (199,0) (200,0) (201,0) (202,0) (203,0) (204,0) (205,0) (206,0) (207,0) (208,0) (209,0) (210,0) (211,0) (212,0) (213,0) (214,0) (215,0) (216,0) (217,0) (218,0) (219,0) (220,0) (221,0) (222,0) (223,0) (224,0) (225,0) (226,0) (227,0) (228,0) (229,0) (230,0) (231,0) (232,0) (233,0) (234,0) (235,0) (236,0) (237,0) (238,0) (239,0) (240,0) (241,0) (242,0) (243,0) (244,0) (245,0) (246,0) (247,0) (248,0) (249,0) (250,0) (251,0) (252,0) (253,0) (254,0) (255,0) (256,0) (257,0) (258,0) (259,0) (260,0) (261,0) (262,0) (263,0) (264,0) (265,0) (266,0) (267,0) (268,0) (269,0) (270,0) (271,0) (272,0) (273,0) (274,0) (275,0) (276,0) (277,0) (278,0) (279,0) (280,0) (281,0) (282,0) (283,0) (284,0) (285,0) (286,0) (287,0) (288,0) (289,0) (290,0) (291,0) (292,0) (293,0) (294,0) (295,0) (296,0) (297,0) (298,0) (299,0) (300,0) (301,0) (302,0) (303,0) (304,0) (305,0) (306,0) (307,0) (308,0) (309,0) (310,0) (311,0) (312,0) (313,0) (314,0) (315,0) (316,0) (317,0) (318,0) (319,0) (320,0) (321,0) (322,0) (323,0) (324,0) (325,0) (326,0) (327,0) (328,0) (329,0) (330,0) (331,0) (332,0) (333,0) (334,0) (335,0) (336,0) (337,0) (338,0) (339,0) (340,0) (341,0) (342,0) (343,0) (344,0) (345,0) (346,0) (347,0) (348,0) (349,0) (350,0) (351,0) (352,0) (353,0) (354,0) (355,0) (356,0) (357,0) (358,0) (359,0) (360,0) (361,0) (362,0) (363,0) (364,0) (365,0) (366,0) (367,0) (368,0) (369,0) (370,0) (371,0) (372,0) (373,0) (374,0) (375,0) (376,0) (377,0) (378,0) (379,0) (380,0) (381,0) (382,0) (383,0) (384,0) (385,0) (386,0) (387,0) (388,0) (389,0) (390,0) (391,0) (392,0) (393,0) (394,0) (395,0) (396,0) (397,0) (398,0) (399,0) (400,0) (401,0) (402,0) (403,0) (404,0) (405,0) (406,0) (407,0) (408,0) (409,0) (410,0) (411,0) (412,0) (413,0) (414,0) (415,0) (416,0) (417,0) (418,0) (419,0) (420,0) (421,0) (422,0) (423,0) (424,0) (425,0) (426,0) (427,0) (428,0) (429,0) (430,0) (431,0) (432,0) (433,0) (434,0) (435,0) (436,0) (437,0) (438,0) (439,0) (440,0) (441,0) (442,0) (443,0) (444,0) (445,0) (446,0) (447,0) (448,0) (449,0) (450,0) (451,0) (452,0) (453,0) (454,0) (455,0) (456,0) (457,0) (458,0) (459,0) (460,0) (461,0) (462,0) (463,0) (464,0) (465,0) (466,0) (467,0) (468,0) (469,0) (470,0) (471,0) (472,0) (473,0) (474,0) (475,0) (476,0) (477,0) (478,0) (479,0) (480,0) (481,0) (482,0) (483,0) (484,0) (485,0) (486,0) (487,0) (488,0) (489,0) (490,0) (491,0) (492,0) (493,0) (494,0) (495,0) (496,0) (497,0) (498,0) (499,0) (500,0) (501,0) (502,0) (503,0) (504,0) (505,0) (506,0) (507,0) (508,0) (509,0) (510,0) (511,0)

一目了然，正是我们定义的512个线程。

现在大家应该知道CUDA是怎么运行了？0 – 0 要是还不清楚，跟贴问我吧 - - ！

好了，现在招兵的任务已经完成了，兵是招来了，下面就改训练了，不过天时已晚，大家还是吃饱喝足，好好的休息一下吧～

明天继续GPU的革命。

且听：《沙家浜》里胡傳魁已經出場了：十幾個人來七八條槍，不是很威風嗎？----我可要说十几个人怎么来分这七八条枪。

士兵们执行任务了，是一哄而上，还是……且听SIMD- -！

PS:上面在输出结果的时候用了一个小技巧，或许有的人知道，就不用看了，有不知道的，接着看下去，或许可以方便你以后的调试：）

再输出的时候，在系统debugging信息里面的command Arguments里面用一个管道输出到一个1.txt文件里面。管道用法大家自己可以到网上查～我就不用讲解了。也很简单，就是在你的哦xx.exe 后面加一个 “>” 管道，对了还有 “<” 管道，呵呵：） 一个是输出，一个是输入。

一、   提要
以前在《 Win32 OpenGL 编程（1）Win32下的OpenGL编程必须步骤》一文中提供了一个较为完整的Win32下编写OpenGL程序的框架，但是仅仅提到的是此Win32框架所需要的一些东西，事实上由于MS为了维护自家的DirectX，所以很早就放弃了对OpenGL的支持，在Windows XP平台上实现的OpenGL还是1.1版(新版的Windows看资料好像会好一点，但是事实上OpenGL最新版已经是3.2了)。本文讲述的就是使用本机显卡支持的最新的OpenGL所需要的东西，顺便还讲了一下GLUT，这样基本形成一个较为完整的Win32下的OpenGL变成学习环境。由于工作很简单，所以本文会很短，最长的可能就是提要^^

另外，事实上参考资料4中的文章已经较为详细的介绍了本文所要讲的大部分内容，大家可以直接去看，写作此文的目的主要还是希望此系列完整，在后面的文章中已经有个完整可用的编程环境。顺面补上glew和glut简单的安装使用说明。

二、   寻找缺失的OpenGL函数
OpenGL是由显卡直接支持并实现的，当时在设计的时候考虑到显卡厂商可能会比操作系统/开发软件制造商先行一步提供新的扩展，所以有提供获取新的扩展的方法，在Windows下就是利用wglGetProcAddress函数：

PROC wglGetProcAddress(

  LPCSTR  lpszProc   // name of the extension function

);

此函数与GetProcAddress功能类似，实现也类似，无非就是从动态库中通过名字直接获取到函数的地址。

FARPROC WINAPI GetProcAddress(

  __in          HMODULE hModule,

  __in          LPCSTR lpProcName

);

因为wglGetProcAddress很明确的知道该从哪个Dll获取函数，所以省略了第一参数模块名，仅仅通过函数名就能获取到对应的函数地址，使用方法也是类似，通过函数名获取到函数地址后，通过强制转换的方式使用。此方法虽然可行，但是用到大量函数的时候比较麻烦。事实上，此方法提供的时候是作为一种补充的手段，是在显卡生厂商走在操作系统和开发工具提供商前面时的一种暂时应付手段，但是，很不幸的是，Windows是在落后太多了。。。。。。。。

三、   Windows下使用OpenGL的救星
MS当年作为OpenGL的发起人之一，仅仅对OpenGL支持到1.1，然后就退出了OpenGL组织，自己发展DirectX系统去了，我们要在Windows上使用OpenGL光靠MS的话，那么OpenGL 1.1后的所有函数都得通过上述手段来完成，那简直是让人吐血的事情-_-!

以前似乎有关官方（SGI）的OpenGL SDK下载，但是现在没有那么好了，OpenGL的官方网站已经明确说明，不再提供SDK下载，叫咱们自己到其他网站上找去-_-!在OpenGL.org上的OpenGL SDK中的Library为我们提供了几个开源项目，其中有两个就是与扩展有关的。

GLee is a free cross-platform extension loading library that takes the burden off your application. GLee makes it easy to check for OpenGL extension and core version availability, automatically setting up the entry points with no effort on your part.

GLEW is an open-source cross-platform extension loading library with thread-safe support for multiple rendering contexts and automatic code generation capability. GLEW provides easy-to-use and efficient methods for checking OpenGL extensions and core functionality.

我选择了GLEW，这个有线程安全支持的项目，大家各取所需吧。不是开源社区的努力，要在Windows下使用OpenGL还真是让人吐血了-_-!说是救星也不为过。

1.      glew的安装
上面的链接地址有下载，glew有提供Windows的二进制版本，使用上可以直接在工程中包含头文件，库，然后在程序的运行目录包含动态库，因为考虑到别人不一定下载了glew，按道理来说不论是源代码分发还是可执行文件的分发最好都是利用此方式使用。但是作为学习来说，要在每个工程中都做这么多额外的工作比较让人受不了，所以推荐讲头文件放在windows SDK的include/GL目录下（可以通过找到gl/gl.h的位置找到），lib放在windows SDK的lib目录下，动态库放到windows/system32这个系统目录下，那么以后的OpenGL工程就不需要额外的配置了，以后文章的工程都是在这个环境下建立的。。。。要运行程序要编译代码请首先搭建好环境。。。。。真是方便了我一个，辛苦了大家了^^-_-!

2.      glew的使用
以后在使用OpenGL的时候，都请通过下述代码

#include <GL/glew.h>

#include <GL/wglew.h>

包含必要的头文件

通过

#pragma comment( lib, "glew32.lib" ) 

#pragma comment( lib, "glew32s.lib" ) 

然后，在初始化OpenGL的时候，不要忘了将glew初始化，可以这样：

       GLenum err = glewInit();

       if(err != GLEW_OK)

       {

              exit(-2);

       }

这样就能使用缺失的函数，就像它们原来就被操作系统支持那样，而不用自己通过wglGetProcAddress去获取了，说是能像原来就被操作系统支持的那样使用其实不完成正确，起码在VS中，由于都是通过强转后的函数指针使用，VS和VA都不能正确的解析其函数原型，使得代码编写时智能提醒函数原型的功能不能正常使用，这也算是个小瑕疵吧。。。。。。。。。

四、   为了红宝书《OpenGL Programming Guide》安装glut
事实上，在《 Win32 OpenGL 编程（1）Win32下的OpenGL编程必须步骤》中我就提到了不想使用glut的窗口管理系统，但是想完整的学习《OpenGL Programming Guide》要完全避过glut好像比较难，就算不想使用其窗口管理系统，但是一些简单易用的快捷函数要是不用，那就太麻烦了。。。。。我想画个立方体都得先去想办法才行-_-!茶壶那就更。。。。。无语了。这里将glut的安装使用顺面也讲一下吧。

glut在http://www.opengl.org/resources/libraries/glut/ 上有介绍，下载页是：http://www.opengl.org/resources/libraries/glut/glut_downloads.php

也有Windows的二进制版本下载，安装方法同glew,拷贝头文件，库，动态库到相应的位置，使用方法嘛。。。见红宝书-_-!呵呵，其实glut包含了平台相关的头文件和gl.h,glu.h，所以在使用的时候包含glut.h一个文件就行了。

包含了glut后，基本上Win32下OpenGL的工程可以以下面的代码开始了：

#include <GL/glew.h>

#include <GL/wglew.h>

#include <GL/glut.h>

#pragma comment( lib, "opengl32.lib" )

#pragma comment( lib, "glu32.lib" ) 

#pragma comment( lib, "glew32.lib" ) 

#pragma comment( lib, "glew32s.lib" ) 

#pragma comment( lib, "glut32.lib" ) 

这些都是必须的。（事实上由于glut32.lib会自动的包含OpenGL核心所需的库，包含glut32.lib时可以省略opengl32.lib和glu32.lib）

五、   完整的示例
这里提供一个使用了glew，glut的完整示例，这个示例在上述两个库都安装完成工程配置完成的时候才能正确运行，也可以用此示例来检验安装，完整的示例从《OpenGL 编程指南》（《OpenGL Programming Guide》）的示例代码mvarray.c改过来的。其中的宏GL_VERSION_1_3是由glew定义用来表示支持OpenGL1.3版本，其中的函数glMultiDrawElementsEXT就是Windows本身不支持，用了glew后才能使用的OpenGL扩展接口（首先要显卡支持啊）。glut的使用那就更不用说了。完整的示例在博客示例代码的2009-9-29\TestOpenGL目录下可以找到，下载方式见本文最后的说明。

/*

 * Copyright (c) 1993-2003, Silicon Graphics, Inc.

 * All Rights Reserved

 *

 * Permission to use, copy, modify, and distribute this software for any

 * purpose and without fee is hereby granted, provided that the above

 * copyright notice appear in all copies and that both the copyright

 * notice and this permission notice appear in supporting documentation,

 * and that the name of Silicon Graphics, Inc. not be used in

 * advertising or publicity pertaining to distribution of the software

 * without specific, written prior permission.

 *

 * THE MATERIAL EMBODIED ON THIS SOFTWARE IS PROVIDED TO YOU "AS-IS" AND

 * WITHOUT WARRANTY OF ANY KIND, EXPRESS, IMPLIED OR OTHERWISE,

 * INCLUDING WITHOUT LIMITATION, ANY WARRANTY OF MERCHANTABILITY OR

 * FITNESS FOR A PARTICULAR PURPOSE.  IN NO EVENT SHALL SILICON

 * GRAPHICS, INC.  BE LIABLE TO YOU OR ANYONE ELSE FOR ANY DIRECT,

 * SPECIAL, INCIDENTAL, INDIRECT OR CONSEQUENTIAL DAMAGES OF ANY KIND,

 * OR ANY DAMAGES WHATSOEVER, INCLUDING WITHOUT LIMITATION, LOSS OF

 * PROFIT, LOSS OF USE, SAVINGS OR REVENUE, OR THE CLAIMS OF THIRD

 * PARTIES, WHETHER OR NOT SILICON GRAPHICS, INC.  HAS BEEN ADVISED OF

 * THE POSSIBILITY OF SUCH LOSS, HOWEVER CAUSED AND ON ANY THEORY OF

 * LIABILITY, ARISING OUT OF OR IN CONNECTION WITH THE POSSESSION, USE

 * OR PERFORMANCE OF THIS SOFTWARE.

 *

 * US Government Users Restricted Rights

 * Use, duplication, or disclosure by the Government is subject to

 * restrictions set forth in FAR 52.227.19(c)(2) or subparagraph

 * (c)(1)(ii) of the Rights in Technical Data and Computer Software

 * clause at DFARS 252.227-7013 and/or in similar or successor clauses

 * in the FAR or the DOD or NASA FAR Supplement.  Unpublished - rights

 * reserved under the copyright laws of the United States.

 *

 * Contractor/manufacturer is:

 *    Silicon Graphics, Inc.

 *    1500 Crittenden Lane

 *    Mountain View, CA  94043

 *    United State of America

 *

 * OpenGL(R) is a registered trademark of Silicon Graphics, Inc.

 */

/*

 *  mvarray.c

 *  This program demonstrates multiple vertex arrays,

 *  specifically the OpenGL routine glMultiDrawElements().

 */

#include <GL/glew.h>

#include <GL/wglew.h>

#include <GL/glut.h>

#include <stdio.h>

#pragma comment( lib, "opengl32.lib" )

#pragma comment( lib, "glu32.lib" ) 

#pragma comment( lib, "glew32.lib" ) 

#pragma comment( lib, "glew32s.lib" ) 

#pragma comment( lib, "glut32.lib" ) 

#ifdef GL_VERSION_1_3

void setupPointer(void)

{

   static GLint vertices[] = {25, 25,

                       75, 75,

                       100, 125,

                       150, 75,

                       200, 175,

                       250, 150,

                       300, 125,

                       100, 200,

                       150, 250,

                       200, 225,

                       250, 300,

                       300, 250};

   glEnableClientState (GL_VERTEX_ARRAY);

   glVertexPointer (2, GL_INT, 0, vertices);

}

void init(void)

{

       GLenum err = glewInit();

       if(err != GLEW_OK)

       {

              exit(-2);

       }

   glClearColor (0.0, 0.0, 0.0, 0.0);

   glShadeModel (GL_SMOOTH);

   setupPointer ();

}

void display(void)

{

   static GLubyte oneIndices[] = {0, 1, 2, 3, 4, 5, 6};

   static const GLubyte twoIndices[] = {1, 7, 8, 9, 10, 11};

   static GLsizei count[] = {7, 6};

   static const GLvoid * indices[2] = {oneIndices, twoIndices};

   glClear (GL_COLOR_BUFFER_BIT);

   glColor3f (1.0, 1.0, 1.0);

   glMultiDrawElementsEXT (GL_LINE_STRIP, count, GL_UNSIGNED_BYTE,

                           indices, 2);

   // 上面的一句相当于下面的两句

   //glDrawElements(GL_LINE_STRIP, count[0], GL_UNSIGNED_BYTE, oneIndices);

   //glDrawElements(GL_LINE_STRIP, count[1], GL_UNSIGNED_BYTE, twoIndices);

   glFlush ();

}

void reshape (int w, int h)

{

   glViewport (0, 0, (GLsizei) w, (GLsizei) h);

   glMatrixMode (GL_PROJECTION);

   glLoadIdentity ();

   gluOrtho2D (0.0, (GLdouble) w, 0.0, (GLdouble) h);

}

void keyboard(unsigned char key, int x, int y)

{

   switch (key) {

      case 27:

         exit(0);

         break;

   }

}

int main(int argc, char** argv)

{

   glutInit(&argc, argv);

   glutInitDisplayMode (GLUT_SINGLE | GLUT_RGB);

   glutInitWindowSize (350, 350);

   glutInitWindowPosition (100, 100);

   glutCreateWindow (argv[0]);

   init ();

   glutDisplayFunc(display);

   glutReshapeFunc(reshape);

   glutKeyboardFunc (keyboard);

   glutMainLoop();

   return 0;

}

#else

int main(int argc, char** argv)

{

    fprintf (stderr, "This program demonstrates a feature which is not in OpenGL Version 1.0.\n");

    fprintf (stderr, "If your implementation of OpenGL Version 1.0 has the right extensions,\n");

    fprintf (stderr, "you may be able to modify this program to make it run.\n");

    return 0;

}

#endif

六、   参考资料
1.       《OpenGL Reference Manual》，OpenGL参考手册

2.       《OpenGL 编程指南》（《OpenGL Programming Guide》），Dave Shreiner，Mason Woo，Jackie Neider，Tom Davis 著，徐波译，机械工业出版社

3.       MSDN，事实上我收回MSDN上有完整的OpenGL参考的话，应该是说MSDN上有OpenGL1.1前版本的较完整资料。

4.       《如何使用OpenGL扩展》，曹添著，已经包含了本文的大部分内容

七、   最后说明
本文中所有代码（如果有的话）都能用Mercurial在Google Code中下载。

文章以博文发表的日期分目录存放，下载地址为：

http://code.google.com/p/jtianling/source/checkout?repo=blog-sample-code

Mercurial使用方法见《分布式的，新一代版本控制系统Mercurial的介绍及简要入门》

本文来自CSDN博客，转载请标明出处：http://blog.csdn.net/vagrxie/archive/2009/09/29/4617556.aspx

http://glew.sourceforg.net

实在不知道Wrangler 应该在这里当什么意思讲，不过他已经被包含到OpenGL SDK了，官方是推荐使用它的。

GLEW有什么用？你可以把它看作是windows上的OpenGL的基本库，他取代了原来的gl.h wgl.h，让你可以方便的调用最新的OpenGL功能，包括OpenGL众多的Extension。

实际上我对这个OpenGL的体系也是稀里糊涂，除了核心的概念和功能，对扩展功能基本上是抓瞎的。这就是为什么我要读GLEW的源码。读源码似乎很傻，不过总比等别人出书要聪明的多。

开始吧。

基本特性：

• 跨操作系统
• 支持动态链接、静态库，你还可以把glew.c复制到工程中一起编译
• 只依赖于OPENGL32.dll和KERNEL32.dll
• 因为不依赖于CRT所有你使用single-threaded, multi-threaded 或者 multi-threaded DLL都没有限制

文件结构：

发布版本包括一个glew32.dll和对应的glew32.lib，头文件glew.h和wglew.h。（windows dll版本）。源文件只有一个glew.c。

glew基本使用：

初始化

#include 放在一切和gl有关的头文件之前，就完成了引用

GLenum err = glewInit();  如果err为真就完成了初始化，应该放在程序的开始部分。现在还是不明白为什么有操作可以早于初始化执行。

检查扩展

if (GLEW_ARB_vertex_program)
{
  /* 通过了ARB_vertex_program extension 的检查. */
  glGenProgramsARB(...);
}

if (GLEW_VERSION_1_3)
{
  /* OpenGL 1.3 是被支持的 */
}

if (glewIsSupported("GL_VERSION_1_4  GL_ARB_point_sprite"))
{
  /* 两个都通过了检查 */
}

if (glewGetExtension("GL_ARB_fragment_program"))
{
  /* 检查ARB_fragment_program 是被支持的. */
}

GLEW为了实现扩展检查的功能，编制了大量的代码。简单的看就是即支持宏标签又支持字符串查询。

glew.h源码结构：

开头就检测有没有和GLEW冲突的库

#if defined(__gl_h_) || defined(__GL_H__)
#error gl.h included before glew.h
#endif
#if defined(__glext_h_) || defined(__GLEXT_H_)
#error glext.h included before glew.h
#endif
#if defined(__gl_ATI_h_)
#error glATI.h included before glew.h
#endif

这些文件功能上和GLEW是重复的，所有开始就屏蔽掉了。

接下来的一堆宏只是为了正确的定义 GLEWAPI GLAPI GLAPIENTRY CALLBACK 等等宏。这些定义本来已经在其他的头文件中定义了，比如WINGDIAPI就在winnt.h中定义过了。这里画蛇添足的做法只是为了保证GLEW的纯洁性，为了不依赖任何不必要的库和头文件。

再下面就是OpenGL1.1的定义了。这里很简单就是类型定义和宏定义，API使用外部函数方式，这就是为什么不需要gl.h了。

#define GLEW_VERSION_1_1 GLEW_GET_VAR(__GLEW_VERSION_1_1)

唯一让人有点不解的是这句，解析了所有的宏之后GLEW_VERSION_1_1被定义为一个GLboolean类型的变量。之后就让人很烦了__GLEW_VERSION_1_1被初始化为GL_FALSE，然后作者又在检测到支持1.1之后把GLEW_VERSION_1_1赋值为真。宏来宏去的让人郁闷。不过想想也是有道理的，GLEW_VERSION_1_1作为检测标志必须是个宏，而宏是不能修改值的，既然他可能为真也可能为假，那就只有一个办法，找一个变量作为他的存储空间，然后这个宏就可以当变量使用了。

再下面#include

再下面是OpenGL1.2的定义，和1.1的定义方式类似，只是函数定义方式不同了。

typedef void (GLAPIENTRY * PFNGLCOPYTEXSUBIMAGE3DPROC) (GLenum target, GLint level, GLint xoffset, GLint yoffset, GLint zoffset, GLint x, GLint y, GLsizei width, GLsizei height);
typedef void (GLAPIENTRY * PFNGLDRAWRANGEELEMENTSPROC) (GLenum mode, GLuint start, GLuint end, GLsizei count, GLenum type, const GLvoid *indices);
typedef void (GLAPIENTRY * PFNGLTEXIMAGE3DPROC) (GLenum target, GLint level, GLint internalFormat, GLsizei width, GLsizei height, GLsizei depth, GLint border, GLenum format, GLenum type, const GLvoid *pixels);
typedef void (GLAPIENTRY * PFNGLTEXSUBIMAGE3DPROC) (GLenum target, GLint level, GLint xoffset, GLint yoffset, GLint zoffset, GLsizei width, GLsizei height, GLsizei depth, GLenum format, GLenum type, const GLvoid *pixels);

#define glCopyTexSubImage3D GLEW_GET_FUN(__glewCopyTexSubImage3D)
#define glDrawRangeElements GLEW_GET_FUN(__glewDrawRangeElements)
#define glTexImage3D GLEW_GET_FUN(__glewTexImage3D)
#define glTexSubImage3D GLEW_GET_FUN(__glewTexSubImage3D)

4个函数指针定义对应4个define，看名字就知道他们是一一对应的，但是就是对不上，郁闷吧。但是在文件最后的部分能找到这样的代码。

GLEW_FUN_EXPORT PFNGLCOPYTEXSUBIMAGE3DPROC __glewCopyTexSubImage3D;
GLEW_FUN_EXPORT PFNGLDRAWRANGEELEMENTSPROC __glewDrawRangeElements;
GLEW_FUN_EXPORT PFNGLTEXIMAGE3DPROC __glewTexImage3D;
GLEW_FUN_EXPORT PFNGLTEXSUBIMAGE3DPROC __glewTexSubImage3D;

也是4个，这下对上号了吧。PFNGLCOPYTEXSUBIMAGE3DPROC 是一个函数指针类型定义 __glewCopyTexSubImage3D是实际被定义的指针函数，#define glCopyTexSubImage3D GLEW_GET_FUN(__glewCopyTexSubImage3D)又指向了glCopyTexSubImage3D 。到最后 glCopyTexSubImage3D  才是真正的API接口。真费劲。

再往下1.3 1.4 1.5 2.0 2.1 3.0结构完全和1.2一致

再往下就是各种扩展了，形式上还是和1.2定义的方式一致

再往下就是GLEW_FUN_EXPORT PFNGLCOPYTEXSUBIMAGE3DPROC __glewCopyTexSubImage3D;等等了。

再往下是GLEW_VAR_EXPORT GLboolean __GLEW_VERSION_1_1;等等。

最后才是GLEW的主要函数接口。这里支持一个GLEW_MX的定义，就是对GLEW MX的支持。

glew.c源码结构：

相对简单的多，除了一些函数的实现以外，全是初始化函数。

全部都是从glewInit开始的，一路下去初始化各个版本和所有扩展。

wglew

看明白上面的，这里就不用说了，比葫芦看瓢吧。

http://glew.sourceforg.net

实在不知道Wrangler 应该在这里当什么意思讲，不过他已经被包含到OpenGL SDK了，官方是推荐使用它的。

GLEW有什么用？你可以把它看作是windows上的OpenGL的基本库，他取代了原来的gl.h wgl.h，让你可以方便的调用最新的OpenGL功能，包括OpenGL众多的Extension。

实际上我对这个OpenGL的体系也是稀里糊涂，除了核心的概念和功能，对扩展功能基本上是抓瞎的。这就是为什么我要读GLEW的源码。读源码似乎很傻，不过总比等别人出书要聪明的多。

开始吧。

基本特性：

• 跨操作系统
• 支持动态链接、静态库，你还可以把glew.c复制到工程中一起编译
• 只依赖于OPENGL32.dll和KERNEL32.dll
• 因为不依赖于CRT所有你使用single-threaded, multi-threaded 或者 multi-threaded DLL都没有限制

文件结构：

发布版本包括一个glew32.dll和对应的glew32.lib，头文件glew.h和wglew.h。（windows dll版本）。源文件只有一个glew.c。

glew基本使用：

初始化

#include <GL/glew.h>放在一切和gl有关的头文件之前，就完成了引用

GLenum err = glewInit();  如果err为真就完成了初始化，应该放在程序的开始部分。现在还是不明白为什么有操作可以早于初始化执行。

检查扩展

if (GLEW_ARB_vertex_program)
{
  /* 通过了ARB_vertex_program extension 的检查. */
  glGenProgramsARB(...);
}

if (GLEW_VERSION_1_3)
{
  /* OpenGL 1.3 是被支持的 */
}

if (glewIsSupported("GL_VERSION_1_4  GL_ARB_point_sprite"))
{
  /* 两个都通过了检查 */
}

if (glewGetExtension("GL_ARB_fragment_program"))
{
  /* 检查ARB_fragment_program 是被支持的. */
}

GLEW为了实现扩展检查的功能，编制了大量的代码。简单的看就是即支持宏标签又支持字符串查询。

glew.h源码结构：

开头就检测有没有和GLEW冲突的库

#if defined(__gl_h_) || defined(__GL_H__)
#error gl.h included before glew.h
#endif
#if defined(__glext_h_) || defined(__GLEXT_H_)
#error glext.h included before glew.h
#endif
#if defined(__gl_ATI_h_)
#error glATI.h included before glew.h
#endif

这些文件功能上和GLEW是重复的，所有开始就屏蔽掉了。

接下来的一堆宏只是为了正确的定义 GLEWAPI GLAPI GLAPIENTRY CALLBACK 等等宏。这些定义本来已经在其他的头文件中定义了，比如WINGDIAPI就在winnt.h中定义过了。这里画蛇添足的做法只是为了保证GLEW的纯洁性，为了不依赖任何不必要的库和头文件。

再下面就是OpenGL1.1的定义了。这里很简单就是类型定义和宏定义，API使用外部函数方式，这就是为什么不需要gl.h了。

#define GLEW_VERSION_1_1 GLEW_GET_VAR(__GLEW_VERSION_1_1)

唯一让人有点不解的是这句，解析了所有的宏之后GLEW_VERSION_1_1被定义为一个GLboolean类型的变量。之后就让人很烦了__GLEW_VERSION_1_1被初始化为GL_FALSE，然后作者又在检测到支持1.1之后把GLEW_VERSION_1_1赋值为真。宏来宏去的让人郁闷。不过想想也是有道理的，GLEW_VERSION_1_1作为检测标志必须是个宏，而宏是不能修改值的，既然他可能为真也可能为假，那就只有一个办法，找一个变量作为他的存储空间，然后这个宏就可以当变量使用了。

再下面#include <GL/glu.h>

再下面是OpenGL1.2的定义，和1.1的定义方式类似，只是函数定义方式不同了。

typedef void (GLAPIENTRY * PFNGLCOPYTEXSUBIMAGE3DPROC) (GLenum target, GLint level, GLint xoffset, GLint yoffset, GLint zoffset, GLint x, GLint y, GLsizei width, GLsizei height);
typedef void (GLAPIENTRY * PFNGLDRAWRANGEELEMENTSPROC) (GLenum mode, GLuint start, GLuint end, GLsizei count, GLenum type, const GLvoid *indices);
typedef void (GLAPIENTRY * PFNGLTEXIMAGE3DPROC) (GLenum target, GLint level, GLint internalFormat, GLsizei width, GLsizei height, GLsizei depth, GLint border, GLenum format, GLenum type, const GLvoid *pixels);
typedef void (GLAPIENTRY * PFNGLTEXSUBIMAGE3DPROC) (GLenum target, GLint level, GLint xoffset, GLint yoffset, GLint zoffset, GLsizei width, GLsizei height, GLsizei depth, GLenum format, GLenum type, const GLvoid *pixels);

#define glCopyTexSubImage3D GLEW_GET_FUN(__glewCopyTexSubImage3D)
#define glDrawRangeElements GLEW_GET_FUN(__glewDrawRangeElements)
#define glTexImage3D GLEW_GET_FUN(__glewTexImage3D)
#define glTexSubImage3D GLEW_GET_FUN(__glewTexSubImage3D)

4个函数指针定义对应4个define，看名字就知道他们是一一对应的，但是就是对不上，郁闷吧。但是在文件最后的部分能找到这样的代码。

GLEW_FUN_EXPORT PFNGLCOPYTEXSUBIMAGE3DPROC __glewCopyTexSubImage3D;
GLEW_FUN_EXPORT PFNGLDRAWRANGEELEMENTSPROC __glewDrawRangeElements;
GLEW_FUN_EXPORT PFNGLTEXIMAGE3DPROC __glewTexImage3D;
GLEW_FUN_EXPORT PFNGLTEXSUBIMAGE3DPROC __glewTexSubImage3D;

也是4个，这下对上号了吧。PFNGLCOPYTEXSUBIMAGE3DPROC 是一个函数指针类型定义 __glewCopyTexSubImage3D是实际被定义的指针函数，#define glCopyTexSubImage3D GLEW_GET_FUN(__glewCopyTexSubImage3D)又指向了glCopyTexSubImage3D 。到最后 glCopyTexSubImage3D  才是真正的API接口。真费劲。

再往下1.3 1.4 1.5 2.0 2.1 3.0结构完全和1.2一致

再往下就是各种扩展了，形式上还是和1.2定义的方式一致

再往下就是GLEW_FUN_EXPORT PFNGLCOPYTEXSUBIMAGE3DPROC __glewCopyTexSubImage3D;等等了。

再往下是GLEW_VAR_EXPORT GLboolean __GLEW_VERSION_1_1;等等。

最后才是GLEW的主要函数接口。这里支持一个GLEW_MX的定义，就是对GLEW MX的支持。

glew.c源码结构：

相对简单的多，除了一些函数的实现以外，全是初始化函数。

全部都是从glewInit开始的，一路下去初始化各个版本和所有扩展。

wglew

看明白上面的，这里就不用说了，比葫芦看瓢吧。

开发基于OpenGL的应用程序，必须先了解OpenGL的库函数。它采用C语言风格，提供大量的函数来进行图形的处理和显示。OpenGL库函数的命名方式非常有规律。所有OpenGL函数采用了以下格式

<库前缀><根命令><可选的参数个数><可选的参数类型>

库前缀有gl、glu、aux、glut、wgl、glx、agl等等，分别表示该函数属于OpenGL那个开发库等，从函数名后面中还可以看出需要多少个参数以及参数的类型。I代表int型，f代表float型，d代表double型，u代表无符号整型。例如glVertex3fv()表示了该函数属于gl库，参数是三个float型参数指针。我们用glVertex*()来表示这一类函数。

OpenGL函数库相关的API有核心库(gl)、实用库(glu)、辅助库(aux)、实用工具库(glut)、窗口库(glx、agl、wgl)和扩展函数库等。从图1可以看出，gl是核心，glu是对gl的部分封装。glx、agl、wgl 是针对不同窗口系统的函数。glut是为跨平台的OpenGL程序的工具包，比aux功能强大。扩展函数库是硬件厂商为实现硬件更新利用OpenGL的扩展机制开发的函数。下面逐一对这些库进行详细介绍。

1．  OpenGL核心库

核心库包含有115个函数，函数名的前缀为gl。

这部分函数用于常规的、核心的图形处理。此函数由gl.dll来负责解释执行。由于许多函数可以接收不同数以下几类。据类型的参数，因此派生出来的函数原形多达300多个。核心库中的函数主要可以分为以下几类函数。

绘制基本几何图元的函数。如绘制图元的函数glBegain()、glEnd()、glNormal*()、glVertex*()。

矩阵操作、几何变换和投影变换的函数。如矩阵入栈函数glPushMatrix()、矩阵出栈 函数glPopMatrix()、装载矩阵函数glLoadMatrix()、矩阵相乘函数glMultMatrix()，当前矩阵函数glMatrixMode()和矩阵标准化函数glLoadIdentity()，几何变换函数glTranslate*()、glRotate*()和glScale*()，投影变换函数glOrtho()、glFrustum()和视口变换函数glViewport()等等。

颜色、光照和材质的函数。如设置颜色模式函数glColor*()、glIndex*()，设置光照效果的函数glLight*() 、glLightModel*()和设置材质效果函数glMaterial()等等。

显示列表函数、主要有创建、结束、生成、删除和调用显示列表的函数glNewList()、 glEndList()、glGenLists()、glCallList()和glDeleteLists()。

纹理映射函数，主要有一维纹理函数glTexImage1D()、二维纹理函数glTexImage2D()、 设置纹理参数、纹理环境和纹理坐标的函数glTexParameter*()、glTexEnv*()和glTetCoord*()等。

特殊效果函数。融合函数glBlendFunc()、反走样函数glHint()和雾化效果glFog*()。

光栅化、象素操作函数。如象素位置glRasterPos*()、线型宽度glLineWidth()、多边形绘制模式glPolygonMode()，读取象素glReadPixel()、复制象素glCopyPixel()等。

选择与反馈函数。主要有渲染模式glRenderMode()、选择缓冲区glSelectBuffer()和反馈缓冲区glFeedbackBuffer()等。

曲线与曲面的绘制函数。生成曲线或曲面的函数glMap*()、glMapGrid*()，求值器的函数glEvalCoord*() glEvalMesh*()。

状态设置与查询函数。主要有glGet*()、glEnable()、glGetError()等。

2．  OpenGL实用库The OpenGL Utility Library (GLU)

包含有43个函数，函数名的前缀为glu。

OpenGL提供了强大的但是为数不多的绘图命令，所有较复杂的绘图都必须从点。线、面开始。Glu 为了减轻繁重的编程工作，封装了OpenGL函数，Glu函数通过调用核心库的函数，为开发者提供相对简单的用法，实现一些较为复杂的操作。此函数由glu.dll来负责解释执行。OpenGL中的核心库和实用库可以在所有的OpenGL平台上运行。主要包括了以下几种。

辅助纹理贴图函数，有gluScaleImage() 、gluBuild1Dmipmaps()、gluBuild2Dmipmaps()。

坐标转换和投影变换函数，定义投影方式函数gluPerspective()、gluOrtho2D() 、gluLookAt()，拾取投影视景体函数gluPickMatrix()，投影矩阵计算gluProject()和 gluUnProject()等等。

多边形镶嵌工具，有gluNewTess()、 gluDeleteTess()、gluTessCallback()、gluBeginPolygon() gluTessVertex()、gluNextContour()、gluEndPolygon()等等。

二次曲面绘制工具，主要有绘制球面、锥面、柱面、圆环面gluNewQuadric()、gluSphere()、gluCylinder()、gluDisk()、gluPartialDisk()、gluDeleteQuadric()等等。

非均匀有理B样条绘制工具，主要用来定义和绘制Nurbs曲线和曲面，包括gluNewNurbsRenderer()、gluNurbsCurve()、gluBeginSurface()、gluEndSurface()、gluBeginCurve()、gluNurbsProperty()等函数。

错误反馈工具，获取出错信息的字符串gluErrorString().

3．  OpenGL辅助库

包含有31个函数，函数名前缀为aux。

这部分函数提供窗口管理、输入输出处理以及绘制一些简单三维物体。此函数由glaux.dll来负责解释执行。创建aux库是为了学习和编写OpenGL程序，它更像是一个用于测试创意的预备基础接管。Aux库在windows实现有很多错误，因此很容易导致频繁的崩溃。在跨平台的编程实例和演示中，aux很大程度上已经被glut库取代。OpenGL中的辅助库不能在所有的OpenGL平台上运行。

辅助库函数主要包括以下几类。

窗口初始化和退出函数，auxInitDisplayMode()和auxInitPosition()。

窗口处理和时间输入函数，auxReshapeFunc()、auxKeyFunc()和auxMouseFunc()。

颜色索引装入函数，auxSetOneColor()。

三维物体绘制函数。包括了两种形式网状体和实心体，如绘制立方体auxWireCube()和 auxSolidCube()。这里以网状体为例，长方体auxWireBox()、环形圆纹面auxWireTorus()、圆柱auxWireCylinder()、二十面体auxWireIcosahedron()、八面体auxWireOctahedron()、四面体auxWireTetrahedron()、十二面体auxWireDodecahedron()、圆锥体auxWireCone()和茶壶auxWireTeapot()。

背景过程管理函数auxIdleFunc()。

程序运行函数auxMainLoop()。

4．  OpenGL工具库 OpenGL Utility Toolkit

包含大约30多个函数，函数名前缀为glut。

glut是不依赖于窗口平台的OpenGL工具包，由Mark KLilgrad在SGI编写（现在在Nvidia），目的是隐藏不同窗口平台API的复杂度。 函数以glut开头，它们作为aux库功能更强的替代品，提供更为复杂的绘制功能，此函数由glut.dll来负责解释执行。由于glut中的窗口管理函数是不依赖于运行环境的，因此OpenGL中的工具库可以在X-Window, Windows NT, OS/2等系统下运行，特别适合于开发不需要复杂界面的OpenGL示例程序。对于有经验的程序员来说，一般先用glut理顺3D图形代码，然后再集成为完整的应用程序。

这部分函数主要包括

窗口操作函数，窗口初始化、窗口大小、窗口位置等函数glutInit() glutInitDisplayMode() glutInitWindowSize() glutInitWindowPosition()等。

回调函数。响应刷新消息、键盘消息、鼠标消息、定时器函数等，GlutDisplayFunc() glutPostRedisplay() glutReshapeFunc() glutTimerFunc() glutKeyboardFunc() glutMouseFunc()。

创建复杂的三维物体。这些和aux库的函数功能相同。创建网状体和实心体。如glutSolidSphere()、glutWireSphere()等。在此不再叙述。

菜单函数。创建添加菜单的函数GlutCreateMenu()、glutSetMenu()、glutAddMenuEntry()、glutAddSubMenu() 和glutAttachMenu()。

程序运行函数，glutMainLoop()。

5．  Windows专用库

针对windows平台的扩展。包含有16个函数，函数名前缀为wgl。

这部分函数主要用于连接OpenGL和Windows ，以弥补OpenGL在文本方面的不足。  Windows专用库只能用于Windows环境中。

这类函数主要包括以下几类

绘图上下文相关函数wglCreateContext(), wglDeleteContext()wglGetCurrentContent() wglGetCurrentDC() wglDeleteContent()等

文字和文本处理函数 wglUseFontBitmaps()、wglUseFontOutlines()。

覆盖层、地层和主平面层处理函数 wglCopyContext()、wglCreateLayerPlane()、  wglDescribeLayerPlane()、wglReakizeLayerPlatte()等

其他函数，wglShareLists()、wglGetProcAddress()等。

6．  Win32 API函数库

 包含有6个函数，函数名无专用前缀。

是win32扩展函数。这部分函数主要用于处理像素存储格式和双帧缓存。这6个函数将替换Windows GDI中原有的同样的函数。Win32API函数库只能用于Windows 95/98/NT环境中。

7         X窗口专用库

是针对Unix和Linux的扩展函数。

包括渲染上下文、绘制图元、显示列表、纹理贴图、等等

初始化 glXQueryExtension()

渲染上下文函数， glXCreateContext() glXDestroyContext() glXCopyContext() glXMakeCurrent() glXCreateGLXPixmap()

执行 glXWaitGL()、glXWaitX()

缓冲区和字体 glXSwapBuffers()、glXUseXFont（）

8 其他扩展库

这些函数可能是新的OpenGL函数，并没有在标准OpenGL库中实现，或者它们是用来扩展已存在的OpenGL函数的功能。和glu、glx和wgl一样，这些OpenGL扩展是由硬件厂商和厂商组织开发的。OpenGL扩展(OpenGL Extention)包含了大量的扩展API函数。

随着硬件的更新，硬件厂商首先向SGI申请登记新的扩展，编写规格说明书（specification）。然后按照说明书进行开发扩展程序。不同的OpenGL实现（OpenGL Implementation）支持的扩展可能不一样，只有随着某一扩展的推广与应用以及硬件技术的提高该扩展才会在所有的OpenGL实现中被给予支持，从而最终成为OpenGL标准库的一部分。扩展由SGI维护，在SGI网站上列出了目前公开的已注册的扩展及其官方说明书。

扩展源由扩展函数的后缀来指明(或使用扩展常量后缀)。例如，后缀WIN表明一个符合Windows规范的扩展，EXT或ARB后缀表明该扩展由多个卖主定义。

下面给出OpenGL官方规定的

命名规则：

ARB

– OpenGL Architecture Review Board 正式核准的扩展，往往由厂商开发的扩展发展而来，如果同时存在厂商开发的扩展和ARB扩展，应该优先使用ARB扩展

EXT

– 多家OpenGL厂商同意支持的扩展

HP

– Hewlett-Packard 惠普

IBM

– International Business Machines

KTX

– Kinetix, maker of 3D Studio Max

INTEL

– Intel 公司

NV

– NVIDIA 公司

MESA

– Brian Paul’s freeware portable OpenGL implementation

SGI

– Silicon Graphics公司开发的扩展

SGIX

– Silicon Graphics (experimental) 公司开发的实验性扩展

SUN

– Sun Microsystems

WIN

– Microsoft

由于OpenGL扩展在针对不同平台和不同驱动，OpenGL不可能把所有的接口程序全部放到gl.h、glx.h、wgl.h中，而是将这些函数头放在了glext.h、glxext.h和wglext.h中。这些扩展被看作时OpenGL核心库规范的增加和修改。

OpenGL扩展也不是没有缺点，正因为各个硬件厂商都可以开发自己的扩展，所以扩展的数目比较大，而且有点混乱，有些扩展实现的相同的功能，可因为是不同厂商开发的，接口却不一样，所以程序中为了实现这个功能，往往要为不同的显卡写不同的程序。这个问题在OpenGL 2.0出来后可能会得到解决，OpenGL 2.0的一个目标就是统一扩展，减少扩展数目。