CUDA 数组点乘

最新推荐文章于 2024-07-02 10:24:12 发布
最新推荐文章于 2024-07-02 10:24:12 发布
阅读量232 收藏
点赞数
文章标签： python
原文链接：https://my.oschina.net/toyandong/blog/228989
版权

欢迎访问小站，阅读原文 http://www.yandong.org/archives/755

Question 1: Dot product

基础函数

基本思路：并行计算乘法，串行将每个BLOCK中的所有线程的结果在thread=0线程中汇总，

 

__global__ void dot( TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    temp[threadIdx.x] = a[index] * b[index];

 

    __syncthreads();

    if( 0 == threadIdx.x ) {

         for( int i= 0; i< threadsPerBlock; i++ )

         {

             c[blockIdx.x] += temp[i];

         }

    }

}

 

 

改进1 树形加法

基本思路：并行乘法，并行加法。每个BLOCK中的数据汇总也可以使用并行计算，采用树形加法。

/*并行乘法，并行加法*/

__global__ void dot3( TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    int offset=1, mask=1;

    temp[threadIdx.x] = a[index] * b[index];

 

     __syncthreads();/*树形加法*/

    while(offset < threadsPerBlock) {

        if((threadIdx.x & mask) == 0) {

            temp[threadIdx.x] += temp[threadIdx.x + offset];

        }

        offset += offset;

        mask = offset + mask;

        __syncthreads();

}

 

    if(threadIdx.x == 0) {

        c[blockIdx.x] = temp[0];  

    }

}

 

 

改进2 浮点数精度

基本思路：浮点运算的精度问题，采用Kahan's Summation Formula 来提高精确度。 并将main函数中的c由float类型改为double类型。

/*并行乘法，并行加法, 浮点数精度优化Kahan's Summation Formula*/

__global__ void dot( TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    int offset=1, mask=1;

    TYPE y=0,r;

    /*Kahan's Summation Formula*/

    temp[threadIdx.x] =0;

    y -= a[index] * b[index];

    r=temp[threadIdx.x]-y;

    y=(r-temp[threadIdx.x])+y;

    temp[threadIdx.x] =r;

 

     __syncthreads();

    while(offset < threadsPerBlock) {

        if((threadIdx.x & mask) == 0) {

            temp[threadIdx.x] += temp[threadIdx.x + offset];

        }

        offset += offset;

        mask = offset + mask;

        __syncthreads();

    }

 

    if(threadIdx.x == 0) {

        c[blockIdx.x] = temp[0];  

    }

}

 

改进3 Bank conflict

基本思路：改进2的树形加法存在Bank conflict的问题。

 

/*并行乘法，并行加法, 浮点数精度优化Kahan's Summation Formula  解决bank conflict*/

__global__ void dot3( TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    int offset=1;

    TYPE y=0,r;

    /*Kahan's Summation Formula*/

    temp[threadIdx.x] =0;

    y -= a[index] * b[index];

    r=temp[threadIdx.x]-y;

    y=(r-temp[threadIdx.x])+y;

    temp[threadIdx.x] =r;

 

     __syncthreads();

    offset = threadsPerBlock / 2;

    while(offset > 0) {

        if(threadIdx.x  < offset) {

           temp[threadIdx.x] += temp[threadIdx.x + offset];

        }

        offset >>= 1;

        __syncthreads();

    }

 

    if(threadIdx.x == 0) {

        c[blockIdx.x] = temp[0];  

    }

}

改进4 循环展开

基本思路：将树形加法进行循环展开

 

/*并行乘法，并行加法, 浮点数精度优化Kahan's Summation Formula  解决bank conflict 循环展开*/

