uzluga.ru
добавить свой файл
1


Архитектура и программирование массивно-параллельных вычислительных систем

Лекторы:

Боресков А.В. (ВМиК МГУ)

Харламов А. (NVIDIA)

План

  • Существующие архитектуры

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

  • CUDA

  • Несколько слов о курсе

  • Дополнительные слайды



План

  • Существующие архитектуры

    • Intel CPU
    • SMP
    • CELL
    • BlueGene
    • NVIDIA Tesla 10
  • Классификация

  • CUDA

  • Несколько слов о курсе

  • Дополнительные слайды



Существующие многоядерные системы

Посмотрим на частоты CPU:
    • 2004 г. - Pentium 4, 3.46 GHz
    • 2005 г. - Pentium 4, 3.8 GHz
    • 2006 г. - Core Duo T2700, 2333 MHz
    • 2007 г. - Core 2 Duo E6700, 2.66 GHz
    • 2007 г. - Core 2 Duo E6800, 3 GHz
    • 2008 г. - Core 2 Duo E8600, 3.33 Ghz
    • 2009 г. - Core i7 950, 3.06 GHz


Существующие многоядерные системы

  • Роста частоты практически нет

    • Энерговыделение ~ четвертой степени частоты
    • Ограничения техпроцесса
    • Одноядерные системы зашли в тупик


Существующие многоядерные системы

  • Повышение быстродействия следует ждать от параллельности.

  • CPU используют параллельную обработку для повышения производительности

    • Конвейер
    • Multithreading
    • SSE


Intel Core 2 Duo

  • 32 Кб L1 кэш для каждого ядра

  • 2/4 Мб общий L2 кэш

  • Единый образ памяти для каждого ядра - необходимость синхронизации кэшей



Intel Core 2 Quad



Intel Core i7



Symmetric Multiprocessor Architecture (SMP)



Symmetric Multiprocessor Architecture (SMP)

Каждый процессор
  • имеет свои L1 и L2 кэши

  • подсоединен к общей шине

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



Cell



Cell

  • Dual-threaded 64-bit PowerPC

  • 8 Synergistic Processing Elements (SPE)

  • 256 Kb on-chip на каждый SPE



BlueGene/L



BlueGene/L

  • 65536 dual-core nodes

  • node

    • 770 Mhz PowerPC
    • Double Hammer FPU (4 Flop/cycle)
    • 4 Mb on-chip L3 кэш
    • 512 Mb off-chip RAM
    • 6 двухсторонних портов для 3D-тора
    • 3 двухсторонних порта для collective network
    • 4 двухсторонних порта для barrier/interrupt


BlueGene/L



Архитектура Tesla 10



Архитектура Tesla Мультипроцессор Tesla 10



Технические детали

  • RTM CUDA Programming Guide

  • Run CUDAHelloWorld

    • Печатает аппаратно зависимые параметры
    • Размер shared памяти
    • Кол-во SM
    • Размер warp’а
    • Кол-во регистров на SM
    • т.д.


План

  • Существующие архитектуры

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

    • Примеры для CPU
  • CUDA

  • Несколько слов о курсе

  • Дополнительные слайды



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



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

  • CPU – SISD

    • Multithreading: позволяет запускать множество потоков – параллелизм на уровне задач (MIMD) или данных (SIMD)
    • SSE: набор 128 битных регистров ЦПУ
    • можно запаковать 4 32битных скаляра и проводить над ними операции одновременно (SIMD)
  • GPU – SIMD*



MultiThreading “Hello World”

#include

#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*)&t1 );

// напечатать из основного потока

mtPrintf( (void*)&t0);

// подождать 100 мс

// создание потока windows требует времени

// если основной поток закончит выполнение

// то и все дочерние потоки будут прерваны

Sleep( 100 );



SSE “Hello World”

#include

#include

struct vec4

{

union

{

float v[4];

__m128 v4;

};

};

int main()

{

vec4 c, a = {5.0f, 2.0f, 1.0f, 3.0f}, b = {5.0f, 3.0f, 9.0f, 7.0f};

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;

}

CPU

  • Параллельное программирование CPU требует специалльных API

    • MPI, OpenMP
  • Программирование ресурсов CPU ограничено

    • Multithreading
    • SSE
    • Ограничивает пропускная способность памяти


SIMD

  • На входе поток однородных элементов, каждый из которых может быть обработан независимо

  • На выходе – однородный поток

  • Обработкой занимается ядро (kernel)



SIMD

  • Каждый элемент может быть обработан независимо от других

    • Их можно обрабатывать параллельно
  • Можно соединять между собой отдельные ядра для получения более сложного конвеера обработки



План

  • Существующие архитектуры

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

  • CUDA

    • Программная модель
    • Связь программной модели с HW
    • SIMT
    • Язык CUDA C
    • Примеры для CUDA
  • Несколько слов о курсе

  • Дополнительные слайды



