Master CUDA: Pre inžinierov strojového učenia

CUDA pre strojové učenie: praktické aplikácie

Teraz, keď sme prebrali základy, poďme preskúmať, ako možno CUDA aplikovať na bežné úlohy strojového učenia.

1. Násobenie matice

Násobenie matice je základnou operáciou v mnohých algoritmoch strojového učenia, najmä v neurónových sieťach. CUDA môže túto operáciu výrazne urýchliť. Tu je jednoduchá implementácia:

__global__ void matrixMulKernel(float *A, float *B, float *C, int N)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;
    
    if (row < N && col < N) {
        for (int i = 0; i < N; i++) {
            sum += A(row * N + i) * B(i * N + col);
        }
        C(row * N + col) = sum;
    }
}
// Host function to set up and launch the kernel
void matrixMul(float *A, float *B, float *C, int N)
{
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                   (N + threadsPerBlock.y - 1) / threadsPerBlock.y);
    
    matrixMulKernelnumBlocks, threadsPerBlock(A, B, C, N);
}

Táto implementácia rozdeľuje výstupnú maticu na bloky, pričom každé vlákno počíta jeden prvok výsledku. Aj keď je táto základná verzia už rýchlejšia ako implementácia CPU pre veľké matice, je tu priestor na optimalizáciu pomocou zdieľanej pamäte a iných techník.

2. Konvolučné operácie

Konvolučné neurónové siete (CNN) silne spoliehať na konvolučné operácie. CUDA môže tieto výpočty výrazne urýchliť. Tu je zjednodušené 2D konvolučné jadro:

__global__ void convolution2DKernel(float *input, float *kernel, float *output, 
                                    int inputWidth, int inputHeight, 
                                    int kernelWidth, int kernelHeight)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x < inputWidth && y < inputHeight) {
        float sum = 0.0f;
        for (int ky = 0; ky < kernelHeight; ky++) {
            for (int kx = 0; kx < kernelWidth; kx++) {
                int inputX = x + kx - kernelWidth / 2;
                int inputY = y + ky - kernelHeight / 2;
                if (inputX >= 0 && inputX < inputWidth && inputY >= 0 && inputY < inputHeight) {
                    sum += input(inputY * inputWidth + inputX) * 
                           kernel(ky * kernelWidth + kx);
                }
            }
        }
        output(y * inputWidth + x) = sum;
    }
}

Toto jadro vykonáva 2D konvolúciu, pričom každé vlákno počíta s jedným výstupným pixelom. V praxi by sofistikovanejšie implementácie využívali zdieľanú pamäť na zníženie prístupu do globálnej pamäte a optimalizáciu pre rôzne veľkosti jadra.

3. Stochastický gradientový zostup (SGD)

SGD je základný optimalizačný algoritmus strojového učenia. CUDA dokáže paralelizovať výpočet gradientov naprieč viacerými dátovými bodmi. Tu je zjednodušený príklad lineárnej regresie:

__global__ void sgdKernel(float *X, float *y, float *weights, float learningRate, int n, int d)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float prediction = 0.0f;
        for (int j = 0; j < d; j++) {
            prediction += X(i * d + j) * weights(j);
        }
        float error = prediction - y(i);
        for (int j = 0; j < d; j++) {
            atomicAdd(&weights(j), -learningRate * error * X(i * d + j));
        }
    }
}
void sgd(float *X, float *y, float *weights, float learningRate, int n, int d, int iterations)
{
    int threadsPerBlock = 256;
    int numBlocks = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    for (int iter = 0; iter < iterations; iter++) {
        sgdKernel<<<numBlocks, threadsPerBlock>>>(X, y, weights, learningRate, n, d);
    }
}

Táto implementácia aktualizuje váhy paralelne pre každý údajový bod. The atomicAdd funkcia sa používa na bezpečné spracovanie súbežných aktualizácií závaží.

Optimalizácia CUDA pre strojové učenie

Zatiaľ čo vyššie uvedené príklady demonštrujú základy používania CUDA pre úlohy strojového učenia, existuje niekoľko optimalizačných techník, ktoré môžu ďalej zvýšiť výkon:

1. Coalesced Memory Access

GPU dosahujú špičkový výkon, keď vlákna v deformácii pristupujú k súvislým pamäťovým miestam. Zabezpečte, aby vaše dátové štruktúry a prístupové vzory podporovali kombinovaný prístup k pamäti.

