前些时间玩WOW在AH倒腾锻造的一些商品时候偶然产生了做这么一个网站的想法。主要功能是指定一些低级原材料的成本价格然后计算出制造的最终产品的成本价，再根据指定的售价得到收益。经过一个礼拜的设计到实现，目前的完成结果如上图。网站地址是： goblin : http://goblin.hapblue.com/

界面糙了些，功能上只能说凑合，非常不实用，添加新的物品的流程还是太繁琐。在计算产品成本时是基于:

1. 将低级原材料的成本提交到服务器上。
2. 根据制造品的配方计算出产品的成本，再由提交的价格计算收益并保存在服务器上。
3. 从服务器调取更新后的成本，价格、收益并显示在页面上。

感觉访问服务器有些过于频繁了，一些计算应该可以独立于本地上进行，但这样又需要将大量的配方和原材料成本数据保存到本地上，相比较而言好像也不够理想，而且实现上比前者复杂。

下面是添加一个新产品并进行关注的流程：

1. 添加物品(insert item)：需要将产品以及所有相关的材料都添加进物品数据库中，否则无法添加相应配方并进行成本计算。 需要输入所要添加物品的名称以及在计算成本时每单位的数量（如矿石一组是20个）。
   

2. 添加配方(insert recipe)：只有产品和原材料全部都在物品数据库中存在时才可添加该配方。

3. 添加关注物品(favorite)：为避免显示过多用户不关心的物品数据，只有该用户关注的物品才显示在成本收益计算页面中。（虽然现在还没有加入用户管理功能= =。) 也可以在items页面下进行添加关注的操作。输入要关注物品的名称即可。

   

4. 指定成本价格(valuate)： 这一步操作没有什么要说的，输入对应的成本和价格，结果就如首页图所示。

已知和未知的BUG很多，还有许多有待完善和补充的地方，待以后有时间再慢慢修改。算起来这还是第一次做网站，可能也算不上是网站，充其量就是一些PHP和HTML的混搭，全当学习练手之用，游戏之余的另外消遣。

The NURBS Model

The NURBS (Non Uniform Rational B-Spline) model proposes the definition of a curve as a piecewise rational polynomial function of the parameter u. See [1] for a complete description of this model.

NURBS Curve Definition

A NURBS curve is defined by control points Pi, i=0..n, which influence is weighted by rational polynomial functions Ri, i=0..n (dependent on the parameter) and weights wi, i=0..n (independent on the parameter). The rational polynomial functions Ri are defined by the means of a basis, called B-Spline basis, set of piecewise polynomial functions Nik, i=0..n, of same degree k. The degree of the NURBS curve is the degree of the polynomial functions.

The definition of the basis Nik is uniquely determined by a knot vector. The pieces of the basis polynomial functions are called arcs. They represent an interval for the parameter values to calculate a segment of shape.

The control points are not, in general, points of the NURBS curve. By convention however, the first and last control points are the begin and end point of the curve respectively, except for the periodic NURBS curves. These control points can be seen as an attracting zone for the curve, which influence is weighted as seen previously.

Knot Vector

The knot vector is the way to state the definition of the basis. In particular, it manages the continuity between the different arcs of the basis functions, and, hence, the curves that use it. The knot vector is a set of non decreasing parameter values (t_0,..,t_m), called knot values or knots. The B-Spline basis is recursively defined as follows, with i=0,..,n:

with the following conventions:

The relation between the number of knots (m+1), the degree (k) of Nik and the number of control points (n+1) is given as follows:

m = (n+1) + k

Knot values are non-decreasing, so a knot vector can have knots with the same value. In this case, the knot is called multiple, and its multiplicity is the number of repetitions of the same value. There are as many arcs as knots of different values plus one. If the increment is always 1, the knot vector is called uniform.

The multiplicity is a way to specify the continuity order between the arcs. Hence, there is a relation between the multiplicity and this continuity order:

