99精品伊人亚洲|最近国产中文炮友|九草在线视频支援|AV网站大全最新|美女黄片免费观看|国产精品资源视频|精彩无码视频一区|91大神在线后入|伊人终合在线播放|久草综合久久中文

0
  • 聊天消息
  • 系統(tǒng)消息
  • 評(píng)論與回復(fù)
登錄后你可以
  • 下載海量資料
  • 學(xué)習(xí)在線課程
  • 觀看技術(shù)視頻
  • 寫(xiě)文章/發(fā)帖/加入社區(qū)
會(huì)員中心
創(chuàng)作中心

完善資料讓更多小伙伴認(rèn)識(shí)你,還能領(lǐng)取20積分哦,立即完善>

3天內(nèi)不再提示

如何有效地從內(nèi)核中訪問(wèn)設(shè)備的全局內(nèi)存

星星科技指導(dǎo)員 ? 來(lái)源:NVIDIA ? 作者:Mark Harris ? 2022-04-11 10:07 ? 次閱讀
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

在前面的兩文章中,我們研究了如何在主機(jī)和設(shè)備之間高效地移動(dòng)數(shù)據(jù)。在我們的 CUDA C / C ++系列的第六篇文章中,我們將討論如何有效地從內(nèi)核中訪問(wèn)設(shè)備存儲(chǔ)器,特別是全局內(nèi)存。

在 CUDA 設(shè)備上有幾種內(nèi)存,每種內(nèi)存的作用域、生存期和緩存行為都不同。到目前為止,在本系列中,我們已經(jīng)使用了駐留在設(shè)備 DRAM 中的全局內(nèi)存,用于主機(jī)和設(shè)備之間的傳輸,以及內(nèi)核的數(shù)據(jù)輸入和輸出。這里的名稱global是指作用域,因?yàn)樗梢詮闹鳈C(jī)和設(shè)備訪問(wèn)和修改。全局內(nèi)存可以像下面代碼片段的第一行那樣使用__device__de Clara 說(shuō)明符在全局(變量)范圍內(nèi)聲明,或者使用cudaMalloc()動(dòng)態(tài)分配并分配給一個(gè)常規(guī)的 C 指針變量,如第 7 行所示。全局內(nèi)存分配可以在應(yīng)用程序的生命周期內(nèi)保持。根據(jù)設(shè)備的計(jì)算能力,全局內(nèi)存可能被緩存在芯片上,也可能不在芯片上緩存。

__device__ int globalArray[256];

void foo()
{
    ...
    int *myDeviceMemory = 0;
    cudaError_t result = cudaMalloc(&myDeviceMemory, 256 * sizeof(int));
    ...

}在討論全局內(nèi)存訪問(wèn)性能之前,我們需要改進(jìn)對(duì) CUDA 執(zhí)行模型的理解。我們已經(jīng)討論了如何將線程被分組為線程塊分配給設(shè)備上的多處理器。在執(zhí)行過(guò)程中,有一個(gè)更精細(xì)的線程分組到warps。 GPU 上的多處理器以 SIMD (單指令多數(shù)據(jù))方式為每個(gè)扭曲執(zhí)行指令。所有當(dāng)前支持 CUDA – 的 GPUs 的翹曲尺寸(實(shí)際上是 SIMD 寬度)是 32 個(gè)線程。

全局內(nèi)存合并

將線程分組為扭曲不僅與計(jì)算有關(guān),而且與全局內(nèi)存訪問(wèn)有關(guān)。設(shè)備coalesces全局內(nèi)存加載并存儲(chǔ)由一個(gè) warp 線程發(fā)出的盡可能少的事務(wù),以最小化 DRAM 帶寬(在計(jì)算能力小于 2 . 0 的舊硬件上,事務(wù)合并在 16 個(gè)線程的一半扭曲內(nèi),而不是整個(gè)扭曲中)。為了弄清楚 CUDA 設(shè)備架構(gòu)中發(fā)生聚結(jié)的條件,我們?cè)谌齻€(gè) Tesla 卡上進(jìn)行了一些簡(jiǎn)單的實(shí)驗(yàn): a Tesla C870 (計(jì)算能力 1 . 0 )、 Tesla C1060 (計(jì)算能力 1 . 3 )和 Tesla C2050 (計(jì)算能力 2 . 0 )。

我們運(yùn)行兩個(gè)實(shí)驗(yàn),使用如下代碼(GitHub 上也有)中所示的增量?jī)?nèi)核的變體,一個(gè)具有數(shù)組偏移量,這可能導(dǎo)致對(duì)輸入數(shù)組的未對(duì)齊訪問(wèn),另一個(gè)是對(duì)輸入數(shù)組的跨步訪問(wèn)。

