CUDA Unity Compute Shader 3

news2024/11/19 9:38:01

计划
这应该是第3章的读书笔记,但是因为第3章读起来比较困难,所以先看了《CUDA并行程序设计编程指南》的第5章和第6章,感觉读起来顺畅多了,《CUDA并行程序设计编程指南》暂定精读第5、6、7章

1.如何生成ptx文件
属性➔CUDA C/C++➔Common➔Keep Preprocess Files➔是(–keep)
在这里插入图片描述2.查看内核使用的寄存器数量;
属性➔CUDA C/C++ ➔Device➔Verbose PTXAS Output➔是 (–ptxas-options=-v)
在这里插入图片描述# NVIDIA_CUDA_Programming_Guide_1.1_chs

4.2.2 同步函数
void __syncthreas();
在一个块内同步所有线程。一旦所有线程到达了这点,恢复正常执行。
__syncthreads()通常用于调整在相同块之间的线程通信。当在一个块内的有些线程访问相同的共享或全局内存时,对于有些内存访问潜在存在read-after-write,write-after-read,或者write-after-write的危险。
这些数据危险可以通过同步线程之间的访问得以避免。
__syncthreads()允许放在条件代码中,但只有当整个线程块有相同的条件贯穿时,否则代码执行可能被挂起或导致没想到的副作用。

注解:这是对__syncthreads()函数的解释

Professional CUDA C Programming

Chapter03 CUDA Execution Model

3.3 并行性的表现

注解:和书中使用的测试程序有所不同,偷个懒,后面看时间再同步

3.3.1 用nvprof检测活跃的线程数

一个内核的可实现占用率被定义为:
每周期内活跃线程束的平均数量与一个SM支持的线程束最大数量的比值。

nvprof --metrics achieved_occupany CUDA.exe

注解:–metrics后面需要跟一个指令
在这里插入图片描述

3.3.2 用nvprof检测内存操作

gld_throughput指标检查内核的内存读取效率
gld_efficiency指标检测全局加载效率,即被请求的全局加载吞吐量占所需的全局加载吞吐量的比值。它衡量了应用程序的加载操作利用设备内存带宽的程度。

3.3.3 增大并行性

指标与性能
▨ 在大部分情况下,一个单独的指标不能产生最佳的性能
▨ 与总体性能最直接相关的指标或事件取决于内核代码的本质
▨ 在相关的指标与事件之间寻求一个好的平衡
▨ 从不同角度查看内核以寻求相关指标间的平衡
▨ 网格/块启发式算法为性能调节提供了一个很好的起点

3.4 避免分支分化

3.4.1 并行规约问题

要对一个有N个元素的整数数组求和。
▨ 相邻配对:元素与它们直接相邻的元素配对
▨ 交错配对:根据给定的跨度配对元素
C语言利用递归实现的一个交错配对方法:

int recursiveReduce(int* data, int const size) {
	if (size == 1)
		return data[0];
	int const stride = size / 2;
	for (size_t i = 0; i < stride; i++)
	{
		data[i] += data[i + stride];
	}
	return recursiveReduce(data, stride);
}

在向量中执行满足交换律和结合律的运算,被称为归约问题。
3.4.2 并行归约中的分化

在这个内核里,有两个全局内存数组:一个大数组用来存放整个数组,进行归约;另一个小数组用来存放每个线程块的部分和。每个线程块在数组的一部分上独立地执行操作。
循环中迭代一次执行一个归约步骤。归约是在就地完成的,这意味着在每一步,全局内存的值都被部分和替代。
两个相邻元素间的距离被称为跨度,初始化均为。在每一次归约循环结束后,这个间隔就被乘以2。在第一次循环结束后,idata(全局数据指针)的偶数元素将被部分和替代。在第二次循环结束后,idata的每四个元素将会被新产生的部分和替代。
因为线程间无法同步,所以每个线程块产生的部分和被赋值回了主机,并且在哪儿进行串行求和。

3.4.3 改善并行归约的分化

3.4.3 交错配对的归约

3.5 展开循环

循环展开是一个尝试通过减少分支出现的频率和循环维护指令来优化循环的技术。
在循环展开中,循环主题在代码中药多次被编写,而不是只编写一次循环主题再使用另一个循环来反复执行的。
任何的封闭循环可将它的迭代次数减少或完全删除。
循环体的复制数量被称为循环展开因子,迭代次数就变为了原始循环迭代次数除以循环展开因此。

for (int i=0;i<100;i++)
	a[i]=b[i]+c[i];
for (int i=0;i<100;i+=2){
	a[i]=b[i]+c[i];
	a[i+1]=b[i+1]+c[i+1];
}
	