1. For an internal knot value (neither the first, nor the last), the continuity order is the degree minus the multiplicity. As a consequence, the multiplicity of a knots cannot be strictly greater then the degree.
2. By convention, for a non periodic basis (open curve), each extreme knot value has a multiplicity equal to the degree plus one.
3. For a periodic basis (they are used to model closed curves), the first point applies for all knots. Moreover, the first and last multiplicities are the same.

The following table summarizes these relations:

Fig. 1: Illustration of Example 1

The green curve has four control points (CP1 to CP4).If you move CP3 to CP3′, the curve is attracted by this new points. Notice two important properties of this kind of NURBS, called Bezier arc: The curve is inside the convex hull of the control points The curve is tangent to the segment joigning the first and second control points at the beginning of the curve, and to the segment joining the last and next to last control points at its end.

Fig. 2: Illustration of Example 2

The green curve is an example of non uniform polynomial B-Spline curve, having the knot vector of the example 2. The weight of each control point is 1.If you assign (1,1,10,20,5,1) to the control points, it gives the red curve. This curve is attracted by the control points CP3 and CP4, that are more weighted than the others. These curves have three arcs: CP1-K1, K1-K2, K2-CP2.

According to the definition, it is possible to create NURBS curves that are only C0 or C1 continuous. Despite of this fact, remember that the CGM geometric operators suppose that the geometry is at least C2 continuous.

Bernstein Basis and Bézier Points

NURBS, NUPS, UPBS have a common feature: the continuity between arcs is managed, with more or less flexibility, by the model with the knot vector. The arcs definitions are linked together. Now, the Bernstein-Bézier model defines a basis for one arc. To define a curve with several arcs, you have to connect arcs, each arc being independent on its neighbors. Hence, the continuity is not managed by the model, you have to put constraints between the different arcs to insure it.

The parametric equation of an arc of degree k and the functions of the basis are as follows:

The control points Pi are called Bézier points.

The transformation of a NURBS curve into a Bernstein-Bézier curve amounts to increase the multiplicity of the knots until having a C0 continuity. Conversely, the transformation of a Bernstein-Bézier curve into a NURBS curve amounts to build a knot vector for which the knots have a multiplicity equal to the degree.

Fig. 1 displays examples of Bézier arcs.

• Vertex shader file

• fragment shader file

这里有关于GLSL的介绍：传送门

地城手册已经列出了现有的和在4.2中将加入的新副本火源之界。

可以在地城手册中查看每个BOSS的掉落列表（看来atlasloot即将被取代了= =）

当然最重要的是可以查看BOSS的详细技能说明（向以前一次一次尝试来获取技能数据的先行者致敬）

顺便一提，在技能列表的右边会有具有辅助提示作用的一些小标记，用于对应职业特别注意，例如盾牌，加号，剑这些图示。上图中的骷髅有说是需要大家注意的杀招，或者是只有在heroic模式中才会出现的技能。当然我更倾向于后者，想象如果整个一个阶段都是骷髅杀招的话……

但是3D角色浏览功能还未正常。在新的英雄榜中可以清晰的浏览装备，宝石，附魔，重铸以及天赋雕纹。

GLSL – An Introduction

By: Florian Rudolf

What Is GLSL?

GLSL (GLslang) is a short term for the official OpenGL Shading Language. GLSL is a C/C++ similar high level programming language for several parts of the graphic card. With GLSL you can code (right up to) short programs, called shaders, which are executed on the GPU.

Why Shaders?

Until DirectX 8 hardware (GeForce 2 and lower, Radoen 7000 and lower) the graphic pipeline could only be configured, but not be programmed. For example there is the OpenGL lighting model with ambient, diffuse, specular and emissive lighting. This model is mainly used but there are many other models for lighting. In fixed-function OpenGL only this lighting model could be used, no other. With Shaders you are able to write your own lighting model. But that’s only one feature of shaders. There are thousands of other really nice possibilities: Shadows, Environment Mapping, Per-Pixel Lighting, Bump Mapping, Parallax Bump Mapping, HDR, and much more!

Why GLSL?

