g u t x .com.tr İpek yolu - Çin'i anlamaya götürürüm

ARM GPU'nun tüm derin öğrenme performansını keşfedin, TVM optimizasyonu 2 kata kadar performans artışı sağlar

Leifeng.com AI Teknolojisi Yorumu: Bu makale, TVM tarafından Shanghai Jiaotong Üniversitesi Apex Lab'den bir lisans öğrencisi olan Lianmin Zheng tarafından yayınlanan bir blogdur.Makale, mobil terminalde ARM GPU'nun derin öğrenimini optimize etmek için TVM'nin nasıl kullanılacağını açıklamaktadır. Leifeng.com AI Technology Review orijinal metni derledi.

Derin öğrenmenin büyük başarısıyla birlikte, mobil cihazlara derin öğrenme sinir ağı modellerini yerleştirme ihtiyacı da hızla artıyor. Masaüstü platformunda yaptığımıza benzer şekilde, mobil cihazlarda GPU kullanmak aynı anda elde edilebilir. Çıkarım hesaplamalarını hızlandırın ve güç tasarrufu yapın . Ancak, mevcut derin öğrenme çerçevelerinin çoğu mobil GPU'ları iyi desteklemiyor. şüpheli Zorluk, mobil GPU ve masaüstü GPU arasındaki mimari farklılıklarda yatmaktadır. Bu, mobil GPU'yu optimize etmek için daha fazla özel çalışma gerektiği anlamına gelir. Çoğu derin öğrenme çerçevesinde mobil GPU'lar için destek eksikliğine yol açan bu ekstra çalışmadır.

TVM geçişi Farklı donanımlardaki dağıtım sorunlarını çözmek için birleşik bir IR yığını oluşturun Bu IR yığını sayesinde, farklı donanımlar için optimizasyonu kolayca tamamlayabilirsiniz. Bu makalede, ARM Mali GPU'lar için verimli çekirdekler oluşturmak ve uçtan uca derleme gerçekleştirmek için TVM / NNVM'nin nasıl kullanılacağını gösteriyoruz. Mali-T860 MP4 tabanlı testimizde, Arm Compute Library'ye kıyasla yöntemimiz VGG-16'da 1,4 kat, MobileNet'te 2,2 kat daha hızlı. Grafik seviyesi (Grafik seviyesi) ve işlem seviyesi (Operatör seviyesi) Eklemin optimizasyonu bu hızlanmayı destekler.

ImageNet'in çıkarım hızını farklı alt katmanlarda test edin

Mali Midgrad GPU

Firefly-RK3399'u test ortamımız olarak Mali-T860 MP4 ile birlikte kullanacağız, bu nedenle esas olarak aşağıdaki Mali T8xx'e odaklanıyoruz.

Mimari

Şekil 1, T860 ve T880'deki Mali mimari diyagramıdır. GPU, 16 bağlı gölgelendirici çekirdeğine (Coherent shader core) genişletilebilir. Her gölgelendirici çekirdeğinde, 2 veya 3 Aritmetik işlem hattı ve 1 yükleme / depolama ardışık düzeni (sözde TriPipe) vardır. Her işlem boru hattındaki ALU, dört 128-bit vektör birimine ve bir skaler birime sahiptir. GPU hesaplama için OpenCL kullanıyoruz. OpenCL modeline eşleme yaparken, her gölgelendirici çekirdeği bir veya daha fazla çalışma grubunu yürütmekten sorumludur. Ve her bir gölgelendirici çekirdeği, 384 adede kadar eşzamanlı yürütme iş parçacığını destekler. OpenCL'deki her iş öğesi genellikle Mali GPU'daki tek bir iş parçacığı ile eşlenir. Mali GPU, VLIW (Very Long Instruction Word) mimarisini kullanır. Her bir komut kelimesi birden fazla işlem içerir. Mali GPU ayrıca SIMD'yi de kullanabilir, bu nedenle çoğu aritmetik talimat aynı anda birden fazla veri öğesi üzerinde çalışacaktır.