3.5.1 展开的归约

3.5.2 展开线程的归约

3.5.3 完全展开的归约

3.5.4 模板函数的归约

#include "CUDA_Header.cuh"

int invokeKernel();

int recursiveReduce(int* data, int const size) {
	if (size == 1)
		return data[0];
	int const stride = size / 2;
	for (size_t i = 0; i < stride; i++)
	{
		data[i] += data[i + stride];
	}
	return recursiveReduce(data, stride);
}

__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n) {

	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
	int* idata = g_idata + blockIdx.x * blockDim.x;

	if (idx >= n)
		return;
	for (size_t stride = 1; stride < blockDim.x; stride *= 2)
	{
		int index = 2 * stride * tid;
		if (index < blockDim.x)
			idata[index] += idata[index + stride];
		//if ((tid % (2 * stride)) == 0)
		//	idata[tid] += idata[tid + stride];
		__syncthreads();
	}
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}

__global__ void reduceInterleaved(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

	int* idata = g_idata + blockIdx.x * blockDim.x;

	if (idx >= n)
		return;

	int stride = blockDim.x / 2;
	for (size_t i = stride; i > 0; i >>= 1)
	{
		if (tid < i)
			idata[tid] += idata[tid + i];
		__syncthreads();
	}
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}

__global__ void reduceUnrolling2(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;
	//printf("tid %d , idx %d , blockIdx.x %d, blockDim.x %d , unrolling target %d \n",tid, idx, blockIdx.x, blockDim.x, idx + blockDim.x);
	int* idata = g_idata + blockIdx.x * blockDim.x * 2;

	if (idx + blockDim.x < n)
		g_idata[idx] += g_idata[idx + blockDim.x];

	__syncthreads();

	//__syncthreads()语句可以保证,线程块中的任一线程在进入下一次迭代之前,在当前迭代里,每个线程的所有部分和都被保存在了全局内存中,进入下一次迭代的所有线程都使用上一步产生的数值。在最后一个循环以后,整个线程块的和被保存进全局内存中
	//注解:__syncthreads()的意思应该是所有线程束中的

	//write result for this block to global mem
	//注解:当迭代结束后,结果保存在了idata的第一个元素里,idata是g_idata偏移后的地址
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}

__global__ void reduceUnrollWarp8(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

	int* idata = g_idata + blockIdx.x * blockDim.x * 8;

	//unrolling 8
	if (idx + 7 * blockDim.x < n) {
		int a1 = g_idata[idx];
		int a2 = g_idata[idx + blockDim.x];
		int a3 = g_idata[idx + 2 * blockDim.x];
		int a4 = g_idata[idx + 3 * blockDim.x];
		int b1 = g_idata[idx + 4 * blockDim.x];
		int b2 = g_idata[idx + 5 * blockDim.x];
		int b3 = g_idata[idx + 6 * blockDim.x];
		int b4 = g_idata[idx + 7 * blockDim.x];
		g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
	}
	__syncthreads();

	for (size_t stride = blockDim.x / 2; stride > 32; stride >>= 1)
	{
		if (tid < stride)
			idata[tid] += idata[tid + stride];

		__syncthreads();
	}
	if (tid < 32) {
		volatile int* vmem = idata;
		vmem[tid] += vmem[tid + 32];
		vmem[tid] += vmem[tid + 16];
		vmem[tid] += vmem[tid + 8];
		vmem[tid] += vmem[tid + 4];
		vmem[tid] += vmem[tid + 2];
		vmem[tid] += vmem[tid + 1];

	}
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}

__global__ void reduceCompleteUnrollWarp8(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

	int* idata = g_idata + blockIdx.x * blockDim.x * 8;
	if (idx + 7 * blockDim.x < n) {
		int a1 = g_idata[idx];
		int a2 = g_idata[idx + blockDim.x];
		int a3 = g_idata[idx + 2 * blockDim.x];
		int a4 = g_idata[idx + 3 * blockDim.x];
		int b1 = g_idata[idx + 4 * blockDim.x];
		int b2 = g_idata[idx + 5 * blockDim.x];
		int b3 = g_idata[idx + 6 * blockDim.x];
		int b4 = g_idata[idx + 7 * blockDim.x];
		g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
	}
	__syncthreads();

	if (blockDim.x >= 1024 && tid < 512)
		idata[tid] += idata[tid + 512];
	__syncthreads();

	if (blockDim.x >= 512 && tid < 256)
		idata[tid] += idata[tid + 256];
	__syncthreads();

	if (blockDim.x >= 256 && tid < 128)
		idata[tid] += idata[tid + 128];
	__syncthreads();

	if (blockDim.x >= 128 && tid < 64)
		idata[tid] += idata[tid + 64];
	__syncthreads();

	if (tid < 32) {
		volatile int* vsmem = idata;
		vsmem[tid] += vsmem[tid + 32];
		vsmem[tid] += vsmem[tid + 16];
		vsmem[tid] += vsmem[tid + 8];
		vsmem[tid] += vsmem[tid + 4];
		vsmem[tid] += vsmem[tid + 2];
		vsmem[tid] += vsmem[tid + 1];
	}

	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}


