GPU Puzzles讲解（二）

 GPU-Puzzles项目是一个很棒的学习cuda编程的项目，可以让你学习到GPU编程和cuda核心并行编程的概念，通过一个个小问题让你理解cuda的编程和调用，创建共享显存空间，实现卷积和矩阵乘法等

https://github.com/srush/GPU-Puzzlesicon-default.png?t=O83Ahttps://github.com/srush/GPU-Puzzles

本文是接续我上一篇文章的讲解，深入分析几个比较困难的puzzles，讲解实现和优化原理

，GPU Puzzles讲解（一）-CSDN博客GPU Puzzles是一个很棒的cuda编程学习仓库，本文是对其中题目介绍和思量讲解https://blog.csdn.net/lijj0304/article/details/142737859GPU Puzzles是一个很棒的cuda编程学习仓库，本文是对其中题目介绍和思量讲解https://blog.csdn.net/lijj0304/article/details/142737859

 Puzzle 11 - 1D Convolution

Implement a kernel that computes a 1D convolution between a and b and stores it in out. You need to handle the general case. You only need 2 global reads and 1 global write per thread.

实现一维卷积的教学，这里要求控制全局读写次数。一开始很自然的会想到把待卷积的向量和卷积核内容分别存下来。同时考虑到共享内存空间一般和每块的线程数直接关联，而卷积需要乘周围的几个元素，我们需要额外存储边界的几个元素。不仅要在共享内存中存储卷积核的内容，又为了控制到2次的全局读写，这里用了一个trick：即利用不需要存卷积核时的索引去存边界的额外元素

def conv_spec(a, b):out = np.zeros(*a.shape)len = b.shape[0]for i in range(a.shape[0]):out[i] = sum([a[i + j] * b[j] for j in range(len) if i + j < a.shape[0]])return outMAX_CONV = 4
TPB = 8
TPB_MAX_CONV = TPB + MAX_CONV
def conv_test(cuda):def call(out, a, b, a_size, b_size) -> None:i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.x# FILL ME IN (roughly 17 lines)shared_a = cuda.shared.array(TPB_MAX_CONV, numba.float32)shared_b = cuda.shared.array(MAX_CONV, numba.float32)if i < a_size:shared_a[local_i] = a[i]if local_i < b_size:shared_b[local_i] = b[local_i]else:local_j = local_i - b_sizeif i-b_size+TPB < a_size and local_j < b_size:shared_a[local_j+TPB] = a[i-b_size+TPB]cuda.syncthreads()total = 0.0for j in range(b_size):if i+j < a_size:total += shared_a[local_i+j] * shared_b[j]if i < a_size:out[i] = totalreturn call# Test 1SIZE = 6
CONV = 3
out = np.zeros(SIZE)
a = np.arange(SIZE)
b = np.arange(CONV)
problem = CudaProblem("1D Conv (Simple)",conv_test,[a, b],out,[SIZE, CONV],Coord(1, 1),Coord(TPB, 1),spec=conv_spec,
)
problem.show()

可视化效果

# 1D Conv (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             6 |             2 |

​

一维卷积Test2：向量长度超过TPB

out = np.zeros(15)
a = np.arange(15)
b = np.arange(4)
problem = CudaProblem("1D Conv (Full)",conv_test,[a, b],out,[15, 4],Coord(2, 1),Coord(TPB, 1),spec=conv_spec,
)
problem.show()

 可视化效果

# 1D Conv (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             8 |             2 |

​

Puzzle 12 - Prefix Sum

Implement a kernel that computes a sum over a and stores it in out. If the size of a is greater than the block size, only store the sum of each block.

​

这个是利用共享内存来计算前缀和，减少访存次数。可以利用两两归并相加的原理，最终结果存储在共享内存块的第1个。利用TPB分割向量长度，输出的out是每一段TPB长度的向量和