Şekil 1. Mali T860 ve T880 (kaynak)

Nvidia GPU ile karşılaştırıldığında farklılıklar

NVIDIA GPU'larla karşılaştırıldığında, Mali GPU'lar için OpenCL kodu yazarken dikkat etmemiz gereken farklardan bazıları şunlardır.

Mali GPU, birleşik küresel bellek kullanır . Nvidia'nın GPU'sunda, verileri genellikle paylaşılan belleğe kopyalarız çünkü Nvidia'nın GPU'su küresel belleği, paylaşılan belleği ve kayıtları fiziksel olarak ayırır. Mali'de, bu kopyalama işlemi hesaplama performansını iyileştirmez, bu nedenle bu işlem kaldırılabilir. Ek olarak, Mali GPU genellikle CPU ile global belleği paylaşır, bu nedenle CPU ve GPU arasında veri aktarımına gerek yoktur.
Mali Midgrad GPU Evet SIMD (Tek Talimat Çoklu Veri) temel alınarak tasarlanmıştır ve açıkça vektörleştirilmesi gerekir . Nvidia'nın CUDA'sında paralellik, açık vektörleştirme olmaksızın SIMT (Single Instruction Multithreading) aracılığıyla elde edilir. Ancak, daha yeni Mali Bitfrost GPU'nun Dörtlü stil vektörleştirmeye dayandığını ve açık vektörleştirme gerektirmediğini de unutmayın.
Mali GPU'daki tüm iş parçacıkları bağımsız program sayaçlarına sahiptir . Bu, çözgü boyutunun 1 olduğu anlamına gelir, bu nedenle dal sapması büyük bir sorun değildir.

Optimizasyon: Evrişim işlemini örnek olarak alın

Evrişim katmanı, en derin sinir ağlarının çekirdeğidir ve hesaplama süresinin çoğunu alır. Bu nedenle, TVM'de paketleme, döşeme, açma ve vektörleştirme gibi yaygın tekniklerin nasıl uygulanacağını göstermek için evrişimi örnek olarak alıyoruz.

Im2Col uygulamak için GEMM kullanın

İyi bilinen evrişimli katman algoritması im2col'dur ve prensibi, küçük bir 3B girdi küpünü bir matrisin sütununa dönüştürmek ve GEMM algoritmasını yürütmektir. Bunun avantajı, yüksek düzeyde optimize edilmiş BLAS kitaplığının matris işlemlerine dönüştürüldükten sonra kullanılabilmesidir. Ancak bellek fazlalığı sorunu (3x3 evrişimin 9 katı bellek fazlalığı) da oldukça korkunç.

Mekansal Ambalaj

Bunun yerine, evrişimi hesaplamak için başka bir yöntem kullanıyoruz ve aşamalı olarak bazı optimizasyon tekniklerini uyguluyoruz. İnce ayar örneği olarak VGG-16'daki evrişimli katmanı kullanın ve yapılandırması aşağıda gösterilmiştir. Burada parti büyüklüğünün 1 olduğunu varsayıyoruz.

Karşılaştırma olarak, bu katmanın performansını Arm Hesaplama Kitaplığı'nda da listeledik.

Açıklama hesaplama süreci: döşeme ve paketleme

Döşeme ve paketleme, daha iyi bellek erişimi için iki yöntemdir. Çini işlemi Tüm hesaplamayı birden çok küçük parçaya bölün Daha iyi olmak için Verilerin yeniden kullanımı verim. Paketleme işlemi giriş matrisini döşemeye göre yeniden düzenler, böylece belleğe sıralı olarak erişebiliriz, böylece Önbellek kaçırma oranını azaltın .

Girdi görüntüsünün genişlik boyutu ve filtre matrisinin CO boyutu üzerinde döşeme işlemleri gerçekleştiriyoruz. Bu, tvm.compute koduyla ilan edilir.

# döşeme faktörünü ayarla

VH = 1VW = VC = 4

# girdi şekli al _, CI, IH, IW = data.shape

CO, CI, KH, KW = kernel.shape

TH = IH + 2 * H_PAD

TW = IW + 2 * W_PAD

# calc çıktı şekli

OH = (IH + 2 * H_PAD-KH) // H_STR + 1

OW = (IW + 2 * W_PAD-KW) // W_STR + 1

# paketlemeden sonraki veri şekli

dvshape = (N, TH // (VH * H_STRIDE), TW // (VW * W_STRIDE), CI, VH * H_STRIDE + HCAT, VW * W_STRIDE + WCAT)

# paketlemeden sonra çekirdek şekli

kvshape = (CO // VC, CI, KH, KW, VC)

ovshape = (N, CO // VC, OH // VH, OW // VW, VH, VW, VC)

oshape = (N, CO, OH, OW)

# paketlemeyi tanımla

data_vec = tvm.compute (dvshape, lambda n, h, w, ci, vh, vw: data_pad, name = 'data_vec')

kernel_vec = tvm.compute (kvshape, lambda co, ci, kh, kw, vc: kernel, name = 'kernel_vec')

# evrişimi tanımla

ci = tvm.reduce_axis ((0, CI), ad = 'ci')

kh = tvm.reduce_axis ((0, KH), ad = 'kh')

kw = tvm.reduce_axis ((0, KW), ad = 'kw')

conv = tvm.compute (ovshape, lambda n, co, h, w, vh, vw, vc:

tvm.sum (data_vec.astype (out_dtype) *

kernel_vec.astype (out_dtype),

eksen =), ad = 'dönüşüm')

# düzeni düzeltmek için paketten çıkarın

output = tvm.compute (oshape, lambda n, co, h, w:

dönş.

name = 'output_unpack', tag = 'direct_conv_output')

Tanımlanan IR'yi aşağıdaki kod aracılığıyla görüntüleyebiliriz.

baskı (tvm.lower (s ,, basit_mod = Doğru))

Burada evrişim kısmını seçtim.

dönüşüm üret {

for (co, 0, 64) {

için (h, 0, 56) {

for (w, 0, 14) {

for (vw.init, 0, 4) {

for (vc.init, 0, 4) {

dönş = 0.000000f

}

(ci, 0, 256) {için

for (kh, 0, 3) {

for (kw, 0, 3) {

for (vw, 0, 4) {

için (vc, 0, 4) {

dönş = (dönüşüm + (veri_vec * kernel_vec))

}

Çekirdek 1: İplik bağlama

TVM'de önce hesaplamaları açıklıyor sonra planlıyoruz. Bu mekanizma, algoritma ve uygulama ayrıntılarını ayırabilir. (Bu fikir Halide'den geliyor)

Aşağıdaki kod, eksenleri GPU iş parçacığına bağlar, böylece kodumuz Mali GPU'da çalışabilir.

# iş parçacığı bağlamak için yardımcı işlev

def tile_and_bind3d (s, tensor, z, y, x, z_factor = 2, y_factor = Yok, x_factor = Yok):

"" "döşeme ve 3d bağlama" ""

y_factor = y_factor veya z_factor

x_factor = x_factor veya y_factor

zo, zi = s.split (z, z_factor)

yo, yi = s.split (y, y_factor)

xo, xi = s.split (x, x_factor)

s.bind (zo, tvm.thread_axis ("blockIdx.z"))

s.bind (zi, tvm.thread_axis ("threadIdx.z"))

s.bind (yo, tvm.thread_axis ("blockIdx.y"))

s.bind (yi, tvm.thread_axis ("threadIdx.y"))

s.bind (xo, tvm.thread_axis ("blockIdx.x"))

s.bind (xi, tvm.thread_axis ("threadIdx.x"))

# ayarlanabilir parametreyi ayarla

num_thread = 8

# veri paketlemeyi programla

_, h, w, ci, vh, vw = s.op.axis

tile_and_bind3d (s, veri_vec, h, w, ci, 1)

# çekirdek paketlemeyi programla

co, ci, kh, kw, vc = s.op.axis

tile_and_bind (s, kernel_vec, co, ci, 1)

# program dönş.

_, c, h, w, vh, vw, vc = s.op.axis

kc, kh, kw = s.op.reduce_axis

s.reorder (_, c, h, w, vh, kc, kh, kw, vw, vc)

tile_and_bind3d (s, dönş., c, h, w, iplik_sayısı, 1, 1)

_, co, oh, ow = s.op.axis

tile_and_bind3d (s, çıktı, co, oh, ow, num_thread, 1, 1)

Bu kodlarla kodumuz çalışabilir, ancak performans çok kötü.

Çekirdek 2: İşlemi genişletme

Döngü açma kutusu Döngü kontrol talimatlarını azaltın, dal cezasını azaltın ve bellek okuma gecikmesini gizleyin . TVM'de bu, s.unroll (eksen) çağrılarak gerçekleştirilebilir.

# ayarlanabilir parametreyi ayarla

num_thread = 8

# veri paketlemeyi programla

_, h, w, ci, vh, vw = s.op.axis

tile_and_bind3d (s, veri_vec, h, w, ci, 1)

"" "!! BURAYA KAYDOLMAYI EKLE !!" ""

s.unroll (vw)

# çekirdek paketlemeyi programla

co, ci, kh, kw, vc = s.op.axis

tile_and_bind (s, kernel_vec, co, ci, 1)

"" "!! BURAYA KAYDOLMAYI EKLE !!" ""

s.unroll (kh)

s.unroll (kw)

s.unroll (vc)

# program dönş.

_, c, h, w, vh, vw, vc = s.op.axis

kc, kh, kw = s.op.reduce_axis

s.reorder (_, c, h, w, vh, kc, kh, kw, vw, vc)

tile_and_bind3d (s, dönş., c, h, w, iplik_sayısı, 1, 1)

"" "!! KAYDIRMAYI BURAYA EKLEYİN !!" ""

s.unroll (kh)

s.unroll (kw)

s.unroll (vw)

s.unroll (vc)

_, co, oh, ow = s.op.axis

tile_and_bind3d (s, çıktı, co, oh, ow, num_thread, 1, 1)

Çekirdek 3: Vektörleştirme

Daha önce de belirtildiği gibi, Mali GPU'larda en iyi performansı elde etmek için açıkça vektörleştirmemiz gerekiyor.

# ayarlanabilir parametre ayarla

ternum_thread = 8

# veri paketlemeyi programla

_, h, w, ci, vh, vw = s.op.axis

tile_and_bind3d (s, veri_vec, h, w, ci, 1)

# unroll

s.unroll (vw)

# çekirdek paketlemeyi programla

co, ci, kh, kw, vc = s.op.axis

tile_and_bind (s, kernel_vec, co, ci, 1)

# unroll

s.unroll (kh)

s.unroll (kw)

"" "!! BURADA VEKTÖRLEŞTİRİN !!" ""

s.vectorize (vc)

# program con

v_, c, h, w, vh, vw, vc = s.op.axis

kc, kh, kw = s.op.reduce_axis

s.reorder (_, c, h, w, vh, kc, kh, kw, vw, vc)

tile_and_bind3d (s, dönş., c, h, w, iplik_sayısı, 1, 1)

# unroll

s.unroll (kh)

s.unroll (kw)

s.unroll (vw)

"" "!! BURADA VEKTÖRLEŞTİRİN !!" ""

s.vectorize (vc)

_, co, oh, ow = s.op.axis

tile_and_bind3d (s, çıktı, co, oh, ow, num_thread, 1, 1)

Ayarlanabilir parametreleri ayarlayın

Yukarıdaki ayarlanabilir parametrelere gelince, bazıları hesaplanabilir. Vektörize edilmiş VC boyutu için 128 bitlik kayıtları doldurmalıyız, böylece float32128/32 = 4 ve float16128/16 = 8 olarak ayarlanabilir.

Ancak operasyonun karmaşıklığından dolayı en iyi hiperparametre değerlerini belirlememiz bizim için zordur. Bu yüzden TVM'de kullanıyoruz Izgara araması . Doğrudan OpenCL kodunu kullanmak yerine TVM'nin üst düzey IR'sine python kodu yazdığımız için çok etkili olabilir.

OpenCL kodu oluşturun

Oluşturulan OpenCL kodunu aşağıdaki kod aracılığıyla görebiliriz.

print (func.imported_modules.get_source)

OpenCL kodu buraya yapıştırmak için çok uzun ve çok fazla genişletme nedeniyle okunması zor. Eğer ilgileniyorsanız, buradan kontrol edebilirsiniz.

Uçtan uca kıyaslama

Bu bölümde, farklı alt katmanlar arasındaki performans farklılıklarını test etmek için bazı popüler derin öğrenme ağlarını kullanacağız. Test ortamımız:

Firefly-RK33994G

CPU: çift çekirdekli Cortex-A72 + dört çekirdekli Cortex-A53

GPU: Mali-T860MP4

Kol Hesaplama Kitaplığı: v17.12

MXNet: v1.0.1

Openblas: v0.2.18

Uçtan uca derlemeyi uygulamak için NNVM ve TVM kullanıyoruz.

verim

Şekil 2. ImageNet'in çıkarım hızını farklı katmanlarda test edin

Şekil 2'de gösterildiği gibi, ImageNet üzerinde çıkarım hızını test ediyoruz. Firefly-RK3399'da Mali GPU, 6 çekirdekli big.LITTLE CPU'dan 2 ila 4 kat daha hızlı olabilir. Uçtan uca ardışık düzenimiz, Arm Hesaplama Kitaplığından 1,4 ila 2,2 kat daha hızlıdır. Arm Hesaplama Kütüphanesi'nde GEMM ve direkt evrişim yöntemlerini kullanmaya çalışıyoruz.Bu test durumlarında GEMM yöntemi her zaman direkt yöntemden daha hızlıdır, bu nedenle sadece GEMM yönteminin sonuçlarını grafiklendiriyoruz.

Şekilde, Arm Hesaplama Kitaplığındaki resnet18 gibi bazı sonuçlar eksik. Bunun nedeni, Arm Hesaplama Kitaplığı'nın grafik çalışma zamanının şimdilik bağlantı atlama işlemlerini desteklememesi ve derinlemesine evrişimin uygulama etkisinin daha iyi olmasıdır. fark. Bu aynı zamanda NNVM yazılım yığınının avantajlarını da yansıtır.

Yarı hassas performans

Derin sinir ağlarının doğruluğu, özellikle mobil cihazların çıkarım süreci için çok önemli değildir. Düşük hassasiyetli aritmetik kullanmak, çıkarımı daha hızlı hale getirebilir. Ayrıca Mali GPU'da yarı hassas kayan nokta sayılarını da test ettik.

Tablo 1. ImageNet'te FP16'nın çıkarım hızı

Teorik olarak, FP16 hem tepe hesaplamasını ikiye katlayabilir hem de bellek ek yükünü yarıya indirebilir, böylece hızı ikiye katlayabilir. Ancak daha uzun vektörleştirme ve bazı parametreleri ayarlamak için daha iyi bir giriş şekli gerektirir.

Mobil cihazlarda daha fazla çalışma

İyileştirme için bazı alanlar olduğunu kabul ediyoruz, bunlar esas olarak grafik seviyesinde. Model sıkıştırma ve ağırlık ön düzeni gibi. NNVM'nin bir sonraki iyileştirmesi bu sorunları çözmeye çalışacaktır.