Shaders are available in OpenGL till 2002 through ARB_vertex_program and ARB_fragment_program extension. But with those extensions you are only able to use assembly shaders. Because of the growing complexity of lighting and shading models assembly shaders are hard to use. GLSL is a high-level shading language, which means that you can write your shader in C/C++ style. This makes shader development much easier!

What Is The Difference Between Fixed Function Pipeline And GLSL?

There are two types of shaders in GLSL: vertex shaders and fragment shaders.

Vertex Shader

A vertex shader operates on every vertex. So if you call glVertex* (or glDrawArrays, …) the vertex shader is executed for each vertex. If you use a vertex shader you have nearly full control over what is happening with each vertex. But if you use a vertex shader ALL Per-Vertex operations of the fixed function OpenGL pipeline are replaced (see Figure 1):

Vertex Transformation

Normal Transformation, Normalization and Rescaling

Lighting

Texture Coordinate Generation and Transformation

…

For a full overview what a vertex shader replaces and what it does not replace please see reference [1], page 41. So if you want to use a vertex shader you HAVE to do all these things above on your own (of course, only if you need it

Fragment Shader

A fragment shader operates on every fragment which is produced by rasterization. With fragment shader you have nearly full control over what is happening with each fragment. But just like a vertex shader, a fragment shader replaces ALL Per-Fragment operations of the fixed function OpenGL pipeline (see Figure 1):

Texture access and application (Texture environments)

Fog

…

For a full overview what a fragment shader replaces and what it does not replace please see reference [1], page 43. Using a fragment shader is just like using a vertex shader, because you HAVE to do all these things above on your own (of course, only if you need it).

Figure 1, OpenGL 1.5 fixed function pipeline [1]

What Does GLSL Look Like?

As mentioned above there are 2 types of shaders, a vertex shader and a fragment shader. Each shader type has other inputs and outputs.

Data Types In GLSL

There are four main types: float, int, bool and sampler. For the first three types, vector types are available:

vec2, vec3, vec4			2D, 3D and 4D floating point vector

ivec2, ivec3, ivec4			2D, 3D and 4D integer vector

bvec2, bvec3, bvec4			2D, 3D and 4D boolean vectors

For floats here are also matrix types:mat2, mat3, mat4 2×2, 3×3, 4×4 floating point matrix

Samplers are types representing textures. They are used for texture sampling. Sampler types have to be uniform. They are not allowed to be declared as a non-uniform type. Here are the different sampler types:

sampler1D, sampler2D, sampler3D		1D, 2D and 3D texture

samplerCube				Cube Map texture

sampler1Dshadow, sampler2Dshadow	1D and 2D depth-component texture

About Attributes, Uniforms And Varyings

There are three types of inputs and outputs in a shader: uniforms, attributes and varyings.

Uniforms are values which do not change during a rendering, for example the light position or the light color. Uniforms are available in vertex and fragment shaders. Uniforms are read-only.

Attributes are only available in vertex shader and they are input values which change every vertex, for example the vertex position or normals. Attributes are read-only.

Varyings are used for passing data from a vertex shader to a fragment shader. Varyings are (perspective correct) interpolated across the primitive. Varyings are read-only in fragment shader but are read- and writeable in vertex shader (but be careful, reading a varying type before writing to it will return an undefined value). If you want to use varyings you have to declare the same varying in your vertex shader and in your fragment shader.

All uniform, attribute and varying types HAVE to be global. You are not allowed to specify a uniform/attribute/varying type in a function or a void.

Built-In Types

GLSL has some built-in attributes in a vertex shader:

gl_Vertex				4D vector representing the vertex position

gl_Normal				3D vector representing the vertex normal

gl_Color				4D vector representing the vertex color

gl_MultiTexCoordX			4D vector representing the texture coordinate of texture unit X

There are some other built-in attributes, see reference [2], page 41 for a full list.

GLSL also has some built-in uniforms:

gl_ModelViewMatrix			4x4 Matrix representing the model-view matrix.

gl_ModelViewProjectionMatrix		4x4 Matrix representing the model-view-projection matrix.

gl_NormalMatrix				3x3 Matrix representing the inverse transpose model-view matrix.

					This matrix is used for normal transformation.

There are some other built-in uniforms, like lighting states. See reference [2], page 42 for a full list.

GLSL Built-In Varyings:

gl_FrontColor				4D vector representing the primitives front color

gl_BackColor				4D vector representing the primitives back color

gl_TexCoord[X]				4D vector representing the Xth texture coordinate

There are some other built-in varyings. See reference [2], page 44 for a full list.

And last but not least there are some built-in types which are used for shader output:

gl_Position				4D vector representing the final processed vertex position. Only 

					available in vertex shader.

gl_FragColor				4D vector representing the final color which is written in the frame 

					buffer. Only available in fragment shader.

gl_FragDepth				float representing the depth which is written in the depth buffer.

					Only available in fragment shader.

The importance of built-in types is that they are mapped to the OpenGL states. For example if you call glLightfv(GL_LIGHT0, GL_POSITION, my_light_position) this value is available as a uniform using gl_LightSource[0].position in a vertex and/or fragment shader.

Generic Types

You are also able to specify your own attributes, uniforms and varyings. For example if you want to pass a 3D tangent vector for each vertex from your application to the vertex shader you can specify a “Tangent” attribute:attribute vec3 Tangent;

See full text.

singleton 是一种对象，它在一个系统中的任何时候只有一个实例。在游戏中，singleton 的一些常见例子是纹理贴图，文件或用户界面的管理程序。它们每一个都是一个子系统，通常假定从游戏开始时可用，一直持续到游戏关闭。

一种通用的解决方法是使用模板来自动定义 singleton 指针，并完成指针设置、查询和清除的工作。它还可以检查（通过 assert() )确保没有将 singleton 实例化多次。最重要的一点是，我们可以免费获得所有这些功能，只需要从以下这个简单的类派生就可以：

