Как я могу обеспечить согласованность глобальной памяти CUDA без объявления указателя как изменчивого?
Я сначала сделаю некоторую контекстуализацию. Я пытаюсь реализовать неблокирующий метод кражи работы с использованием deques в CUDA. Deques (aDeques) находятся в массиве с сегментированными блоками в глобальной памяти, а функция устройства popWork() имеет целью перенести работу на подачу потоков. В дополнение к глобальным запросам каждый блок имеет стек в общей памяти (aLocalStack), где он может работать локально. Поп происходит в 3 уровня. Первая попытка находится в разделяемом стеке, вторая попытка находится в деке, принадлежащем блоку, а третья попытка - украсть другие запросы. Каждый deque имеет глобальные нижние и поп-указатели, которые лежат в глобальных массивах памяти (aiDequesBottoms и auiDequesAges). Моя проблема заключается в том, что когда блок меняет глобальный указатель deque, эти изменения не видны другим блокам, когда я тестирую код в GTS450. Кажется, что кэш не обновляется. Я также проверил на карте GT520, где проблема не возникает. У меня были похожие проблемы с массивом aiDequeFlags. Эти проблемы решаются путем объявления его изменчивым. К сожалению, я не могу сделать то же самое с массивами указателей deque, так как позже мне нужно использовать атомарные функции для них. Я сожалею, что не поместил проблему в более простой пример, но я не смог воспроизвести это поведение. Этот первый фрагмент имеет объяснение интерфейса popWork().
template <int iDequeSize> //Size of each segment in aDeques
bool __inline__ __device__ popWork(
volatile int *aiDequeFlags , //Flags that indicates if a deque is active (has work)
int *aiDequesBottoms , //Deque bottom pointers
unsigned int *auiDequesAges , //Deque top pointers (29 higher bits) +
//Tag bits(3 lower bits).
const Int2Array *aDeques , //Deques (Int2Array is an interface for 2 int arrays)
int &uiStackBot , //Shared memory stack pointer
int2 *aLocalStack , //Shared memory local stack
const int &iTid , //threadIdx.x
const int &iBid , //blockIdx.x
//All other parameters are output
unsigned int &uiPopDequeIdx , //Choosen deque for pop
int2 *popStartIdxAndSize , //Arrays of pop start index and sizes
bool *bPopFlag , //Array of flags for pop in each level
unsigned int &uiActiveDequesIdx , //Flag to indicate pop failed (no more work)
int2 &work //Actual acquired thread work)
Этот второй фрагмент имеет всю функцию. Ядро, которое использует функцию, было запущено с 8 блоками, 64 потоками, и в начале только у deque 0 была 1 работа, в то время как все другие запросы пусты. Существует несколько отладочных вызовов printf для создания журнала, который будет показан в следующем фрагменте.
template <int iDequeSize>
bool __inline__ __device__ popWork(volatile int *aiDequeFlags , int *aiDequesBottoms , unsigned int *auiDequesAges ,
const Int2Array *aDeques , int &uiStackBot , int2 *aLocalStack , const int &iTid , const int &iBid ,
unsigned int &uiPopDequeIdx , int2 *popStartIdxAndSize , bool *bPopFlag , unsigned int &uiActiveDequesIdx , int2 &work)
{
//Pop from local stack
if(iTid == 0)
{
unsigned int uiAge = 0;
bPopFlag[0] = popBottom(uiStackBot , uiAge , popStartIdxAndSize[iBid]);
bPopFlag[3] = bPopFlag[0];
}
__syncthreads();
if(bPopFlag[0])
{
if(iTid < popStartIdxAndSize[iBid].y)
{
work = aLocalStack[popStartIdxAndSize[iBid].x + iTid];
}
}
else
{
if(iTid == 0)
{ //Try to pop from block deque
bPopFlag[1] = popBottom(aiDequesBottoms[iBid] , auiDequesAges[iBid] , popStartIdxAndSize[iBid]);
if(bPopFlag[1])
{
uiPopDequeIdx = iBid;
//Debug
if(iBid == 0)
{
printf("Block %d pop global deque. Bottom=%d\n" , iBid , aiDequesBottoms[iBid]);
}
//
}
else
{
aiDequeFlags[iBid] = 0;
popStartIdxAndSize[iBid].x = INFTY;
uiPopDequeIdx = INFTY;
}
bPopFlag[3] = bPopFlag[1];
bPopFlag[2] = false;
}
__syncthreads();
if(!bPopFlag[1])
{
//Verify if lazy steal can be done.
if(iTid < NDEQUES)
{
if(popStartIdxAndSize[iTid].x != INFTY && iTid != iBid)
{
atomicMin(&uiPopDequeIdx , iTid);
bPopFlag[2] = true;
bPopFlag[3] = true;
}
}
__syncthreads();
if(iTid == uiPopDequeIdx)
{
popStartIdxAndSize[iBid] = popStartIdxAndSize[iTid];
popStartIdxAndSize[iTid].x = INFTY;
}
while(!bPopFlag[3])
{ //No more work, try to steal some!
__syncthreads();
if(iTid == 0)
{
uiActiveDequesIdx = 0;
}
__syncthreads();
if(iTid < NDEQUES)
{
if(aiDequeFlags[iTid] == 1)
{
uiActiveDequesIdx = 1;
//Debug
printf("Block %d steal attempt on block %d. Victim bottom=%d\n" , blockIdx.x , threadIdx.x , aiDequesBottoms[iTid]);
//
if(popTop(aiDequesBottoms , auiDequesAges , iTid , popStartIdxAndSize[iTid]))
{
aiDequeFlags[iBid] = 1;
atomicMin(&uiPopDequeIdx , iTid);
bPopFlag[3] = true;
//Debug
//printf("%d ss %d %d %d\n" , iBid , iTid , popStartIdxAndSize[iTid].x , popStartIdxAndSize[iTid].y);
//
}
}
}
__syncthreads();
if(uiActiveDequesIdx == 0)
{ //No more work to steal. End.
break;
}
if(iTid == uiPopDequeIdx)
{
popStartIdxAndSize[iBid] = popStartIdxAndSize[iTid];
popStartIdxAndSize[iTid].x = INFTY;
}
__syncthreads();
}
}
__syncthreads();
if(bPopFlag[3] && iTid < popStartIdxAndSize[iBid].y) //assuming number of threads >= WORK_SIZE
{
aDeques->getElement(work , uiPopDequeIdx*iDequeSize + popStartIdxAndSize[iBid].x + iTid);
}
}
return bPopFlag[3];
}
Этот последний фрагмент является сгенерированным журналом. Строки push ("Блок X push. Bottom=Y") были созданы функцией push, которая здесь не показана. Помните, что в начале только у блока 0 есть 1 работа.
Block 0 pop global deque. Bottom=0
Block 4 steal attempt on block 0. Victim bottom=0
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 7 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 4 steal attempt on block 0. Victim bottom=0
Block 7 steal attempt on block 0. Victim bottom=1
Block 0 push. Bottom=448
Block 2 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 7 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 4. Victim bottom=0
Block 1 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 4. Victim bottom=0
Block 5 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 4. Victim bottom=0
Block 4 push. Bottom=384
Как видно, только блок 4 может видеть изменения в блоке 0 нижнего указателя deque. Я пытался добавить некоторые вызовы __threadfence() после любого изменения указателей, но безуспешно. Спасибо за внимание!
2 ответа
Судя по комментариям, кажется, что единственным работающим решением является отключение кэширования L1. Это может быть достигнуто в рамках всей программы путем передачи следующего параметра в nvcc при компиляции:
–Xptxas –dlcm=cg
Кэши L1 являются свойством / ресурсом SM, а не устройством в целом. Поскольку потоковые блоки выполняются на определенных SM, активность одного потокового блока в его кэше L1 может быть не связана с активностью другого потокового блока и его кэша L1 (при условии, что он работает на другом SM), даже если они оба ссылаются на одну и ту же места в глобальной памяти. Кэши L1 в разных SM не имеют никакой связи друг с другом и не гарантированно согласованы друг с другом.
Обратите внимание, что кэш L2 распространяется на все устройство и, следовательно, является "связным" с точки зрения отдельных потоковых блоков. Отключение кэширования L1 не влияет на кэширование L2, поэтому все еще существует вероятность некоторого преимущества кэширования, однако время, необходимое для удовлетворения запроса из L2, больше, чем время, необходимое для удовлетворения запроса из L1, поэтому отключение Кэширование L1 по всей программе - довольно большой молоток, чтобы попытаться заставить все работать.
volatile Ключевое слово перед определением переменной должно указывать компилятору пропускать кеширование L1 при загрузке (согласно моему пониманию). Но volatile сама по себе не обращается к пути записи, поэтому один потоковый блок в одном SM может выполнить volatile прочитайте, извлекая значение из L2, измените это значение, а затем запишите его обратно, где оно заканчивается в L1 (до тех пор, пока оно не будет выселено). Если другой потоковый блок считывает то же глобальное значение, он может не увидеть эффект обновления.
Прилежное использование __threadfence (), хотя и утомительно, должно заставить любые такие обновления из L1 в L2, чтобы их могли прочитать другие блоки потоков. Однако это все еще оставляет промежуток синхронизации от того, когда значение было записано, до того момента, когда оно наблюдается другими SM / потоковыми блоками.
(Global) Atomics также должен иметь эффект прямого перехода к "глобальной памяти" для чтения и записи используемых значений.
Также может быть полезно пройти через код, чтобы гарантировать, что каждое возможное чтение из глобально синхронизированного местоположения обрабатывается должным образом (например, с volatile или используя атомику) и что каждая возможная запись в глобально синхронизированное местоположение обрабатывается должным образом (например, с __threadfence() или атомарность), а также проверьте состояние гонки между различными блоками.
Как выяснилось, процесс создания стабильной глобально синхронизированной среды в графическом процессоре нетривиален. Эти другие вопросы также могут представлять интерес (например, в отношении Кеплера) (и, например, обсуждение глобальных семафоров).
Изменить: Чтобы ответить на вопрос, размещенный в комментариях, я бы сказал так:
Возможно, нет проблем. тем не мение __threadfence() не дает никаких гарантий (насколько я знаю) в течение максимального времени завершения. Поэтому в момент обновления глобального местоположения обновляется только L1, связанный с исполняющим блоком потоков /SM. Тогда мы попали в __threadfence(), Предположительно, выполнение threadfence занимает некоторое время, и в течение этого времени другой потоковый блок может быть резидентным на том же SM, доставлен для выполнения (в то время как предыдущий поток / деформация / блок остановлен в threadfence) и "увидит" обновленное глобальное значение. в (местном) L1, связанном с этим SM. Другие потоковые блоки, выполняющиеся в других SM, будут видеть устаревшее значение до __threadfence() завершается. Это то, что я называю возможным "разрывом синхронизации". Два разных блока все еще могут видеть два разных значения в течение короткого периода времени. Будет ли это важно или нет, будет зависеть от того, как глобальное значение используется для синхронизации между блоками (так как это обсуждаемая тема). Поэтому atomics + volatile может быть лучшим выбором, чем volatile + threadfence, чтобы попытаться охватить оба чтения и напишите пути для синхронизации.
Редактировать #2: Из комментариев видно, что сочетание использования атомных плюс volatile также решил проблему.
Честно говоря, я нахожу ваш код слишком сложным с индексами и, что более важно, неполным. Как popBottom а также popTop функционировать? Кроме того, как это push операция выполнена? Эти два должны быть тщательно разработаны, чтобы работать правильно и гарантировать, что некоторые проблемы синхронизации не происходят.
Например: что произойдет, когда один блок попытается вставить что-то в свою очередь глобальной памяти, а другой блок попытается прочитать из нее в тот же самый момент? Это очень важно, и если это не сделано правильно, это может привести к сбою в некоторых очень редких случаях, например, вы можете выскочить из ячейки данных, в которую еще не было записано.
Когда я реализовывал аналогичную вещь - один дуэк глобальной памяти, совместно используемый всеми блоками, я дополнительно отмечал каждую ячейку данных как: пустую, занятую и мертвую. В псевдокоде алгоритм работал примерно так:
/* Objects of this class should reside in CUDA global memory */
template <typename T, size_t size>
class WorkQueue {
private:
size_t head, tail;
size_t status[size];
T data[size];
enum {
FieldFree = 0,
FieldDead = 1,
FieldTaken = 2
};
public:
/*
This construction should actually be done by host on the device,
before the actual kernel using it is launched!
Zeroing the memory should suffice.
*/
WorkQueue() : head(0), tail(0) {
for (size_t i=0; i<size; ++i)
status[i]=FieldFree;
}
__device__ bool isEmpty() { return head==tail; }
/* single thread of a block should call this */
__device__ bool push(const T& val) {
size_t oldFieldStatus;
do {
size_t cell = atomicInc(&tail,size-1);
data[cell]=val;
__threadfence(); //wait untill all blocks see the above change
oldFieldStatus=atomicCAS(&status[cell],FieldFree,FieldTaken); //mark the cell as occupied
} while (oldFieldStatus!=FieldFree);
return true;
}
/* single thread of a block should call this */
__device__ bool pop(T& out) {
size_t cellStatus;
size_t cell;
do {
cell=atomicInc(&head,size-1);
cellStatus=atomicCAS(&status[cell],FieldFree,FieldDead);
//If cell was free, make it dead - any data stored there will not be processed. Ever.
} while (cellStatus==FieldDead);
if (cellStatus!=FieldTaken)
return false;
out = data[cell];
status[cell]=FieldFree;
return true;
}
};
Я не вижу надежного способа реализовать его без статуса ячейки - иначе будут плохие вещи, если два потока из двух разных блоков попытаются протолкнуть / вытолкнуть одну и ту же ячейку очереди. При описанном выше подходе в худшем случае может случиться так, что всплывающий поток не сможет всплыть, вернуть false и пометить ячейку dead, и нить подталкивания будет повторять толкание в следующую ячейку. Идея заключается в том, что если всплывающий поток не может всплыть, то в любом случае, вероятно, не так много работы, и блок может завершиться. При таком подходе вы "убьете" только столько ячеек, сколько параллельных блоков.
Обратите внимание, в приведенном выше коде я не проверяю на переполнение!