日B视频 亚洲,啪啪啪网站一区二区,91色情精品久久,日日噜狠狠色综合久,超碰人妻少妇97在线,999青青视频,亚洲一区二卡,让本一区二区视频,日韩网站推荐

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

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

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

通過GPU內(nèi)存訪問調(diào)整提高應(yīng)用程序性能

星星科技指導(dǎo)員 ? 來源:NVIDIA ? 作者:NVIDIA ? 2022-08-15 16:24 ? 次閱讀
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

NVIDIA GPU 具有強(qiáng)大的計算能力,通常需要高速傳輸數(shù)據(jù)才能部署這種能力。原則上,這是可能的,因為 GPU 也有很高的內(nèi)存帶寬,但有時他們需要程序員的幫助來飽和帶寬。在這篇博文中,我們研究了一種實現(xiàn)這一點的方法,并將其應(yīng)用于金融計算中的一個示例。我們將解釋在什么情況下這種方法可以很好地工作,以及如何找出這些情況是否適用于您的工作負(fù)載。

上下文

NVIDIA GPU 的力量來自大規(guī)模并行??梢詫?32 個線程的許多扭曲放置在流式多處理器( SM )上,等待輪到它們執(zhí)行。當(dāng)一個 warp 因任何原因暫停時, warp 調(diào)度程序?qū)⑶袚Q到另一個,開銷為零,確保 SM 始終有工作要做。在高性能 NVIDIA Ampere 100 ( A100 ) GPU 上,多達(dá) 64 個活動經(jīng)線可以共享一個 SM ,每個都有自己的資源。除此之外, A100 還有許多 SMs-108 ,它們都可以同時執(zhí)行 warp 指令。大多數(shù)指令都必須對數(shù)據(jù)進(jìn)行操作,而這些數(shù)據(jù)幾乎總是源自連接到 GPU 的設(shè)備內(nèi)存( DRAM )。 SM 上大量的翹曲也可能無法工作的一個主要原因是,它們正在等待來自內(nèi)存的數(shù)據(jù)。如果發(fā)生這種情況,并且內(nèi)存帶寬沒有得到充分利用,則可以重新組織程序以改進(jìn)內(nèi)存訪問并減少扭曲暫停,從而使程序更快完成。

第一步:寬負(fù)載

在之前的博客文章中,我們檢查了一個工作負(fù)載,該工作負(fù)載沒有充分利用 GPU 的可用計算和內(nèi)存帶寬資源。我們確定,在需要之前從內(nèi)存中預(yù)取數(shù)據(jù)可以大大減少內(nèi)存暫停并提高性能。當(dāng)預(yù)取不適用時,需要確定哪些其他因素可能會限制內(nèi)存子系統(tǒng)的性能。一種可能性是,向該子系統(tǒng)發(fā)出請求的速率太高。直觀地說,我們可以通過在每個加載指令中提取多個單詞來降低請求速率。最好用一個例子來說明這一點。

在本文的所有代碼示例中,大寫變量都是編譯時常量。 BLOCKDIMX 采用預(yù)定義變量 blockDim 的值。 x 、 出于某些目的,它必須是編譯時已知的常量,而出于其他目的,它有助于避免在運行時進(jìn)行計算。
原始代碼如下所示,index是計算數(shù)組索引的輔助函數(shù)。它隱式地假設(shè)只使用了一個一維線程塊,而派生它的激勵應(yīng)用程序則不是這樣。但是,它減少了代碼混亂,并且不會更改參數(shù)。

for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; ++k) { double c = big_array[index(pt, k)]; c += small_array[k] ; best = max(c, best); } final[pt] = best;
}

請注意,每個線程從建議命名的small_array中加載kmax個連續(xù)值。此陣列足夠小,完全適合一級緩存,但要求它以非常高的速率返回數(shù)據(jù)可能會出現(xiàn)問題。下面的更改表明,如果我們稍微重新構(gòu)造代碼并引入 double2 數(shù)據(jù)類型,則每個線程可以在同一條指令中發(fā)出兩個雙精度字的請求,這在 NVIDIA GPU 上本機(jī)支持;它將兩個雙精度字存儲在相鄰的內(nèi)存位置,可以使用字段選擇器“ x ”和“ y ”訪問這些位置。之所以這樣做,是因為每個線程都訪問small_array的連續(xù)元素。我們稱這種技術(shù)為 VZX28 。請注意,索引“k”上的內(nèi)部循環(huán)現(xiàn)在增加了 2 ,而不是 1 。

for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double c = big_array[index(pt, k)]; double2 val = *(double2 *) &small_array[k]; c += val.x; best = max(c, best); c = big_array[index(pt, k+1)]; c += val.y; best = max(c, best); } final[pt] = best;
}

有幾個注意事項。首先,我們沒有檢查kmax是否為偶數(shù)。如果沒有,修改后的k循環(huán)將執(zhí)行額外的迭代,我們需要編寫一些特殊代碼來防止這種情況發(fā)生。其次,我們沒有確認(rèn)small_array是否在 16 字節(jié)邊界上正確對齊。否則,寬荷載將失效。如果它是使用cudaMalloc分配的,它將自動在 256 字節(jié)的邊界上對齊。但是,如果使用指針算法將其傳遞給內(nèi)核,則需要執(zhí)行一些檢查。

接下來,我們檢查輔助函數(shù)指數(shù),發(fā)現(xiàn)它在 pt 中與系數(shù) 1 呈線性關(guān)系。因此,通過在一條指令中請求兩個雙精度值,我們可以對從 big \ U 數(shù)組獲取的值應(yīng)用類似的寬負(fù)載方法。對big_arraysmall_array的訪問之間的區(qū)別在于,現(xiàn)在 warp 中的連續(xù)線程訪問相鄰的數(shù)組元素。下面重構(gòu)的代碼將數(shù)組元素上的循環(huán)增量加倍big_array,現(xiàn)在每個線程在每次迭代中處理兩個數(shù)組元素。

for (pt = 2*threadIdx.x; pt < ptmax ; pt += 2*BLOCKDIMX ) { double best1 = 0.0, best2 = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double2 c1 = *(double2 *) &big_array[index(pt, k)]; double2 c2 = *(double2 *) &big_array[index(pt, k+1)]; double2 val = *(double2 *) &small_array[k]; c1.x += val.x; best1 = max(c1.x, best1); c2.x += val.y; best1 = max(c2.x, best1); c1.y += val.x; best2 = max(c1.y, best2); c2.y += val.y; best2 = max(c2.y, best2); } final[pt] = best1; final[pt+1] = best2;
}

與之前相同的注意事項也適用,現(xiàn)在應(yīng)該擴(kuò)展到ptmax的奇偶校驗和big_array的對齊。幸運的是,從中派生此示例的應(yīng)用程序滿足所有要求。下圖顯示了在應(yīng)用程序中重復(fù)多次的一組內(nèi)核的持續(xù)時間(以納秒為單位)。對于寬負(fù)載組合,內(nèi)核的平均加速比為 1.63 倍。

圖 1 :由于負(fù)載較寬,內(nèi)核持續(xù)時間減少

第二步:寄存器使用

我們可能想到此為止并宣布成功,但使用 NVIDIA Nsight Compute 對程序執(zhí)行的深入分析表明,即使我們將加載指令的數(shù)量減少了一半,我們也沒有從根本上改變對內(nèi)存子系統(tǒng)的請求速率。原因是一條扭曲加載指令(即 32 個線程同時發(fā)出加載指令)會導(dǎo)致一個或多個扇區(qū)請求,這是硬件處理的實際內(nèi)存訪問單元。每個扇區(qū)是 32 字節(jié),因此每個線程一條 8 字節(jié)雙精度字的扭曲加載指令會導(dǎo)致 8 個扇區(qū)請求(訪問以單位跨距進(jìn)行),而一條雙精度字的扭曲加載指令會導(dǎo)致 16 個扇區(qū)請求。普通負(fù)載和寬負(fù)載的扇區(qū)請求總數(shù)相同。那么,是什么導(dǎo)致了性能的提高呢?