#include <cassert>
 
template <typename T> class Singleton
{
	static T* ms_Singleton;
 
public:
	Singleton(void)
	{
		assert(!ms_Singleton);
		int offset = (int)(T*)1 - (int)(Singleton<t>*)(T*)1;
		ms_Singleton = (T*)((int)this + offset);
	}
	~Singleton(void)
	{
		assert(ms_Singleton);
		ms_Singleton = 0;
	}
	static T& GetSingleton(void)
	{
		assert(ms_Singleton);
		return (*ms_Singleton);
	}
	static T* GetSingletonPtr(void)
	{
		return (ms_Singleton);
	}
};
 
template <typename T> T* Singleton<T>::ms_Singleton = 0;

将任何类转换成 singleton ，只需要按以下3个简单步骤来做：

1. 从Singleton<MyClass> 公开派生你的类 MyClass。
2. 确保使用前在系统中创建了 MyClass 的实例。如果实例化并不重要，可以通过将它设为全局或局部静态，让编译程序来做；也可以通过一个属主类，用 new 和 delete 来自己做。不论实例在何时如何创建，它将被追踪，并通过公共接口为系统其他部分用作 singleton。
3. 在系统的任何地方调用 MyClass::GetSingleton() 来使用对象。也可以 #define g_MyClass 为 MyClass::GetSingleton() ，并将其作为全局对象来对待，以方便书写。

以下是使用该类的例子：

//sample code
class TextureMgr : public Singleton<TextureMgr>
{
public:
	Texture*  GetTexture( const char* name);
	// ...
};
 
#define g_TextureMgr TextureMgr::GetSingleton()
 
void Func(void)
{
	Texture* stone1 = TextureMgr::GetSingleton().GetTexture("stone1");
	Texture* wood6 = g_TextureMgr.GetTexture("wood6");
	// ...
}