template <unsigned int iBlockSize>
__global__ void reduceCompleteUnroll(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;
	int* idata = g_idata + blockIdx.x * blockDim.x*8;

	if (idx + 7 * blockDim.x < n) {
		int a1 = g_idata[idx];
		int a2 = g_idata[idx + blockDim.x];
		int a3 = g_idata[idx + 2 * blockDim.x];
		int a4 = g_idata[idx + 3 * blockDim.x];
		int b1 = g_idata[idx + 4 * blockDim.x];
		int b2 = g_idata[idx + 5 * blockDim.x];
		int b3 = g_idata[idx + 6 * blockDim.x];
		int b4 = g_idata[idx + 7 * blockDim.x];
		g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
	}
	__syncthreads();

	if (blockDim.x >= 1024 && tid < 512)	
		idata[tid] += idata[tid + 512];		
	__syncthreads();
	if (blockDim.x >= 512 && tid < 256)		
		idata[tid] += idata[tid + 256];		
	__syncthreads();
	if (blockDim.x >= 256 && tid < 128)		
		idata[tid] += idata[tid + 128];		
	__syncthreads();
	if (blockDim.x >= 128 && tid < 64)		
		idata[tid] += idata[tid + 64];		
	__syncthreads();


	if (tid < 32) {
		volatile int* vsmem = idata;
		vsmem[tid] += vsmem[tid + 32];
		vsmem[tid] += vsmem[tid + 16];
		vsmem[tid] += vsmem[tid + 8];
		vsmem[tid] += vsmem[tid + 4];
		vsmem[tid] += vsmem[tid + 2];
		vsmem[tid] += vsmem[tid + 1];
	}

	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}
int main() {
	invokeKernel();

}

static int invokeKernel() {

	int dev = 0;
	cudaDeviceProp deviceProp;
	CHECK(cudaGetDeviceProperties(&deviceProp, dev));
	printf("device %d: %s \n", dev, deviceProp.name);

	CHECK(cudaSetDevice(dev));

	bool bResult = false;
	long   size = 1 << 24;
	printf("	with array size %d \n", size);


	int blockSize = 1024;

	dim3 block(blockSize, 1);

	dim3 grid((size + block.x - 1) / block.x, 1);
	grid.x /= 8;
	printf("grid %d block %d \n", grid.x, block.x);

	size_t bytes = size * sizeof(int);


	int* h_idata = (int*)malloc(bytes);
	int* h_odata = (int*)malloc(grid.x * sizeof(int));
	int* tmp = (int*)malloc(bytes);

	for (size_t i = 0; i < size; i++)
	{
		h_idata[i] = (int)(rand() & 0xFF);
	}
	memcpy(tmp, h_idata, bytes);

	clock_t iStart, iElaps;
	int gpu_sum = 0;

	int* d_idata = NULL;
	int* d_odata = NULL;

	CHECK(cudaMalloc((void**)&d_idata, bytes));
	CHECK(cudaMalloc((void**)&d_odata, grid.x * sizeof(int)));

	iStart = cpuSeconds();

	int cpu_sum = recursiveReduce(tmp, size);

	iElaps = cpuSeconds() - iStart;

	printf("cpu reduce elapsed %d sec cpu_sum: %d \n", iElaps, cpu_sum);
	//kernel 1: reduceNeighbored

	CHECK(cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice));
	CHECK(cudaDeviceSynchronize());
	iStart = cpuSeconds();

	reduceCompleteUnroll<1024> << <grid.x, block >> > (d_idata, d_odata, size);
	CHECK(cudaDeviceSynchronize());

	iElaps = cpuSeconds() - iStart;
	CHECK(cudaMemcpy(h_odata, d_odata, grid.x * sizeof(int), cudaMemcpyDeviceToHost));
	gpu_sum = 0;

	for (size_t i = 0; i < grid.x; i++)
	{
		gpu_sum += h_odata[i];
	}
	printf("gpu Kernel elapsed %d sec gpu_sum: %d <<<grid %d, block %d>>> \n", iElaps, gpu_sum, grid.x, block.x);

	free(h_idata);
	free(h_odata);

	cudaFree(d_idata);
	cudaFree(d_odata);
}

