CUDA Unity Compute Shader 3

计划
这应该是第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


在这里插入图片描述

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

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

鸿蒙轻内核M核源码分析系列六 任务及任务调度(3)任务调度模块

调度&#xff0c;Schedule也称为Dispatch&#xff0c;是操作系统的一个重要模块&#xff0c;它负责选择系统要处理的下一个任务。调度模块需要协调处于就绪状态的任务对资源的竞争&#xff0c;按优先级策略从就绪队列中获取高优先级的任务&#xff0c;给予资源使用权。本文我们…

java多线程原理

1.线程创建与启动&#xff1a;通过继承Thread类或实现Runnable接口创建线程&#xff0c;并调用start()方法启动线程。 1.线程状态&#xff1a;线程在其生命周期中有多种状态&#xff0c;包括新建、运行、阻塞、死亡等。了解这些状态以及如何在它们之间转换对于管理线程至关重要…

完美解决 mysql 报错ERROR 1524 (HY000): Plugin ‘mysql_native_password‘ is not loaded

文章目录 错误描述错误原因解决步骤 跟着我下面的步骤走&#xff0c;解决你的问题&#xff0c;如果解决不了 私信我来给你解决 错误描述 执行ALTER USER root% IDENTIFIED WITH mysql_native_password BY 123456;报错ERROR 1524 (HY000): Plugin mysql_native_password is not …

RPA实战演练UiBot6.0校园学生教评机器人

前言 校园学生教评机器人&#xff0c;也称为全自动校园教评RPA&#xff08;Robotic Process Automation&#xff0c;机器人流程自动化&#xff09;机器人&#xff0c;是一种利用软件机器人技术来模拟和执行学生教评流程中的各项任务和操作的智能化系统。以下是关于校园学生教评…

Mamba v2诞生:1 儒(Transformers)释(SSD)道(Mamba)本是一家?!

大模型技术论文不断&#xff0c;每个月总会新增上千篇。本专栏精选论文重点解读&#xff0c;主题还是围绕着行业实践和工程量产。若在某个环节出现卡点&#xff0c;可以回到大模型必备腔调或者LLM背后的基础模型新阅读。而最新科技&#xff08;Mamba,xLSTM,KAN&#xff09;则提…

因子区间[牛客周赛44]

思路分析: 我们可以发现125是因子个数的极限了,所以我们可以用二维数组来维护第几个数有几个因子,然后用前缀和算出来每个区间合法个数,通过一个排列和从num里面选2个 ,c num 2 来计算即可 #include<iostream> #include<cstring> #include<string> #include…

数据库中锁的机制和MVCC协议以及隔离级别

文章目录 数据库中的锁锁与索引的关系释放锁的时机乐观锁与悲观锁行锁与表锁共享锁与排它锁意向锁记录锁、间隙锁和临键锁记录锁间隙锁临键锁 锁优化方案 MVCC协议MySQL的隔离级别脏读和幻读快照读和当前读 版本链Read ViewRead View 与已提交读Read View 与可重复读m_up_limit…

8. C#多线程基础概念

文章目录 一. 目标二. 技能介绍① 进程和线程② 为什么需要多线程③ C#实现多线程的方式④ 线程的操作(创建_终止_挂起_恢复) 一. 目标 进程和线程基本概念为什么需要多线程?C#实现多线程的方式?线程Thread的创建,终止,挂起和恢复? 二. 技能介绍 ① 进程和线程 什么是进程…

F5G城市光网,助力“一网通城”筑基数字中国

《淮南子》中说&#xff0c;“临河而羡鱼&#xff0c;不如归家织网”。 这句话在后世比喻为做任何事情都需要提前做好准备&#xff0c;有了合适的工具&#xff0c;牢固的基础&#xff0c;各种难题也会迎刃而解。 如今&#xff0c;数字中国发展建设如火如荼&#xff0c;各项任务…

C语言 | Leetcode C语言题解之第119题杨辉三角II

题目&#xff1a; 题解&#xff1a; int* getRow(int rowIndex, int* returnSize) {*returnSize rowIndex 1;int* row malloc(sizeof(int) * (*returnSize));row[0] 1;for (int i 1; i < rowIndex; i) {row[i] 1LL * row[i - 1] * (rowIndex - i 1) / i;}return row…

