Оптимизация работы с глобальной памятью.

Глобальная память обладает высокой латентностью, поэтому для нее важна оптимизация доступа. Обращение к глобальной памяти происходит через чтение и запись 32, 64 или 128-битовых слов, поэтому важно, чтобы адрес, по которому происходит доступ, был выровнен по размеру слова. Пусть имеется массив, выделенный в глобальной памяти: struct vec3 { float x,y,z;} Хотя каждый элемент этого массива размером 12 байт, полностью помещается в 16 байтах. И даже если первый элемент массива выровнен по 16 байтам, то адрес следующего уже не будет выровнен и его чтение потребует двух обращений. Это можно исправить, если обеспечить выравнивание элементов массива: struct _ _align_ _(16) vec3 {float x,y,z;). В результате этой записи все элементы массива будут находится на адресах, кратных 16, что обеспечит чтение одного элемента за раз. Важным для оптимизации работы является объединение нескольких запросов к глобальной памяти в один. При использовании этой возможности получается 16х ускорение. Все обращения мультипроцессора к памяти происходят независимо от каждой половины варпа, поэтому максимальным является объединение, когда все запросы одного полуварпа удается объединить в один. Для объединенного запроса необходимо выполнение ряда условий:

1. все нити обращаются к 32-битовым словам, давая в результате один 64-байтовый блок, или же все нити обращаются к 64-битовым словам, давая один 128-байтовый блок.
2. полученный блок выровнен по своему размеру, т.е. адрес 64-байтового блока кратен 64.
3. все 16 слов, к которым обращаются нити, лежат в пределах данного блока.
4. нити обращаются к словам последовательно. К-тая нить должна обращаться к к-тому слову. Допускается, что нить пропустит свои обращения.

Если нити полуварпа не удовлетворяют какому-либо из условий, то каждое обращение к памяти производится как отдельная транзакция. Гораздо эффективнее с точки зрения использования памяти употребление не массивов структур, а структуры массивов.

struct A_align_(16) {float a,b; uint c; }; A array [1024]; A.a = array[threadIdx.x];

float a[1024], b[1024]; uint c[1024]; float fa = a[threadeIdx.x]; float fb = b[threadIdx.x]; float fc = c[threadedx.x];

В результате подобного разбиения каждый из 3 запросов приведет к объединению и понадобится всего по 3 транзакции на полуварп вместо 16.

Задача об N телах.

Дано несколько тел (N) со своими положениями и скоростями. Необходимо просчитать их движение под действием сил взаимного притяжения. Известна формула, описывающая полную силу, действующую на i–тое тело.

Fi = Σc/(abs(pj-pi)^3)*(pj-pi), j = 0…N, pj – координаты тела.

Для реализации понадобятся 4 массива: положение и скорости в текущий момент времени, положение и скорости устройства.