Compute Unified Device Architecture

  • CUDA – программно-аппаратный стек для программирования GPU



Программная модель CUDA

  • Код состоит из последовательных и параллельных частей

  • Последовательные части кода выполняются на CPU (host)

  • Массивно-параллельные части кода выполняются на GPU (device)

    • Является сопроцессором к CPU (host)
    • Имеет собственную память (DRAM)
    • Выполняет одновременно очень много нитей


Программная модель CUDA

  • Параллельная часть кода выполняется как большое количество нитей (threads)

  • Нити группируются в блоки (blocks) фиксированного размера

  • Блоки объединяются в сеть блоков (grid)

  • Ядро выполняется на сетке из блоков

  • Каждая нить и блок имеют свой уникальный идентификатор



Программная модель CUDA

  • Десятки тысяч нитей



Программная модель CUDA

  • Нити в CUDA объединяются в блоки:

    • 1D топология блока
    • 2D топология блока
    • 3D топология блока
  • Общее кол-во нитей в блоке ограничено

  • В текущем HW это 512* нитей



Программная модель CUDA

  • Блоки могут использовать shared память

    • Нити могут обмениваться общими данными
  • Внутри блока потоки могут синхронизоваться



Программная модель CUDA

  • Блоки потоков объединяются в сеть (grid) блоков потоков

    • 1D топология сетки блоков потоков
    • 2D топология сетки блоков потоков
  • Блоки в сети выполняются независимо друг от друга



Связь программной модели с HW

  • Блоки могут использовать shared память

    • Т.к. блок целиком выполняется на одном SM
    • Объем shared памяти ограничен и зависит от HW
  • Внутри блока нити могут синхронизоваться

    • Т.к. блок целиком выполняется на одном SM
  • Масштабирование архитектуры и производительности



Масштабирование архитектуры Tesla



Масштабирование мультипроцессора Tesla 10



Масштабирование производительности



Связь программной модели с HW

  • Очень высокая степень параллелизма

    • Десятки тысяч потоков на чипе
    • Потоки на GPU очень «легкие»
    • HW планировщик задач
  • Основная часть чипа занята логикой, а не кэшем

  • Для полноценной загрузки GPU нужны тысячи потоков

    • Для покрытия латентностей операций чтения / записи
    • Для покрытия латентностей sfu инструкций


Блоки и warp’ы?

  • Блоки – абстракция программной модели

  • Warp – реальная единица исполнения HW



Single Instruction Multiple Threads (SIMT)

  • Параллельно на каждом SM выполняется большое число отдельных нитей (threads)

  • Нити в пределах одного warp’а выполняются физически параллельно (SIMD)

  • Разные warp’ы могут исполнять разные команды

  • Большое число warp’ов покрывает латентность



Язык CUDA C

  • CUDA С – это расширение языка C/C++

    • спецификаторы для функций и переменных
    • новые встроенные типы
    • встроенные переменные (внутри ядра)
    • директива для запуска ядра из C кода
  • Как скомпилировать CUDA код

    • nvcc компилятор
    • .cu расширение файла


Язык CUDA С Спецификаторы



Язык CUDA С Спецификаторы

  • Спецификатор __global__ соответствует ядру

    • Может возвращать только void
  • Спецификаторы __host__ и __device__ могут использоваться одновременно

    • Компилятор сам создаст версии для CPU и GPU
  • Спецификаторы __global__ и __host__ не могут быть использованы одновременно



Язык CUDA С Ограничения

  • Ограничения на функции, выполняемые на GPU:

    • Нельзя брать адрес (за исключением __global__)
    • Не поддерживается рекурсия
    • Не поддерживаются static-переменные внутри функции
    • Не поддерживается переменное число входных аргументов


Язык CUDA С Ограничения

  • Ограничения на спецификаторы переменных:

    • Нельзя применять к полям структуры или union
    • Не могут быть extern
    • Запись в __constant__ может выполнять только CPU через специальные функции
    • __shared__ - переменные не могут инициализироваться при объявлении


Язык CUDA С Типы данных

  • Новые типы данных:

    • 1/2/3/4-мерные вектора из базовых типов
    • (u)char, (u)int, (u)short, (u)long, longlong
    • float, double
    • dim3 – uint3 с нормальным конструкторов, позволяющим задавать не все компоненты
    • Не заданные инициализируются единицей


Язык CUDA С Встроенные переменные

Сравним CPU vs CUDA С код:

Язык CUDA С Встроенные переменные

  • В любом CUDA kernel’e доступны:

    • dim3 gridDim;
    • uint3 blockIdx;
    • dim3 blockDim;
    • uint3 threadIdx;
    • int warpSize;


Язык CUDA С Директивы запуска ядра

  • Как запустить ядро с общим кол-во тредов равным nx?



Язык CUDA С Директивы запуска ядра

  • Общий вид команды для запуска ядра

