草庐IT

memory - CUDA 常量内存分配是如何工作的？

coder 2023-06-04 原文

我想了解一下如何分配常量内存(使用 CUDA 4.2)。我知道总可用的常量内存是 64KB。但是这个内存是什么时候在设备上实际分配的呢？此限制适用于每个内核、cuda 上下文还是整个应用程序？

假设 .cu 文件中有多个内核，每个内核使用的常量内存都少于 64K。但总的恒定内存使用量超过 64K。是否可以按顺序调用这些内核？如果使用不同的流同时调用它们会发生什么？

如果有一个大型 CUDA 动态库，其中包含许多内核，每个内核都使用不同数量的常量内存？

如果有两个应用程序都需要一半以上的可用常量内存会怎样？第一个应用程序运行良好，但第二个应用程序何时会失败？在应用启动、cudaMemcpyToSymbol() 调用或内核执行时？

最佳答案

Parallel Thread Execution ISA Version 3.1第 5.1.3 节讨论常量库。

Constant memory is restricted in size, currently limited to 64KB which can be used to hold statically-sized constant variables. There is an additional 640KB of constant memory, organized as ten independent 64KB regions. The driver may allocate and initialize constant buffers in these regions and pass pointers to the buffers as kernel function parameters. Since the ten regions are not contiguous, the driver must ensure that constant buffers are allocated so that each buffer fits entirely within a 64KB region and does not span a region boundary.

可以用一个简单的程序来说明常量内存的使用。

__constant__ int    kd_p1;
__constant__ short  kd_p2;
__constant__ char   kd_p3;
__constant__ double kd_p4;

__constant__ float kd_floats[8];

__global__ void parameters(int p1, short p2, char p3, double p4, int* pp1, short* pp2, char* pp3,     double* pp4)
{
    *pp1 = p1;
    *pp2 = p2;
    *pp3 = p3;
    *pp4 = p4;
    return;
}

__global__ void constants(int* pp1, short* pp2, char* pp3, double* pp4)
{
    *pp1 = kd_p1;
    *pp2 = kd_p2;
    *pp3 = kd_p3;
    *pp4 = kd_p4;
    return;
}

为compute_30、sm_30编译并执行cuobjdump -sass <executable or obj>拆卸你应该看到