#include
#include

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %sn", cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
#endif
  return result;
}

template
__global__ void offset(T* a, int s)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x + s;
  a[i] = a[i] + 1;
}

template
__global__ void stride(T* a, int s)
{
  int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
  a[i] = a[i] + 1;
}

template
void runTest(int deviceId, int nMB)
{
  int blockSize = 256;
  float ms;

  T *d_a;
  cudaEvent_t startEvent, stopEvent;

  int n = nMB*1024*1024/sizeof(T);

  // NB:  d_a(33*nMB) for stride case
  checkCuda( cudaMalloc(&d_a, n * 33 * sizeof(T)) );

  checkCuda( cudaEventCreate(&startEvent) );
  checkCuda( cudaEventCreate(&stopEvent) );

  printf("Offset, Bandwidth (GB/s):n");

  offset<<>>(d_a, 0); // warm up

  for (int i = 0; i <= 32; i++) {
    checkCuda( cudaMemset(d_a, 0.0, n * sizeof(T)) );

    checkCuda( cudaEventRecord(startEvent,0) );
    offset<<>>(d_a, i);
    checkCuda( cudaEventRecord(stopEvent,0) );
    checkCuda( cudaEventSynchronize(stopEvent) );

    checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
    printf("%d, %fn", i, 2*nMB/ms);
  }

  printf("n");
  printf("Stride, Bandwidth (GB/s):n");

  stride<<>>(d_a, 1); // warm up
  for (int i = 1; i <= 32; i++) {
    checkCuda( cudaMemset(d_a, 0.0, n * sizeof(T)) );

    checkCuda( cudaEventRecord(startEvent,0) );
    stride<<>>(d_a, i);
    checkCuda( cudaEventRecord(stopEvent,0) );
    checkCuda( cudaEventSynchronize(stopEvent) );

    checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
    printf("%d, %fn", i, 2*nMB/ms);
  }

  checkCuda( cudaEventDestroy(startEvent) );
  checkCuda( cudaEventDestroy(stopEvent) );
  cudaFree(d_a);
}

int main(int argc, char **argv)
{
  int nMB = 4;
  int deviceId = 0;
  bool bFp64 = false;

  for (int i = 1; i < argc; i++) {
    if (!strncmp(argv[i], "dev=", 4))
      deviceId = atoi((char*)(&argv[i][4]));
    else if (!strcmp(argv[i], "fp64"))
      bFp64 = true;
  }

  cudaDeviceProp prop;

  checkCuda( cudaSetDevice(deviceId) )
  ;
  checkCuda( cudaGetDeviceProperties(&prop, deviceId) );
  printf("Device: %sn", prop.name);
  printf("Transfer size (MB): %dn", nMB);

  printf("%s Precisionn", bFp64 ? "Double" : "Single");

  if (bFp64) runTest(deviceId, nMB);
  else       runTest(deviceId, nMB);?

}此代碼可以通過(guò)傳遞“ fp64 ”命令行選項(xiàng)以單精度(默認(rèn)值)或雙精度運(yùn)行偏移量?jī)?nèi)核和跨步內(nèi)核。每個(gè)內(nèi)核接受兩個(gè)參數(shù),一個(gè)輸入數(shù)組和一個(gè)表示訪問(wèn)數(shù)組元素的偏移量或步長(zhǎng)的整數(shù)。內(nèi)核在一系列偏移和跨距的循環(huán)中被稱為。

未對(duì)齊的數(shù)據(jù)訪問(wèn)

下圖顯示了 Tesla C870 、 C1060 和 C2050 上的偏移內(nèi)核的結(jié)果。

設(shè)備內(nèi)存中分配的數(shù)組由 CUDA 驅(qū)動(dòng)程序與 256 字節(jié)內(nèi)存段對(duì)齊。該設(shè)備可以通過(guò) 32 字節(jié)、 64 字節(jié)或 128 字節(jié)的事務(wù)來(lái)訪問(wèn)全局內(nèi)存。對(duì)于 C870 或計(jì)算能力為 1 . 0 的任何其他設(shè)備,半線程的任何未對(duì)齊訪問(wèn)(或半扭曲線程不按順序訪問(wèn)內(nèi)存的對(duì)齊訪問(wèn))將導(dǎo)致 16 個(gè)獨(dú)立的 32 字節(jié)事務(wù)。由于每個(gè) 32 字節(jié)事務(wù)只請(qǐng)求 4 個(gè)字節(jié),因此可以預(yù)期有效帶寬將減少 8 倍,這與上圖(棕色線)中看到的偏移量(不是 16 個(gè)元素的倍數(shù))大致相同,對(duì)應(yīng)于線程的一半扭曲。

