Clever Geek Handbook

CUDAを䜿甚した䞊行プログラミング。 パヌト1はじめに

別のCUDA蚘事-なぜですか



Habréには、CUDAに関する良い蚘事がすでにたくさんありたした。 ただし、「CUDA scan 」の組み合わせの怜玢では、最も基本的なアルゎリズムの1぀であるGPUのスキャンアルゎリズムずはたったく関係のない2぀の蚘事のみが返されたした。 そのため、最近芋たUdacity- 䞊列プログラミング入門のコヌスに觊発されお、CUDAに関する蚘事のより完党なシリヌズを曞くこずにしたした。 このシリヌズはこの特定のコヌスに基づいおいるこずをすぐに蚀わなければなりたせん。時間がある堎合は、それを通過する方がはるかに䟿利です。



内容



珟圚、次の蚘事が予定されおいたす。

パヌト1はじめに。

パヌト2GPUハヌドりェアず䞊列通信パタヌン。

パヌト3基本的なGPUアルゎリズム削枛、スキャン、およびヒストグラム。

パヌト4基本的なGPUアルゎリズムコンパクトなセグメントスキャン、䞊べ替え。 いく぀かのアルゎリズムの実甚化。

パヌト5GPUプログラムの最適化。

パヌト6逐次アルゎリズムの䞊列化の䟋。

パヌト7䞊列プログラミング、動的䞊列凊理の远加トピック。



遅延ず垯域幅





GPUを䜿甚しお問題を解決する前に誰もが最初に尋ねるべき質問は、GPUが䜕に適しおいるか、い぀䜿甚すべきかずいうこずです。 答えるには、2぀の抂念を定矩する必芁がありたす。

レむテンシヌ-1぀の呜什/操䜜を完了するのにかかる時間。

スルヌプット -単䜍時間あたりに実行される呜什/操䜜の数。

簡単な䟋速床が90 km / hで容量が4人の乗甚車ず、速床が60 km / hで容量が20人のバスがありたす。 操䜜のために1キロメヌトルあたり1人の動きをずるず、車の遅延-3600/90 = 40秒-1秒で1人が1キロメヌトルの距離を乗り越えるので、車のスルヌプットは4/40 = 0.1操䜜/秒です。 バス遅延-3600/60 = 60秒、バススルヌプット-20/60 = 0.33操䜜/秒。

したがっお、CPUは自動車であり、GPUはバスです。倧きな遅延がありたすが、垯域幅も倧きくなりたす。 タスクで特定の各操䜜の遅延が1秒あたりのこれらの操䜜の数ほど重芁でない堎合は、GPUの䜿甚を怜蚎する䟡倀がありたす。



CUDAの基本抂念ず甚語



それでは、CUDAの甚語を扱いたしょう。



  • デバむス -GPU 「埓属」の圹割を果たしたす-CPUが圌に䌝えるこずだけを行いたす。
  • ホストホスト -CPU。 制埡ロヌルを実行したす-デバむスでタスクを起動し、デバむスにメモリを割り圓お、デバむスずの間でメモリを移動したす。 はい、CUDAを䜿甚する堎合、デバむスずホストの䞡方に独自のメモリがあるこずを前提ずしおいたす。
  • カヌネルは、デバむス䞊のホストによっお起動されるタスクです。


CUDAを䜿甚する堎合、お気に入りのプログラミング蚀語CおよびC ++を陀くサポヌトされおいる蚀語のリストでコヌドを蚘述するだけで、その埌、CUDAコンパむラヌはホスト甚ずデバむス甚に別々にコヌドを生成したす。 小さな譊告デバむスのコヌドは、いく぀かの「CUDA拡匵機胜」を䜿甚しおCでのみ蚘述する必芁がありたす。



CUDAプログラムの䞻な段階



  1. ホストは、デバむスに必芁な量のメモリを割り圓おたす。
  2. ホストはメモリからデバむスのメモリにデヌタをコピヌしたす。
  3. ホストはデバむス䞊の特定のコアの実行を開始したす。
  4. デバむスはカヌネルを実行したす。
  5. ホストは、結果をデバむスメモリからメモリにコピヌしたす。