Fatbin elf code:
================
arch = sm_30
code version = [1,6]
producer = cuda
host = windows
compile_size = 32bit
identifier = c:/dev/constant_banks/kernel.cu

    code for sm_30
            Function : _Z10parametersiscdPiPsPcPd
    /*0008*/     /*0x10005de428004001*/     MOV R1, c [0x0] [0x44];        // stack pointer
    /*0010*/     /*0x40001de428004005*/     MOV R0, c [0x0] [0x150];       // pp1
    /*0018*/     /*0x50009de428004005*/     MOV R2, c [0x0] [0x154];       // pp2
    /*0020*/     /*0x0001dde428004005*/     MOV R7, c [0x0] [0x140];       // p1
    /*0028*/     /*0x13f0dc4614000005*/     LDC.U16 R3, c [0x0] [0x144];   // p2
    /*0030*/     /*0x60011de428004005*/     MOV R4, c [0x0] [0x158];       // pp3
    /*0038*/     /*0x70019de428004005*/     MOV R6, c [0x0] [0x15c];       // pp4
    /*0048*/     /*0x20021de428004005*/     MOV R8, c [0x0] [0x148];       // p4
    /*0050*/     /*0x30025de428004005*/     MOV R9, c [0x0] [0x14c];       // p4
    /*0058*/     /*0x1bf15c0614000005*/     LDC.U8 R5, c [0x0] [0x146];    // p3
    /*0060*/     /*0x0001dc8590000000*/     ST [R0], R7;                   // *pp1 = p1
    /*0068*/     /*0x0020dc4590000000*/     ST.U16 [R2], R3;               // *pp2 = p2
    /*0070*/     /*0x00415c0590000000*/     ST.U8 [R4], R5;                // *pp3 = p3
    /*0078*/     /*0x00621ca590000000*/     ST.64 [R6], R8;                // *pp4 = p4
    /*0088*/     /*0x00001de780000000*/     EXIT;
    /*0090*/     /*0xe0001de74003ffff*/     BRA 0x90;
    /*0098*/     /*0x00001de440000000*/     NOP CC.T;
    /*00a0*/     /*0x00001de440000000*/     NOP CC.T;
    /*00a8*/     /*0x00001de440000000*/     NOP CC.T;
    /*00b0*/     /*0x00001de440000000*/     NOP CC.T;
    /*00b8*/     /*0x00001de440000000*/     NOP CC.T;
            ...........................................


            Function : _Z9constantsPiPsPcPd
    /*0008*/     /*0x10005de428004001*/     MOV R1, c [0x0] [0x44];       // stack pointer
    /*0010*/     /*0x00001de428004005*/     MOV R0, c [0x0] [0x140];      // p1
    /*0018*/     /*0x10009de428004005*/     MOV R2, c [0x0] [0x144];      // p2
    /*0020*/     /*0x0001dde428004c00*/     MOV R7, c [0x3] [0x0];        // kd_p1
    /*0028*/     /*0x13f0dc4614000c00*/     LDC.U16 R3, c [0x3] [0x4];    // kd_p2
    /*0030*/     /*0x20011de428004005*/     MOV R4, c [0x0] [0x148];      // p3
    /*0038*/     /*0x30019de428004005*/     MOV R6, c [0x0] [0x14c];      // p4
    /*0048*/     /*0x20021de428004c00*/     MOV R8, c [0x3] [0x8];        // kd_p4
    /*0050*/     /*0x30025de428004c00*/     MOV R9, c [0x3] [0xc];        // kd_p4
    /*0058*/     /*0x1bf15c0614000c00*/     LDC.U8 R5, c [0x3] [0x6];     // kd_p3
    /*0060*/     /*0x0001dc8590000000*/     ST [R0], R7;
    /*0068*/     /*0x0020dc4590000000*/     ST.U16 [R2], R3;
    /*0070*/     /*0x00415c0590000000*/     ST.U8 [R4], R5;
    /*0078*/     /*0x00621ca590000000*/     ST.64 [R6], R8;
    /*0088*/     /*0x00001de780000000*/     EXIT;
    /*0090*/     /*0xe0001de74003ffff*/     BRA 0x90;
    /*0098*/     /*0x00001de440000000*/     NOP CC.T;
    /*00a0*/     /*0x00001de440000000*/     NOP CC.T;
    /*00a8*/     /*0x00001de440000000*/     NOP CC.T;
    /*00b0*/     /*0x00001de440000000*/     NOP CC.T;
    /*00b8*/     /*0x00001de440000000*/     NOP CC.T;
            .....................................

我在 SASS 的右侧进行了注释。

在 sm30 上，您可以看到参数在从偏移量 0x140 开始的常量 bank 0 中传递。

用户定义__constant__变量在常量库 3 中定义。

如果你执行 cuobjdump --dump-elf <executable or obj>您可以找到其他有趣的常量信息。

32bit elf: abi=6, sm=30, flags = 0x1e011e
Sections:
Index Offset   Size ES Align   Type   Flags Link     Info Name
    1     34    142  0  1    STRTAB       0    0        0 .shstrtab
    2    176    19b  0  1    STRTAB       0    0        0 .strtab
    3    314     d0 10  4    SYMTAB       0    2        a .symtab
    4    3e4     50  0  4 CUDA_INFO       0    3        b .nv.info._Z9constantsPiPsPcPd
    5    434     30  0  4 CUDA_INFO       0    3        0 .nv.info
    6    464     90  0  4 CUDA_INFO       0    3        a .nv.info._Z10parametersiscdPiPsPcPd
    7    4f4    160  0  4  PROGBITS       2    0        a .nv.constant0._Z10parametersiscdPiPsPcPd
    8    654    150  0  4  PROGBITS       2    0        b .nv.constant0._Z9constantsPiPsPcPd
    9    7a8     30  0  8  PROGBITS       2    0        0 .nv.constant3
    a    7d8     c0  0  4  PROGBITS       6    3  a00000b .text._Z10parametersiscdPiPsPcPd
    b    898     c0  0  4  PROGBITS       6    3  a00000c .text._Z9constantsPiPsPcPd

