CUDA学习笔记

最新推荐文章于 2021-12-08 23:54:25 发布

溪边的小石头

最新推荐文章于 2021-12-08 23:54:25 发布

阅读量1.2k

点赞数 1

本文链接：https://blog.csdn.net/yesterdail/article/details/8178742

版权

1.一些基本概念

CPU及其系统的内存称为主机，GPU及其内存称为设备；函数前的 __global__告诉编译器，函数应该编译为在设备上运行；

kernel<<<M,N>>>，运行时将创建M个线程块(block)，如何知道当前核函数运行的是哪个线程块？ int tid=blockIdx.x;

以上M个线程块的集合叫线程格(grid)，在启动线程块数组时，数组每一维的最大数量不能超过65535

如何操作二维线程格？

	int x = blockIdx.x;
	int y = blockIdx.y;
	int offset = x + y*gridDim.x;

最大的线程数量不能超过该设备属性结构中maxThreadsPerBlock域的值(一般是512)

	int tid=threadIdx.x+blockIdx.x*blockDim.x
	if(tid<N)
	...

	add<<<(N+127)/128, 128>>>(...)

GPU上实现任意长度的计算

	int tid=threadIdx.x+blockIdx.x*blockDim.x
	while(tid<N){
		//...
		tid+=blockDim.x+gridDim.x;
	}

使用二维线程块、二维线程

	dim3 blocks(DIM/16, DIM/16);
	dim3 threads(16, 16);
	kernel<<<blocks, threads>>>(...)

__global__ void kernel(...){
	int x=threadIdx.x+blockIdx.x*blockDim.x;
	int y=threadIdx.y+blockIdx.y*blockDim.y;
	int offset=x+y*blockDim.x*gridDim.x;
	//...
}

2.共享内存与同步

__share__ 添加到变量声明中，使这个变量驻留在共享内存中。对于GPU启动的每个线程块，CUDA C编译器都将创建该变量的一个副本，线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。

__syncthreads(); 这个函数将确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。

当某些线程需要执行一条指令，而其他线程不需要执行时，这种情况称为线程发散。

使用__syncthreads()时一定要注意线程发散问题。CUDA架构确保，除非线程块中的第个线程都执行了__syncthreads()，否则没有任何线程能执行__syncthreads()之后的指令。

3.常量内存

和全局内存的使用只有2处不同：使用__constant__在全局声明变量；当从主机内存复制到GPU上的内存时，使用cudaMemcpyToSymbol()函数，会复制到常量内存，而使用参数为cudaMemcpyHostToDevice的cudaMemcpy()会复制到全局内存。

#define SPHERES 20
__constant__ Sphere s[SPHERES];

//...
cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES)

4.使用事件来测量性能

	cudaEvent_t     start, stop;
	cudaEventCreate( &start );
	cudaEventCreate( &stop );
	cudaEventRecord( start, 0 );

	// 在GPU上执行一些工作(核函数、设备内存复制)

	cudaEventRecord( stop, 0 );
	cudaEventSynchronize( stop );
	float   elapsedTime;
	cudaEventElapsedTime( &elapsedTime, start, stop );
	printf( "Time to generate:  %3.1f ms\n", elapsedTime );

	cudaEventDestroy( start );
	cudaEventDestroy( stop );

5.纹理内存

一维纹理

	// 1.申明纹理引用
	float* src;
	texture<float>  tex;

	// 2.在设备上分配内存，将内存绑定到之前声明的纹理引用
	cudaMalloc((void**)&src, imageSize);
	cudaBindTexture(NULL, tex, src, imageSize);

	// 3.在核函数中采样纹理
	float c = tex1Dfetch(tex,offset);

	// 4.解绑定并释放内存
	cudaUnbindTexture(tex);
	cudaFree(src);

二维纹理

	// 1.申明纹理引用
	float* src;
	texture<float, 2>  tex;

	// 2.在设备上分配内存，将内存绑定到之前声明的纹理引用
	cudaMalloc((void**)&src, imageSize);
	cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
	cudaBindTexture2D(NULL, tex, src, desc, DIM, DIM, sizeof(float)*DIM);

	// 3.在核函数中采样纹理
	float c = tex2D(tex,offset);

	// 4.解绑定并释放内存
	cudaUnbindTexture(tex);
	cudaFree(src);

当使用tex2D()时，不用担心发生溢出问题，当采样坐标小于0时，tex2D()会返回0处的值，当采样坐标大于宽度或高度时，tex2D()将返回位于宽度或高度处的值。

6.图形互操作性

通用计算与渲染模式之间的互操作：能否在一个应用程序中GPU既执行渲染计算，又执行通用计算？如果要渲染的图像依赖通用计算的结果，该如何处理？如果想要在已经渲染的帧上执行某种图像处理或者统计，又该如何实现？这些都是图形互操作性可以解决的问题。

#include <gl/glew.h>
#include <gl/glut.h>
#include "cuda.h"
#include "cuda_gl_interop.h"

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


#define DIM 512


// 指向同一个缓冲区的不同句柄，因为OpenGL和CUDA对于这个缓冲区各自有着不同的"命名"
GLuint imgId;
cudaGraphicsResource* resource;

__global__ void kernel(uchar4* 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 fx = x/(float)DIM - 0.5f;
	float fy = y/(float)DIM - 0.5f;
	unsigned char green = 128 + 127*sin(abs(fx*100)-abs(fy*100));

	ptr[offset].x = 0;
	ptr[offset].y = green;
	ptr[offset].z = 0;
	ptr[offset].w = 255;
}