為了理解代碼行為,我們需要考慮一個尚未討論的資源,即寄存器。這些用于存儲從內(nèi)存加載的數(shù)據(jù),并用作算術(shù)指令的輸入。寄存器是一種有限的資源。如果流式多處理器( SM )在 A100 GPU 上承載盡可能多的扭曲,則每個線程可以使用 32 個 4 字節(jié)寄存器,這些寄存器總共可以容納 16 個雙精度字。將代碼翻譯成機(jī)器語言的編譯器知道這一點,并將限制每個線程的寄存器數(shù)量。我們?nèi)绾未_定代碼的寄存器使用及其在性能中所起的作用?我們使用 Nsight Compute 中的“ source ”視圖來并排查看匯編代碼(“ SASS ”)和 C 源代碼。

代碼的最內(nèi)層循環(huán)是執(zhí)行次數(shù)最多的循環(huán),因此,如果我們在導(dǎo)航菜單中選擇“已執(zhí)行的指令”,然后要求轉(zhuǎn)到 SASS 代碼中數(shù)量最多的那一行,我們會自動進(jìn)入內(nèi)部循環(huán)。如果不確定,可以將 SASS 與突出顯示的相應(yīng)源代碼進(jìn)行比較以確認(rèn)。接下來,我們在內(nèi)環(huán)的 SASS 代碼中識別從內(nèi)存( LDG )加載數(shù)據(jù)的所有指令。圖 2 顯示了 SASS 的一個片段,我們在其中搜索以找到內(nèi)部循環(huán)的開始;在第 166 行,指令的執(zhí)行次數(shù)突然跳到其最大值。

圖 2 :演示內(nèi)部循環(huán)開始的 SASS 代碼段(第 166 行)

LDG 。 E 、 64 是我們所追求的指令。它從全局內(nèi)存( DRAM )加載一個具有擴(kuò)展地址的 64 位字。寬單詞的負(fù)載對應(yīng)于 LDG 。 E 、 128 。加載指令名稱后的第一個參數(shù)(圖 2 中的 R34 )是接收該值的寄存器。由于雙精度值占用兩個相鄰寄存器,因此加載指令中隱含 R35 。接下來,我們比較三個版本的代碼( 1.基線, 2.寬負(fù)載的small_array, 3.寬負(fù)載的small_array和big_array)在內(nèi)部循環(huán)中使用寄存器的方式。回想一下,編譯器試圖保持在限制范圍內(nèi),有時需要對寄存器進(jìn)行處理。也就是說,如果沒有足夠的寄存器可用于從內(nèi)存接收每個唯一值,它將重用以前在內(nèi)部循環(huán)中使用的寄存器。

這樣做的結(jié)果是,算術(shù)指令需要使用以前的值,以便新值可以覆蓋它。此時,從內(nèi)存加載需要等待該指令完成:內(nèi)存延遲暴露。在所有現(xiàn)代計算機(jī)體系結(jié)構(gòu)上,此延遲構(gòu)成了一個顯著的延遲。在 GPU 上,可以通過切換到另一個扭曲來隱藏部分扭曲,但通常不是全部扭曲。因此,寄存器在內(nèi)環(huán)中被重用的次數(shù)可以表示代碼的速度變慢。

有了這一見解,我們分析了代碼的三個版本,發(fā)現(xiàn)它們在每個內(nèi)部循環(huán)中分別經(jīng)歷了 8 、 6 和 3 個內(nèi)存延遲,這解釋了圖 1 所示的性能差異。不同寄存器重用模式背后的主要原因是,當(dāng)兩個普通加載融合為單個寬加載時,通常需要更少的地址計算,并且地址計算的結(jié)果也會進(jìn)入寄存器。隨著持有地址的寄存器越來越多,剩下來充當(dāng)從內(nèi)存中提取的值的“著陸區(qū)”的地址越來越少,我們在 Music chairs 游戲中失去了席位;寄存器壓力增大。

第三步:啟動邊界

我們還沒有完成?,F(xiàn)在我們知道了寄存器在程序性能中所起的關(guān)鍵作用,我們將查看三個版本的代碼使用的寄存器總數(shù)。最簡單的方法是再次檢查 Nsight Compute 報告。我們發(fā)現(xiàn)使用的寄存器數(shù)量分別為 40 、 36 和 44 。

編譯器確定這些數(shù)字的方法是使用復(fù)雜的啟發(fā)式算法,該算法考慮了大量因素,包括 SM 上可能存在多少活動扭曲、在忙循環(huán)中加載的唯一值的數(shù)量以及每個操作所需的寄存器數(shù)量。如果編譯器不知道 SM 上可能存在的扭曲數(shù),它將嘗試將每個線程的寄存器數(shù)限制為 32 ,因為如果存在硬件允許的絕對最大同時扭曲數(shù)( 64 ),那么這就是可用的數(shù)字。在我們的例子中,我們沒有告訴編譯器期望的是什么,所以它盡了最大努力,但顯然確定僅使用 32 個寄存器生成的代碼效率太低。

然而,內(nèi)核的 launch 語句中指定的線程塊的實際大小是 1024 個線程,因此有 32 個扭曲。這意味著,如果 SM 上只存在一個線程塊,則每個線程最多可以使用 64 個線程。在實際使用的每個線程中有 40 、 36 和 44 個寄存器時,沒有足夠的寄存器可用于支持每個 SM 的兩個或多個線程塊,因此將只啟動一個,每個線程分別保留 24 、 28 和 20 個未使用的寄存器。

通過使用 launch bounds 將我們的意圖告知編譯器,我們可以做得更好。通過告訴編譯器一個線程塊中的最大線程數(shù)( 1024 )和同時支持的最小塊數(shù)( 1 ),編譯器可以放松,并且很高興每個線程分別使用 63 、 56 和 64 個寄存器。

有趣的是,最快的代碼版本現(xiàn)在是基線版本,沒有任何廣泛的負(fù)載。雖然組合寬負(fù)載 without 啟動邊界的加速比為 1.64 倍,但寬負(fù)載 with 啟動邊界的加速比為 1.76 倍,而基線代碼的加速比為 1.77 倍。這意味著我們不必費心修改內(nèi)核定義;在這種情況下,僅提供啟動邊界就足以獲得這種特定線程塊大小的最佳性能。

通過對 SM 上的線程塊大小和預(yù)期的最小線程塊數(shù)進(jìn)行更多的實驗,我們在每個 SM 有 512 個線程的 2 個線程塊的情況下達(dá)到了 1.79 倍的加速,對于沒有寬負(fù)載的基線版本也是如此。

結(jié)論

寄存器的有效使用對于獲得良好的 GPU 內(nèi)核性能至關(guān)重要。有時,一種稱為“寬負(fù)載”的技術(shù)可以帶來顯著的好處。它減少了計算并需要存儲在寄存器中的內(nèi)存地址的數(shù)量,留下更多的寄存器來接收來自內(nèi)存的數(shù)據(jù)。然而,向編譯器提示在應(yīng)用程序中啟動內(nèi)核的方式可能會帶來同樣的好處,而無需更改內(nèi)核本身。

關(guān)于作者

Rob Van der Wijngaart 是 NVIDIA 的高級高性能計算( HPC )架構(gòu)師。他在各種工業(yè)和政府實驗室從事 HPC 領(lǐng)域的研究超過三十年,是廣泛使用的 NAS 并行基準(zhǔn)測試的共同開發(fā)者。

Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級產(chǎn)品營銷經(jīng)理。弗雷德?lián)碛屑又荽髮W(xué)戴維斯分校計算機(jī)科學(xué)和數(shù)學(xué)學(xué)士學(xué)位。他的職業(yè)生涯開始于一名 UNIX 軟件工程師,負(fù)責(zé)將內(nèi)核服務(wù)和設(shè)備驅(qū)動程序移植到 x86 體系結(jié)構(gòu)。