.section .strtab

.section .shstrtab

.section .symtab
 index     value     size      info    other  shndx    name
   0          0        0        0        0      0     (null)
   1          0        0        3        0      a     .text._Z10parametersiscdPiPsPcPd
   2          0        0        3        0      7     .nv.constant0._Z10parametersiscdPiPsPcPd
   3          0        0        3        0      b     .text._Z9constantsPiPsPcPd
   4          0        0        3        0      8     .nv.constant0._Z9constantsPiPsPcPd
   5          0        0        3        0      9     .nv.constant3
   6          0        4        1        0      9     kd_p1
   7          4        2        1        0      9     kd_p2
   8          6        1        1        0      9     kd_p3
   9          8        8        1        0      9     kd_p4
  10         16       32        1        0      9     kd_floats
  11          0      192       12       10      a     _Z10parametersiscdPiPsPcPd
  12          0      192       12       10      b     _Z9constantsPiPsPcPd

每次启动都会对内核参数常量库进行版本控制，以便可以执行并发内核。编译器和用户常量是每个 CUmodule。开发人员有责任管理这些数据的一致性。例如，开发人员必须确保 cudaMemcpyToSymbol 以安全的方式更新。

关于memory - CUDA 常量内存分配是如何工作的？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/14483077/

memoryCUDA0xconstantcodeconstantsnvidiagpu-constant-memory