void display()
{
	glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, imgId);
	glDrawPixels(DIM, DIM, GL_RGBA, GL_UNSIGNED_BYTE, 0);
	glutSwapBuffers();
}

void key(unsigned char key, int x, int y)
{
	switch(key)
	{
	case 27:
		cudaGraphicsUnregisterResource(resource);
		glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
		glDeleteBuffers(1, &imgId);
		exit(0);
	}
}

int main(int argc, char* argv[])
{
	// 选择执行程序的CUDA设备
	cudaDeviceProp prop;
	int dev;
	memset(&prop, 0, sizeof(cudaDeviceProp));
	prop.major = 1;
	prop.minor = 0;
	cudaChooseDevice(&dev, &prop);

	// 选择执行OpenGL的设备
	cudaGLSetGLDevice(dev);

	// 初始化OpenGL窗口
	glutInit(&argc, argv);
	glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA);
	glutInitWindowSize(DIM, DIM);
	glutCreateWindow("图形互操作性");
	glewInit();

	// 创建像素缓冲区对象
	glGenBuffers(1, &imgId);
	glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, imgId);
	glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, DIM*DIM*4, NULL, GL_DYNAMIC_DRAW_ARB);

	// imgId运行时将在CUDA和OpenGL间共享，通过把imgId注册为一个图形资源
	cudaGraphicsGLRegisterBuffer(&resource, imgId, cudaGraphicsMapFlagsNone);

	// 映射该共享资源
	cudaGraphicsMapResources(1, &resource, NULL); 

	// 请求一个指向映射资源的指针
	uchar4* devPtr;
	size_t size;
	cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource);

	dim3 grids(DIM/16, DIM/16);
	dim3 threads(16, 16);
	kernel<<<grids, threads>>>(devPtr);

	// 取消映射，确保cudaGraphicsUnmapResource()之前的所有CUDA操作完成
	cudaGraphicsUnmapResources(1, &resource, NULL);

	// 设置好GLUT并启动循环
	glutKeyboardFunc(key);
	glutDisplayFunc(display);
	glutMainLoop();

	return 0;
}

7.原子性

GPU支持的各种功能统称为计算功能集。只有1.1或更高版本的计算功能集才支持全局内存上的原子操作，只有1.2或更高版本的计算功能集才支持共享内存上的原子操作。计算功能集的版本为1.3或者更高的显卡才支持双精度浮点数学计算，prop.major, prop.minor。

由于一些操作的执行过程不能分解为更小的部分，因此将满足这种条件限制的操作称为原子操作。

计算直方图，atomicAdd()保证在CUDA C中使用原子操作的方式。

__global__ void histo_kernel( unsigned char *buffer,long size,unsigned int *histo ) 
{					
	 int i = threadIdx.x + blockIdx.x * blockDim.x;
	 int stride = blockDim.x * gridDim.x;
	 while(i < size){
		 atomicAdd(&histo[buffer[i]], 1);
		 i += stride;
	 }
}

对100MB的输入数据使用以上进行直方图统计需要1.752秒，与基于CPU的直图计算相比(只需要0.416秒)，性能更糟！这是因为核函数中只包含了非常少的计算工作，全局内存上的原子操作导致了性能的降低。

__global__ void histo_kernel( unsigned char *buffer,long size,unsigned int *histo ) 
{
	 __shared__  unsigned int temp[256];
	 temp[threadIdx.x] = 0;
	 __syncthreads();

	 int i = threadIdx.x + blockIdx.x * blockDim.x;
	 int stride = blockDim.x * gridDim.x;
	 while (i < size) {
		 atomicAdd( &temp[buffer[i]], 1 );
		 i += stride;
	 }

	 __syncthreads();
	 atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );
}

首先分配一个共享内存缓冲区并进行初始化，使用__syncthreads()确保每个线程的写入操作在线程继续前进之前完成。下一个步骤与之前的计算非常类似，只是这里使用了共享内存缓冲区temp[]而不是全局内存缓冲区histo[]，交调用__syncthreads()确保提交最后的写入操作。最后将每个线程块的临时直方图合并到全局缓冲区中。

通过使用共享内存，程序执行时间为0.057秒。

有时候依赖原子操作会带来性能问题，这些问题只能通过对算法的某些部分进行重构来加以解决。如上直方图示例中，我们使用一种两阶段算法，该算法降低了在全局内存俯冲上竞争程度。通常，这种降低内存竞争程度的策略总能带来不错的效果，因此在使用原子操作时，要记住这种策略。

补充：

cudaMemset(dev_data, 0, size);与标准C函数的memset()是相似的。

当线程块的数量为GPU中处理器数量的2倍时，性能最佳。

cudaDeviceProp  prop;
HANDLE_ERROR( cudaGetDeviceProperties( &prop, 0 ) );
int blocks = prop.multiProcessorCount;
histo_kernel<<<blocks*2,256>>>(...);

溪边的小石头

关注

1
点赞
踩
1

收藏

觉得还不错? 一键收藏
0
评论
CUDA学习笔记

1.一些基本概念CPU及其系统的内存称为主机，GPU及其内存称为设备；函数前的 __global__告诉编译器，函数应该编译为在设备上运行；kernel>>，运行时将创建M个线程块(block)，如何知道当前核函数运行的是哪个线程块？ int tid=blockIdx.x;以上M个线程块的集合叫线程格(grid)，在启动线程块数组时，数组每一维的最大数量不能超过65535如何操作二
复制链接

扫一扫