对于以上所讲的内容,使用Excel模拟了下内核的运算过程,应该更加容易理解了
带颜色的部分并非计算结果,而是处理的数据索引
在这里插入图片描述

严重性	代码	说明	项目	文件	行	禁止显示状态
错误		kernel launch from __device__ or __global__ functions requires separate compilation mode	CUDA	H:\C_CPP_CUDA\CUDA\CUDA\nestedHelloWorld.cu	15	

属性➔CUDA C/C++ ➔Common➔Generate Relocatable Device Code➔ 是(-rdc=true)
在这里插入图片描述
使用nvvp查看代码,NVIDIA Visual Profile在

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin

路径下有个nvvp.bat,双击可以打开这个窗口
从这个bat中也可以看到,nvvp.exe的路径是在

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\libnvvp


在这里插入图片描述

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.coloradmin.cn/o/1791855.html

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈,一经查实,立即删除!

相关文章

类的特殊成员函数

使用类的嵌套&#xff0c;并自定义析构函数 #include <iostream>using namespace std; class Per{ private:string name;int age;double hight;double weight; public:Per(string name,int age,double hight,double weight):name(name),age(age),hight(hight),weight(we…

WPF Binding对象、数据校验、数据转换

在WinForm中&#xff0c;我们要想对控件赋值&#xff0c;需要在后台代码中拿到控件对象进行操作&#xff0c;这种赋值形式&#xff0c;从根本上是无法实现界面与逻辑分离的。 在WPF中&#xff0c;微软引入了Binding对象&#xff0c;通过Binding&#xff0c;我们可以直接将控件与…

企业(园区)智慧能源双碳平台解决方案

园区作为工业企业集聚区&#xff0c;在提供了大量基础设施和公共服务的同时也成为了碳排放的主要源头。工业园区的耗能约占全社会总耗能的69%&#xff0c;碳排放占全国总排放约31%。工业园区节能、减耗、提质、减碳工作的落实&#xff0c;是我国实现碳达峰碳中和目标的必然要求…

Excel行列条件转换问题,怎么实现如图一到图二的效果?

图一 图二 如果数据比较&#xff0c;不建议一上来就用公式&#xff0c;风速值那一列的数据可以确定都是数值型数字&#xff0c;可以先试试用数据透视表做转换工具&#xff1a; 1.创建数据透视表 将采集时间放在行字段&#xff0c;测风放在列字段&#xff0c;风速放在值字段 2.…

RA8D1-Vision Board上OSPI-Flash实践

Vision-Board 开发板是 RT-Thread 推出基于瑞萨 Cortex-M85 架构 RA8D1 芯片,拥有Helium和TrustZone技术的加持,性能非常强大。 内核:480 MHz Arm Cortex-M85,包含Helium和TrustZone技术 存储:集成2MB/1MB闪存和1MB SRAM(包括TCM,512KB ECC保护) 外设:兼容xSPI的四线O…

Siemens-NXUG二次开发-创建倒斜角特征、边倒圆角特征、设置对象颜色、获取面信息[Python UF][20240605]

Siemens-NXUG二次开发-创建倒斜角特征、边倒圆角特征、设置对象颜色、获取面信息[Python UF][20240605] 1.python uf函数1.1 NXOpen.UF.Modeling.AskFaceData1.2 NXOpen.UF.Modeling.CreateChamfer1.3 NXOpen.UF.ModlFeatures.CreateBlend1.4 NXOpen.UF.Obj.SetColor 2.实体目标…

逐步掌握最佳Ai Agents框架-AutoGen 三 写新闻稿新方式

前言 今天带来的仍然是AutoGen基于AssistantAgent和UserProxyAgent的例子&#xff0c;以帮助大家一起消化目前最前卫的AI应用框架。这是一个AIGC最擅长&#xff0c;因为生成新闻稿嘛&#xff0c;同时又需要利用Agent的一个常规Demo。了解LangChain的同学&#xff0c;会通过对比…

C++基础编程100题-004 OpenJudge-1.1-06 空格分隔输出

更多资源请关注纽扣编程微信公众号 http://noi.openjudge.cn/ch0101/06/ 描述 读入一个字符&#xff0c;一个整数&#xff0c;一个单精度浮点数&#xff0c;一个双精度浮点数&#xff0c;然后按顺序输出它们&#xff0c;并且要求在他们之间用一个空格分隔。输出浮点数时保留…

v1.2.70-FastJson的AutoType机制研究

