Как правильно смоделировать `atomicAdd` на `u64`, используя два буфера `u32`?
Я пытаюсь выполнить атомарные операции над u64. Но поскольку он не поддерживается, номер хранится в ДВА u32 буферах.
Проблема в том, что я не знаю, как правильно сделать atomicAdd, чтобы имитировать эффект, который бы это оказало на u64. При этом избегая модификации памяти другими потоками между загрузкой и сохранением значений.
моя текущая идея такова:
fn tou64(value: u32) -> vec2u {
return vec2u(u32(value / BASE), value % BASE);
}
fn add(a: vec2u, b: vec2u) -> vec2u {
let x = a.x + b.x + u32((a.y + b.y) / BASE);
let y = (a.y + b.y) % BASE;
return vec2u(x, y);
}
fn main() {
// .....
// convert the value from u32 to 2-buffer representation of u64
let b: vec2u = tou64(value);
// fetch the old value from the 2 buffers
var a = vec2u(0);
a.x = atomicLoad(&buffer[index]);
a.y = atomicLoad(&buffer[index+1]);
// add the value to the buffer value
let result = add(a, b);
// store back the buffer results
atomicStore(&buffer[index], result.x);
atomicStore(&buffer[index+1], result.y);
}
Это работает только тогда, когда никакой другой поток одновременно не изменяет буфер index. Но в остальном это очень слабая реализация. Поток 1 может изменить значение buffer[index+1], а поток 2 просто прочитать старое значение buffer[index] и новое значение buffer[index+1].
Обновлено: в руководстве CUDA отмечено, что:
Обратите внимание, что любая атомарная операция может быть реализована на основеatomicCAS() (Сравнение и замена).
и представлен этот пример AtomicAdd для double
#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull =
(unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
__double_as_longlong(val +
__longlong_as_double(assumed)));
// Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
} while (assumed != old);
return __longlong_as_double(old);
}
#endif
можно ли это применить и к WebGPU, используя atomicExchange? этот ответ показывает, как выполнять атомарные операции над пользовательскими типами. Как я могу сделать что-то подобное, но для WebGPU?
Если все, что вы можете сделать, это записать отдельные элементы u32 по одному, вы не сможете выполнить атомарный RMW. Вы даже не выполняете атомарную RMW ни для одного элемента отдельно, а просто разделяете загрузку и сохранение, поэтому даже упрощение этого алгоритма до одного u32 не является потокобезопасным атомарным приращением. Два потока могут прочитать одно и то же начальное значение и оба сохранить одно и то же конечное значение, теряя счетчик. Или, что еще хуже, поток может загрузиться, зависнуть или быть исключен из расписания и выполнить сохранение намного позже, эффективно перезаписывая многие приращения.
Если вы сможете выполнить атомарное сложение на нижнем уровне u32 и обнаружить перенос, который можно передать атомарному RMW в старшей половине, тогда можно будет получить правильное окончательное значение счетчика, как только уляжется пыль. Несмотря на то, что две половины могут быть не синхронизированы в любой момент во время их изменения. Перенос распространяется только от младшего к старшему, а сложение является коммутативным + ассоциативным. Таким образом, добавление переносов из нижней половины в верхнюю можно выполнить независимо от RMW в нижней половине. (С помощью fetch_add вы можете обнаружить перенос с помощью sum = a+b/carry = sum < a, где a — старое значение
@PeterCordes, к сожалению, webgpu пока ВООБЩЕ не поддерживает. Даже определение типа внутри шейдера не поддерживается, не говоря уже об атомарности. Я пытался найти в шейдере механизм блокировки/барьера, но безрезультатно. Однако я нашел атомарную биржу сравнения для u64. Теперь я понимаю, как правильно увеличивать верхнюю половину (спасибо за отличное объяснение!). а как насчет нижней половины и переполнения? как предотвратить переполнение буфера?
Переполнение (переполнение u32) — это выход, который вы обнаруживаете с помощью carry = sum < a. Математика u32 в webgpu не работает так же, как C uint32_t, где сложение автоматически оборачивается? В C вы буквально делаете sum_low += input_low;sum_high += input_high + (sum_low < input_low). (Сравнение может проводиться с любым входным значением младшего сложения, поскольку оно коммутативно. Выполнение атомарного RMW с точки зрения CAS означает, что вам уже нужны как старые, так и новые sum_low в переменных, поэтому можно реализовать fetch_add или add_fetch.)
@PeterCordes, все сработало отлично !! большое спасибо!
Ответы 1
Вот решение для atomicAddU64, спасибо @PeterCordes!
Вручную выполните перенос между младшей половиной u32 и старшей половиной u32, причем каждое добавление u32 представляет собой атомарный RMW. Для этого требуется возвращаемое значение fetch_add/atomicAdd из младшей половины, чтобы мы могли проверить выполнение.
Поскольку у нас есть только две половинки, а не более широкое целое число, нам не нужно обрабатывать перенос и перенос из одного и того же сложения. Просто sum = a+b;carry_out = sum<b; трюк.
Общее количество в конце будет правильным, но невозможно прочитать правильный снимок текущего счетчика, пока другие потоки добавляются. Например, несколько потоков могли добавиться к нижней половине и все еще ожидать добавления к старшей половине.
(Возможно, было бы лучше, чтобы каждый поток добавлял в локальные аккумуляторы каждого потока для суммирования один раз в конце, вместо того, чтобы все потоки постоянно обращались к одной общей переменной.)
fn main() {
// .....
// convert the value from u32 or f32 to 2-buffer representation of u64
let b: vec2u = tou64(value);
// low: no need for modulu since it will wrap by itself
// take old value on the buffer to check for carry
let oldValue = atomicAdd(&buffer[low_index], b.y);
// high: add high part + carry
// check if the sum cause value to wrap
atomicAdd(&buffer[high_index], b.x + u32((oldValue + b.y) < b.y));
}
Чтобы преобразовать u64 в vec2u и обратно:
// this is just a pseudocode! wgsl doesn't support u64 yet
// do this operation in c++/js or while data processing
fn u64ToVec2u(value: u64) -> vec2u {
let low = u32(value);
let high = u32(value >> 32);
return vec2u(high, low);
}
fn vec2uToU64(value: vec2u) -> u64 {
return (value.x << 32) + value.y; // x = high, y = low
}
и вот решение для вычислений с плавающей запятой, можно сделать в wgsl
// to save 2 ^ 10 = 1024 (ie 3 digits after the decimal point)
const DEGREE_TO_SAVE = 10;
fn tou64(value: f32) -> vec2u {
// modulu is important here!!
// because converting from float to u32 will not automatically wrap
let low = u32((value * pow(2., DEGREE_TO_SAVE)) % pow(2., 32));
let high = u32(value / pow(2., 32 - DEGREE_TO_SAVE));
return vec2u(high, low);
}
fn tof32(value: vec2u) -> f32 {
return f32(value.x) * pow(2., 32 - DEGREE_TO_SAVE) + f32(value.y) / pow(2., DEGREE_TO_SAVE);
}
Небольшое предостережение. Расширение этого кода с 64 бит до 96 бит имеет некоторую сложность. Приведенный выше код не может отличить добавление 0xFFFFFFFF с переносом 1 от добавления 0 без переноса. При выполнении третьего (и последующих) сложений проверка исходящего переноса выполняется так: result <= addend был ли входящий перенос, и result < addend — если его не было.
Есть ли в webgpu какие-либо атомарные операции с 64-битным значением? Типа сравнить-обменять? В противном случае вам понадобится какая-то блокировка. Если есть только один писатель, это может быть секлок, где читатели смогут обнаружить разрыв и повторить попытку чтения. Если модификации происходят нечасто, это может работать хорошо.