PPMP_char3
PMPP char3 – Multidimensional grids and data
五一过后,有些工作要赶,抽出时间更新一下。这一章基本都熟练掌握,在做习题过程中有一些思考。这里涉及到了一点点GEMM(矩阵乘),GEMM有太多可深挖的了,推荐一篇博客How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog (siboehm.com)。另外,我还发现上一篇博客,有写错的地方,这篇博客的末尾做了一下勘误。这里记录我的个人理解,有不正确的地方,欢迎留言或者私信讨论。
课后习题
- In this chapter we implemented a matrix multiplication kernel that has each
thread produce one output matrix element. In this question, you will
implement different matrix-matrix multiplication kernels and compare them.
a. Write a kernel that has each thread produce one output matrix row. Fill in
the execution configuration parameters for the design.
b. Write a kernel that has each thread produce one output matrix column. Fill
in the execution configuration parameters for the design.
c. Analyze the pros and cons of each of the two kernel designs.
答案:
a 部分 (来自大模型)
__global__ void matrixMulRow(float *A, float *B, float *C, int m, int n, int k) {int row = blockIdx.y * blockDim.y + threadIdx.y;if (row < m) {for (int col = 0; col < k; ++col) {float sum = 0;for (int i = 0; i < n; ++i) {sum += A[row * n + i] * B[i * k + col];}C[row * k + col] = sum;}}
}
dim3 threadsPerBlock(1, 256);
dim3 blocksPerGrid(1, (m + threadsPerBlock.y - 1) / threadsPerBlock.y);
matrixMulRow<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, m, n, k);
b部分 (来自大模型)
__global__ void matrixMulCol(float *A, float *B, float *C, int m, int n, int k) {int col = blockIdx.x * blockDim.x + threadIdx.x;if (col < k) {for (int row = 0; row < m; ++row) {float sum = 0;for (int i = 0; i < n; ++i) {sum += A[row * n + i] * B[i * k + col];}C[row * k + col] = sum;}}
}
dim3 threadsPerBlock(256, 1);
dim3 blocksPerGrid((k + threadsPerBlock.x - 1) / threadsPerBlock.x, 1);
matrixMulCol<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, m, n, k);
c部分
假设A、B、C都是行主序;这里也不使用共享内存。一次读入一个缓存行,访问一行数据的话,存在访存局部性,需要用到的数据就在缓存中;访问一列数据的话,不存在访存局部性,需要用到的数据就不在缓存中,需要再读一个缓存行。假设一个缓存行时64B,对应16个float32。
访问A | 访问B | 访问C | |
---|---|---|---|
一个线程处理C中的一行 | 连续访问,只访问一行。访存次数=K/16 | 不连续访问,一次访问一列,访存次数=K;总共要访问N列,访问次数=N*K | 连续访问,只访问一行。访存次数=K/16 |
一个线程处理C中的一列 | 不连续访问,只访问一列。访存次数=M | 不连续访问,一次访问一列,访存次数=K;总共要访问N列,访问次数=N*K | 不连续访问,只访问一列。访存次数=M |
如果访存数据量较小,也就是a那种,可以直接放到寄存器中,这样访存cycle更短。
如果访存数据量较大,reg的容量就不够用,需要一部分存到L1上面,甚至L1也装不下,要从global里读取。涉及多个存储数据一致性、data hazard的问题,效率就很低。
- A matrix-vector multiplication takes an input matrix B and a vector C and
produces one output vector A. Each element of the output vector A is the dot
product of one row of the input matrix B and C, that is, A[i] = ΣB[i][j] * C[j].
For simplicity we will handle only square matrices whose elements are single-
precision floating-point numbers. Write a matrix-vector multiplication kernel and
the host stub function that can be called with four parameters: pointer to the output
matrix, pointer to the input matrix, pointer to the input vector, and the number of
elements in each dimension. Use one thread to calculate an output vector element.
答案:
其实,这道题就是参考gemm,实现gemv。就是把输出C的一个变量(C[i])映射到一个线程上,这个线程遍历j个变量,再各自点乘+求和(也就是j维度上做规约)。
以下是大模型写出来的程序,我看了下没啥问题,测试了下也是ok的。输入为A和B,输出为C。
#include <iostream>
#include <vector>
#include <random>
#include <cuda_runtime.h>#define CHECK(call) \
{ \const cudaError_t error = call; \if (error != cudaSuccess) { \std::cout << "Error: " << __FILE__ << ":" << __LINE__ << ", " << cudaGetErrorString(error) << std::endl; \exit(1); \} \
}__global__ void matrixVectorMultiply(const float* A, const float* B, float* C, int rows, int cols) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < rows) {float sum = 0.0f;for (int j = 0; j < cols; ++j) {sum += A[j] * B[i * cols + j];}C[i] = sum;}
}int main() {int rows = 1024;int cols = 768;std::vector<float> hostA(cols);std::vector<float> hostB(rows * cols);std::vector<float> hostC(rows);std::vector<float> resultC(rows);// 初始化输入数据std::random_device rd;std::mt19937 gen(rd());std::uniform_real_distribution<float> dis(-1.0, 1.0);for (int j = 0; j < cols; ++j) {hostA[j] = dis(gen);}for (int i = 0; i < rows; ++i) {for (int j = 0; j < cols; ++j) {hostB[i * cols + j] = dis(gen);}}// 分配设备内存float* deviceA;float* deviceB;float* deviceC;CHECK(cudaMalloc(&deviceA, cols * sizeof(float)));CHECK(cudaMalloc(&deviceB, rows * cols * sizeof(float)));CHECK(cudaMalloc(&deviceC, rows * sizeof(float)));// 将输入数据从主机复制到设备CHECK(cudaMemcpy(deviceA, hostA.data(), cols * sizeof(float), cudaMemcpyHostToDevice));CHECK(cudaMemcpy(deviceB, hostB.data(), rows * cols * sizeof(float), cudaMemcpyHostToDevice));// 启动内核int threadsPerBlock = 256;int blocksPerGrid = (rows + threadsPerBlock - 1) / threadsPerBlock;matrixVectorMultiply<<<blocksPerGrid, threadsPerBlock>>>(deviceA, deviceB, deviceC, rows, cols);CHECK(cudaGetLastError());// 将输出数据从设备复制到主机CHECK(cudaMemcpy(hostC.data(), deviceC, rows * sizeof(float), cudaMemcpyDeviceToHost));// 验证结果正确性for (int i = 0; i < rows; ++i) {float sum = 0.0f;for (int j = 0; j < cols; ++j) {sum += hostA[j] * hostB[i * cols + j];}resultC[i] = sum;}for (int i = 0; i < rows; ++i) {if (std::abs(hostC[i] - resultC[i]) > 1e-5) {std::cout << "Result verification failed at index " << i << std::endl;return 1;}}std::cout << "Result verification passed" << std::endl;// 释放设备内存CHECK(cudaFree(deviceA));CHECK(cudaFree(deviceB));CHECK(cudaFree(deviceC));return 0;
}
- Consider the following CUDA kernel and the corresponding host function that
calls it:a. What is the number of threads per block?
b. What is the number of threads in the grid?c. What is the number of blocks in the grid?
d. What is the number of threads that execute the code on line 05?
答案:a、16*32=512;b、48640;c、95;d、45000
- Consider a 2D matrix with a width of 400 and a height of 500. The matrix is
stored as a one-dimensional array. Specify the array index of the matrix
element at row 20 and column 10:
a. If the matrix is stored in row-major order.
b. If the matrix is stored in column-major order.
答案:a、20*400+10=8010
b、10*500+20=5020
Consider a 3D tensor with a width of 400, a height of 500, and a depth of 300
The tensor is stored as a one-dimensional array in row-major order.
Specify the array index of the tensor element at x 5 10, y 5 20, and z 5 5.
答案: 5 * (400 * 500) + 10 * 400 + 5 = 1004005
勘误
上一篇博客中,这个图里面的d_a的类型写错了,应该是int* d_a,而不是void* d_a。传入函数的是(void**)&d_a,函数结束后,d_a还是int*类型。
相关文章:

PPMP_char3
PMPP char3 – Multidimensional grids and data 五一过后,有些工作要赶,抽出时间更新一下。这一章基本都熟练掌握,在做习题过程中有一些思考。这里涉及到了一点点GEMM(矩阵乘),GEMM有太多可深挖的了&a…...
VulkanSDK Demos vkcube 编译失败
操作系统: Windows 11 23H2 Vulkan 版本: 1.3.2.280.0 Visual Studio 版本: 2022 在VulkanSDK/Demos目录下存在一个demo solution,其中包含两个project, vkcube和vkcubepp,两个分别为C语言和C写的示例程序, 但是直接编译这两个project时会编译失败,报了以下错误: fatal err…...

(二)Jetpack Compose 布局模型
前文回顾 (一)Jetpack Compose 从入门到会写-CSDN博客 首先让我们回顾一下上一篇文章中里提到过几个问题: ComposeView的层级关系,互相嵌套存在的问题? 为什么Compose可以实现只测量一次? ComposeView和…...
【Oracle impdp导入dmp文件(windows)】
Oracle impdp导入dmp文件(windows) 1、连接数据库2、创建与导出的模式相同名称的用户WIRELESS2,并赋予权限3、创建directory 的物理目录f:\radio\dmp,并把.dmp文件放进去4、连接新用户WIRELESS25、创建表空间的物理目录F:\radio\t…...