TPB = 8
def sum_spec(a):out = np.zeros((a.shape[0] + TPB - 1) // TPB)for j, i in enumerate(range(0, a.shape[-1], TPB)):out[j] = a[i : i + TPB].sum()return outdef sum_test(cuda):def call(out, a, size: int) -> None:cache = cuda.shared.array(TPB, numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.x# FILL ME IN (roughly 12 lines)if i < size:cache[local_i] = a[i]else:cache[local_i] = 0cuda.syncthreads()p = 2while (p <= TPB):if local_i % p == 0:cache[local_i] += cache[local_i+p/2]cuda.syncthreads()p *= 2if local_i == 0:out[cuda.blockIdx.x] = cache[local_i]return call# Test 1SIZE = 8
out = np.zeros(1)
inp = np.arange(SIZE)
problem = CudaProblem("Sum (Simple)",sum_test,[inp],out,[SIZE],Coord(1, 1),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Sum (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             1 |             1 |             7 |             4 |

​

前缀和Test2：

SIZE = 15
out = np.zeros(2)
inp = np.arange(SIZE)
problem = CudaProblem("Sum (Full)",sum_test,[inp],out,[SIZE],Coord(2, 1),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Sum (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             1 |             1 |             7 |             4 |

 ​​

 Puzzle 13 - Axis Sum

Implement a kernel that computes a sum over each column of a and stores it in out.

其实就是前缀和的运用，这里加入了batch作为列标计，每个batch计算和，out尺寸和bath大小一致

TPB = 8
def sum_spec(a):out = np.zeros((a.shape[0], (a.shape[1] + TPB - 1) // TPB))for j, i in enumerate(range(0, a.shape[-1], TPB)):out[..., j] = a[..., i : i + TPB].sum(-1)return outdef axis_sum_test(cuda):def call(out, a, size: int) -> None:cache = cuda.shared.array(TPB, numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.xbatch = cuda.blockIdx.y# FILL ME IN (roughly 12 lines)if i < size:cache[local_i] = a[batch, i]else:cache[local_i] = 0cuda.syncthreads() p = 2while (p <= TPB):if local_i % p == 0:cache[local_i] += cache[local_i+p/2]cuda.syncthreads()p *= 2if local_i == 0:out[batch, 0] += cache[local_i]return callBATCH = 4
SIZE = 6
out = np.zeros((BATCH, 1))
inp = np.arange(BATCH * SIZE).reshape((BATCH, SIZE))
problem = CudaProblem("Axis Sum",axis_sum_test,[inp],out,[SIZE],Coord(1, BATCH),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Axis SumScore (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             7 |             4 |

​

Puzzle 14 - Matrix Multiply!

Implement a kernel that multiplies square matrices a and b and stores the result in out.

矩阵乘法实现教学。第一直觉肯定就是利用二维的线程矩阵，循环行列遍历乘积求和。但是会有一个致命的问题就是线程矩阵比实际的矩阵小，而且想要控制全局读取的次数。这时就要利用好二维的共享空间，可以从对应行列出发，按行横向，列竖向顺序取部分的（TPB，TPB）矩阵块然后求乘积。从Test2的可视化图可以很直观看出这个思想

def matmul_spec(a, b):return a @ bTPB = 3
def mm_oneblock_test(cuda):def call(out, a, b, size: int) -> None:a_shared = cuda.shared.array((TPB, TPB), numba.float32)b_shared = cuda.shared.array((TPB, TPB), numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xj = cuda.blockIdx.y * cuda.blockDim.y + cuda.threadIdx.ylocal_i = cuda.threadIdx.xlocal_j = cuda.threadIdx.y# FILL ME IN (roughly 14 lines)total = 0.0for k in range(0, size, TPB): if local_j + k < size and i < size:a_shared[local_i, local_j] = a[i, local_j+k]if local_i + k < size and j < size:b_shared[local_i, local_j] = b[local_i+k, j]cuda.syncthreads()for local_k in range(min(TPB, size-k)):total += a_shared[local_i, local_k] * b_shared[local_k, local_j]if i < size and j < size:out[i, j] = totalreturn call# Test 1SIZE = 2
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).Tproblem = CudaProblem("Matmul (Simple)",mm_oneblock_test,[inp1, inp2],out,[SIZE],Coord(1, 1),Coord(TPB, TPB),spec=matmul_spec,
)
problem.show(sparse=True)

可视化效果

# Matmul (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             4 |             2 |

​

Test2大于线程矩阵的矩阵乘法：

SIZE = 8
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).Tproblem = CudaProblem("Matmul (Full)",mm_oneblock_test,[inp1, inp2],out,[SIZE],Coord(3, 3),Coord(TPB, TPB),spec=matmul_spec,
)
problem.show(sparse=True)

可视化效果

# Matmul (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             6 |             1 |            16 |             6 |

​

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

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

NVIDIA Hopper 架构深入

NVIDIA Hopper 架构深入

在 2022 年 NVIDIA GTC 主题演讲中，NVIDIA 首席执行官黄仁勋介绍了基于全新 NVIDIA Hopper GPU 架构的全新 NVIDIA H100 Tensor Core GPU。 文章目录 前言一、NVIDIA H100 Tensor Core GPU 简介二、NVIDIA H100 GPU 主要功能概述1. 新的流式多处理器 （SM） 具有许多性能和效率…
阅读更多...
leetcode58:最后一个单词的长度

leetcode58:最后一个单词的长度

给你一个字符串 s&#xff0c;由若干单词组成&#xff0c;单词前后用一些空格字符隔开。返回字符串中 最后一个 单词的长度。 单词 是指仅由字母组成、不包含任何空格字符的最大 子字符串 。 示例 1&#xff1a; 输入&#xff1a;s "Hello World" 输出&#xff…
阅读更多...
浅谈汽车智能座舱如何实现多通道音频

浅谈汽车智能座舱如何实现多通道音频

一、引言 随着汽车智能座舱的功能迭代发展&#xff0c;传统的 4 通道、6 通道、8 通道等音响系统难以在满足驾驶场景的需求&#xff0c;未来对于智能座舱音频质量和通道数会越来越高。接下来本文将浅析目前智能座舱如何实现音频功放&#xff0c;以及如何实现多路音频功放方案。…
阅读更多...
C语言文件操作(上)(27)

C语言文件操作(上)(27)

文章目录 前言一、为什么要用文件&#xff1f;二、什么是文件&#xff1f;程序文件数据文件文件名文件类型文件缓冲区文件指针 三、流流的概念标准流 总结 前言 C语言可以直接操作文件&#xff0c;如果你是第一次听说这个特性&#xff0c;可能会眼前一亮&#xff0c;感到惊奇  …
阅读更多...
MongoDB的安装与增删改查基本操作

MongoDB的安装与增删改查基本操作

MongoDB是一种非关系型数据库，是NoSQL语言，但是又是最接近关系型数据库的。内部存储不是表结构，但是可以对数据进行表结构的操作。 一、安装 在官网：Download MongoDB Community Server | MongoDB下载系统对应的版本进行安装即可 二、编辑器 在安装MongoDB后会自带一个编…
阅读更多...
图片格式入门

图片格式入门

主要参考资料&#xff1a; 常见的图片格式介绍: https://blog.csdn.net/cnds123/article/details/127165291 目录 像素图与矢量图像素图&#xff08;pixel image&#xff09;矢量图&#xff08;Vector graphics&#xff09; 像素图与矢量图 像素图&#xff08;pixel image&…
阅读更多...
D3.js中国地图可视化

D3.js中国地图可视化

1、项目介绍 该项目来自Github&#xff0c;基于D3.js中国地图可视化。 D3.js is a JavaScript library for manipulating documents based on data. It uses HTML, SVG, and CSS to display data. The full name of D3 is "Data-Driven Documents," which means it a…
阅读更多...
话术挂断之后是否处理事件

话术挂断之后是否处理事件

文章目录 前言联系我们解决方案方案一方案二 前言 流程&#xff1a;自动外呼进入机器人话术。问题&#xff1a;在机器人放音时用户挂断后&#xff0c;话术还会继续匹配流程&#xff0c;如果匹配上的是放音节点&#xff0c;还会进行放音&#xff0c;那么在数据库表conversation…
阅读更多...
Redis 缓存策略详解：提升性能的四种常见模式

Redis 缓存策略详解：提升性能的四种常见模式

在现代分布式系统中&#xff0c;缓存是提升性能和减轻数据库负载的关键组件。Redis 作为一种高性能的内存数据库&#xff0c;被广泛应用于缓存层。本文将深入探讨几种常用的 Redis 缓存策略&#xff0c;包括旁路缓存模式&#xff08;Cache-Aside Pattern&#xff09;、读穿透模…
阅读更多...
阿里云 SAE Web：百毫秒高弹性的实时事件中心的架构和挑战

阿里云 SAE Web：百毫秒高弹性的实时事件中心的架构和挑战

作者&#xff1a;胡志广(独鳌) 背景 Serverless 应用引擎 SAE 事件中心主要面向早期的 SAE 控制台只有针对于应用维度的事件&#xff0c;这个事件是 K8s 原生的事件&#xff0c;其实绝大多数的用户并不会关心&#xff0c;同时也可能看不懂。而事件中心&#xff0c;是希望能够…
阅读更多...
SpringBoot MyBatis连接数据库设置了encoding=utf-8还是不能用中文来查询

SpringBoot MyBatis连接数据库设置了encoding=utf-8还是不能用中文来查询

properties的MySQL连接时已经指定了字符编码格式&#xff1a; url: jdbc:mysql://localhost:3306/sky_take_out?useUnicodetrue&characterEncodingutf-8使用MyBatis查询&#xff0c;带有中文参数&#xff0c;查询出的内容为空。 执行的语句为&#xff1a; <select id&…
阅读更多...
Tensorflow2.0

Tensorflow2.0

Tensorflow2.0 有深度学习基础的建议直接看class3 class1 介绍 人工智能3学派 行为主义:基于控制论&#xff0c;构建感知-动作控制系统。(控制论&#xff0c;如平衡、行走、避障等自适应控制系统) 符号主义:基于算数逻辑表达式&#xff0c;求解问题时先把问题描述为表达式…
阅读更多...
【Kubernetes】常见面试题汇总（五十三）

【Kubernetes】常见面试题汇总（五十三）

目录 118. pod 状态为 ErrlmagePull &#xff1f; 119.探测存活 pod 状态为 CrashLoopBackOff &#xff1f; 特别说明&#xff1a; 题目 1-68 属于【Kubernetes】的常规概念题&#xff0c;即 “ 汇总&#xff08;一&#xff09;~&#xff08;二十二&#xff09;” 。…
阅读更多...
使用NumPy进行线性代数的快速指南

使用NumPy进行线性代数的快速指南

介绍 NumPy 是 Python 中用于数值计算的基础包。它提供了处理数组和矩阵的高效操作&#xff0c;这对于数据分析和科学计算至关重要。在本指南中&#xff0c;我们将探讨 NumPy 中可用的一些基本线性代数操作&#xff0c;展示如何通过运算符重载和内置函数执行这些操作。 元素级…
阅读更多...
【汇编语言】寄存器（CPU工作原理）（一）—— 寄存器的基础知识及存储

【汇编语言】寄存器（CPU工作原理）（一）—— 寄存器的基础知识及存储

文章目录 前言1. 寄存器2. 通用寄存器3. 字在寄存器中的存储结语 前言 &#x1f4cc; 汇编语言是很多相关课程&#xff08;如数据结构、操作系统、微机原理&#xff09;的重要基础。但仅仅从课程的角度出发就太片面了&#xff0c;其实学习汇编语言可以深入理解计算机底层工作原…
阅读更多...
docker拉取镜像推送到阿里云镜像仓库

docker拉取镜像推送到阿里云镜像仓库

文章目录 个人GitHub仓库地址镜像源地址 Docker拉取失败&#xff0c;利用github将镜像推送到阿里云 docker_image_pusher hub-mirror仓库 1、windows没有升级&#xff0c;用不了WSL。可以使用wsl&#xff0c;配合docker desktop直接拉取镜像&#xff0c;windows10安装WSL2及使…
阅读更多...
Linux和指令初识

Linux和指令初识

前言 Linux是我们在服务器中常用的操作系统&#xff0c;我们有必要对这个操作系统有足够的认识&#xff0c;并且能够使相关的指令操作。今天我们就来简单的认识一下这个操作的前世今生&#xff0c;并且介绍一些基础的指令操作 Linux的前世今生 要说Linux&#xff0c;还得从U…
阅读更多...
【C++11】右值引用和移动语义

【C++11】右值引用和移动语义

文章目录 左值和右值的概念左值右值 左值与右值引用移动语义的概念std::move 的作用使用std::move的注意事项 右值引用的使用场景右值引用的其他概念万能引用完美转发std::forward万能引用和右值引用的区别 新的类功能默认成员函数 左值和右值的概念 在C中&#xff0c; 左值 和…
阅读更多...
MATLAB下的RSSI定位程序，二维平面上的定位，基站数量可自适应

MATLAB下的RSSI定位程序，二维平面上的定位，基站数量可自适应

文章目录 引言程序概述程序代码运行结果待定位点、锚点、计算结果显示待定位点和计算结果坐标 引言 随着无线通信技术的发展&#xff0c;基于 R S S I RSSI RSSI&#xff08;接收信号强度指示&#xff09;的方法在定位系统中变得越来越流行。 R S S I RSSI RSSI定位技术特别适…
阅读更多...
面试题之- null和undefined的区别

面试题之- null和undefined的区别

前言 首先undefined和null都是基本数据类型&#xff0c;这两个基本数据类型分别都只有一个值&#xff0c;就是undefined和null。 undefined代表的含义是未定义&#xff0c;null代表的的含义是空对象&#xff0c;一般变量声明了但是还有没有定义的时候会返回undefined&#xf…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023