審核編輯:郭婷


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

    關(guān)注

    68

    文章

    20339

    瀏覽量

    255355
  • NVIDIA
    +關(guān)注

    關(guān)注

    14

    文章

    5696

    瀏覽量

    110142
  • gpu
    gpu
    +關(guān)注

    關(guān)注

    28

    文章

    5283

    瀏覽量

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

掃碼添加小助手

加入工程師交流群

    評論

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

    內(nèi)存要取代GPU?HBM之父警告:以英偉達(dá)GPU為核心的架構(gòu)要被顛覆

    主板和CPU成為了主角。 ? 而最近“HBM之父”金正浩教授也語出驚人,提出未來內(nèi)存將成為主角:“GPU和CPU將會被集成到內(nèi)存(HBM和HBF)里,淪為內(nèi)存中的一個組件”。 ? 倒反
    的頭像 發(fā)表于 04-03 09:54 ?7242次閱讀
    <b class='flag-5'>內(nèi)存</b>要取代<b class='flag-5'>GPU</b>?HBM之父警告:以英偉達(dá)<b class='flag-5'>GPU</b>為核心的架構(gòu)要被顛覆

    運行測試程序以讀取通過受信任應(yīng)用程序 (TA) 存儲的安全 blob 時,內(nèi)存不足怎么解決?

    當(dāng)我運行測試程序以讀取通過受信任應(yīng)用程序 (TA) 存儲的安全 blob 時,我遇到了內(nèi)存不足 (OOM) 問題。 我仔細(xì)觀察了代碼,但沒有發(fā)現(xiàn)任何
    發(fā)表于 04-10 10:52

    C語言訪問某特定內(nèi)存位置

    嵌入式系統(tǒng)經(jīng)常具有要求程序員去訪問某特定的內(nèi)存位置的特點。在某工程中,要求設(shè)置一絕對地址為0x67a9的整型變量的值為0xaa66。編譯器是一個純粹的ANSI編譯器。寫代碼去完成這一任務(wù)。 考察點
    發(fā)表于 12-22 15:42

    內(nèi)存與數(shù)據(jù)處理優(yōu)化藝術(shù)

    內(nèi)存訪問程序運行的瓶頸之一。減少內(nèi)存訪問次數(shù)可以顯著提高程序的運行速度。 在C語言中,指針是直
    發(fā)表于 11-14 07:46

    通過sysmem接口擴(kuò)展內(nèi)存空間

    和0x90000000起始的64k空間范圍內(nèi)時,內(nèi)核會訪問ITCM和DTCM;如果不在上述空間范圍內(nèi),內(nèi)核會通過sysmem接口訪問外部存儲器。這里通過sysmem接口擴(kuò)展
    發(fā)表于 10-24 08:12

    提高RISC-V在Drystone測試中得分的方法

    性能內(nèi)存的讀寫速度、延遲和帶寬等都會影響到 Drystone 的性能。 指令集優(yōu)化:對RISC-V指令集的優(yōu)化也會影響性能。例如,對于特定的應(yīng)用或計算任務(wù),可以
    發(fā)表于 10-21 13:58

    蜂鳥E203內(nèi)核優(yōu)化方法

    。 修改內(nèi)核參數(shù):對蜂鳥E203的內(nèi)核參數(shù)進(jìn)行相應(yīng)修改,可以優(yōu)化內(nèi)核運行效率,提高系統(tǒng)性能,比如調(diào)整緩存大小、內(nèi)存分配策略等。 資源管理:進(jìn)行有針對的資源管理,例如調(diào)度算法的修改,
    發(fā)表于 10-21 07:55

    學(xué)生適合使用的SOLIDWORKS 云應(yīng)用程序

    隨著科技的不斷發(fā)展,計算機(jī)輔助設(shè)計(CAD)技術(shù)已經(jīng)成為現(xiàn)代工程教育的重要組成部分。SOLIDWORKS作為一款CAD軟件,其教育版云應(yīng)用程序為學(xué)生提供了強(qiáng)大而靈活的設(shè)計平臺。本文將探討
    的頭像 發(fā)表于 09-15 10:39 ?968次閱讀
    學(xué)生適合使用的SOLIDWORKS 云<b class='flag-5'>應(yīng)用程序</b>

    樹莓派5超頻指南:安全高效地提升性能!

    為什么要對樹莓派5進(jìn)行超頻?對樹莓派進(jìn)行超頻,可通過提高CPU和GPU的時鐘頻率來釋放額外的性能。在需要額外處理能力以提高響應(yīng)速度、減少延遲
    的頭像 發(fā)表于 08-14 17:45 ?2942次閱讀
    樹莓派5超頻指南:安全高效地提升<b class='flag-5'>性能</b>!

    Linux系統(tǒng)性能指南

    Linux服務(wù)器運行了很多應(yīng)用,在高負(fù)載下,服務(wù)器可能會出現(xiàn)性能瓶頸,例如CPU利用率過高、內(nèi)存不足、磁盤I/O瓶頸等,從而導(dǎo)致系統(tǒng)卡頓,服務(wù)無法正常運行等問題。所以針對以上問題,可以通過調(diào)整
    的頭像 發(fā)表于 06-23 14:12 ?1959次閱讀
    Linux系統(tǒng)<b class='flag-5'>性能</b>指南

    【「算力芯片 | 高性能 CPU/GPU/NPU 微架構(gòu)分析」閱讀體驗】+NVlink技術(shù)從應(yīng)用到原理

    帶來了總雙向帶寬160GB/s的通訊速率,遠(yuǎn)高于當(dāng)時的PCIe接口(實際比現(xiàn)在的PCIe5.0也還要快)。首代的NVlink主要是增強(qiáng)了GPUGPU的通信性能GPU對系統(tǒng)
    發(fā)表于 06-18 19:31

    HarmonyOS優(yōu)化應(yīng)用內(nèi)存占用問題性能優(yōu)化四

    的尺寸大小,使其與組件的大小保持一致。這樣可以避免不必要的內(nèi)存浪費,并提高應(yīng)用程序性能和效率。開發(fā)者可以使用圖像處理工具來調(diào)整圖像的尺寸大
    發(fā)表于 05-24 17:20

    HarmonyOS優(yōu)化應(yīng)用內(nèi)存占用問題性能優(yōu)化一

    :開發(fā)者可通過該接口監(jiān)聽系統(tǒng)內(nèi)存的變化,并根據(jù)系統(tǒng)內(nèi)存的實時情況,動態(tài)地調(diào)整應(yīng)用程序內(nèi)存,以避
    發(fā)表于 05-21 11:27

    如何使用USB中斷傳輸方法訪問FPGA?

    我目前正在設(shè)計一個可以通過 CY7C65216 從 Windows PC 訪問 FPGA 的單元。 我正在考慮使用USB中斷傳輸方法訪問FPGA。 這可能嗎? 如果有,是否有任何示例軟件程序
    發(fā)表于 05-19 06:04

    如何使用CYUSB3KIT-003使用GPIO訪問SRAM的應(yīng)用程序?

    你好。我是CYUSB3的初學(xué)者。 我想創(chuàng)建一個使用 CYUSB3KIT-003 使用 GPIO 訪問 SRAM 的應(yīng)用程序。 目前我已經(jīng)在我的電腦上安裝了SDK,但是有什么參考資料嗎?
    發(fā)表于 05-14 06:51
    河东区| 孟津县| 和田市| 长乐市| 蒲江县| 应城市| 弋阳县| 巴林右旗| 秦皇岛市| 鹤岗市| 台南市| 怀安县| 崇州市| 调兵山市| 怀仁县| 新闻| 五大连池市| 轮台县| 伊春市| 宣恩县| 石河子市| 河池市| 长春市| 峡江县| 永新县| 湟源县| 博客| 九龙坡区| 江安县| 金堂县| 纳雍县| 尼玛县| 冷水江市| 景泰县| 舟山市| 迭部县| 安陆市| 磴口县| 育儿| 高要市| 延边|