__global__ void dot4( TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    int offset=1;

    TYPE y=0,r;

    /*Kahan's Summation Formula*/

    temp[threadIdx.x] =0;

    y -= a[index] * b[index];

    r=temp[threadIdx.x]-y;

    y=(r-temp[threadIdx.x])+y;

    temp[threadIdx.x] =r;

 

     __syncthreads();

    if(threadIdx.x < 256) {   temp[threadIdx.x] += temp[threadIdx.x +  256]; }

    __syncthreads();

    if(threadIdx.x < 128) {   temp[threadIdx.x] += temp[threadIdx.x +  128]; }

    __syncthreads();

    if(threadIdx.x < 64) {   temp[threadIdx.x] += temp[threadIdx.x +  64]; }

    __syncthreads();

    if(threadIdx.x < 32) {   temp[threadIdx.x] += temp[threadIdx.x + 32]; }

    __syncthreads();

    if(threadIdx.x < 16) {   temp[threadIdx.x] += temp[threadIdx.x +  16]; }

    __syncthreads();

    if(threadIdx.x < 8) {   temp[threadIdx.x] += temp[threadIdx.x +  8]; }

    __syncthreads();

    if(threadIdx.x < 4) {   temp[threadIdx.x] += temp[threadIdx.x + 4]; }

    __syncthreads();

    if(threadIdx.x < 2) {   temp[threadIdx.x] += temp[threadIdx.x +  2]; }

    __syncthreads();

    if(threadIdx.x < 1) {   temp[threadIdx.x] += temp[threadIdx.x +  1]; }

    __syncthreads();

 

    if(threadIdx.x == 0) {

        c[blockIdx.x] = temp[0];  

    }

}

改进5 Ideal Thread

基本思路：解决空闲线程问题，即<32不再展开

/*并行乘法，并行加法, 浮点数精度优化Kahan's Summation Formula  解决bank conflict 循环展开， 空闲线程*/

__global__ void dot5(TYPE *a, TYPE *b, TYPE *c)

{

    __shared__ TYPE temp[threadsPerBlock];

    int index = threadIdx.x + blockIdx.x * blockDim.x;

    int offset=1;

    TYPE y=0,r;

    /*Kahan's Summation Formula*/

    temp[threadIdx.x] =0;

    y -= a[index] * b[index];

    r=temp[threadIdx.x]-y;

    y=(r-temp[threadIdx.x])+y;

    temp[threadIdx.x] =r;

 

     __syncthreads();

    if(threadIdx.x < 256) {   temp[threadIdx.x] += temp[threadIdx.x +  256]; }

    __syncthreads();

    if(threadIdx.x < 128) {   temp[threadIdx.x] += temp[threadIdx.x +  128]; }

    __syncthreads();

    if(threadIdx.x < 64) {   temp[threadIdx.x] += temp[threadIdx.x +  64]; }

    __syncthreads();

    if (threadIdx.x < 32)

    {

         temp[threadIdx.x] += temp[threadIdx.x + 32];

         temp[threadIdx.x] += temp[threadIdx.x + 16];

         temp[threadIdx.x] += temp[threadIdx.x +8];

         temp[threadIdx.x] += temp[threadIdx.x +4];

         temp[threadIdx.x] += temp[threadIdx.x + 2];

         temp[threadIdx.x] += temp[threadIdx.x + 1];

    }

     __syncthreads();

 

    if(threadIdx.x == 0) {

        c[blockIdx.x] = temp[0];  

    }

}

 

性能总结

改进

性能us

dot0

1699738.016

Dot1

114682.144

Dot2

120819.072

Dot3

71126.24

Dot4

57975.072

Dot5

49489.856

 

 


 

 

 

 

 

 

 

 

 


转载于:https://my.oschina.net/toyandong/blog/228989

chiminmang7244
关注
  • 0
    点赞
  • 踩
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
CUDA 分块矩阵乘法
weixin_43906500的博客
01-11 1085
cpp文件 #include "stdafx.h" #include <stdio.h> #include <stdlib.h> //为rand()及srand()提供函数声明 #include <time.h> extern "C" int mulWithCuda(float *c, const float *a, const float *b, int size); int _tmain(int argc, _TCHAR* argv[]...
CUDA三维数组
weixin_30924087的博客
01-02 1042
http://hpcbbs.it168.com/forum.php?mod=viewthread&tid=1643 根据上面链接的帖子研究了下三维数组，就像他自己说的一样是有问题的，我自己修改了下，结果终于正确了。大家有兴趣的可以对照着看看。 整个过程关键参考了这篇文章http://www.xuebuyuan.com/685353.html 1 #include <std...
Numpy数组相乘、数组点乘，矩阵元素相乘
qq_61177769的博客
08-29 2950
Numpy数组相乘、数组点乘，矩阵元素相乘
Numpy多维数组的点乘规律总结
daimashiren的博客
12-10 7474
Numpy中数组的点乘用函数dot()来实现，而单个*则表示的是数组对应元素的乘法。 一维情况: In: a = np.array([1,2,3]) b = np.array([1,2,3]) print(a.dot(b)) Out: 14 一维数组点乘相当于向量乘法，计算结果为一个数(标量)，注意要是等长的一维数组才能点乘 二维情况: In: a = np.array([[1,2,3], [1,2,3]]) b = np.array([[1,2,3],
numpy 点乘_细说NumPy数组的四种乘法，带你走进向量运算的奇妙世界
edward_zcl的博客
06-08 7304
当孔乙己说回字有四样写法的时候，相信各位都是这样的表情吧？ 但是，如果孔乙己说NumPy数组有四种乘法的时候，各位大约就是这样的表情了吧？ 实际上，NumPy数组乘法远不止四种。为了在写作和阅读时保持清晰的逻辑和清醒的头脑，本文仅对四种最常见的数组乘法给出详细说明，并用一道数学题来演示向量点乘和叉乘的用法。 1. 星乘（*） 先声明一下：星乘这个说法，是我自己创造的，...
CUDA编程实现矩阵乘法（代码解析）
ing100的博客
10-13 1136
算法思路+代码，CUDA编程之矩阵乘法全面解析，涉及到共享内存与全局内存的使用以及矩阵分块问题等
CUDA实现任意行列大小常数矩阵相乘
Hughie999的博客
05-17 309
实现CUDA加速的任意大小常数矩阵乘法计算
cuda 高性能计算作业 点乘 dot
05-24
中科院 高性能课程 cuda优化并行点乘
matlab代码中向量的点乘-matlube:用于Scala的类似于Matlab的线性代数DSL。因为Scala中的线性代数不应该受到损害
05-23
matlab代码中向量的点乘Matlube 用于矩阵库的Scala DSL。 Matlube是Scala的类似于Matlab的线性代数DSL。 该库仅支持密集双矩阵。 它旨在用作其他Java Matrix库的前端。 当前，有两个后端：一个用于Jama（），另一个...
CUDA计算矢量点乘并和MKL作性能比较的VS工程
02-22
CUDA计算矢量点乘并和MKL作性能比较的VS工程是一个示例项目，它旨在展示如何使用NVIDIA的CUDA技术进行并行计算，并与Intel的Math Kernel Library (MKL)进行性能对比。CUDA是一种由NVIDIA开发的编程模型，允许开发者...
C.rar_点乘
09-24
在C语言中，实现点乘通常是通过遍历两个相同长度的数组，将对应元素相乘然后求和来完成的。例如，对于两个长度为n的一维数组a和b，点乘的结果可以通过以下方式计算： ```c int result = 0; for(int i = 0; i ; i++)...
蒙哥马利梯子点乘.pdf
08-01
蒙哥马利梯子点乘算法。[k]P：椭圆曲线上点P的k倍点，即：[k]P = {P+P+……+P}，k个P相加，即点乘。其中k是正整数。
细说NumPy数组的四种乘法，带你走进向量运算的奇妙世界
热门推荐
Python作业辅导员 - 天元浪子
12-07 2万+
当孔乙己说回字有四样写法的时候，相信各位都是这样的表情吧？但是，如果孔乙己说NumPy数组有四种乘法的时候，各位大约就是这样的表情了吧？实际上，NumPy数组乘法远不止四种。为了在写作和阅读时保持清晰的逻辑和清醒的头脑，本文仅对四种最常见的数组乘法给出详细说明，并用一道数学题来演示向量点乘和叉乘的用法。
PyCUDA数组点乘
taifyang的博客
01-07 186
两个向量各有100个元素，用随机整数初始化以计算点乘，用Python 的time模块来计算点乘所需的时间。*运算符用于计算两个向量的元素相乘，其结果相加以计算整体点乘。请注意，numpy.dot方法只用于矩阵乘法，而不能用于点乘。使用to_gpu方法将两个向量上传给GPU执行点乘计算。gpuarray类提供了一个点乘方法可以直接计算点乘，需要两个GPU数组作为参数，使用get()方法将计算结果下载回主机。计算结果和使用CUDA事件测量的时间显示在控制台。
CUDA:矩阵乘法的实现(Share Memory)
m0_46197553的博客
07-07 4027
记录了使用Share Memory优化的CUDA矩阵乘法
Python numpy中的一维数组和矩阵之间的点乘和矩阵乘法计算
txh3093的博客
09-30 8995
在numpy中，一维数组是一个很奇葩的存在，在数组与矩阵计算时，很容易产生混淆，到底数组在矩阵计算时是行向量还是列向量？ 经过测试，本文预先给出如下结论： 1.数组点乘矩阵时，将数组看做是行向量，按行依次与矩阵最内层方括号的内容（最后一个维度）点乘。 2.数组左乘矩阵时，将数组作为行向量。 3.数组右乘矩阵时将数组作为列向量。 4.注意： 按照上述规则，遇到长度不同时，将报错 1.数组与数组和矩阵之间的点乘 数组之间维度相同时，可以任意进行点乘和点除计算，其他的计算如，sqrt、power、l.
numpy 矩阵与向量相乘_Python学习第115课--numpy点乘余弦定理推导及其空间意义
weixin_39837041的博客
11-27 452
【每天几分钟，从零入门python编程的世界！】●numpy中的对位相乘我们之前学习的numpy中的相乘有两种方法：①直接用*号对两个数组进行相乘。②用numpy.multiply(x,y)方法对两个数组相乘。以上两种相乘的方法叫做element wise，即对位相乘。假如有x和y两个数组，则数组x中的第一个元素和数组y中的第一个元素相乘，数组x中的第二个元素和数组y中的第二个元素相乘，...，数...
python数组（矩阵）乘法（点乘、叉乘）
qq_32649321的博客
02-18 1万+
转载：https://blog.csdn.net/haiziccc/article/details/101361583 总结： （1）tf.matmul(A,C)=np.dot(A,C)= A@C都属于叉乘，而tf.multiply(A,C)= A*C=A∙C属于点乘。 （2）叉乘称向量积；点乘指对应元素相乘，点乘的结果表示 在 方向上的投影与 的乘积，反映了两个向量的相似度，结果越大越相似。 （3）python中的数组和矩阵： python中的numpy包下有array(obj)和mat(obj)这两个方
Ansys Zemax｜场曲跟畸变图的前世今生
最新发布
ueotek的博客
07-02 385
这里的y和z坐标和方向余弦是（Hx, Hy, Px, Py）=（X, X, 0 ,0）和（Hx, Hy, Px, Py）=（X, X, 0 ,0.001），在光线追迹像面的前一个平面的 z坐标和方向余弦。OpticStudio通过在X和Y方向（弧矢和子午方向）的傍轴光线追踪确定近轴图像平面的Z坐标，并测量该近轴焦平面与系统图像平面的Z坐标之间的距离。使用附件中的样本文件，近轴像面的全局Z坐标与视场0处的像面位置之差与场曲和畸变图中的的Tan Shift相同。近轴像面的定义是以下两个光线交点的全局Z坐标。
python数组点乘
03-15
在Python中，可以使用NumPy库来进行数组的点乘操作。NumPy是一个强大的数值计算库，提供了高效的数组操作和数学函数。 要进行数组的点乘，首先需要导入NumPy库，并创建两个数组。然后，可以使用NumPy提供的dot函数来进行点乘操作。点乘操作会将两个数组对应位置的元素相乘，并将结果相加得到一个标量值。 下面是一个示例代码： ```python import numpy as np # 创建两个数组 array1 = np.array([1, 2, 3]) array2 = np.array([4, 5, 6]) # 进行点乘操作 result = np.dot(array1, array2) print("点乘结果：", result) ``` 输出为： ``` 点乘结果： 32 ```

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值