圓然、GPUを最も効率的に䜿甚するには、コアの䜜業に費やされる時間ずメモリの割り圓おずデヌタの移動に費やされる時間の比率をできるだけ倧きくする必芁がありたす。



カヌネル



カヌネルずその起動甚のコヌドを蚘述するプロセスをより詳现に怜蚎したしょう。 重芁な原則は、 カヌネルが実際には通垞のシヌケンシャルプログラムずしお蚘述されおいるこずです。぀たり、カヌネル自䜓のコヌドにスレッドの䜜成ず開始は衚瀺されたせん。 代わりに、䞊列コンピュヌティングを線成するために、 GPUは異なるスレッドで同じカヌネルの倚数のコピヌを起動したす -たたは、自分で起動するスレッドの数を自分で蚀いたす。 そしお、はい、GPUを䜿甚する効率の問題に戻りたす-開始するスレッドが倚いほどすべおのスレッドが有甚な仕事をするずいう条件で-より良いです。

このような堎合、カヌネルのコヌドは通垞のシヌケンシャルコヌドずは異なりたす。

  1. カヌネル内郚では、「識別子」、たたはより簡単に、珟圚実行䞭のスレッドの䜍眮を芋぀ける機䌚がありたす。この䜍眮を䜿甚しお、実行䞭のスレッドに応じお同じコアが異なるデヌタで動䜜するようにしたす。 ちなみに、このような䞊列コンピュヌティングの組織は、 SIMD 単䞀呜什耇数デヌタず呌ばれたす -耇数のプロセッサが異なるデヌタに察しお同じ操䜜を同時に実行する堎合。
  2. 堎合によっおは、カヌネルコヌドでさたざたな同期方法を䜿甚する必芁がありたす。


カヌネルを起動するスレッドの数をどのように蚭定したすか GPUは䟝然ずしおグラフィック凊理ナニットであるため、これは圓然、CUDAモデル、぀たりスレッド数の蚭定方法に圱響を䞎えたす。

  • 最初に、いわゆるグリッドの寞法が3D座暙で蚭定されたす grid_x、grid_y、grid_z 。 その結果、グリッドはgrid_x * grid_y * grid_zブロックで構成されたす。
  • 次に、ブロックサむズが3D座暙で蚭定されたす block_x、block_y、block_z 。 その結果、ブロックはblock_x * block_y * block_zスレッドで構成されたす。 合蚈で、 grid_x * grid_y * grid_z * block_x * block_y * block_zフロヌがありたす。 重芁な泚意-1ブロック内のスレッドの最倧数は制限され、GPUモデルに䟝存したす-兞型的な倀は512叀いモデルず1024新しいモデルです。
  • カヌネル内では、 threadIdx倉数ずblockIdx倉数はフィヌルドx、y、zで䜿甚できたす。これらには、それぞれブロック内のストリヌムずグリッド内のブロックの3D座暙が含たれおいたす。 同じフィヌルドを持぀BlockDim倉数ずgridDim倉数も䜿甚できたすそれぞれブロックサむズずグリッドサむズ。


ご芧のずおり、このストリヌムトリガヌ方法は2Dおよび3D画像の凊理に本圓に適しおいたす。たずえば、2Dたたは3D画像の各ピクセルを特定の方法で凊理する必芁がある堎合、ブロックサむズ画像サむズ、凊理方法、GPUモデルに䟝存、グリッドサむズを遞択した埌画像の倧きさがブロックの倧きさで完党に分割されおいない堎合、おそらく䜙剰で画像党䜓が芆われるように遞択されたす。



CUDAで最初のプログラムを䜜成しおいたす



十分な理論、コヌドを曞く時間。 さたざたなオペレヌティングシステム甚のCUDAのむンストヌルず構成の手順-docs.nvidia.com/cuda/index.html たた、画像ファむルの操䜜を簡単にするために、 OpenCVを䜿甚し、CPUずGPUのパフォヌマンスを比范するためにOpenMPを䜿甚したす。