對(duì)于計(jì)算能力為 1 . 2 或 1 . 3 的 Tesla C1060 或其他設(shè)備,未對(duì)準(zhǔn)訪問(wèn)的問(wèn)題較少?;旧希ㄟ^(guò)半個(gè)線程對(duì)連續(xù)數(shù)據(jù)的未對(duì)齊訪問(wèn)在幾個(gè)“覆蓋”請(qǐng)求的數(shù)據(jù)的事務(wù)中提供服務(wù)。由于未請(qǐng)求的數(shù)據(jù)正在傳輸,以及不同的半翹曲所請(qǐng)求的數(shù)據(jù)有些重疊,因此相對(duì)于對(duì)齊的情況仍然存在性能損失,但是這種損失遠(yuǎn)遠(yuǎn)小于 C870 。

計(jì)算能力為 2 . 0 的設(shè)備,如 Tesla C250 ,在每個(gè)多處理器中都有一個(gè) L1 緩存,其行大小為 128 字節(jié)。該設(shè)備將線程的訪問(wèn)合并到盡可能少的緩存線中,從而導(dǎo)致對(duì)齊對(duì)跨線程順序內(nèi)存訪問(wèn)吞吐量的影響可以忽略不計(jì)。

快速內(nèi)存訪問(wèn)

步幅內(nèi)核的結(jié)果如下圖所示。

對(duì)于快速的全局內(nèi)存訪問(wèn),我們有不同的看法。對(duì)于大步進(jìn),無(wú)論架構(gòu)版本如何,有效帶寬都很差。這并不奇怪:當(dāng)并發(fā)線程同時(shí)訪問(wèn)物理內(nèi)存中相距很遠(yuǎn)的內(nèi)存地址時(shí),硬件就沒(méi)有機(jī)會(huì)合并這些訪問(wèn)。從上圖中可以看出,在 Tesla C870 上,除 1 以外的任何步幅都會(huì)導(dǎo)致有效帶寬大幅降低。這是因?yàn)?compute capability 1 . 0 和 1 . 1 硬件需要跨線程進(jìn)行線性、對(duì)齊的訪問(wèn)以進(jìn)行合并,因此我們?cè)?offset 內(nèi)核中看到了熟悉的 1 / 8 帶寬。 Compute capability 1 . 2 及更高版本的硬件可以將訪問(wèn)合并為對(duì)齊的段( CC 1 . 2 / 1 . 3 上為 32 、 64 或 128 字節(jié)段,在 CC 2 . 0 及更高版本上為 128 字節(jié)緩存線),因此該硬件可以產(chǎn)生平滑的帶寬曲線。

當(dāng)訪問(wèn)多維數(shù)組時(shí),線程通常需要索引數(shù)組的更高維,因此快速訪問(wèn)是不可避免的。我們可以使用一種名為共享內(nèi)存的 CUDA 內(nèi)存來(lái)處理這些情況。共享內(nèi)存是一個(gè)線程塊中所有線程共享的片上內(nèi)存。共享內(nèi)存的一個(gè)用途是將多維數(shù)組的 2D 塊以合并的方式從全局內(nèi)存提取到共享內(nèi)存中,然后讓連續(xù)的線程跨過(guò)共享內(nèi)存塊。與全局內(nèi)存不同,對(duì)共享內(nèi)存的快速訪問(wèn)沒(méi)有懲罰。我們將在下一篇文章中詳細(xì)介紹共享內(nèi)存。

概括

在這篇文章中,我們討論了如何從 CUDA 內(nèi)核代碼中有效地訪問(wèn)全局內(nèi)存的一些方面。設(shè)備上的全局內(nèi)存訪問(wèn)與主機(jī)上的數(shù)據(jù)訪問(wèn)具有相同的性能特征,即數(shù)據(jù)局部性非常重要。在早期的 CUDA 硬件中,內(nèi)存訪問(wèn)對(duì)齊和跨線程的局部性一樣重要,但在最近的硬件上,對(duì)齊并不是什么大問(wèn)題。另一方面,快速的內(nèi)存訪問(wèn)會(huì)損害性能,使用片上共享內(nèi)存可以減輕這種影響。在下一篇文章中,我們將詳細(xì)探討共享內(nèi)存,之后的文章中,我們將展示如何使用共享內(nèi)存來(lái)避免在矩陣轉(zhuǎn)置過(guò)程中出現(xiàn)跨步全局內(nèi)存訪問(wèn)。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師,致力于 RAPIDS 。 Mark 擁有超過(guò) 20 年的 GPUs 軟件開(kāi)發(fā)經(jīng)驗(yàn),從圖形和游戲到基于物理的模擬,到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來(lái)納大學(xué)的博士生時(shí),他意識(shí)到了一種新生的趨勢(shì),并為此創(chuàng)造了一個(gè)名字: GPGPU (圖形處理單元上的通用計(jì)算)。

