草庐IT

c++ - 在 CUDA 中用小 M 对两个 MxN 矩阵执行逐 vector 点积的最快方法是什么？

coder 2024-02-17 原文

我有两个矩阵，每个都是 MxN，其中 M = 16 和 N 大得多(比如 n = 262144，例如)。我的目标是生成一个长度为 N 的 vector ，其中每个元素对应于每个矩阵中的 nth vector 的点积。

我尝试了以下方法，其中 cIdx 对应于每个矩阵中列 vector 的列索引。毫不奇怪，NVIDIA Visual Profiler 告诉我这种方法主要受内存带宽限制。

    public static void MatrixDotProduct(
        float* matrix1,
        float* matrix2,
        float* dotProduct,
        int2 matrixDimensions)
    {
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        int stride = gridDim.x * blockDim.x;
        float sum;

        for (int cIdx = i; cIdx < matrixDimensions.y; cIdx += stride)
        {
            int ci = cIdx * matrixDimensions.x;
            sum = 0f;

            for (int j = 0; j < matrixDimensions.x; j++)
            {
                sum += matrix1[ci + j] * matrix2[ci + j];
            }

            dotProduct[cIdx] = sum;
        }
    }

我还找到了一个 vector-by-vector 的版本点积，我也尝试将其并行化。不幸的是，这比上面的实现慢了 20%(也许是因为我的 M = 16？)。有没有更好的方法来解决我在这里缺少的这个问题？

最佳答案

这不是一个容易优化的案例，因为缺乏数据重用和小 vector 大小(小于扭曲)。代码应受内存限制，缓存性能至关重要。

可以尝试的一件事是通过对负载使用 vector 类型来减少容量并提高内存事务的效率。我希望是这样的:

__device__ float vdot(float2 v1, float2 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y);
}

__device__ float vdot(float4 v1, float4 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w);
}

template<typename VT, int NT>
__device__ float vector_dotprod(const VT* v1, const VT* v2) {
    float sum = 0.f;
#pragma unroll
    for (int j = 0; j < NT; j++) {
        sum += vdot(v1[j], v2[j]);
    }
    return sum;
}

template<typename VT, int Nrows>
__global__
void MatrixDotProductPlus(float* matrix1, float* matrix2, float* dotProduct, int2 matrixDimensions)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    int stride2 = stride * Nrows;

    VT* m1 = reinterpret_cast<VT*>(matrix1) + i * Nrows;
    VT* m2 = reinterpret_cast<VT*>(matrix2) + i * Nrows;
    for (; i < matrixDimensions.y; i += stride, m1 += stride2, m2 += stride2) {
        dotProduct[i] = vector_dotprod<VT,Nrows>(m1, m2);
    }
}

[警告:仅经过非常轻微的测试——使用风险自负]

在 float2 情况下比您的代码快两倍，在 Maxwell 或 Pascal 架构上的长度 vector 的 float4 情况下大约快四倍16. 这假设您在编译时知道 vector 的长度并且它们是 2 或 4 的整数倍，尽管我怀疑循环展开与 vector 类型本身一样重要。

关于c++ - 在 CUDA 中用小 M 对两个 MxN 矩阵执行逐 vector 点积的最快方法是什么？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/44253005/

中用ampcodefloatvectorc++matrixcudadot-product

