Архитектура и программирование массивно-параллельных вычислительных систем zЛекторы: yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) yХарламов А. (NVidia)Харламов А. (NVidia)

Существующие многоядерные системы Посмотрим на частоты CPU: y2004 г. - Pentium 4, 3.46 GHz y2005 г. - Pentium 4, 3.8 GHz y2006 г. - Core Duo T2700, 2333 MHz y2007 г. - Core 2 Duo E6700, 2.66 GHz y2007 г. - Core 2 Duo E6800, 3 GHz y2008 г. - Core 2 Duo E8600, 3.33 Ghz y2009 г. - Core i7 950, 3.06 GHz

Существующие многоядерные системы Легко видно, что роста частоты практически нет yЭнерговыделение ~ четвертой степени частоты yОграничения техпроцесса

Существующие многоядерные системы zТаким образом, повышение быстродействия следует ждать именно от параллельности. zУже давно CPU используют параллельную обработку для повышения производительности y Конвейер y Multithreading y SSE

Intel Core 2 Duo z32 Кб L1 кэш для каждого ядра z2/4 Мб общий L2 кэш zЕдиный образ памяти для каждого ядра - необходимость синхронизации кэшей

Symmetric Multiprocessor Architecture (SMP)

Каждый процессор zимеет свои L1 и L2 кэши zподсоединен к общей шине zотслеживает доступ других процессоров к памяти для обеспечения единого образа памяти (например, один процессор хочет изменить данные, кэшированные другим процессором)

Cell

zDual-threaded 64-bit PowerPC z8 Synergistic Processing Elements (SPE) z256 Kb on-chip на каждый SPE

BlueGene/L

z dual-core nodes znode y770 Mhz PowerPC yDouble Hammer FPU (4 Flop/cycle) y4 Mb on-chip L3 кэш y512 Mb off-chip RAM y6 двухсторонних портов для 3D-тора y3 двухсторонних порта для collective network y4 двухсторонних порта для barrier/interrupt

Архитектура G80 Массив из потоковых мультипроцессоров

Архитектура G80

Классификация

zCPU – SISD yMultithreading: позволяет запускать множество потоков – параллелизм на уровне задач (MIMD) или данных (SIMD) ySSE: набор 128 битных регистров ЦПУ xможно запаковать 4 32битных скаляра и проводить над ними операции одновременно (SIMD) zGPU – SIMD*

MultiThreading Hello World #include #include // для beginthread() void mtPrintf( void * pArg); int main() { int t0 = 0; int t1 = 1; _beginthread(mtPrintf, 0, (void*)&t0 ); mtPrintf( (void*)&t1); Sleep( 100 ); return 0; } void mtPrintf( void * pArg ) { int * pIntArg = (int *) pArg; printf( "The function was passed %d\n", (*pIntArg) ); }

MultiThreading Hello World // создание нового потока // необходимо указать: // entry point функцию, // размер стека, при 0 – OS выберет сама // (void *) – указатель на аргументы функции _beginthread(mtPrintf, 0, (void*)1 ); // напечатать из основного потока mtPrintf( (void*)0); // подождать 100 мс // создание потока windows требует времени // если основной поток закончит выполнение // то и все дочерние потоки будут прерваны Sleep( 100 );

SSE Hello World #include struct vec4 { union { float v[4]; __m128 v4; }; int main() { vec4 a = {5.0f, 2.0f, 1.0f, 3.0f}; vec4 b = {5.0f, 3.0f, 9.0f, 7.0f}; vec4 c; c.v4 = _mm_add_ps(a.v4, b.v4); printf("c = {%.3f, %.3f, %.3f, %.3f}\n", c.v[0], c.v[1], c.v[2], c.v[3]); return 0; }

CUDA (Compute Unified Device Architecture) zПрограммирование массивно-параллельных систем требует специалльных систем/языков. zПрограммирование ресурсов CPU ограничено yMultithreading ySSE yЧасто bottleneck – в пропускной способности памяти zCUDA - система (библиотеки и расширенный C) для программирования GPU

CUDA Hello World #defineN(1024*1024) __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; float x = 2.0f * f * (float) idx / (float) N; data [idx] = sinf ( sqrtf ( x ) ); } int main ( int argc, char * argv [] ) { float a [N]; float * dev = NULL; cudaMalloc ( (void**)&dev, N * sizeof ( float ) ); kernel >> ( dev ); cudaMemcpy ( a, dev, N * sizeof ( float ), cudaMemcpyDeviceToHost ); cudaFree ( dev ); for (int idx = 0; idx < N; idx++) printf("a[%d] = %.5f\n", idx, a[idx]); return 0; }

CUDA Hello World __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; // номер текущей нити float x = 2.0f * f * (float) idx / (float) N; // значение аргумента data [idx] = sinf ( sqrtf ( x ) ); // найти значение и } // записать его в массив zДля каждого элемента массива (всего N) запускается отдельная нить, вычисляющая требуемое значение. zКаждая нить обладает уникальным id

CUDA Hello World float a [N]; float * dev = NULL; // выделить память на GPU под N элементов cudaMalloc ( (void**)&dev, N * sizeof ( float ) ); // запустить N нитей блоками по 512 нитей // выполняемая на нити функция - kernel // массив данных - dev kernel >> ( dev ); // скопировать результаты из памяти GPU (DRAM) в // память CPU (N элементов) cudaMemcpy ( a, dev, N * sizeof ( float ), cudaMemcpyDeviceToHost ); // освободить память GPU cudaFree ( dev );

Ресуры нашего курса zCUDA.CS.MSU.SUCUDA.CS.MSU.SU yМесто для вопросов и дискуссий yМесто для материалов нашего курса yМесто для ваших статей! xЕсли вы нашли какой-то интересный подход! xИли исследовали производительность разных подходов и знаете, какой из них самый быстрый! xИли знаете способы сделать работу с CUDA проще! z Steps3d Steps3d z

Ресурсы нашего курса zК той лекции: yCUDA / MT / SSE hello world проекты yCUDA / MT / SSE Hello World проекты xЧуть более насыщенные чем маленькие hello world ы y SVN ?

Несколько слов о курсе z Математический спецкурс z 11 лекций z 5 семинарский занятий yРаз в две недели yЦель занятий: x Начать быстро программировать на CUDA x Написать и сдать практические задания z 5 практических заданий

Несколько слов о курсе zОтчетность по курсу y 5 практических заданий x Задания сдаются на семинаре x Либо по почте В течении недели со дня семинара, на котором задание выдано Если у вас не получается – дайте нам знать yАльтернатива x Дайте нам знать