这个 Singleton 类的惟一目的是在它的派生 (MyClass) 类型的任何实例创建和销毁时自动注册和注销它们。所有重要的工作在 Singleton 的构造函数中完成，在此它计算出派生实例的相对位置，并将结果存储到 singleton 指针 (ms_Singleton) 。注意，派生类可能不仅仅从 Singleton 派生，在这种情况下 MyClass 的 “this” 不同。这种解决方法假设一个不存在的对象在内存的 0×1 位置上，将此对象强制转换为两种类型，并得到其偏移量的差值。这个差值可以有效的作为 Singleton<MyClass> 和它派生类型 MyClass 的距离，可用于计算 singleton 指针。

Reference

Meyers, Scott, More Effective C++, Addison-Wesley Publishing Co., 1995

在用GDI函数 SetDIBitsToDevice 显示DIB位图时，需要一个类型为 BITMAPINFO * 的参数 lpbmi。这个参数一般需要自己赋值。

int SetDIBitsToDevice(
  __in  HDC hdc,
  __in  int XDest,
  __in  int YDest,
  __in  DWORD dwWidth,
  __in  DWORD dwHeight,
  __in  int XSrc,
  __in  int YSrc,
  __in  UINT uStartScan,
  __in  UINT cScanLines,
  __in  const VOID *lpvBits,
  __in  const BITMAPINFO *lpbmi,
  __in  UINT fuColorUse
);

BITMAPINFO结构体的定义如下：

typedef struct tagBITMAPINFO {
  BITMAPINFOHEADER bmiHeader;
  RGBQUAD          bmiColors[1];
} BITMAPINFO, *PBITMAPINFO;

其中 BITMAPINFOHEADER 的结构定义如下：

typedef struct tagBITMAPINFOHEADER {
DWORD biSize; 
LONG biWidth; 
LONG biHeight; 
WORD biPlanes; 
WORD biBitCount;
DWORD biCompression; 
DWORD biSizeImage; 
LONG biXPelsPerMeter;
LONG biYPelsPerMeter;
DWORD biClrUsed;
DWORD biClrImportant;
} BITMAPINFOHEADER, *PBITMAPINFOHEADER;

bmiColors[1] 是一个只有一个项的数组，其结构 RGBQUAD 定义如下：

typedef struct tagRGBQUAD {
    BYTE    rgbBlue;
    BYTE    rgbGreen;
    BYTE    rgbRed;
    BYTE    rgbReserved;
} RGBQUAD

bmiColors[1] 的这个设计是为了照顾bmp文件，bmp文件主要有4部分组成：

1. BITMAPFILEHEADER bmp头文件
2. BITMAPINFOHEADER bmp文件信息头
3. RGBQUAD                 bmp文件调色板
4. BITMAP                     bmp文件数据

而对于不同BPP (bits per pixel) 的 bitmap 来说，palette （调色板）的大小是不一样的。对于8位灰度图像而言，palette 有256 个RGBQUAD 。因为 palette 是连续存储的，在BITMAPINFO 里只存放palette的第一项。

下面代码为一个8位灰度图像的BITMAPINFO 赋值：

BITMAPINFOHEADER bmih;
memset(&bmih, 0, sizeof(BITMAPINFOHEADER));
bmih.biSize = sizeof(BITMAPINFOHEADER);
bmih.biWidth = width;
bmih.biHeight = -height;     //biHeight should be negative
bmih.biPlanes = 1;           //biPlanes must be 1
bmih.biBitCount = 8;         //bpp
bmih.biCompression = BI_RGB; //no compression
 
RGBQUAD colorTable[256];
DWORD color = 0;
for (int i = 0; i < 256; ++i, color += 0x010101)
	*((DWORD *)&colorTable[i]) = color;
 
BITMAPINFO * pbmi;
pbmi = (BITMAPINFO *)malloc(sizeof(BITMAPINFOHEADER) + sizeof(RGBQUAD) * 256);
pbmi->bmiHeader = bmih;
memcpy((void*)pbmi->bmiColors, (void*)colorTable, sizeof(RGBQUAD) * 256);

The WSAWaitForMultipleEvents function returns when one or all of the specified event objects are in the signaled state, when the time-out interval expires, or when an I/O completion routine has executed.