排序-快速排序

前言 本期主角 是这个小老头 图灵奖得主&#xff0c; 美国国家科学院外籍院士&#xff0c; 美国国家工程院外籍院士&#xff0c; 英国皇家工程院院士&#xff0c; 英国皇家学会院士 鼓掌&#x1f44f;&#x1f44f;&#x1f44f; 感觉这个小老头很叼噢(确实很叼) 从标…

MQTT.FX的使用

背景 在如今物联网的时代下&#xff0c;诞生了许多的物联网产品&#xff0c;这些产品通过BLE、WIFI、4G等各种各样的通信方式讲数据传输到各种各样的平台。 除了各个公司私有的云平台外&#xff0c;更多的初学者会接触到腾讯云、阿里云之类的平台。设备接入方式也有着多种多样…

react基础学习 JSX

JSX的测试网站 Babel Babel 可以测试代码的效果 JSX实现map列表 注意 key不一样&#xff08;使用遍历的时候&#xff09; 简单条件渲染 复杂条件渲染 绑定事件 function App() {const colorse (e)>{console.log("测试点击",e);}const colorse1 (name)>{…

数仓建模—指标体系指标拆解和选取

数仓建模—指标拆解和选取 第一节指标体系初识介绍了什么是指标体系 第二节指标体系分类分级和评价管理介绍了指标体系管理相关的,也就是指标体系的分级分类 这一节我们看一下指标体系的拆解和指标选取,这里我们先说指标选取,其实在整个企业的数字化建设过程中我们其实最…

vuInhub靶场实战系列-DC-6实战

免责声明 本文档仅供学习和研究使用,请勿使用文中的技术源码用于非法用途,任何人造成的任何负面影响,与本人无关。 目录 免责声明前言一、环境配置二、信息收集2.1 主机发现2.1.1 nmap扫描存活主机2.1.2 arp-scan扫描存活主机 2.2 端口扫描2.3 指纹识别2.3.1 尝试指纹识别2.3.…

2024050302-重学 Java 设计模式《实战享元模式》

重学 Java 设计模式&#xff1a;实战享元模式「基于Redis秒杀&#xff0c;提供活动与库存信息查询场景」 一、前言 程序员&#x1f468;‍&#x1f4bb;‍的上下文是什么&#xff1f; 很多时候一大部分编程开发的人员都只是关注于功能的实现&#xff0c;只要自己把这部分需求…

现代控制中可控性的Gramian判据

知乎三角猫frank对于这块内容写的非常好&#xff0c;但这个输入的构造还是很难过于没头没尾 数学好的人&#xff0c;可能看一眼根据形式就能推出gramian的构造&#xff0c;但对我这种比较钻牛角尖的人&#xff0c;我就想有一个逻辑链条——gramian是怎么被构造出来的&#xff1…

eNSP学习——配置RIPv2认证

目录 主要命令 原理概述 实验目的 实验内容 实验拓扑 实验编址 实验步骤 1、基本配置 2、搭建RIP网络 3、模拟网络攻击 4、配置RIPv2简单验证 5、配置RIPv2 MD5密文验证 需要eNSP各种配置命令的点击链接自取&#xff1a;华为&#xff45;NSP各种设备配置命令大全PD…

区块链游戏(链游)安全防御:抵御攻击的策略与实践

一、引言 区块链游戏&#xff0c;或称为链游&#xff0c;近年来随着区块链技术的普及而迅速崛起。然而&#xff0c;如同其他任何在线平台一样&#xff0c;链游也面临着各种安全威胁。本文将探讨链游可能遭遇的攻击类型以及如何通过有效的策略和技术手段进行防御。 二、链游可…

如何手动批准内核扩展 Tuxera NTFS for mac内核扩展需要批准 内核扩展怎么打开

在了解如何手动批准内核扩展之前&#xff0c;我们应该先了解什么叫做内核扩展。内核扩展又被称为KEXT&#xff0c;通过它可以实现macOS系统与软件组件之间的交互&#xff0c;例如磁盘管理、任务管理和内存管理等等。 kext 是内核扩展&#xff08;Kernel Extension&#xff09;…