CUDA 数组点乘

最新推荐文章于 2023-10-13 12:23:20 发布
最新推荐文章于 2023-10-13 12:23:20 发布
阅读量231 收藏
点赞数
文章标签： 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 高性能计算作业 点乘 dot
05-24
中科院 高性能课程 cuda优化并行点乘
Numpy数组相乘、数组点乘，矩阵元素相乘
qq_61177769的博客
08-29 2949
Numpy数组相乘、数组点乘，矩阵元素相乘
Numpy多维数组的点乘规律总结
daimashiren的博客
12-10 7440
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 7280
当孔乙己说回字有四样写法的时候，相信各位都是这样的表情吧？ 但是，如果孔乙己说NumPy数组有四种乘法的时候，各位大约就是这样的表情了吧？ 实际上，NumPy数组乘法远不止四种。为了在写作和阅读时保持清晰的逻辑和清醒的头脑，本文仅对四种最常见的数组乘法给出详细说明，并用一道数学题来演示向量点乘和叉乘的用法。 1. 星乘（*） 先声明一下：星乘这个说法，是我自己创造的，...
CUDA编程实现矩阵乘法（代码解析）
ing100的博客
10-13 1114
算法思路+代码，CUDA编程之矩阵乘法全面解析，涉及到共享内存与全局内存的使用以及矩阵分块问题等
CUDA实现任意行列大小常数矩阵相乘
Hughie999的博客
05-17 309
实现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程序，包含了矩阵-矢量乘积，矢量-矢量点积的计算，主要用于说明cuda的使用和一些特性，并与intel mkl库的结果进行了比较。
C.rar_点乘
09-24
C语言基本代码，计算点乘相加的和，打印星号图形，还有税收计算，代码比较精简，符合初学者的水平
蒙哥马利梯子点乘.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 184
两个向量各有100个元素，用随机整数初始化以计算点乘，用Python 的time模块来计算点乘所需的时间。*运算符用于计算两个向量的元素相乘，其结果相加以计算整体点乘。请注意，numpy.dot方法只用于矩阵乘法，而不能用于点乘。使用to_gpu方法将两个向量上传给GPU执行点乘计算。gpuarray类提供了一个点乘方法可以直接计算点乘，需要两个GPU数组作为参数，使用get()方法将计算结果下载回主机。计算结果和使用CUDA事件测量的时间显示在控制台。
CUDA:矩阵乘法的实现(Share Memory)
m0_46197553的博客
07-07 4008
记录了使用Share Memory优化的CUDA矩阵乘法
Python numpy中的一维数组和矩阵之间的点乘和矩阵乘法计算
txh3093的博客
09-30 8976
在numpy中，一维数组是一个很奇葩的存在，在数组与矩阵计算时，很容易产生混淆，到底数组在矩阵计算时是行向量还是列向量？ 经过测试，本文预先给出如下结论： 1.数组点乘矩阵时，将数组看做是行向量，按行依次与矩阵最内层方括号的内容（最后一个维度）点乘。 2.数组左乘矩阵时，将数组作为行向量。 3.数组右乘矩阵时将数组作为列向量。 4.注意： 按照上述规则，遇到长度不同时，将报错 1.数组与数组和矩阵之间的点乘 数组之间维度相同时，可以任意进行点乘和点除计算，其他的计算如，sqrt、power、l.
numpy 矩阵与向量相乘_Python学习第115课--numpy点乘余弦定理推导及其空间意义
weixin_39837041的博客
11-27 448
【每天几分钟，从零入门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)这两个方
cuda做点乘
test
09-22 945
#include "cuda_runtime.h" #include "cublas_v2.h" #include "time.h" #include using namespace std; #define imin(a,b) (a const int N = 31 * 1024; const int threadsPerBlock = 256; const int block
使用cuda C完成矩阵相乘算法详解
sinat_38368658的博客
03-26 4941
矩阵相乘大家应该都不陌生。 设有两个矩阵M和N，假设M和N都是方阵，维度均为width × width 如果M和N均为1000 × 1000的矩阵，总共要进行1000000次点乘。其中，每次点乘有1000次乘法和1000次加法。 Matirx Multiply：CPU实现 先来看看使用普通的c代码在CPU端如何实现 void MatrixMulOnHost(float* M,float* N,...
CUDA矩阵乘法
DreamerZhang的专栏
06-07 4492
一、使用全局内存 矩阵乘法，即用矩阵A每行与矩阵B的每列，依次做乘积累加就可以得到各个元素的值。在CPU上用三层循环实现。这里是将二维数组用一维的形式表示，即按行存储。 size_t size = WIDTH*WIDTH * sizeof(int); int *h_A = (int *)malloc(size); int *h_B = (int*)malloc(size); int *
python数组点乘
最新发布
03-15
在Python中，可以使用NumPy库来进行数组的点乘操作。NumPy是一个强大的数值计算库，提供了高效的数组操作和数学函数。 要进行数组的点乘，首先需要导入NumPy库，并创建两个数组。然后，可以使用NumPy提供的dot函数...

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值