Ошибка запуска ядра, если объем разделяемой памяти, выделенной для всей сетки, превышает 48 КБ

Я работаю над проблемой N-тела, требующей большого количества общей памяти.

В основном, есть N независимые задачи, каждая из которых использует 4 двойных переменных, то есть 32 байта. И одна задача выполняется потоком.

Ради быстроты я использовал разделяемую память для этих переменных (учитывая, что регистры также используются потоками). Поскольку число N задач не известно во время компиляции, разделяемая память выделяется динамически.

  • Размерность сетки и общей памяти вычисляются в зависимости от N и размер блока:

    const size_t BLOCK_SIZE = 512;
const size_t GRID_SIZE = (N % BLOCK_SIZE) ? (int) N/BLOCK_SIZE : (int) N/BLOCK_SIZE +1;
const size_t SHARED_MEM_SIZE = BLOCK_SIZE * 4 * sizeof(double);

  • Затем ядро ​​запускается с использованием этих 3 переменных.

    kernel_function<<<GRID_SIZE, BLOCK_SIZE, SHARED_MEM_SIZE>>>(N, ...);

Для маленьких N, это работает нормально, и ядро ​​выполняется без ошибок.

Но если превышать N = 1500 запуск ядра завершается неудачно (следующие сообщения появляются несколько раз):

========= Invalid __global__ write of size 8
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaLaunch.

Насколько я понимаю, это связано с попыткой выписывания за пределы выделенной разделяемой памяти. Это происходит, когда в ядре глобальная память копируется в общую память:

__global__ void kernel_function(const size_t N, double *pN, ...)
{
    unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if(idx<N)
    {
        extern __shared__ double pN_shared[];
        for(int i=0; i < 4; i++)
        {
            pN_shared[4*idx + i] = pN[4*idx + i];
        }
        ...
    }
}

Эта ошибка происходит только если N > 1500 следовательно, когда общий объем разделяемой памяти превышает 48 КБ (1500 * 4 * sizeof(double) = 1500 * 32 = 48000).
Это ограничение одинаково независимо от сетки и размера блока.

Если я правильно понял, как работает CUDA, накопленный объем разделяемой памяти, которую использует сетка, не ограничен 48 КБ, и это только предел разделяемой памяти, который может использоваться одним блоком потока.

Эта ошибка не имеет смысла для меня, поскольку накопленный объем разделяемой памяти должен влиять только на то, как сетка распределяется между потоковыми мультипроцессорами (и, кроме того, на устройстве с графическим процессором имеется 15 SM в распоряжении).

Источник

3 ответа

Решение

Вы получаете доступ к общему массиву по индексу idx*4+0:3. Программа некорректна, начиная с N > BLOCK_SIZE. К счастью, кажется, что он работает до 1500. Но использование cuda mem-check должно указать на проблему. Обратите внимание, что в связанной теме статически распределенная общая память в другом месте может использовать общую память. Распечатка значения указателя поможет выяснить.

Источник

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

kernel_function<<<GRID_SIZE, BLOCK_SIZE, SHARED_MEM_SIZE>>>(N, ...);
                                         ^^^^^^^^^^^^^^^

сумма за блок потока, и эта сумма ограничена 48 КБ (что составляет 49152, а не 48000). Поэтому, если вы попытаетесь выделить более 48 КБ, вы должны получить ошибку, если проверяете ее.

Однако из этого можно сделать два вывода:

========= Invalid __global__ write of size 8
  1. Ядро действительно запустилось.
  2. Сообщаемая ошибка связана с внеплановой индексацией в глобальную память, при записи в глобальную память, а не в разделяемую память. (Таким образом, это не может происходить при чтении из глобальной памяти для заполнения общей памяти, как предполагает ваша гипотеза.)

В общем, я думаю, что ваши выводы неверны, и вам, вероятно, нужно больше отлаживать, чем приходить к выводам об общей памяти.

Если вы хотите отследить источник недопустимой глобальной записи в конкретную строку кода в вашем ядре, этот вопрос / ответ может представлять интерес.

Источник

Я думаю, что проблема здесь в том, что все потоки внутри блока должны работать в одном SM. Поэтому каждый блок по-прежнему имеет жесткий предел 48 КБ общей памяти. Не имеет значения, сколько потоков запущено в этом блоке. Планирование не имеет значения, поскольку графический процессор не может разделить потоки в блоке между несколькими SM. Я бы попытался уменьшить BLOCK_SIZE, если вы можете, так как это будет непосредственно определять объем разделяемой памяти на блок. Однако, если вы уменьшите его слишком сильно, вы можете столкнуться с проблемами, когда вы не используете полностью вычислительные ресурсы в SM. Это уравновешивающее действие, и, исходя из моего опыта, архитектура CUDA представляет множество интересных компромиссов, подобных этому.

Также в вашем случае я даже не уверен, что вам нужна общая память. Я бы просто использовал локальную переменную. Я думаю, что локальные переменные хранятся в глобальной памяти, но доступ к ним выровнен, так что это очень быстро. Если вы хотите сделать что-то аккуратное с общей памятью, чтобы улучшить производительность, вот ядро ​​OpenCL моего симулятора N-Body. Использование общей памяти для создания кэша для каждого потока в блоке дает мне примерно 10-кратное ускорение.

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

__kernel void acceleration_kernel
(
    __global const double* masses, 
    __global const double3* positions,
    __global double3* accelerations,
    const double G,
    const int N,
    __local double4* cache //shared memory cache (local means shared memory in OpenCL)
)
{
    int idx = get_global_id(0);
    int lid = get_local_id(0);
    int lsz = get_local_size(0);

    if(idx >= N)
        return;

    double3 pos = positions[idx];
    double3 a = { };

    //number of loads required to compute accelerating on Body(idx) from all other bodies
    int loads = (N + (lsz - 1)) / lsz;

    for(int load = 0; load < loads; load++)
    {
        barrier(CLK_LOCAL_MEM_FENCE);

        //compute which body this thread is responsible for loading into the cache
        int load_index = load * lsz + lid;
        if(load_index < N)
            cache[lid] = (double4)(positions[load_index], masses[load_index]);

        barrier(CLK_LOCAL_MEM_FENCE);

        //now compute the acceleration from every body added to the cache
        for(int i = load * lsz, j = 0; i < N && j < lsz; i++, j++)
        {
            if(i == idx)
                continue;
            double3 r_hat = cache[j].xyz - pos; 
            double over_r = rsqrt(0.0001 + r_hat.x * r_hat.x + r_hat.y * r_hat.y + r_hat.z * r_hat.z);
            a += r_hat * G * cache[j].w * over_r * over_r * over_r;
        }
    }

    accelerations[idx] = a;
}
double3 pos = positions[idx];
double3 a = { };

int loads = (N + (lsz - 1)) / lsz;

for(int load = 0; load < loads; load++)
{
    barrier(CLK_LOCAL_MEM_FENCE);
    int load_index = load * lsz + lid;
    if(load_index < N)
        cache[lid] = (double4)(positions[load_index], masses[load_index]);
    barrier(CLK_LOCAL_MEM_FENCE);

    for(int i = load * lsz, j = 0; i < N && j < lsz; i++, j++)
    {
        if(i == idx)
            continue;
        double3 r_hat = cache[j].xyz - pos; 
        double over_r = rsqrt(0.0001 + r_hat.x * r_hat.x + r_hat.y * r_hat.y + r_hat.z * r_hat.z);
        a += r_hat * G * cache[j].w * over_r * over_r * over_r;
    }
}

accelerations[idx] = a;

}

Источник
Другие вопросы по тегам