понедельник, 23 марта 2009 г.

Компилятор Фортрана и Си для nVidia GPU

Если кто-то пробовал писать на CUDA, то он знает, что без познания архитектуры вычислителя, вникания в SIMD-ы и варпы и даже, прости господи, "сеток к которым применяется ядро" не обойтись. Даже если вы собрались перенести на GPU код умножения матриц вида:


do i = 1,n
do j = 1,m
do k = 1,p
a(i,j) = a(i,j) + b(i,k)*c(k,j)
enddo
enddo
enddo


На CUDA реализация такого алгоритма будет выглядеть как-то так:

__global__ void 
matmulKernel( float* C, float* A, float* B, int N2, int N3 ){
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int aFirst = 16 * by * N2;
int bFirst = 16 * bx;
float Csub = 0;

for( int j = 0; j < N2; j += 16 ) {
__shared__ float Atile[16][16], Btile[16][16];
Atile[ty][tx] = A[aFirst + j + N2 * ty + tx];
Btile[ty][tx] = B[bFirst + j*N3 + b + N3 * ty + tx];

__syncthreads();

for( int k = 0; k < 16; ++k )
Csub += Atile[ty][k] * Btile[k][tx];

__syncthreads();
}

int c = N3 * 16 * by + 16 * bx;
C[c + N3 * ty + tx] = Csub;
}

void
matmul( float* A, float* B, float* C,
size_t N1, size_t N2, size_t N3 ){
void *devA, *devB, *devC;
cudaSetDevice(0);

cudaMalloc( &devA, N1*N2*sizeof(float) );
cudaMalloc( &devB, N2*N3*sizeof(float) );
cudaMalloc( &devC, N1*N3*sizeof(float) );

cudaMemcpy( devA, A, N1*N2*sizeof(float), cudaMemcpyHostToDevice );
cudaMemcpy( devB, B, N2*N3*sizeof(float), cudaMemcpyHostToDevice );

dim3 threads( 16, 16 );
dim3 grid( N1 / threads.x, N3 / threads.y);

matmulKernel<<< grid, threads >>>( devC, devA, devB, N2, N3 );

cudaMemcpy( C, devC, N1*N3*sizeof(float), cudaMemcpyDeviceToHost );
cudaFree( devA );
cudaFree( devB );
cudaFree( devC );
}

Вот если бы был ключик у cl.exe или gcc... специальный такой gpu-ключик, так бы и назывался "-gpu"... :)
А пока ИИ по распознаванию хорошо параллелящихся кусков в коде не вышел за пределы лабораторий, всем сочувствующим GPU-оптимизации кода предлагается посмотреть на анонс PGI-акселератора.

PGI является компилятором Си и Фортрана, с возможностью раскидать на карточке обозначенные циклы. Выглядит это, как анонсируется, следующим образом:

 !$acc region
!$acc do parallel
do j=1,m
do k=1,p
!$acc do parallel, vector(2)
do i=1,n
a(i,j) = a(i,j) + b(i,k)*c(k,j)
enddo
enddo
enddo
!$acc end region


Структуры и данные будут автоматически проанализированны и при компиляции переведены в SIMD инструкции. Конечно, не без нашей подсказки о том, какой код будем параллелить, но все же. Для Си, кстати, подсказка будет выглядеть так:

#pragma acc region
{
/* user loops and code for compiler GPU acceleration go here */
}

Если так пойдет и дальше, то можно будет помечтать о первых портах под GPU... :)

А на данный момент компилятор поддерживает кучу ограничений и совсем мало функционала.
В частности, пока только PGF95 Fortran и PGCC ANSI C99 и только под Linux 64.
Со временем планируется поддержка и C++ и Windows, Mac OS, но конкретных сроков пока не обозначают. Наверное, все будет зависеть от востребованности на рынке.

Из хороших новостей - если на компьютере, где выполняется скомпилированный код не будет обнаружен GPU-акселератор, то код будет выполнен на отдельном процессоре.
От CUDA унаследована прозрачность для разработчика кол-ва GPU и модели видеокарты (поддерживаются только nVidia).
Вычисления с плавающей точкой на видеокарте выполняются с точностью single precision. В будущем также обещают double, т.к. у nVidia уже есть такие GPU.

По всей видимости такой компилятор можно по праву счтать первым в новом поколении высокоуровневых средств для программирования на GPU.

Комментариев нет: