Особенности разделяемой памяти.

Разделяемая память находится на каждом мультипроцессоре и используется совместно всеми потоками внутри блока. Разделяемая память, как правило, требует использования команды явной синхронизации __syncthreads() после операций копирования и изменения. Для повышения пропускной способности разделяемая память разбита на 16 банков, доступ к которым идет независимо. Банки состоят из 32-битовых слов, причем подряд идущие слова расположены в идущих подряд блоках. Если возникает ситуация, что несколько нитей обращаются к одному и тому же банку (не элементу, а именно банку), то их обращения обрабатываются последовательно (происходит конфликт банков), однако если все нити обращаются к одному и тому же элементу, то конфликта банков не происходит.

Переменные, объявленные спецификатором __shared__, размещаются в разделяемой памяти:

__shared__ float mem_shared[N_THREADS];

Во время исполнения блока на мультипроцессоре содержимое разделяемой памяти сохраняется. Однако, после того как на мультипроцессоре сменится блок, не гарантируется, что содержимое старого блока сохранится. Поэтому не рекомендуется синхронизировать задачи между блоками и оставлять в разделяемой памяти какие-либо данные, надеясь на их сохранность.

Приведем несколько простых примеров:

__global__ void kern(float* data)

{

//Создается массив float в разделяемой памяти:

__shared__ float a[256];

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

//Копируем из глобальной памяти в разделяемую

a[idx] = data[idx];

//Синхронизируем.

__syncthreads();

//Используем

data[idx] = a[idx]+ a[(idx + 1) % 256];

}

Конфликтов не возникает, так как float занимают 32 бита и последовательные float располагаются в последовательных банках.

__global__ void kern(short* data)

{

//Создается массив short в разделяемой памяти:

__shared__ short a[256];

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

//Копируем из глобальной памяти в разделяемую

a[idx] = data[idx];

// Синхронизируем.

__syncthreads();

//Используем

data[idx] = a[idx]+ a[(idx + 1) % 256];

}

Возникают конфликты 2-го порядка, так как short занимают 16 бит, и два последовательных short располагаются в одном банке. Аналогично если объявлять массив char, возникнет конфликт 4-го порядка, так как char использует 8 бит.

4. Отчет по практической работе

Отчет по практической работе должен соответствовать порядку выполнения работы и содержать конспект изученного материала, данные о вычислительных возможностях видеокарты учебного компьютера:

• объем глобальной памяти;

• объем разделяемой памяти;

• максимальную размерность текстуры (x,y,z);

• КЭШ текстурной памяти на мультипроцессор.

3. Практическая работа №3. Программный пакет «Sarcis»

   1. Цель работы

Изучение интерфейса программного пакета «Sarcis».

Закрепление теоретического материала учебной дисциплины «Цифровая обработка изображений».

2. Порядок выполнения работы

1. Открыть любой графический файл (.bmp) в пакете «Sarcis»..

2. Построить гистограмму изображения.

3. Улучшить изображения с помощью выравнивания гистограммы.

4. Произвести пороговую сегментацию.

5. Показать функцию яркости изображения.

6. Обработать графический файл любым имеющимся в библиотеке методом.

7. Изучить механизм подключения внешних библиотек.

8. Изучить пример создания внешней библиотеки.

9. Создать внешнюю подключаемую библиотеку, проверить корректность ее работы.