[8] CUDA之向量点乘和矩阵乘法

CUDA之向量点乘和矩阵乘法

  • 计算类似矩阵乘法的数学运算

1. 向量点乘

  • 两个向量点乘运算定义如下:
#真正的向量可能很长,两个向量里边可能有多个元素
(X1,Y1,Z1) * (Y1,Y2,Y3) = X1Y1 + X2Y2 + X3Y3
  • 这种原始输入是两个数组而输出却缩减为一个(单一值)的运算,在CUDA里边叫规约运算
  • 该运算对应的内核函数如下:
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 1024
#define threadsPerBlock 512__global__ void gpu_dot(float* d_a, float* d_b, float* d_c) {//Declare shared memory__shared__ float partial_sum[threadsPerBlock];int tid = threadIdx.x + blockIdx.x * blockDim.x;//Calculate index for shared memory int index = threadIdx.x;//Calculate Partial Sumfloat sum = 0;while (tid < N){sum += d_a[tid] * d_b[tid];tid += blockDim.x * gridDim.x;}// Store partial sum in shared memorypartial_sum[index] = sum;// synchronize threads __syncthreads();// Calculating partial sum for whole block in reduce operationint i = blockDim.x / 2;while (i != 0) {if (index < i)partial_sum[index] += partial_sum[index + i];__syncthreads();i /= 2;}//Store block partial sum in global memoryif (index == 0)d_c[blockIdx.x] = partial_sum[0];
}
  • 每个块都有单独的一份共享内存副本,所以每个线程ID索引到的共享内存只能是当前块自己的那个副本

  • 当线线程总数小于元素数量的时候,它也会循环将tid 索引累加偏移到当前线程总数,继续索引下一对元素,并进行计算。每个线程得到的部分和结果被写入到共享内存。我们将继续使用共享内存上的这些线程的部分和计算出当前块的总体部分和

  • 在对共享内存中的数据读取之前,必须确保每个线程都已经完成了对共享内存的写入,可以通过 __syncthreads() 同步函数做到这一点

  • 计算当前块部分和的方法:

    • 1.让一个线程串行循环将这些所有的线程的部分和进行累加
    • 2.并行化:每个线程累加2个数的操作,并将每个线程的得到的1个结果覆盖写入这两个数中第一个数的位置,因为每个线程都累加了2个数,因此可以在第一个数中完成操作(此时第一个数就是两个数的和),后边对剩余的部分重复这个过程,类似将所有数对半分组相加吧,一组两个数,加完算出的新的结果作为新的被加数
    • 上述并行化的方法是通过条件为(i!=0)while循环进行的,后边的计算类似二分法,重复计算中间值与下一值的和,知道总数为0
  • main函数如下:

int main(void) 
{//Declare Host Arrayfloat *h_a, *h_b, h_c, *partial_sum;//Declare device Arrayfloat *d_a, *d_b, *d_partial_sum;//Calculate total number of blocks per gridint block_calc = (N + threadsPerBlock - 1) / threadsPerBlock;int blocksPerGrid = (32 < block_calc ? 32 : block_calc);// allocate memory on the host sideh_a = (float*)malloc(N * sizeof(float));h_b = (float*)malloc(N * sizeof(float));partial_sum = (float*)malloc(blocksPerGrid * sizeof(float));// allocate the memory on the devicecudaMalloc((void**)&d_a, N * sizeof(float));cudaMalloc((void**)&d_b, N * sizeof(float));cudaMalloc((void**)&d_partial_sum, blocksPerGrid * sizeof(float));// fill the host array with datafor (int i = 0; i<N; i++) {h_a[i] = i;h_b[i] = 2;}cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);//Call kernel gpu_dot << <blocksPerGrid, threadsPerBlock >> >(d_a, d_b, d_partial_sum);// copy the array back to host memorycudaMemcpy(partial_sum, d_partial_sum, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);// Calculate final dot product on hosth_c = 0;for (int i = 0; i<blocksPerGrid; i++) {h_c += partial_sum[i];}printf("The computed dot product is: %f\n", h_c);
}
  • 在main函数中添加如下代码,检查该点乘结果是否正确:
#define cpu_sum(x) (x*(x+1))if (h_c == cpu_sum((float)(N - 1))){printf("The dot product computed by GPU is correct\n");}else{printf("Error in dot product computation");}// free memory on host and devicecudaFree(d_a);cudaFree(d_b);cudaFree(d_partial_sum);free(h_a);free(h_b);free(partial_sum);

在这里插入图片描述

矩阵乘法

  • 矩阵乘法A*B,A的行数需等于B的列数,将A的某行与B的所有的列进行点乘,然后对A的每一行以此类推
  • 下面将给出不使用共享内存和使用共享内存的内核函数来计算矩阵乘法
  • 先给出不使用共享内存的内核:
//Matrix multiplication using shared and non shared kernal
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#define TILE_SIZE 2//Matrix multiplication using non shared kernel
__global__ void gpu_Matrix_Mul_nonshared(float* d_a, float* d_b, float* d_c, const int size)
{int row, col;col = TILE_SIZE * blockIdx.x + threadIdx.x;row = TILE_SIZE * blockIdx.y + threadIdx.y;for (int k = 0; k < size; k++){d_c[row * size + col] += d_a[row * size + k] * d_b[k * size + col];}
}
  • 每个元素的线性索引可以这样计算:用它的行号乘以矩阵的宽度,再加上它的列号即可
  • 使用共享内存的内核函数如下:
// Matrix multiplication using shared kernel
__global__ void gpu_Matrix_Mul_shared(float *d_a, float *d_b, float *d_c, const int size)
{int row, col;//Defining Shared Memory,共享内存数量=块数__shared__ float shared_a[TILE_SIZE][TILE_SIZE];__shared__ float shared_b[TILE_SIZE][TILE_SIZE];col = TILE_SIZE * blockIdx.x + threadIdx.x;row = TILE_SIZE * blockIdx.y + threadIdx.y;for (int i = 0; i< size / TILE_SIZE; i++) {shared_a[threadIdx.y][threadIdx.x] = d_a[row* size + (i*TILE_SIZE + threadIdx.x)];shared_b[threadIdx.y][threadIdx.x] = d_b[(i*TILE_SIZE + threadIdx.y) * size + col];__syncthreads(); for (int j = 0; j<TILE_SIZE; j++)d_c[row*size + col] += shared_a[threadIdx.y][j] * shared_b[j][threadIdx.x];__syncthreads(); }
}
  • 主函数代码如下:
int main()
{const int size = 4;//Define Host Arrayfloat h_a[size][size], h_b[size][size],h_result[size][size];//Defining device Arrayfloat *d_a, *d_b, *d_result; //Initialize host Arrayfor (int i = 0; i<size; i++){for (int j = 0; j<size; j++){h_a[i][j] = i;h_b[i][j] = j;}}cudaMalloc((void **)&d_a, size*size*sizeof(int));cudaMalloc((void **)&d_b, size*size * sizeof(int));cudaMalloc((void **)&d_result, size*size* sizeof(int));//copy host array to device arraycudaMemcpy(d_a, h_a, size*size* sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, size*size* sizeof(int), cudaMemcpyHostToDevice);//Define grid and block dimensionsdim3 dimGrid(size / TILE_SIZE, size / TILE_SIZE, 1);dim3 dimBlock(TILE_SIZE, TILE_SIZE, 1);//gpu_Matrix_Mul_nonshared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);gpu_Matrix_Mul_shared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);cudaMemcpy(h_result, d_result, size*size * sizeof(int),	cudaMemcpyDeviceToHost);printf("The result of Matrix multiplication is: \n");for (int i = 0; i< size; i++){for (int j = 0; j < size; j++){printf("%f   ", h_result[i][j]);}printf("\n");}cudaFree(d_a);cudaFree(d_b);cudaFree(d_result);return 0;
}

在这里插入图片描述

  • ——————END——————

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

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

相关文章

linux 查看csv文件,按指定列聚合 排序