DWORD WSAWaitForMultipleEvents(
  __in  DWORD cEvents,
  __in  const WSAEVENT *lphEvents,
  __in  BOOL fWaitAll,
  __in  DWORD dwTimeout,
  __in  BOOL fAlertable
);

WSAWaitForMultipleEvents 基本使用模式如下。

int ret = WSAWaitForMultipleEvents(EventTotal, EventArray, FALSE, WSA_INFINITE, FALSE);
int index = ret - WSA_WAIT_EVENT_0;
for (; index < EventTotal; ++index)
{
    ret = WSAWaitForMultipleEvents(1, EventArray[i], TRUE, 1000, FALSE);
    if ((ret == WSA_WAIT_TIMEOUT) || (ret == WSA_WAIT_FAILED))
        continue;
    //do something ...
    //if it is a socket event
    WSAEnumNetworkEvents(socket, EventArray[index], &NetworkEvents);
    if (NetworkEvents.lNetworkEvents & FD_READ)
        //do something
    if (NetworkEvents.lNetworkEvents & FD_CLOSE)
        //do something
    //if (...)
}

相关结构和函数。

WSANETWORKEVENTS    NetworkEvents;
WSAOVERLAPPED       AcceptOverlapped;
WSABUF              DataBuf;

EventArray[EventTotal] = WSACreateEvent();

WSAEventSelect(socket, EventArray[EventTotal], FD_READ | FD_CLOSE);

WSAResetEvent(EventArray[index]);

WSACloseEvent(EventArray[index]);

WSARecv(socket, &RecvDataBuf, 1, &RecvBytes, &Flags, &AcceptOverlapped, NULL);

WSASend(socket, &SendDataBuf, 1, &SendBytes, Flags, &AcceptOverlapped, NULL);

HANDLE event;
event = CreateEvent(NULL, TRUE, FALSE, (LPCTSTR)("event name"));
SetEvent(event);
ResetEvent(event);
CloseHandle(event);

Requirements

Header: Winsock2.h

Library: Ws2_32.lib

DLL: Ws2_32.dll

• 共享存储器

示例：共享存储器的动态与静态分配与初始化

int main(int argc, char** argv){
    testKernel<<<1, 10, mem_size>>>(d_idata, d_odata);
    CUT_EXIT(argc, argv);
}
 
__global__ void testKernel(float* g_idata, float* g_odata){
    extern __shared__ float sdata_dynamic[]; //extern 声明，大小由主机端程序决定。动态声明
    __shared__ int sdata_static[16]; //静态声明数组大小
 
    sdata_static[tid] = 0; //shared memory 不能在定义时初始化
}

将共享存储器中的变量声明为外部数组时，数组的大小将在Kernel 启动时确定，通过其执行参数确定。通过这种方式定义的所有变量都开始于相同的地址，因此数组中的变量的布局必须通过偏移量显式管理。例：如果希望在动态分配的共享存储器内获得与以下代码对应的内容：

short array0[128];
float array1[64];
int array2[256];

应该按照下面的方式定义：

extern __shared__ char array[];
__device__ void func()
{
	short* array0 = (short*)array;
	float* array1 = (float*)&array0[128];
	int* array2 = (int*)&array1[64];
}

• 全局存储器

显存中的全局存储器也称为线性内存。线性内存通常使用 cudaMalloc() 函数分配， cudaFree() 函数释放，并由 cudaMemcpy() 进行主机端与设备端的数据传输。通过CUDA API分配的空间未经过初始化，初始化共享存储器需要调用 cudaMemset 函数。

对于二维、三维数组，我们使用 cudaMallocPitch() 和 cudaMalloc3D() 分配线性存储空间。这些函数能够确保分配满足对齐要求。

例：分配一个尺寸为 width * height 的 float 型2D 数组，以及遍历数组元素。

//主机端代码
float* devPtr;
int pitch;
cudaMallocPitch((void**)&devPtr, &pitch,
		width * sizeof(float), height);
myKernel<<<100, 512>>>(devPtr, pitch);
//设备端代码
__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];
		}
	}
}

