CUDA中的常量内存constant

2019-11-11 04:27:28

字体：大中小

来源：转载

供稿：网友

GPU包含数百个数学计算单元，具有强大的处理运算能力，可以强大到计算速率高于输入数据的速率，即充分利用带宽，满负荷向GPU传输数据还不够它计算的。CUDA C除全局内存和共享内存外，还支持常量内存，常量内存用于保存在核函数执行期间不会发生变化的数据，使用常量内存在一些情况下，能有效减少内存带宽，降低GPU运算单元的空闲等待。

使用常量内存提升性能

使用常量内存可以提升运算性能的原因如下：

对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程，这将节约15次读取操作；高速缓存。常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量；在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量数据，那么这种方式产生的内存流量只是使用全局内存时的1/16。

常量内存的声明

为普通变量分配内存时是先声明一个指针，然后通过cudaMalloc()来为指针分配GPU内存。而当我们将其改为常量内存时，则要将这个声明修改为在常量内存中静态地分配空间。我们不再需要对变量指针调用cudaMalloc()或者cudaFree()，而是在编译时为这个变量（如一个数组）提交固定的大小。首先用“___constant_”声明一个常量内存变量，然后使用cudaMemcpyToSymbol（而不是cudaMemcpy）把数据从主机拷贝到设备GPU中。

常量内存使用示例

以下程序用CUDA+OpenCv实现一个简单场景的光线跟踪，光线跟踪是从三维场景生成二维图像的一种方式。主要思想为：在场景中选择一个位置放上一台假想的相机，该相机包含一个光传感器来生成图像，需要判断那些光将接触到这个传感器。图像中每个像素与命中传感器的光线有相同的颜色和强度。传感器中命中的光线可能来自场景中的任意位置，想象从该像素发出一道射线进入场景中，跟踪该光线穿过场景，直到光线命中某个物体。代码实现：

#include "cuda_runtime.h"#include <highgui/highgui.hpp>#include <time.h>using namespace cv;#define INF 2e10f  #define rnd(x) (x*rand()/RAND_MAX)  #define SPHERES 100 //球体数量#define DIM 1024    //图像尺寸struct Sphere{	float r, g, b;	float radius;	float x, y, z;	__device__ float hit(float ox, float oy, float *n)	{		float dx = ox - x;		float dy = oy - y;		if (dx*dx + dy*dy < radius*radius)		{			float dz = sqrt(radius*radius - dx*dx - dy*dy);			*n = dz / sqrt(radius*radius);			return dz + z;		}		return -INF;	}};// Sphere *s;__constant__ Sphere s[SPHERES];/************************************************************************///__global__ void rayTracing(unsigned char* ptr, Sphere* s)  __global__ void rayTracing(unsigned char* ptr){	int x = threadIdx.x + blockIdx.x * blockDim.x;	int y = threadIdx.y + blockIdx.y * blockDim.y;	int offset = x + y  * blockDim.x * gridDim.x;	float ox = (x - DIM / 2);	float oy = (y - DIM / 2);	float r = 0, g = 0, b = 0;	float maxz = -INF;	for (int i = 0; i < SPHERES; i++)	{		float n;		float t = s[i].hit(ox, oy, &n);		if (t > maxz)		{			float fscale = n;			r = s[i].r * fscale;			g = s[i].g * fscale;			b = s[i].b * fscale;			maxz = t;		}	}	ptr[offset * 3 + 2] = (int)(r * 255);	ptr[offset * 3 + 1] = (int)(g * 255);	ptr[offset * 3 + 0] = (int)(b * 255);}/************************************************************************/int main(int argc, char* argv[]){	cudaEvent_t start, stop;	cudaEventCreate(&start);	cudaEventCreate(&stop);	cudaEventRecord(start, 0);	Mat bitmap = Mat(Size(DIM, DIM), CV_8UC3, Scalar::all(0));	unsigned char *devBitmap;	(cudaMalloc((void**)&devBitmap, 3 * bitmap.rows*bitmap.cols));	//  cudaMalloc((void**)&s, sizeof(Sphere)*SPHERES);  	Sphere *temps = (Sphere*)malloc(sizeof(Sphere)*SPHERES);	srand(time(0));  //随机数种子	for (int i = 0; i < SPHERES; i++)	{		temps[i].r = rnd(1.0f);		temps[i].g = rnd(1.0f);		temps[i].b = rnd(1.0f);		temps[i].x = rnd(1000.0f) - 500;		temps[i].y = rnd(1000.0f) - 500;		temps[i].z = rnd(1000.0f) - 500;		temps[i].radius = rnd(100.0f) + 20;	}	//  cudaMemcpy(s, temps, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice);  	cudaMemcpyToSymbol(s, temps, sizeof(Sphere)*SPHERES);	free(temps);	dim3 grids(DIM / 16, DIM / 16);	dim3 threads(16, 16);	//  rayTracing<<<grids, threads>>>(devBitmap, s);  	rayTracing << <grids, threads >> > (devBitmap);	cudaMemcpy(bitmap.data, devBitmap, 3 * bitmap.rows*bitmap.cols, cudaMemcpyDeviceToHost);	cudaEventRecord(stop, 0);	cudaEventSynchronize(stop);	float elapsedTime;	cudaEventElapsedTime(&elapsedTime, start, stop);	PRintf("Processing time: %3.1f ms/n", elapsedTime);	imshow("CUDA常量内存使用示例", bitmap);	waitKey();	cudaFree(devBitmap);	//  cudaFree(s);  	return 0;}程序里生成球体的大小和位置是随机的，为了产生随机数，加入了随机数种子srand（）。运行效果：