在Linux中&#xff0c;你可以使用awk工具来查看CSV文件的内容&#xff0c;并按照指定的列进行聚合。awk是一种强大的文本处理工具&#xff0c;它可以处理文本文件中的数据&#xff0c;并根据条件执行相应的操作。 以下是一个示例&#xff0c;假设你有一个名为data.csv的CSV文件…

单点登录(JWT实现)

单点登陆的英文名是&#xff1a;Single Sign On&#xff08;简称SSO&#xff09;&#xff0c;只需要登陆一次&#xff0c;就可以访问所有信任的应用系统。 在单体项目中&#xff0c;我们登陆之后可以把验证用户信息的值放入session中&#xff0c;单个tomcat中的session是可以共…

C++ 数据结构算法 学习笔记(33) -查找算法及企业级应用

C 数据结构算法 学习笔记(33) -查找算法及企业级应用 数组和索引 日常生活中&#xff0c;我们经常会在电话号码簿中查阅“某人”的电话号码&#xff0c;按姓查询或者按字母排 序查询&#xff1b;在字典中查阅“某个词”的读音和含义等等。在这里&#xff0c;“电话号码簿”和…

【FPGA】Verilog:2-bit 二进制比较器的实现(2-bit binary comparator)

解释 2-bit 二进制比较器仿真结果及过程说明(包括真值表和卡诺图) 真值表和卡洛图如下: 2-bit Binary Comparator A1 A2 B1

写好的文章怎样联系媒体投稿?

作为单位信息宣传的桥梁,我肩负着将单位的每一次活动、每一项成就转化为社会认可与赞美的重任。初涉此职,我满腔热血,以为凭借扎实的文字功底与不懈的努力,便能在各大媒体平台上为单位赢得一席之地。然而,现实很快就给了我一记响亮的耳光。 我最初采取的是最直接的方式——邮箱…

QT 使用QLsitView 实现多个子项选中取消效果

文章目录 效果图概述部分代码总结 效果图 概述 整个界面的布局介绍请看这篇博客想要的到这种自由选择中的Item效果&#xff0c;需要使用到Model-view的思想&#xff0c;每个item中都要存放一个标志位&#xff0c;用在Paint函数去判断是否绘制为按下的状态。每次item被点击时&a…

记录下所遇到远程桌面连接方法winSCP跟mstsc

之前公司使用过连接远程桌面&#xff0c;今天又遇到要使用远程桌面问题&#xff0c;来记录下。 之前公司使用的是winR 然后回车弹出 后面按照用户名密码就能登陆了 今天后台给了我一张图片准备接着用这个方法&#xff0c;后台就说这个东西要下载winSCP 后台发给我图片 然后去…

[笔试强训day09]

文章目录 BC146 添加逗号DP2 跳台阶JZ61 扑克牌顺子解法一&#xff1a;排序模拟解法二&#xff1a;规律哈希 BC146 添加逗号 BC146 添加逗号 #include<iostream> #include<string>using namespace std;int main() {string s;cin>>s;string ans;for(int i0;i…

LLaMa系列模型详解(原理介绍、代码解读):LLaMA 2

LLaMA 2 大型语言模型&#xff08;LLMs&#xff09;作为高度能力的人工智能助手&#xff0c;在需要跨多个领域专家知识的复杂推理任务中表现出巨大潜力&#xff0c;包括编程和创意写作等专业领域。它们通过直观的聊天界面与人类互动&#xff0c;这导致了快速和广泛的公众采用。…

军工单位如何做到安全跨网文件交换与导出的

在现代信息化战争中&#xff0c;军工单位在信息安全方面的需求尤为突出。跨网文件交换与导出作为军工单位日常运营的重要环节&#xff0c;面临着网络带宽限制、数据安全风险、合规性要求和传输稳定性等挑战。下面&#xff0c;我们将从以下几个方面探讨军工单位如何实现安全、高…

【C++】开源:RabbitMQ安装与配置使用(SimpleAmqpClient)