代数结构:5、格与布尔代数
16.1 偏序与格 偏序集:设P是集合,P上的二元关系“≤”满足以下三个条件,则称“≤”是P上的偏序关系(或部分序关系) (1)自反性:a≤a,∀a∈P; (2…...

如何使用DEEPL免费翻译PDF
如何使用DEEPL免费翻译PDF 安装DEEPL取消PDF限制 安装DEEPL 安装教程比较多,这里不重复。 把英文pdf拖进去,点翻译,在下面的框中有已经翻译完毕的文档。 但是存在两个问题 问题1:这些文档是加密的。 问题2:带有DeepL标…...
Spring-全面详解
Spring,就像是软件开发界的一个超级英雄,它让编写Java程序变得更简单、更灵活。想象一下,如果你要盖一栋大楼,Spring就是那个提供各种工具、框架和最佳实践的建筑大师,帮助你高效、优雅地搭建起整个项目。 Spring是啥&…...

QT自适应界面 处理高DPI 缩放比界面乱问题
1.pro文件添加 必须添加要不找不到 QT版本需要 5。4 以上才支持 QT widgets 2.main界面提前处理 // 1. 全局缩放使能QApplication::setAttribute(Qt::AA_EnableHighDpiScaling, true);// 2. 适配非整数倍缩放QGuiApplication::setHighDpiScaleFactorRoundingPolicy(Qt::High…...

序列到序列模型在语言识别Speech Applications中的应用 Transformer应用于TTS Transformer应用于ASR 端到端RNN
序列到序列模型在语言识别Speech Applications中的应用 A Comparative Study on Transformer vs RNN in Speech Applications 序列到序列(Seq2Seq)模型在语音识别(Speech Applications)中有重要的应用。虽然Seq2Seq模型最初是为了解决自然语言处理中的序列生成问题而设计的…...

【Linux】- Linux环境变量[8]
目录 环境变量 $符号 自行设置环境变量 环境变量 环境变量是操作系统(Windows、Linux、Mac)在运行的时候,记录的一些关键性信息,用以辅助系统运行。在Linux系统中执行:env命令即可查看当前系统中记录的环境变量。 …...

前端笔记-day04
文章目录 01-后代选择器02-子代选择器03-并集选择器04-交集选择器05-伪类选择器06-拓展-超链接伪类07-CSS特性-继承性08-CSS特性-层叠性09-CSS特性-优先级11-Emmet写法12-背景图13-背景图平铺方式14-背景图位置15-背景图缩放16-背景图固定17-background属性18-显示模式19-显示模…...

计算机字符集产生的历史与乱码
你好,我是 shengjk1,多年大厂经验,努力构建 通俗易懂的、好玩的编程语言教程。 欢迎关注!你会有如下收益: 了解大厂经验拥有和大厂相匹配的技术等 希望看什么,评论或者私信告诉我! 文章目录 一…...

Rerank进一步提升RAG效果
RAG & Rerank 目前大模型应用中,RAG(Retrieval Augmented Generation,检索增强生成)是一种在对话(QA)场景下最主要的应用形式,它主要解决大模型的知识存储和更新问题。 简述RAG without R…...

使用train.py----yolov7
准备工作 在训练之前,数据集的工作和配置环境的工作要做好 数据集:看这里划分数据集,训练自己的数据集。_划分数据集后如何训练-CSDN博客 划分数据集2,详细说明-CSDN博客 配置环境看这里 从0开始配置环境-yolov7_gpu0是inter g…...

机器学习第37周周报 GGNN
文章目录 week37 GGNN摘要Abstract一、文献阅读1. 题目2. abstract3. 网络架构3.1 数据处理部分3.2 门控图神经网络3.3 掩码操作 4. 文献解读4.1 Introduction4.2 创新点4.3 实验过程4.3.1 传感器设置策略4.3.2 数据集4.3.3 实验设置4.3.4 模型参数设置4.3.5 实验结果 5. 结论 …...

Baidu Comate:释放编码潜能,革新软件开发
Baidu Comate Baidu Comate,智能代码助手,凭借着文心大模型的强大支撑,结合了百度多年的编程实战数据和丰富的开源资源,形成了一款崭新的编码辅助利器。它不仅具备着高智能、多场景、价值创造的特质,更可广泛应用于各…...

MATLAB的Bar3函数调节渐变色(内附渐变色库.mat及.m文件免费下载链接)
一. colormap函数 可以使用colormap函数: t1[281.1,584.6, 884.3,1182.9,1485.2; 291.6,592.6,896,1197.75,1497.33; 293.8,596.4,898.6,1204.4,1506.4; 295.8,598,904.4,1209.0,1514.6];bar3(t1,1) set(gca,XTickLabel,{300,600,900,1200,1500},FontSize,10) set…...
使用 TensorFlow.js 和 OffscreenCanvas 实现实时防挡脸弹幕
首先,要理解我们的目标,我们将实时获取视频中的面部区域并将其周围的内容转为不透明以制造出弹幕的“遮挡效应”。 步骤一:环境准备 我们将使用 TensorFlow.js 的 Body-segmentation 库来完成面部识别部分,并使用 OffscreenCanv…...

【计算机网络篇】数据链路层(10)在物理层扩展以太网
文章目录 🍔扩展站点与集线器之间的距离🛸扩展共享式以太网的覆盖范围和站点数量 🍔扩展站点与集线器之间的距离 🛸扩展共享式以太网的覆盖范围和站点数量 以太网集线器一般具有8~32个接口,如果要连接的站点数量超过了…...

conan2 基础入门(03)-使用(msvc为例)
conan2 基础入门(03)-使用(msvc为例) 文章目录 conan2 基础入门(03)-使用(msvc为例)⭐准备生成profile文件预备文件和Code ⭐使用指令预览正确执行结果可能出现的问题 ⭐具体讲解conanconanfile.txt执行 install cmakeCMakeLists.txt生成项目构建 END ⭐准备 在阅读和学习本文…...

网络编程(Modbus进阶)
思维导图 Modbus RTU(先学一点理论) 概念 Modbus RTU 是工业自动化领域 最广泛应用的串行通信协议,由 Modicon 公司(现施耐德电气)于 1979 年推出。它以 高效率、强健性、易实现的特点成为工业控制系统的通信标准。 包…...

LBE-LEX系列工业语音播放器|预警播报器|喇叭蜂鸣器的上位机配置操作说明
LBE-LEX系列工业语音播放器|预警播报器|喇叭蜂鸣器专为工业环境精心打造,完美适配AGV和无人叉车。同时,集成以太网与语音合成技术,为各类高级系统(如MES、调度系统、库位管理、立库等)提供高效便捷的语音交互体验。 L…...

linux之kylin系统nginx的安装
一、nginx的作用 1.可做高性能的web服务器 直接处理静态资源(HTML/CSS/图片等),响应速度远超传统服务器类似apache支持高并发连接 2.反向代理服务器 隐藏后端服务器IP地址,提高安全性 3.负载均衡服务器 支持多种策略分发流量…...
【AI学习】三、AI算法中的向量
在人工智能(AI)算法中,向量(Vector)是一种将现实世界中的数据(如图像、文本、音频等)转化为计算机可处理的数值型特征表示的工具。它是连接人类认知(如语义、视觉特征)与…...
反射获取方法和属性
Java反射获取方法 在Java中,反射(Reflection)是一种强大的机制,允许程序在运行时访问和操作类的内部属性和方法。通过反射,可以动态地创建对象、调用方法、改变属性值,这在很多Java框架中如Spring和Hiberna…...

QT: `long long` 类型转换为 `QString` 2025.6.5
在 Qt 中,将 long long 类型转换为 QString 可以通过以下两种常用方法实现: 方法 1:使用 QString::number() 直接调用 QString 的静态方法 number(),将数值转换为字符串: long long value 1234567890123456789LL; …...

Unity | AmplifyShaderEditor插件基础(第七集:平面波动shader)
目录 一、👋🏻前言 二、😈sinx波动的基本原理 三、😈波动起来 1.sinx节点介绍 2.vertexPosition 3.集成Vector3 a.节点Append b.连起来 4.波动起来 a.波动的原理 b.时间节点 c.sinx的处理 四、🌊波动优化…...
Webpack性能优化:构建速度与体积优化策略
一、构建速度优化 1、升级Webpack和Node.js 优化效果:Webpack 4比Webpack 3构建时间降低60%-98%。原因: V8引擎优化(for of替代forEach、Map/Set替代Object)。默认使用更快的md4哈希算法。AST直接从Loa…...

uniapp 开发ios, xcode 提交app store connect 和 testflight内测
uniapp 中配置 配置manifest 文档:manifest.json 应用配置 | uni-app官网 hbuilderx中本地打包 下载IOS最新SDK 开发环境 | uni小程序SDK hbulderx 版本号:4.66 对应的sdk版本 4.66 两者必须一致 本地打包的资源导入到SDK 导入资源 | uni小程序SDK …...

Ubuntu Cursor升级成v1.0
0. 当前版本低 使用当前 Cursor v0.50时 GitHub Copilot Chat 打不开,快捷键也不好用,当看到 Cursor 升级后,还是蛮高兴的 1. 下载 Cursor 下载地址:https://www.cursor.com/cn/downloads 点击下载 Linux (x64) ,…...