CUDA编程记之一基本使用及线程、同步、存储器

一.基本使用

这里我礼貌性的介绍一下最基本的操作：

cuda编程最基础的模式就是创建一个用__global__修饰的核函数，之后在主程序中创建设备数据指针和主机数据变量，设备指针先通过cudaMalloc开辟空间，再通过cudaMemcpy进行赋值，调用核函数后再通过cudaMemcpy拷贝结果。

示例代码解析

#include <stdio.h>
#include <cuda_runtime.h>
#include <opencv2/opencv.hpp>

#define NUM_THREADS 2000000000
#define SIZE  10
#define BLOCK_WIDTH 1000

__global__ void gpu_increment_atomic(int *d_a)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	__shared__ int sh_arr[10];
	if (threadIdx.x < 10) {
		sh_arr[threadIdx.x] = 0;
	}
	__syncthreads();

	tid = tid % SIZE;
	atomicAdd(&sh_arr[tid], 1);
	__syncthreads();
	if (threadIdx.x < 10) {
		atomicAdd(&d_a[threadIdx.x], sh_arr[threadIdx.x]);
	}
}

首先就是一个叫gpu_increment_atomic的核函数，内部先获取当前调用这个核函数的block和thread，然后通过__sharde__修饰一个共享内存数组数组，

__syncthreads()表示等待同步，所有block和thread会等到都执行到此处再一起往下走，之后就是累加，由于共享内存属于block内通用，所以相当于在第二个__syncthreads()时，每个block都得到了共享内存数组累加的结果。在此同步，再将block内的结果累加到输出中。

这里使用atomicAdd()是为了防止线程同时占用，类似于线程锁，不然用i++之类的方式会有问题。（不理解的就自己再思考下，不太想写）

个人感觉CUDA编程最麻烦的就是共享的东西或者同步问题，反正调的很辛酸。

int main(int argc, char **argv)
{
	printf("%d total threads in %d blocks writing into %d array elements\n",
		NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);

	// declare and allocate host memory
	int h_a[SIZE];
	const int ARRAY_BYTES = SIZE * sizeof(int);

	// declare and allocate GPU memory
	int * d_a;

	clock_t start = clock();
	cudaMalloc((void **)&d_a, ARRAY_BYTES);
	// Initialize GPU memory to zero
	cudaMemset((void *)d_a, 0, ARRAY_BYTES);

	gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
	// cudaDeviceSynchronize();

	// copy back the array to host memory
	
	cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);
	double time = (double)(clock() - start) / CLOCKS_PER_SEC;
	
	std::cout << "time" << time << std::endl;
	printf("Number of times a particular Array index has been incremented is: \n");

	for (int i = 0; i < SIZE; i++)
	{
		printf("index: %d --> %d times\n ", i, h_a[i]);
	}

	cudaFree(d_a);
	system("pause");
	return 0;
}

主程序中创建设备数据指针的d_a和主机数据变量h_a，通过cudaMalloc为的d_a开辟空间，再通过cudaMemcpy进行赋值，调用核函数后再通过cudaMemcpy拷贝结果。

调用核函数时<<<NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH>>>前面表示使用的block数量，后面表示单个block内的thread数量，你可以理解为总共启动了block_num×thread_num数量的线程，一方面block内有些内存能公用，比较方便，一方面显卡对单个block内的thread数量有限制。稍新的显卡大约是1024。

运算完成之后需要通过cudaFree()释放资源。

存储器

之前提到了共享内存，其实是因为属于不同内存的资源使用时速度不一致，为了追求极限的性能，还是需要注意的。

全局内存可以被所有BLOCK访问，但是速度最慢。共享内存只能被块中线程访问，但是速度比较快。
存储器访问是程序快速执行的最大瓶颈。（看到没有，这句话岁我从书上抄的，不过说明了什么？说明写的好和写的菜，程序性能是不一样的，有可能你确实使用了cuda，但没达到应有的加速）
全局内存主要是通过cudaMalloc分配的存储器。
寄存器是每个线程的最快存储器，主要是在线程内声明的变量，当寄存器空间不足，或某些不能在存放在寄存器中的情况，会使用本地内存，本地内存性能较慢
L1缓存和共享内存公用64kb内存，可由工程师分配比例
共享内存访问速度大约是全局内存的100倍。但经常需要搭配同步函数__syncthread() 使用。
cuda编程中同步是个巨大问题，很多操作需要搭配atomic原子操作，如++ --,相当于线程锁，防止访问冲突。
常量内存从过__constant__修饰，在核函数以外声明并通过cudaMemcpyToSymbol赋值，有助于节省全局内存的访问带宽。

总结陈词

通过水了一个博客证明偶尔还是有学习的，反正虽然会的技能都是大家几乎生而知之的，可能这就是人生吧。（尝试谦虚）