CUDA CPP Unity Compute Shader 1

2024/05/12

Professional CUDA C Programming

Chapter02 CUDA Programming Model

2.1 CUDA 编程模型概述

CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁。
CUDA编程模型利用GPU架构的计算能力提供的几个特有功能

• 一种通过层次结构在GPU中组织线程的方法
• 一种通过层次结构在GPU中组织内存的方法

CUDA 编程结构
主机：CPU及其内存（主机内存）
设备：GPU及其内存（设备内存）
代码规范
主机内存中的变量名以h_为前缀，设备内存中的变量名以d_为前缀。
一个典型的CUDA程序实现流程遵循以下模式：

1. 把数据从CPU内存拷贝到GPU内存；
2. 调用核函数对存储在GPU内存中的数据进行操作；
3. 将数据从GPU内存传送回CPU内存；

内存管理
表 2-1 主机和设备内存函数

用于执行GPU内存分配的cudaMalloc函数，其函数原型为：

cudaError_t cudaMalloc(void** devPtr,size_t size)

该函数负责向设备分配一定字节的线性内存，并以devPtr的形式返回指向所分配内存的指针。
cudaMemcpy函数负责主机和设备之间的数据传输，其函数原型为：

cudaError_t cudaMemcpy(void* dst,void* src,size_t count,cudaMemcpyKind lind)

此函数从src指向的源存储区复制一定数量的字节到dst指向的目标存储区。复制方向由kind执行，其中kind有以下几种

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

这个函数以同步方式执行，因为在cudaMemcpuy函数返回以及完成传输操作之前主机应用程序是阻塞的。除了内核启动之外的CUDA调用都会返回一个错误的枚举类型cudaError_t。如果GPU内存分配成功。函数返回：

cudaSuccess

否则返回

cudaErrorMemoryAllocation

可使用以下CUDA运行时函数将错误信息转化为可读的错误信息：

char* cudaGetErrorString(cudaError_t error)

内存层次结构
CUDA编程模型最显著的一个特点就是揭示了内存层次结构。每一个GPU设备都有用于不同用途的存储类型。
在GPU内存层次中，最主要的两种内存是全局内存和共享内存。全局类似于CPU的系统内存，而共享内存类似于CPU的缓存。然而GPU的共享内存可以由CUDA C的内核直接控制。

//第2章代码公用函数头文件
//ch02_Header.cuh
#if !defined __CH02_HEADER_H__

#define __CH02_HEADER_H__
#include <cuda_runtime.h>

#include <time.h>
#include <stdlib.h>
#include <stdio.h>

#define CHECK(call){															\
	const cudaError_t error = call;												\
	if (error != cudaSuccess) {													\
		printf("Error: %s:%d ", __FILE__, __LINE__);							\
		printf("code:%d, reason: %s\n", error, cudaGetErrorString(error));		\
		exit(1);															\
	}																			\
}	

static clock_t cpuSecond() {
	return clock();
}

static void initialData(float* p_A, const int p_size) {
	for (size_t i = 0; i < p_size; i++)
	{
		p_A[i] = i;
	}
}

static void printResult(float *p_A,float *p_B, float*p_hostRef,float* p_gpuRef, const int p_size) {

	for (size_t i = 0; i < p_size; i++)
	{
		printf("A %5f  B %5f	C %5f	SUM %5f \n", p_A[i], p_B[i], p_hostRef[i], p_gpuRef[i]);
	}

}