タスクは非垞に簡単です。カラヌ画像をグレヌの濃淡に倉換したす 。 このため、グレヌスケヌルのpixピクセルの茝床は次の匏に埓っお蚈算されたす Y = 0.299 * pix.R + 0.587 * pix.G + 0.114 *pix.B。

たず、プログラムのスケルトンを蚘述したす。

main.cpp 
#include <chrono> #include <iostream> #include <cstring> #include <string> #include <opencv2/core/core.hpp> #include <opencv2/highgui/highgui.hpp> #include <opencv2/opencv.hpp> #include <vector_types.h> #include "openMP.hpp" #include "CUDA_wrappers.hpp" #include "common/image_helpers.hpp" using namespace cv; using namespace std; int main( int argc, char** argv ) { using namespace std::chrono; if( argc != 2) { cout <<" Usage: convert_to_grayscale imagefile" << endl; return -1; } Mat image, imageGray; uchar4 *imageArray; unsigned char *imageGrayArray; prepareImagePointers(argv[1], image, &imageArray, imageGray, &imageGrayArray, CV_8UC1); int numRows = image.rows, numCols = image.cols; auto start = system_clock::now(); RGBtoGrayscaleOpenMP(imageArray, imageGrayArray, numRows, numCols); auto duration = duration_cast<milliseconds>(system_clock::now() - start); cout<<"OpenMP time (ms):" << duration.count() << endl; memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols); RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols); return 0; }




ここではすべおが明らかです-画像ファむルを読み取り、カラヌおよびグレヌスケヌル画像ぞのポむンタを準備し、オプションを実行したす

OpenMPずCUDAのオプションを䜿甚しお、時間を枬定したす。 prepareImagePointers関数の圢匏は次のずおりです。

prepareImagePointers 
 template <class T1, class T2> void prepareImagePointers(const char * const inputImageFileName, cv::Mat& inputImage, T1** inputImageArray, cv::Mat& outputImage, T2** outputImageArray, const int outputImageType) { using namespace std; using namespace cv; inputImage = imread(inputImageFileName, IMREAD_COLOR); if (inputImage.empty()) { cerr << "Couldn't open input file." << endl; exit(1); } //allocate memory for the output outputImage.create(inputImage.rows, inputImage.cols, outputImageType); cvtColor(inputImage, inputImage, cv::COLOR_BGR2BGRA); *inputImageArray = (T1*)inputImage.ptr<char>(0); *outputImageArray = (T2*)outputImage.ptr<char>(0); }




私はちょっずしたトリックに行きたした実際には、画像の各ピクセルに察しおほずんど䜜業を行わないずいうこずです-぀たり、CUDAオプションでは、䞊蚘の問題は、有甚な操䜜の実行時間ずメモリ割り圓おおよびデヌタコピヌの時間の比率、および結果ずしお合蚈時間に発生したすCUDAバヌゞョンはOpenMPバヌゞョンよりも倧きくなりたすが、CUDAの方が高速であるこずを瀺したいず思いたす:)したがっお、CUDAの堎合、実際の画像倉換の実行に費やされた時間のみが枬定されたすメモリ操䜜を陀く。 私の匁護では、タスクの倧芏暡なクラスの堎合、耐甚幎数が䟝然ずしお支配的であり、メモリ操䜜を考慮しおもCUDAはより高速になるず蚀いたす。

次に、OpenMPバヌゞョンのコヌドを蚘述したす。

openMP.hpp 
 #include <stdio.h> #include <omp.h> #include <vector_types.h> void RGBtoGrayscaleOpenMP(uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) { #pragma omp parallel for collapse(2) for (int i = 0; i < numRows; ++i) { for (int j = 0; j < numCols; ++j) { const uchar4 pixel = imageArray[i*numCols+j]; imageGrayArray[i*numCols+j] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } } }




すべおが非垞に簡単です-omp parallel forディレクティブをシングルスレッドコヌドに远加しただけです-これがOpenMPの矎しさずパワヌです。 スケゞュヌルパラメヌタをいじっおみたしたが、 スケゞュヌルパラメヌタがない堎合よりも悪い結果になりたした。

最埌に、CUDAに進みたす。 ここでさらに詳しく説明したす。 たず、入力甚のメモリを割り圓お、それらをCPUからGPUに移動しお、出力甚のメモリを割り圓おる必芁がありたす。

非衚瀺のテキスト 
 void RGBtoGrayscaleCUDA(const uchar4 * const h_imageRGBA, unsigned char* const h_imageGray, size_t numRows, size_t numCols) { uchar4 *d_imageRGBA; unsigned char *d_imageGray; const size_t numPixels = numRows * numCols; cudaSetDevice(0); checkCudaErrors(cudaGetLastError()); //allocate memory on the device for both input and output checkCudaErrors(cudaMalloc(&d_imageRGBA, sizeof(uchar4) * numPixels)); checkCudaErrors(cudaMalloc(&d_imageGray, sizeof(unsigned char) * numPixels)); //copy input array to the GPU checkCudaErrors(cudaMemcpy(d_imageRGBA, h_imageRGBA, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));




CPU䞊のデヌタはh_  h ostで始たり、デヌタずGPUはd_  d eviceで始たる-CUDAの倉数の呜名基準に泚意を払う䟡倀がありたす。 checkCudaErrors-コヌスUdacity githubリポゞトリから取埗したマクロ。 次の圢匏がありたす。

非衚瀺のテキスト 
 #include <cuda.h> #define checkCudaErrors(val) check( (val), #val, __FILE__, __LINE__) template<typename T> void check(T err, const char* const func, const char* const file, const int line) { if (err != cudaSuccess) { std::cerr << "CUDA error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString(err) << " " << func << std::endl; exit(1); } }




cudaMalloc -GPUのmallocの類䌌䜓、 cudaMemcpy - memcpyの類䌌䜓には、コピヌのタむプを瀺す列挙型の远加パラメヌタヌがありたすcudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice。

次に、時間の枬定を忘れずに、グリッドずブロックの次元を蚭定し、カヌネルを呌び出す必芁がありたす。

非衚瀺のテキスト 
  dim3 blockSize; dim3 gridSize; int threadNum; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); threadNum = 1024; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols/threadNum+1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_simple<<<gridSize, blockSize>>>(d_imageRGBA, d_imageGray, numRows, numCols); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout << "CUDA time simple (ms): " << milliseconds << std::endl;




カヌネル呌び出し圢匏-kernel_name <<< gridSize、blockSize >>>に泚意しおください。 カヌネルコヌド自䜓もそれほど耇雑ではありたせん。

rgba_to_grayscale_simple 
 __global__ void rgba_to_grayscale_simple(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols) { int y = blockDim.y*blockIdx.y + threadIdx.y; int x = blockDim.x*blockIdx.x + threadIdx.x; if (x>=numCols || y>=numRows) return; const int offset = y*numCols+x; const uchar4 pixel = d_imageRGBA[offset]; d_imageGray[offset] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; }




ここで、前述の倉数threadIdx 、 blockIdx 、 blockDimを䜿甚しお、凊理されたピクセルのy座暙ずx座暙を蚈算し、倉換を行いたす。 x> = numCols || y> = numRowsのチェックに泚意しおください-画像のサむズは必ずしもブロックのサむズで完党に分割されるずは限らないため、䞀郚のブロックは画像​​の「範囲を超える」こずがありたす。したがっお、このチェックが必芁です。 たた、カヌネル関数は__global__指定子でマヌクする必芁がありたす 。

最埌のステップは、結果をGPUからCPUにコピヌしお、割り圓おられたメモリを解攟するこずです。

非衚瀺のテキスト 
  checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost)); cudaFree(d_imageGray); cudaFree(d_imageRGBA);




ずころで、CUDAを䜿甚するず、ホストコヌドにC ++コンパむラを䜿甚できたす。したがっお、自動的にメモリを解攟するラッパヌを簡単に䜜成できたす。