v1.2.70-FastJson的AutoType机制研究 最近在对接Alexa亚马逊语音技能&#xff0c;Smart Home Skill Apis时&#xff0c;有一个配置的JSON字符串是这样的&#xff1a; { "capabilityResources": {"friendlyNames": [{"type": "asset",…

智慧农业灌溉系统的主要工作原理

一、概述   智慧农业灌溉系统是一种基于传感器技术和智能控制技术的灌溉系统。它能够根据土壤湿度、气象条件、作物需水量等多种因素&#xff0c;自动控制灌溉水量和灌溉时间&#xff0c;实现精准灌溉。相比传统的手动灌溉和定时灌溉&#xff0c;智慧农业灌溉系统更加高效、准…

【Kubernetes】9-Pod控制器

一、什么是 pod 的控制器 Pod控制器&#xff0c;又称之为工作负载&#xff08;workload&#xff09;&#xff0c;是用于实现管理pod的中间层 确保pod资源符合预期状态&#xff1b;pod的资源故障时会进行重启&#xff1b; 当重启策略无效时&#xff0c;则会重新新建pod的资源 二…

通过ssr-echarts,服务端生成echarts图

ssr-echarts &#xff1a;一个开源项目&#xff0c;它能够服务端渲染 ECharts 图表&#xff0c;并直接生成 PNG 图片返回。该项目提供了多种主题&#xff0c;并且支持 GET 和 POST 请求。如果参数较多&#xff0c;建议使用 POST 方法。用户可以自己部署这个服务。 1. 服务端安装…

uniapp打包网页版配置页面窗口标题

问题描述&#xff1a; 项目场景&#xff1a;在uniapp打包网页的时候会出现项目标题为&#xff1a;uni-app 原因分析&#xff1a; 当时是直接在这里面修改的&#xff0c;但是打包之后发现标题还是没有改过来 解决方案&#xff1a; 后来才发现这里修改是没有用的&#xff0…

最新OpenAI免费API-openai api key获取方式

最近又开始准备LLM 应用开发&#xff0c;要用到api key&#xff0c;才发现过我之前免费发放的额度没了&#xff01;我都没咋用过&#xff0c;痛心&#x1f62d;&#x1f62d;&#x1f62d;&#xff01; 现在 OpenAI 有要求必须充值 5 刀才能使用&#xff0c;问就是没钱&#x…

人工智能芯片封装技术及应用趋势分析

简介人工智能&#xff08;AI&#xff09;、物联网&#xff08;IoT&#xff09;和大数据的融合正在开创全新的智能时代&#xff0c;以智能解决方案改变各行各业。人工智能芯片在支持人工智能学习和推理计算方面发挥着非常重要的作用&#xff0c;可实现各行各业的多样化应用。 本…

Arduino网页服务器:如何将Arduino开发板用作Web服务器

大家好&#xff0c;我是咕噜铁蛋&#xff01;今天&#xff0c;我将和大家分享一个有趣且实用的项目——如何使用Arduino开发板搭建一个简易的网页服务器。通过这个项目&#xff0c;你可以将Arduino连接到互联网&#xff0c;并通过网页控制或查询Arduino的状态。 一、项目背景与…

C#——随机类Random类

Random类 C#的Random类是用于生成随机数的类&#xff0c;属于System命名空间&#xff0c;可以生成各种类型的随机数&#xff0c;例如整型、双精度浮点型、布尔型等。 使用方法&#xff1a; 使用random数据类型关键字 声明一个random的变量 值使用new random 来实例化这个变量…

基于STM32开发的智能农业灌溉控制系统

目录 引言环境准备智能农业灌溉控制系统基础代码实现&#xff1a;实现智能农业灌溉控制系统 4.1 土壤湿度传感器数据读取4.2 水泵控制4.3 环境监测与数据记录4.4 用户界面与多功能显示应用场景&#xff1a;农业灌溉与环境监测问题解决方案与优化收尾与总结 1. 引言 随着农业…

大模型+编程,未来程序员躺平还是失业?

前言 随着科技的飞速发展&#xff0c;大模型与编程技术的结合正在逐步改变着我们的世界。**在这样的背景下&#xff0c;很多程序员开始担忧&#xff1a;未来的我们&#xff0c;是会“躺平”享受技术的红利&#xff0c;还是会因为技术变革而面临失业的风险&#xff1f;**今天&a…

Linuxftp服务001匿名登入

在Linux系统中搭建FTP&#xff08;File Transfer Protocol&#xff09;服务&#xff0c;可以让用户通过网络在服务器与其他客户端之间传输文件。它有几种登入模式&#xff0c;今天我们讲一下匿名登入。 操作系统 CentOS Stream9 操作步骤 首先我们先下载ftp [rootlocalhost…