有关memory - CUDA 常量内存分配是如何工作的？的更多相关文章

  1. ruby - 如何使用 Nokogiri 的 xpath 和 at_xpath 方法 - 2

    我正在学习如何使用Nokogiri，根据这段代码我遇到了一些问题:require'rubygems'require'mechanize'post_agent=WWW::Mechanize.newpost_page=post_agent.get('http://www.vbulletin.org/forum/showthread.php?t=230708')puts"\nabsolutepathwithtbodygivesnil"putspost_page.parser.xpath('/html/body/div/div/div/div/div/table/tbody/tr/td/div

  2. ruby - 如何从 ruby​​ 中的字符串运行任意对象方法？ - 2

    总的来说，我对ruby​​还比较陌生，我正在为我正在创建的对象编写一些rspec测试用例。许多测试用例都非常基础，我只是想确保正确填充和返回值。我想知道是否有办法使用循环结构来执行此操作。不必为我要测试的每个方法都设置一个assertEquals。例如:describeitem,"TestingtheItem"doit"willhaveanullvaluetostart"doitem=Item.new#HereIcoulddotheitem.name.shouldbe_nil#thenIcoulddoitem.category.shouldbe_nilendend但我想要一些方法来使用

  3. ruby-on-rails - Ruby net/ldap 模块中的内存泄漏 - 2

    作为我的Rails应用程序的一部分，我编写了一个小导入程序，它从我们的LDAP系统中吸取数据并将其塞入一个用户表中。不幸的是，与LDAP相关的代码在遍历我们的32K用户时泄漏了大量内存，我一直无法弄清楚如何解决这个问题。这个问题似乎在某种程度上与LDAP库有关，因为当我删除对LDAP内容的调用时，内存使用情况会很好地稳定下来。此外，不断增加的对象是Net::BER::BerIdentifiedString和Net::BER::BerIdentifiedArray，它们都是LDAP库的一部分。当我运行导入时，内存使用量最终达到超过1GB的峰值。如果问题存在，我需要找到一些方法来更正我的代

  4. python - 如何使用 Ruby 或 Python 创建一系列高音调和低音调的蜂鸣声？ - 2

    关闭。这个问题是opinion-based.它目前不接受答案。想要改进这个问题？更新问题，以便editingthispost可以用事实和引用来回答它.关闭4年前。Improvethisquestion我想在固定时间创建一系列低音和高音调的哔哔声。例如:在150毫秒时发出高音调的蜂鸣声在151毫秒时发出低音调的蜂鸣声200毫秒时发出低音调的蜂鸣声250毫秒的高音调蜂鸣声有没有办法在Ruby或Python中做到这一点？我真的不在乎输出编码是什么(.wav、.mp3、.ogg等等)，但我确实想创建一个输出文件。

  5. ruby-on-rails - 由于 "wkhtmltopdf"，PDFKIT 显然无法正常工作 - 2

    我在从html页面生成PDF时遇到问题。我正在使用PDFkit。在安装它的过程中，我注意到我需要wkhtmltopdf。所以我也安装了它。我做了PDFkit的文档所说的一切......现在我在尝试加载PDF时遇到了这个错误。这里是错误:commandfailed:"/usr/local/bin/wkhtmltopdf""--margin-right""0.75in""--page-size""Letter""--margin-top""0.75in""--margin-bottom""0.75in""--encoding""UTF-8""--margin-left""0.75in""-

  6. ruby-on-rails - 如何验证 update_all 是否实际在 Rails 中更新 - 2

    给定这段代码defcreate@upgrades=User.update_all(["role=?","upgraded"],:id=>params[:upgrade])redirect_toadmin_upgrades_path,:notice=>"Successfullyupgradeduser."end我如何在该操作中实际验证它们是否已保存或未重定向到适当的页面和消息？ 最佳答案 在Rails3中，update_all不返回任何有意义的信息，除了已更新的记录数(这可能取决于您的DBMS是否返回该信息)。http://ar.ru

  7. ruby-on-rails - 'compass watch' 是如何工作的/它是如何与 rails 一起使用的 - 2

    我在我的项目目录中完成了compasscreate.和compassinitrails。几个问题:我已将我的.sass文件放在public/stylesheets中。这是放置它们的正确位置吗？当我运行compasswatch时，它不会自动编译这些.sass文件。我必须手动指定文件:compasswatchpublic/stylesheets/myfile.sass等。如何让它自动运行？文件ie.css、print.css和screen.css已放在stylesheets/compiled。如何在编译后不让它们重新出现的情况下删除它们？我自己编译的.sass文件编译成compiled/t

  8. ruby - 如何将脚本文件的末尾读取为数据文件(Perl 或任何其他语言) - 2

    我正在寻找执行以下操作的正确语法(在Perl、Shell或Ruby中):#variabletoaccessthedatalinesappendedasafileEND_OF_SCRIPT_MARKERrawdatastartshereanditcontinues. 最佳答案 Perl用__DATA__做这个:#!/usr/bin/perlusestrict;usewarnings;while(){print;}__DATA__Texttoprintgoeshere 关于ruby-如何将脚

  9. ruby-on-rails - 未初始化的常量 Psych::Syck (NameError) - 2

    在我的gem中，我需要yaml并且在我的本地计算机上运行良好。但是在将我的gem推送到ruby​​gems.org之后，当我尝试使用我的gem时，我收到一条错误消息=>"uninitializedconstantPsych::Syck(NameError)"谁能帮我解决这个问题？附言RubyVersion=>ruby1.9.2,GemVersion=>1.6.2,Bundlerversion=>1.0.15 最佳答案 经过几个小时的研究，我发现=>“YAML使用未维护的Syck库，而Psych使用现代的LibYAML”因此，为了解决

  10. ruby - 如何指定 Rack 处理程序 - 2

    Rackup通过Rack的默认处理程序成功运行任何Rack应用程序。例如:classRackAppdefcall(environment)['200',{'Content-Type'=>'text/html'},["Helloworld"]]endendrunRackApp.new但是当最后一行更改为使用Rack的内置CGI处理程序时，rackup给出“NoMethodErrorat/undefinedmethod`call'fornil:NilClass”:Rack::Handler::CGI.runRackApp.newRack的其他内置处理程序也提出了同样的反对意见。例如Rack

随机推荐