[CUDA编程]基础入门例程2

例程地址：https://github.com/TycoonL/awesome-cuda.git/

3.basic_add

这个例程是一个基础的CUDA程序结构，使用"单指令-多线程"的方式编写代码，通过分配内存、调用核函数、数据传递，实现了并行计算加法。核函数：

__global__
void add(double *x, double *y, double *r)
{
	int i = blockIdx.x*blockDim.x + threadIdx.x;   // 线程索引计算公式
	r[i] = x[i] + y[i];   // 计算结果存储在 r 数组中
	printf("threadId:%d,%f\n", i, r[i]);   // 打印每个线程的结果
}

add<<<1, 128>>>(cuda_x,cuda_y,cuda_r)结果：
threadId:0,5.790000
threadId:1,5.790000
threadId:2,5.790000
threadId:3,5.790000
threadId:4,5.790000
threadId:5,5.790000
threadId:6,5.790000
threadId:7,5.790000
threadId:8,5.790000
threadId:9,5.790000
threadId:10,0.000000
threadId:11,0.000000
threadId:12,0.000000
threadId:13,0.000000...
注意：例程中N为10，而块线程块大小为128，剩余的118个线程也会运行

为了防止线程大于数组元素个数的非法内存，需要加入if语句：

__global__
void add(double *x, double *y, double *r, int N)
{
	int i = blockIdx.x*blockDim.x + threadIdx.x;   // 线程索引计算公式
	if (i >= N) return;
	r[i] = x[i] + y[i];   // 计算结果存储在 r 数组中
	printf("threadId:%d,%f\n", i, r[i]);   // 打印每个线程的结果
}

数据传递cudaMemcpy(dst,src,count,kind)

• dst:目标地址
• src：源地址
• count：复制数据的字节数
• kind：传递方向
  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToHost
  • cudaMemcpyHostToHost
  • cudaMemcpyDeviceToDevice
  • cudaMemcpyDefault：自动判断传递的方向

4.basic_device

函数执行空间标识符：

• __global__：修饰的函数被称为核函数，由主机调用，设备执行
• __device__：修饰的函数被称为设备函数，由核函数和其它设备函数调用，设备执行
• __host__：修饰的函数就是主机端普通的函数，一般省略

注意：可以同时用__device__和__host__修饰一个函数，但其他的组合不行。

该例程将例程3的加法核函数加入设备函数写法：

//返回值写法
__device__
double add_1(double x, double y)
{
	return (x + y);
}
//指针写法
__device__
double add_2(double x, double y,double *z)
{
	*z=x + y;
}
//引用写法
__device__
double add_3(double x, double y, double &z)
{
	z = x + y;
}
__global__
void add(double *x, double *y, double *r, int N)
{
	int i = blockIdx.x*blockDim.x + threadIdx.x;   // 线程索引计算公式
	if (i >= N) return;
	//r[i] = add_1(x[i] , y[i]);   // 计算结果存储在 r 数组中
	add_2(x[i], y[i],&r[i]);
	//add_3(x[i], y[i], r[i]);
	printf("threadId:%d,%f\n", i, r[i]);   // 打印每个线程的结果
}