static void initialData_random(float* p_ip, int p_size) {
	time_t t;
	srand((unsigned int)time(&t));
	for (size_t i = 0; i < p_size; i++)
	{
		p_ip[i] = (float)(rand() & 0xFF) / 10.0F;
	}
}
static void initilaInt(int* p_ip, int p_size) {
	for (size_t i = 0; i < p_size; i++)
	{
		p_ip[i] = i;
	}
}
static void sumArrayOnHost(float* p_A, float* p_B, float* p_C, const int p_N) {
	for (size_t i = 0; i < p_N; i++)
	{
		p_C[i] = p_A[i] + p_B[i];
	}
}
static void checkResult(float* p_hostRef, float* p_gpuRef, const int p_N) {
	double epsilon = 1.0E-8;
	int match = 1;
	for (size_t i = 0; i < p_N; i++)
	{
		if (abs(p_hostRef[i] - p_gpuRef[i]) > epsilon) {
			match = 0;
			printf("Array do not match\n");
			printf("host %.2f , gpu %.2f , at current %3d\n", p_hostRef[i], p_gpuRef[i],i);
			break;
		}


	}
	if (match)
		printf("Array match.\n\n");
	return;
}

#endif 

//sumArrayOnHost_GPU.cu
//两个数组相加
#include "ch02_Header.cuh"
int invokeKernel();
//int main() {
//	return invokeKernel();
//}

__global__ void sumArraysOnGPU(float* p_A, float* p_B, float* p_C, const int p_N) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < p_N)
		p_C[i] = p_A[i] + p_B[i];
	printf("%f", p_C[i]);
}


static int invokeKernel() {
	printf("%s Starting...\n");
	int dev = 0;

	cudaDeviceProp deviceProp;
	CHECK(cudaGetDeviceProperties(&deviceProp, dev));
	printf("Using Device %d: %s\n", dev, deviceProp.name);
	CHECK(cudaSetDevice(0));
	//不是可以无限大
	int nElem = 1 << 24;
	printf("Vector size %d\n", nElem);

	size_t nBytes = nElem * sizeof(float);
	float* h_A, * h_B, * hostRef, * gpuRef;
	h_A = (float*)malloc(nBytes);
	h_B = (float*)malloc(nBytes);
	hostRef = (float*)malloc(nBytes);
	gpuRef = (float*)malloc(nBytes);


	long iStart, iElaps;

	iStart = cpuSecond();
	printf("iStart %d\n", iStart);

	initialData_random(h_A, nElem);
	initialData_random(h_B, nElem);
	iElaps = cpuSecond() - iStart;
	memset(hostRef, 0, nBytes);
	memset(gpuRef, 0, nBytes);

	iElaps = cpuSecond() - iStart;
	printf("memset Time elapsed %d ms\n", iElaps);
	iStart = cpuSecond();
	sumArrayOnHost(h_A, h_B, hostRef, nElem);
	printf("iStart %d\n", cpuSecond());
	iElaps = cpuSecond() - iStart;
	printf("sumArrayOnHost Time elapsed %20d ms\n", iElaps);

	float* d_A, * d_B, * d_C;
	//用cudaMalloc在GPU上申请内存
	cudaMalloc((float**)&d_A, nBytes);
	cudaMalloc((float**)&d_B, nBytes);
	cudaMalloc((float**)&d_C, nBytes);

	//使用cudaMemcpy函数把数据从主机内存拷贝到GPU的全局内存中，参数cudaMemcpyHostToDevice指定数据拷贝方向。
	cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);


	dim3 block(nElem);
	dim3 grid(nElem / block.x);
	iStart = cpuSecond();
	sumArraysOnGPU << <grid, block >> > (d_A, d_B, d_C, nElem);

	iElaps = cpuSecond() - iStart;
	printf("sumArraysOnGPU <<<%d, %d>>> Time elapsed %d ms\n", grid.x, block.x, iElaps / CLOCKS_PER_SEC);
	printf("Execution configuration <<<%d, %d>>>\n", grid.x, block.x);

	//使用cudaMemcpy函数把结果从GPU复制到主机的数组gpuRef中
	CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

	checkResult(hostRef, gpuRef, nElem);
	printResult(h_A, h_B, hostRef, gpuRef, nElem);

	//调用cudaFree释放GPU的内存
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);

	free(h_A);
	free(h_B);
	free(hostRef);
	free(gpuRef);


	return 0;
}