2. Využitie zdieľanej pamäte

Zdieľaná pamäť je oveľa rýchlejšia ako globálna pamäť. Použite ho na vyrovnávanie často používaných údajov v rámci bloku vlákien.

Pochopenie hierarchie pamäte pomocou CUDA

Tento diagram znázorňuje architektúru viacprocesorového systému so zdieľanou pamäťou. Každý procesor má vlastnú vyrovnávaciu pamäť, ktorá umožňuje rýchly prístup k často používaným údajom. Procesory komunikujú prostredníctvom zdieľanej zbernice, ktorá ich spája s väčším priestorom zdieľanej pamäte.

Napríklad pri maticovom násobení:

__global__ void matrixMulSharedKernel(float *A, float *B, float *C, int N)
{
    __shared__ float sharedA(TILE_SIZE)(TILE_SIZE);
    __shared__ float sharedB(TILE_SIZE)(TILE_SIZE);
    
    int bx = blockIdx.x; int by = blockIdx.y;
    int tx = threadIdx.x; int ty = threadIdx.y;
    
    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;
    
    float sum = 0.0f;
    
    for (int tile = 0; tile < (N + TILE_SIZE - 1) / TILE_SIZE; tile++) {
        if (row < N && tile * TILE_SIZE + tx < N)
            sharedA(ty)(tx) = A(row * N + tile * TILE_SIZE + tx);
        else
            sharedA(ty)(tx) = 0.0f;
        
        if (col < N && tile * TILE_SIZE + ty < N)
            sharedB(ty)(tx) = B((tile * TILE_SIZE + ty) * N + col);
        else
            sharedB(ty)(tx) = 0.0f;
        
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++)
            sum += sharedA(ty)(k) * sharedB(k)(tx);
        
        __syncthreads();
    }
    
    if (row < N && col < N)
        C(row * N + col) = sum;
}

Táto optimalizovaná verzia využíva zdieľanú pamäť na zníženie prístupu do globálnej pamäte, čím sa výrazne zvyšuje výkon pre veľké matice.

3. Asynchrónne operácie

CUDA podporuje asynchrónne operácie, čo vám umožňuje prekrývať výpočty s prenosom dát. To je užitočné najmä v kanáloch strojového učenia, kde môžete pripraviť ďalšiu dávku údajov, zatiaľ čo sa aktuálna dávka spracováva.

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// Asynchronous memory transfers and kernel launches
cudaMemcpyAsync(d_data1, h_data1, size, cudaMemcpyHostToDevice, stream1);
myKernel<<<grid, block, 0, stream1>>>(d_data1, ...);
cudaMemcpyAsync(d_data2, h_data2, size, cudaMemcpyHostToDevice, stream2);
myKernel<<<grid, block, 0, stream2>>>(d_data2, ...);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

4. Tenzorové jadrá

Pre pracovné zaťaženie strojového učenia, Tensor Cores NVIDIA (dostupné v novších architektúrach GPU) môže poskytnúť výrazné zrýchlenie operácií násobenia matíc a konvolúcie. Knižnice ako cuDNN a cuBLAS automaticky využívajú jadrá Tensor, keď sú k dispozícii.

Výzvy a úvahy

Zatiaľ čo CUDA ponúka obrovské výhody pre strojové učenie, je dôležité si uvedomiť potenciálne výzvy:

1. Správa pamäte: Pamäť GPU je v porovnaní so systémovou pamäťou obmedzená. Efektívna správa pamäte je kľúčová, najmä pri práci s veľkými súbormi údajov alebo modelmi.
2. Réžia prenosu dát: Prenos údajov medzi CPU a GPU môže byť prekážkou. Minimalizujte prenosy a ak je to možné, používajte asynchrónne operácie.
3. Presnosť: GPU tradične vynikajú vo výpočtoch s jednou presnosťou (FP32). Podpora pre dvojitú presnosť (FP64) sa síce zlepšila, no často je pomalšia. Mnoho úloh strojového učenia môže fungovať dobre s nižšou presnosťou (napr. FP16), ktorú moderné GPU zvládajú veľmi efektívne.
4. Zložitosť kódu: Zápis efektívneho kódu CUDA môže byť zložitejší ako kód CPU. Využitie knižníc ako cuDNNcuBLAS a rámce ako TensorFlow alebo PyTorch môžu pomôcť abstrahovať časť tejto zložitosti.