&#x1f60f;★,:.☆(&#xffe3;▽&#xffe3;)/$:.★ &#x1f60f; 这篇文章主要介绍。 无专精则不能成&#xff0c;无涉猎则不能通。——梁启超 欢迎来到我的博客&#xff0c;一起学习&#xff0c;共同进步。 喜欢的朋友可以关注一下&#xff0c;下次更新不迷路&#x1…

用队列实现栈 用栈实现队列 设计循环队列

用队列实现栈 思路 栈的特点&#xff1a;后进先出 队列的特点&#xff1a;先进先出 使用两个队列实现栈&#xff1a; 我们可以使用两个队列&#xff0c;一个队列为&#xff1a;空队列&#xff0c;一个队列为&#xff1a;非空队列 当我们要出队列时&#xff1a; 将 size - …

多线程JUC 第2季 BlockingQueue 阻塞队列

一 阻塞队列 1.1 阻塞队列介绍 阻塞队列&#xff08;BlockingQueue&#xff09;是一个在队列基础上又支持了两个附加操作的队列&#xff1a; put方法&#xff1a;当队列装满时&#xff0c;添加的线程则被阻塞&#xff0c;直到队列不满&#xff0c;则可用。 take方法&#x…

力扣2028. 找出缺失的观测数据

题目&#xff1a; 现有一份 n m 次投掷单个 六面 骰子的观测数据&#xff0c;骰子的每个面从 1 到 6 编号。观测数据中缺失了 n 份&#xff0c;你手上只拿到剩余 m 次投掷的数据。幸好你有之前计算过的这 n m 次投掷数据的 平均值 。 给你一个长度为 m 的整数数组 rolls &a…

【Linux 网络】网络基础(三)(数据链路层协议:以太网协议、ARP 协议)

一、以太网 两个不同局域网的主机传递数据并不是直接传递的&#xff0c;而是通过路由器 “一跳一跳” 的传递过去。 跨网络传输的本质&#xff1a;由无数个局域网&#xff08;子网&#xff09;转发的结果。 所以&#xff0c;要理解数据跨网络转发原理就要先理解一个局域网中数…

安全厂商第一站!OASA 走进绿盟科技圆满结束

近日&#xff0c;龙蜥社区安全联盟&#xff08;OASA&#xff09;走进联盟成员单位绿盟科技集团股份有限公司&#xff08;以下简称“绿盟科技”&#xff09;&#xff0c;就未来合作方向&#xff0c;双方进行了一次深入的合作交流。该会议共有 11 位人员出席&#xff0c;有来自绿…

视频监控平台AS-V1000产品介绍:账户或用户数据的导入和导出功能介绍

目录 一、功能描述 &#xff08;一&#xff09;导入功能定义 &#xff08;二&#xff09;导出功能定义 二、用户数据的导入导出的作用 三、AS-V1000新版本的导出和导入功能介绍 &#xff08;一&#xff09;功能主界面 &#xff08;二&#xff09;导出功能 1、导出操作 …

从零开始构建 Vision Transformer(ViT) 模型

Transformer 模型最早由 Vaswani 等人在 2017 年论文 Attention Is All You Need 中提出&#xff0c;并已广泛应用于自然语言处理。 2021年&#xff0c;Dosovitsky 等人在论文An Image is Worth 16x16 Words: Transformers for Image Recognition at Scale中提出将 Transforme…

第十二届蓝桥杯物联网试题(国赛)

不得不说国赛相比较省赛而言确实&#xff0c;功能变得更加复杂&#xff0c;更加繁琐&#xff0c;特别是串口LORA通信相结合的更加频繁&#xff0c;且对收取的字符处理要求要更加复杂&#xff0c;处理判别起来会更加复杂。 对于收发数据本身来说&#xff0c;收发的数据本身是以…

每日一题——Python实现PAT甲级1029 Median(举一反三+思想解读+逐步优化)

一个认为一切根源都是“自己不够强”的INTJ 个人主页&#xff1a;用哲学编程-CSDN博客专栏&#xff1a;每日一题——举一反三Python编程学习Python内置函数 Python-3.12.0文档解读 目录 我的方法 代码功能和结构点评 时间复杂度分析 空间复杂度分析 优化建议 我要更强…