incKernel<<<bl, th, ns, st>>> ( data );
  • bl – число блоков в сетке

  • th – число нитей в сетке

  • ns – количество дополнительной shared-памяти, выделяемое блоку

  • st – поток, в котором нужно запустить ядро



Как скомпилировать CUDA код

  • NVCC – компилятор для CUDA

    • Основными опциями команды nvcc являются:
    • --use_fast_math - заменить все вызовы стандартных математических функций на их быстрые (но менее точные) аналоги
    • -o  - задать имя выходного файла
  • CUDA файлы обычно носят расширение .cu



CUDA “Hello World”

#define N (1024*1024)

__global__ void kernel ( float * data )

{

int idx = blockIdx.x * blockDim.x + threadIdx.x;

float x = 2.0f * 3.1415926f * (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 * 3.1415926f * idx / N; // значение аргумента

data [idx] = sinf ( sqrtf ( x ) ); // найти значение и записать в массив

}



CUDA “Hello World”



План

  • Существующие архитектуры

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

  • CUDA

  • Несколько слов о курсе

    • Отчетность по курсу
    • Ресурсы нашего курса
  • Дополнительные слайды



Несколько слов о курсе

  • Математический спецкурс

  • 5 семинарский занятий

    • Раз в две недели
    • Цель занятий:
    • Начать быстро программировать на CUDA
    • Написать и сдать практические задания
    • На удаленной машине *nix
    • Тренеруйте shell-skill
  • 5 практических заданий



Отчетность по курсу

  • 5 практических заданий

    • Задания сдаются на семинаре
    • Либо по почте
    • с темой CUDA Assignment #
    • В течении недели с момента публикации
    • Если у вас не получается – дайте нам знать
    • Заранее
  • Альтернатива

    • Дайте нам знать
    • Заранее


Отчетность по курсу

  • Если тема email отличается от CUDA Assignment #

    • MINOR FAIL
  • Если ваш код не собирается или не запускается

    • MAJOR FAIL
  • Если обнаруживается дубликат

    • EPIC FAIL


Отчетность по курсу

  • Если вы не сдадите задания и не предупредите заранее

    • FAIL
  • Если вы выбрали альтернативу, но нас не предупредлили заранее

    • EPIC FAIL


Ресурсы нашего курса





План

  • Существующие архитектуры

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

  • CUDA

  • Несколько слов о курсе

  • Дополнительные слайды

    • Эволюция GPU
    • Архитектура Tesla 8
    • Архитектура Tesla 20


Эволюция GPU

  • Voodoo - растеризация треугольников, наложение текстуры и буфер глубины

  • Очень легко распараллеливается

  • На своих задачах легко обходил CPU



Эволюция GPU

  • Быстрый рост производительности

  • Добавление новых возможностей

    • Мультитекстурирование (RivaTNT2)
    • T&L
    • Вершинные программы (шейдеры)
    • Фрагментные программы (GeForceFX)
    • Текстуры с floating point-значениями


Эволюция GPU: Шейдеры

  • Работают с 4D float-векторами

  • Специальный ассемблер

  • Компилируется драйвером устройства

  • Отсутствие переходов и ветвления

    • Вводились как vendor-расширения


GPGPU

  • Использование GPU для решения не графических задач

  • Вся работа с GPU идет через графический API (OpenGL, D3D)

  • Программы используют сразу два языка – один традиционный (С++) и один шейдерный

  • Ограничения, присущие графическим API



Эволюция GPU: Шейдеры

  • Появление шейдерных языков высокого уровня (Cg, GLSL, HLSL)

  • Поддержка ветвлений и циклов (GeForce 6xxx)

  • Появление GPU, превосходящие CPU в 10 и более раз по Flop’ам



Архитектура Tesla 8



Архитектура Tesla: Мультипроцессор Tesla 8



Архитектура Tesla 20

  • Объединенный L2 кэш (768 Kb)

  • До 1 Tb памяти (64-битная адресация)

  • Общее адресное пространство памяти

  • ККО (DRAM, регистры, разделяемая память, кэш)

  • Одновременное исполнение ядер, копирования памяти (CPU->GPU, GPU->CPU)

  • Быстрая смена контекста (10x)

  • Одновременное исполнение ядер (до 16)



Архитектура Tesla 20 Потоковый мультипроцессор



Архитектура Tesla 20

  • 32 ядра на SM

  • Одновременное исполнение 2х варпов.

  • 48 Kb разделяемой памяти

  • 16 Kb кэш

    • или 16 Kb разделяемй + 48 Kb кэш
  • Дешевые атомарные операции



Intellisence для CUDA

  • Start  Run  Regedit

  • HKEY_LOCAL_MACHINE

    • Software
    • Microsoft
    • Visual Studio
    • 9.0 MSVS 2008 или
    • 8.0 MSVS 2005
    • Languages
    • Language Services
    • C/C++
    • NCB Default C/C++ Extensions
    • Добавить .cu;
  • Перезапустить VS