有关c++ - 在 CUDA 中用小 M 对两个 MxN 矩阵执行逐 vector 点积的最快方法是什么？的更多相关文章

  1. ruby - 为什么我可以在 Ruby 中使用 Object#send 访问私有(private)/ protected 方法？ - 2

    类classAprivatedeffooputs:fooendpublicdefbarputs:barendprivatedefzimputs:zimendprotecteddefdibputs:dibendendA的实例a=A.new测试a.foorescueputs:faila.barrescueputs:faila.zimrescueputs:faila.dibrescueputs:faila.gazrescueputs:fail测试输出failbarfailfailfail.发送测试[:foo,:bar,:zim,:dib,:gaz].each{|m|a.send(m)resc

  2. ruby-on-rails - Rails - 子类化模型的设计模式是什么？ - 2

    我有一个模型:classItem项目有一个属性“商店”基于存储的值，我希望Item对象对特定方法具有不同的行为。Rails中是否有针对此的通用设计模式？如果方法中没有大的if-else语句，这是如何干净利落地完成的？ 最佳答案 通常通过Single-TableInheritance. 关于ruby-on-rails-Rails-子类化模型的设计模式是什么？，我们在StackOverflow上找到一个类似的问题： https://stackoverflow.co

  3. ruby - 什么是填充的 Base64 编码字符串以及如何在 ruby​​ 中生成它们？ - 2

    我正在使用的第三方API的文档状态:"[O]urAPIonlyacceptspaddedBase64encodedstrings."什么是“填充的Base64编码字符串”以及如何在Ruby中生成它们。下面的代码是我第一次尝试创建转换为Base64的JSON格式数据。xa=Base64.encode64(a.to_json) 最佳答案 他们说的padding其实就是Base64本身的一部分。它是末尾的“=”和“==”。Base64将3个字节的数据包编码为4个编码字符。所以如果你的输入数据有长度n和n%3=1=>"=="末尾用于填充n%

  4. ruby - 解析 RDFa、微数据等的最佳方式是什么，使用统一的模式/词汇(例如 schema.org)存储和显示信息 - 2

    我主要使用Ruby来执行此操作，但到目前为止我的攻击计划如下:使用gemsrdf、rdf-rdfa和rdf-microdata或mida来解析给定任何URI的数据。我认为最好映射到像schema.org这样的统一模式，例如使用这个yaml文件，它试图描述数据词汇表和opengraph到schema.org之间的转换:#SchemaXtoschema.orgconversion#data-vocabularyDV:name:namestreet-address:streetAddressregion:addressRegionlocality:addressLocalityphoto:i

  5. ruby - 为什么 4.1%2 使用 Ruby 返回 0.0999999999999996？但是 4.2%2==0.2 - 2

    为什么4.1%2返回0.0999999999999996？但是4.2%2==0.2。 最佳答案 参见此处:WhatEveryProgrammerShouldKnowAboutFloating-PointArithmetic实数是无限的。计算机使用的位数有限(今天是32位、64位)。因此计算机进行的浮点运算不能代表所有的实数。0.1是这些数字之一。请注意，这不是与Ruby相关的问题，而是与所有编程语言相关的问题，因为它来自计算机表示实数的方式。 关于ruby-为什么4.1%2使用Ruby返

  6. ruby-on-rails - 如何在 ruby​​ 中使用两个参数异步运行 exe？ - 2

    exe应该在我打开页面时运行。异步进程需要运行。有什么方法可以在ruby​​中使用两个参数异步运行exe吗？我已经尝试过ruby​​命令-system()、exec()但它正在等待过程完成。我需要用参数启动exe，无需等待进程完成是否有任何ruby​​gems会支持我的问题？ 最佳答案 您可以使用Process.spawn和Process.wait2:pid=Process.spawn'your.exe','--option'#Later...pid,status=Process.wait2pid您的程序将作为解释器的子进程执行。除

  7. ruby-on-rails - 如何优雅地重启 thin + nginx？ - 2

    我的瘦服务器配置了nginx，我的ROR应用程序正在它们上运行。在我发布代码更新时运行thinrestart会给我的应用程序带来一些停机时间。我试图弄清楚如何优雅地重启正在运行的Thin实例，但找不到好的解决方案。有没有人能做到这一点？ 最佳答案 #Restartjustthethinserverdescribedbythatconfigsudothin-C/etc/thin/mysite.ymlrestartNginx将继续运行并代理请求。如果您将Nginx设置为使用多个上游服务器，例如server{listen80;server

  8. ruby - ruby 中的 TOPLEVEL_BINDING 是什么？ - 2

    它不等于主线程的binding，这个toplevel作用域是什么？此作用域与主线程中的binding有何不同？>ruby-e'putsTOPLEVEL_BINDING===binding'false 最佳答案 事实是，TOPLEVEL_BINDING始终引用Binding的预定义全局实例，而Kernel#binding创建的新实例>Binding每次封装当前执行上下文。在顶层，它们都包含相同的绑定(bind)，但它们不是同一个对象，您无法使用==或===测试它们的绑定(bind)相等性。putsTOPLEVEL_BINDINGput

  9. ruby - Infinity 和 NaN 的类型是什么？ - 2

    我可以得到Infinity和NaNn=9.0/0#=>Infinityn.class#=>Floatm=0/0.0#=>NaNm.class#=>Float但是当我想直接访问Infinity或NaN时:Infinity#=>uninitializedconstantInfinity(NameError)NaN#=>uninitializedconstantNaN(NameError)什么是Infinity和NaN？它们是对象、关键字还是其他东西？ 最佳答案 您看到打印为Infinity和NaN的只是Float类的两个特殊实例的字符串

  10. ruby-on-rails - 如果 Object::try 被发送到一个 nil 对象，为什么它会起作用？ - 2

    如果您尝试在Ruby中的nil对象上调用方法，则会出现NoMethodError异常并显示消息:"undefinedmethod‘...’fornil:NilClass"然而，有一个tryRails中的方法，如果它被发送到一个nil对象，它只返回nil:require'rubygems'require'active_support/all'nil.try(:nonexisting_method)#noNoMethodErrorexceptionanymore那么try如何在内部工作以防止该异常？ 最佳答案 像Ruby中的所有其他对象

随机推荐