審核編輯:郭婷

聲明:本文內(nèi)容及配圖由入駐作者撰寫(xiě)或者入駐合作網(wǎng)站授權(quán)轉(zhuǎn)載。文章觀點(diǎn)僅代表作者本人,不代表電子發(fā)燒友網(wǎng)立場(chǎng)。文章及其配圖僅供工程師學(xué)習(xí)之用,如有內(nèi)容侵權(quán)或者其他違規(guī)問(wèn)題,請(qǐng)聯(lián)系本站處理。 舉報(bào)投訴
  • 存儲(chǔ)器
    +關(guān)注

    關(guān)注

    38

    文章

    7653

    瀏覽量

    167423
  • NVIDIA
    +關(guān)注

    關(guān)注

    14

    文章

    5309

    瀏覽量

    106422
收藏 人收藏
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

    評(píng)論

    相關(guān)推薦
    熱點(diǎn)推薦

    請(qǐng)問(wèn)STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet?

    STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet? 能否有現(xiàn)成的例程Samples借鑒?
    發(fā)表于 05-14 08:23

    請(qǐng)問(wèn)STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet?

    STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet? 能否有現(xiàn)成的例程Samples借鑒?
    發(fā)表于 04-27 07:06

    請(qǐng)問(wèn)STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet?

    STM32MP257CM33內(nèi)核能否訪問(wèn)以太網(wǎng)口Ethernet? 能否有現(xiàn)成的例程Samples借鑒?
    發(fā)表于 04-25 06:12

    Linux系統(tǒng)通過(guò)預(yù)留物理內(nèi)存實(shí)現(xiàn)ARM與FPGA高效通信的方法

    在嵌入式系統(tǒng)開(kāi)發(fā),ARM 和 FPGA 之間的通信可以使用 ARM 側(cè)的 DDR 作為通道。由于 FPGA 也可以直接訪問(wèn)到 ARM 側(cè) DDR,但 DDR 作為共享通信時(shí),就不能被操作系統(tǒng)的內(nèi)存
    的頭像 發(fā)表于 04-16 13:42 ?683次閱讀
    Linux系統(tǒng)<b class='flag-5'>中</b>通過(guò)預(yù)留物理<b class='flag-5'>內(nèi)存</b>實(shí)現(xiàn)ARM與FPGA高效通信的方法

    如何在Linux內(nèi)核5.18版本之后和64位架構(gòu)內(nèi)核空間調(diào)用ioctl?

    我嘗試在最近的內(nèi)核重新構(gòu)建以前版本 (4.19) 的 Linux 設(shè)備驅(qū)動(dòng)程序,即嵌入式平臺(tái)上的 6.1.22,ARM64 架構(gòu)。 驅(qū)動(dòng)程序管理 tty 設(shè)備。 當(dāng)我調(diào)用類似于用戶空
    發(fā)表于 04-02 06:06

    嵌入式學(xué)習(xí)-飛凌嵌入式ElfBoard ELF 1板卡-內(nèi)核空間與用戶空間的數(shù)據(jù)拷貝之?dāng)?shù)據(jù)拷貝介紹

    本帖最后由 jf_13411809 于 2025-3-20 14:10 編輯 在Linux系統(tǒng),內(nèi)核空間和用戶空間是兩個(gè)獨(dú)立的地址空間,它們有不同的訪問(wèn)權(quán)限和內(nèi)存保護(hù)機(jī)制。在
    發(fā)表于 03-20 11:50

    飛凌嵌入式ElfBoard ELF 1板卡-內(nèi)核空間與用戶空間的數(shù)據(jù)拷貝之?dāng)?shù)據(jù)拷貝介紹

    在Linux系統(tǒng),內(nèi)核空間和用戶空間是兩個(gè)獨(dú)立的地址空間,它們有不同的訪問(wèn)權(quán)限和內(nèi)存保護(hù)機(jī)制。在內(nèi)核空間和用戶空間之間進(jìn)行數(shù)據(jù)傳輸時(shí),需要
    發(fā)表于 03-19 08:55

    利用eBPF程序繞過(guò)內(nèi)核以加速存儲(chǔ)訪問(wèn)

    隨著微秒級(jí)NVMe存儲(chǔ)的蓬勃發(fā)展,Linux內(nèi)核存儲(chǔ)棧的開(kāi)銷(xiāo)幾乎是存儲(chǔ)訪問(wèn)時(shí)間的兩倍,已經(jīng)成為性能瓶頸。
    的頭像 發(fā)表于 03-01 16:09 ?531次閱讀
    利用eBPF程序繞過(guò)<b class='flag-5'>內(nèi)核</b>以加速存儲(chǔ)<b class='flag-5'>訪問(wèn)</b>

    如何有效地提高傳感器的測(cè)試精度

    問(wèn)題描述 如何有效地提高傳感器的測(cè)試精度是行業(yè)的發(fā)展趨勢(shì);近來(lái),對(duì)傳感器進(jìn)行實(shí)驗(yàn)測(cè)試過(guò)程中發(fā)現(xiàn)結(jié)果存在明顯的工頻干擾,信號(hào)夾雜有明顯噪音,具體頻率為50hz,因此,近來(lái)以解決實(shí)際問(wèn)題為出發(fā)點(diǎn)
    的頭像 發(fā)表于 01-24 10:55 ?1041次閱讀
    如何<b class='flag-5'>有效地</b>提高傳感器的測(cè)試精度

    如何有效地安裝孔隙水壓力計(jì)

    孔隙水壓力計(jì)作為一種重要的監(jiān)測(cè)工具,廣泛應(yīng)用于多種工程場(chǎng)景,包括士方填筑、混凝土澆筑、測(cè)壓管式埋設(shè)等。下面,峟思工程儀器將和大家詳細(xì)探討在這些具體應(yīng)用場(chǎng)景,如何有效地安裝孔隙水壓力計(jì)。在士方填筑
    的頭像 發(fā)表于 01-21 17:02 ?412次閱讀
    如何<b class='flag-5'>有效地</b>安裝孔隙水壓力計(jì)

    如何使用內(nèi)存加速存儲(chǔ)訪問(wèn)速度

    本篇文章是首爾大學(xué)發(fā)表在FAST 2023上的文章。隨著閃存容量的增加,邏輯地址到物理地址的映射表項(xiàng)也相應(yīng)增加。映射表項(xiàng)通常存放在設(shè)備控制器的SRAM來(lái)加速訪問(wèn)。然而由于成本問(wèn)題SRAM一直無(wú)法
    的頭像 發(fā)表于 12-19 10:54 ?720次閱讀
    如何使用<b class='flag-5'>內(nèi)存</b>加速存儲(chǔ)<b class='flag-5'>訪問(wèn)</b>速度

    嵌入式工程師都在找的【Linux內(nèi)核調(diào)試技術(shù)】建議收藏!

    在嵌入式系統(tǒng)的開(kāi)發(fā),Linux內(nèi)核調(diào)試是一個(gè)至關(guān)重要的環(huán)節(jié)。 隨著處理器技術(shù)的不斷進(jìn)步和嵌入式領(lǐng)域的蓬勃發(fā)展,掌握有效內(nèi)核調(diào)試技術(shù)成為了開(kāi)發(fā)者們的一項(xiàng)必備技能。本文將介紹幾種常見(jiàn)
    發(fā)表于 11-28 15:37

    內(nèi)存管理的硬件結(jié)構(gòu)

    常見(jiàn)的內(nèi)存分配函數(shù)有malloc,mmap等,但大家有沒(méi)有想過(guò),這些函數(shù)在內(nèi)核是怎么實(shí)現(xiàn)的?換句話說(shuō),Linux內(nèi)核內(nèi)存管理是怎么實(shí)現(xiàn)的
    的頭像 發(fā)表于 09-04 14:28 ?661次閱讀
    <b class='flag-5'>內(nèi)存</b>管理的硬件結(jié)構(gòu)

    影響內(nèi)存延遲的因素有哪些

    內(nèi)存延遲是指等待對(duì)系統(tǒng)內(nèi)存存儲(chǔ)數(shù)據(jù)的訪問(wèn)完成時(shí)引起的延期,它是衡量內(nèi)存響應(yīng)速度的重要指標(biāo)。影響內(nèi)存
    的頭像 發(fā)表于 09-04 11:46 ?3568次閱讀

    這個(gè)機(jī)子內(nèi)存是不是集成到板子上了,哪里可以插個(gè)接口訪問(wèn)內(nèi)存

    我這個(gè)機(jī)子內(nèi)存是不是集成到板子上了呀,哪里可以插個(gè)接口訪問(wèn)內(nèi)存呀。
    發(fā)表于 08-30 09:04