例：分配一个 width * height * depth 的 float 型3D 数组，以及遍历数组元素。

//主机端代码
cudaPitchedPtr devPitchedPtr;
cudaExtent extent = make_cudaExtent(64, 64, 64);
cudaMalloc3D(&devPitchedPtr, extent);
myKernel<<<100, 512>>>(devPitchedPtr, extent);
 
//设备端代码
__global__ void myKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent){
	char* devPtr = devPitchedPtr.ptr;
	size_t pitch = devPitchedPtr.pitch;
	size_t slicePitch = pitch * extent.height;
	for (int z = 0; z < extent.depth; ++z){
		char* slice = devPtr + z * slicePitch;
		for (int y = 0; y < extent.height; ++y){
			float* row = (float*)(slice + y * pitch);
			for (int x = 0; x < extent.width; ++x){
				float element = row[x];
			}
		}
	}
}

例：二维数组和CUDA数组间的数据拷贝。

cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch,
		    width * sizeof(float), height,
		    cudaMemcpyDeviceToDevice);

• 主机端页锁定内存

通过 cudaHostAlloc() 和 cudaFreeHost() 来分配和释放 pinned memory。

• 常数存储器

定义常数存储器时，需要将其定义在所有函数之外，作用范围为整个文件，并且对主机端和设备端函数都可见。下面两段代码说明了两种常数存储器的使用方法。

第一种方法是直接在定义时直接初始化常熟存储器，然后再Kernel里面直接使用就可以了。

__constant__ int t_HelloCUDA[11] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
__constant__ int num = 11;

第二种方法是定义一个costant 数组，然后使用函数进行赋值。

__constant__ char p_HelloCUDA[11];
CUDA_SAFE_CALL(cudaMemcpyToSymbol(p_HelloCUDA, helloCUDA, sizeof(char) * 11));

• 纹理存储器

在Kernel 中访问纹理存储器的操作称为纹理拾取（texture fetching）。纹理拾取使用的坐标与数据在显存中的位置可以不同，我们通过纹理参照系（texture reference）约定二者的映射方式。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定（texture binding）。显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器（Linear Memory）和 CUDA 数组（CUDA Array）。

纹理存储器的使用

1. 声明CUDA数组，分配空间
2. 声明纹理参照系
3. 设置运行时纹理参照系属性
4. 纹理绑定
5. 纹理拾取

例：简单的纹理使用。

/*
声明纹理参照系
texture<Type, Dim, ReadMode> texRef;
*/
//2D float texture
texture<float, 2, cudaReadModeElementType> texRef;
 
//设备端代码，一个简单的转换kernel
__global__ void transformKernel(float* output, int width, int height, float theta){
	//计算归一化的纹理坐标
	unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
	float u = x / (float)width;
	float v = y / (float)height;
	//坐标转换
	u -= 0.5f;
	v -= 0.5f;
	float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
	float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;
	//读纹理并向全局存储器写会
	output[y * width + x] = tex2D(tex, tu, tv);
}
 
//主机端代码
int main()
{
	//在先存上为CUDA array 分配空间
	cudaChannelFormatDesc channelDesc = cudaCreatChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
	cudaArray* cuArray;
	cudaMallocArray(&cuArray, &channelDesc, width, height);
	//内存中h_data地址处的数据向显存进行拷贝
	cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
	//设置纹理参数
	texRef.addressMode[0] = cudaAddressModeWrap;
	texRef.addressMode[1] = cudaAddressModeWrap;
	texRef.filterMode = cudaFilterModeLinear;
	texRef.normalized = true;
	//数组绑定到纹理
	cudaBindTextureToArray(texRef, cuArray, &channelDesc);
	//问转换结果分配显存空间
	float* output;
	cudaMalloc((void**)&output, width * height * sizeof(float));
	//启动 kernel
	dim3 dimBlock(16, 16);
	dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);
	transformKernel<<<dimGrid, dimBlock>>>(output, width, height, angle);
	//释放显存空间
	cudaFreeArray(cuArray);
	cudaFree(output);
}