常量内存和纹理内存的问题

加速计算 CUDA
1

（1）常量内存的使用是先声明__constant__变量，然后使用cudaMemcpyToSymbol()设置常量，之后就可以在核函数中使用了，但最好是所有线程都访问相同的地址，而硬件装置会将该地址中的数据每次广播到半个线程束（16个线程）中。（2）常量内存的访问方式是直接用数组访问。
（3）常量内存不可以被修改。

不知对于常量内存我这样的理解是否全面？
但对于纹理内存，我觉得它是可以被修改的，只不过是可以修改它的绑定内存缓冲区，比如：

global void BlendKernel(float* outSrc,const float* inSrc)
{
int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockIdx.yblockDim.y;
int offset=x+yblockDim.xgridDim.x;

int top,bottom,left,right;

top=offset-DIM;
bottom=offset+DIM;
left=offset-1;
right=offset+1;

if (y==0)
{
	top+=DIM;
}

if (y==(DIM-1))
{
	bottom-=DIM;
}

if (x==0)
{
	++left;
}

if (x==(DIM-1))
{
	--right;
}

float t,b,l,r,c;

t=tex1Dfetch(textIn,top);
b=tex1Dfetch(textIn,bottom);
l=tex1Dfetch(textIn,left);
r=tex1Dfetch(textIn,right);
c=tex1Dfetch(textIn,offset);

outSrc[offset]=c+SPEED*(t+b+l+r-4*c);

return;

}

void AnimGpu(DataBlock* db,int ticks)
{
float elapsedTimes;

cudaEventRecord(db->start,0);
dim3 blocks(DIM/16,DIM/16);
dim3 threads(16,16);
CPUAnimBitmap* bitmap=db->bitmap;

for (int i=0;i<90;++i)
{
	CopyConstKernel<<<blocks,threads>>>(db->dev_inSrc);
	BlendKernel<<<blocks,threads>>>(db->dev_outSrc,db->dev_inSrc);
	swap(db->dev_inSrc,db->dev_outSrc);
}

float_to_color<<<blocks,threads>>>(db->output_bitmap,db->dev_inSrc);
cudaMemcpy(bitmap->get_ptr(),db->output_bitmap,bitmap->image_size(),cudaMemcpyDeviceToHost);

cudaEventRecord(db->stop,0);
cudaEventSynchronize(db->stop);
cudaEventElapsedTime(&elapsedTimes,db->start,db->stop);

db->totolTime+=elapsedTimes;
++db->frames;

return;

}

不知这种理解对不对？

2

关于您的第一点理解，是不对的：
（1）对constant cache / constant memory的访问可以使用__constant__开头的变量
（2）以及，最常见的使用方式是直接写上常数，如果可以，编译器会自动放置到constant memory中。
（3）每次是按warp访问的，而不是半个warp(您那个是1.x时代的老资料了）
（4）内容是可以修改了，您上文第一段就给出了一种修改手段。

关于您的第二个问题，
实际上您没有给出您的outSrc和您的texture的绑定关系，因为无法直接对您的代码下结论，
但是一般的说：
（1）texture后备存储可以是cuda array或者是普通的线性内存。
（2）前者可以通过surface写入修改（下次读取有效）
（3）后者可以直接在kernel内通过普通指针写入修改（下次读取有效）
（4）前者和后者都可以使用cudaMemcpy修改

因为他们的确是可以修改的，但不能评价您的代码（因为您的代码没给出texIn和outSrc的绑定关系）

3

不好意思，我为了简化代码突出重点没有写main()中的内容：
int main()
{
DataBlock data;
CPUAnimBitmap bitmap(DIM,DIM,&data);
data.bitmap=&bitmap;
data.totolTime=0;
data.frames=0;

cudaEventCreate(&data.start);
cudaEventCreate(&data.stop);

cudaMalloc((void**)&data.output_bitmap,bitmap.image_size());
cudaMalloc((void**)&data.dev_inSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_outSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_constSrc,bitmap.image_size());

cudaBindTexture(NULL,textIn,data.dev_inSrc,bitmap.image_size());
cudaBindTexture(NULL,textOut,data.dev_outSrc,bitmap.image_size());
cudaBindTexture(NULL,textConst,data.dev_constSrc,bitmap.image_size());

float* temp=(float*)malloc(bitmap.image_size());
for (int i=0;i<DIM*DIM;++i)
{
	temp[i]=0;
	int x=i%DIM;
	int y=i/DIM;

	if ((x>300)&&(x<600)&&(y>310)&&(y<601))
	{
		temp[i]=MAX_TEMP;
	}
}

temp[DIM*100+100]=(MAX_TEMP+MIN_TEMP)/2;
temp[DIM*700+100]=MIN_TEMP;
temp[DIM*300+300]=MIN_TEMP;
temp[DIM*200+700]=MIN_TEMP;

for (int y=800;y<900;++y)
{
	for (int x=400;x<500;++x)
	{
		temp[x+y*DIM]=MIN_TEMP;
	}
}

cudaMemcpy(data.dev_constSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice);

for (int y=800;y<DIM;++y)
{
	for (int x=0;x<200;++x)
	{
		temp[x+y*DIM]=MAX_TEMP;
	}
}

cudaMemcpy(data.dev_inSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice);

free(temp);

bitmap.anim_and_exit( (void (*)(void*,int))AnimGpu,(void (*)(void*))AnimExit);

return 0;
4

嗯嗯。看到您的补充了。

如同我前文所说：
“（3）后者（普通线性显存）可以直接在kernel内通过普通指针写入修改”

以及，建议楼主用的时候小心，此写入(对texture的后备的普通显存的写入）和纹理读取不维持一致性的。一般可以认为，下次kernel启动再使用对应的texture是安全的行为。

5

感谢您的指点