それで、枬定を開始したす入力画像のサむズは10.109×4.542です  

 OpenMP time (ms):45 CUDA time simple (ms): 43.1941


テストが実行されたマシンの構成

非衚瀺のテキスト
プロセッサヌIntel®CoreTMi7-3615QM CPU @ 2.30GHz。

GPUNVIDIA GeForce GT 650M、1024 MB、900 MHz。

RAMDD3、2x4GB、1600 MHz。

OSOS X 10.9.5。

コンパむラg ++GCC4.9.2 20141029。

CUDAコンパむラCudaコンパむルツヌル、リリヌス6.0、V6.0.1。

サポヌトされおいるOpenMPのバヌゞョンOpenMP 4.0。



どういうわけかあたり印象的ではありたせんでした:)しかし、問題は同じです-各ピクセルで行われる䜜業が少なすぎる-数千のスレッドを実行し、それぞれがほが瞬時に動䜜したす。 CPUの堎合、この問題は発生したせん-OpenMPは比范的少数のスレッド私の堎合は8を起動し、スレッド間で均等に䜜業を分割したす-このようにしお、プロセッサはほが100占有されたすが、GPUでは実際に、そのすべおの力を䜿甚しないでください。 解決策はかなり明癜です-カヌネル内のいく぀かのピクセルを凊理したす。 新しい、最適化されたカヌネルは次のようになりたす。

rgba_to_grayscale_optimized 
 #define WARP_SIZE 32 __global__ void rgba_to_grayscale_optimized(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols, int elemsPerThread) { int y = blockDim.y*blockIdx.y + threadIdx.y; int x = blockDim.x*blockIdx.x + threadIdx.x; const int loop_start = (x/WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x; for (int i=loop_start, j=0; j<elemsPerThread && i<numCols; i+=WARP_SIZE, ++j) { const int offset = y*numCols+i; const uchar4 pixel = d_imageRGBA[offset]; d_imageGray[offset] = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } }




すべおが以前のカヌネルほど簡単ではありたせん。 ご芧のずおり 、各スレッドはelemsPerThreadピクセルを凊理し、連続ではなく、それらの間のWARP_SIZEの距離で凊理したす。 WARP_SIZEずは䜕か、32である理由、ピクセルを自由な方法で凊理する必芁がある理由に぀いおは、次の郚分でさらに詳しく説明したす。メモリでより効率的な䜜業を行うこずができたす。 各スレッドはelemsPerThreadピクセルをWARP_SIZEの距離で凊理するようになったため、ブロック内の䜍眮に基づくこのスレッドの最初のピクセルのx座暙は、以前よりもやや耇雑な数匏を䜿甚しお蚈算されるようになりたした。

このカヌネルは次のように起動したす。

非衚瀺のテキスト 
  threadNum=128; const int elemsPerThread = 16; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols / (threadNum*elemsPerThread) + 1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_optimized<<<gridSize, blockSize>>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout << "CUDA time optimized (ms): " << milliseconds << std::endl;




x座暙によるブロック数は、 numCols / threadNum + 1ではなくnumCols /threadNum * elemsPerThread+ 1ずしお蚈算されるようになりたした。 そうでなければ、すべおが同じたたでした。

以䞋を開始したす。 

 OpenMP time (ms):44 CUDA time simple (ms): 53.1625 CUDA time optimized (ms): 15.9273


速床が2.76倍になりたしたここでも、メモリ操䜜の時間を考慮しおいたせん-このような単玔な問題の堎合、これはかなり良いです。 はい、このタスクは単玔すぎたす-CPUも非垞にうたく察凊したす。 2番目のテストからわかるように、GPUでの単玔な実装でも、CPUでの実装速床が䜎䞋する可胜性がありたす。

今日はこれで終わりです。次のパヌトでは、GPUハヌドりェアず基本的な䞊列通信パタヌンに぀いお説明したす。

すべおの゜ヌスコヌドはbitbucketで入手できたす。


More articles:


All Articles