Exponenta Banner

gpucoder.atomicAdd

Атомарно добавьте заданное значение в переменную в глобальной памяти или общей памяти

свернуть все на странице

    Синтаксис

    [A,oldA] = gpucoder.atomicAdd(A,B)

    Описание

    пример

    [A,oldA] = gpucoder.atomicAdd(A,B) добавляет B к значению A в глобальной памяти или общей памяти и записывает результат обратно в A. Операция является атомарной в некотором смысле, что целая операция "чтение изменяет запись", как, гарантируют, будет выполнен без интерференции от других потоков. Порядок аргументов ввода и вывода должен совпадать с обеспеченным синтаксисом.

    Примеры

    свернуть все

    Выполните простую атомарную операцию сложения при помощи gpucoder.atomicAdd функционируйте и сгенерируйте CUDA® код, который вызывает соответствующий atomicAdd() CUDA API.

    В одном файле запишите функции точки входа myAtomicAdd это признает, что матрица вводит a и b.

    function [a,oldA] = myAtomicAdd(a,b)

oldA = coder.nullcopy(a);

for i = 1:size(a,1)
    for j = 1:size(a,2)
        for k = 1:size(a,3)
            [a(i,j,k),oldA(i,j,k)] = gpucoder.atomicAdd(a(i,j,k),b(i,j,k));
        end
    end
end

    Создать тип для матрицы удваивается для использования в генерации кода, используйте coder.newtype функция.

    A = coder.newtype('single', [30 30 20], [1 0 1]);
B = coder.newtype('single', [30 30 20], [1 0 1]);
inputArgs = {A,B};

    Чтобы сгенерировать библиотеку CUDA, используйте codegen функция.

    cfg = coder.gpuConfig('lib');
cfg.GenerateReport = true;

codegen -config cfg -args inputArgs myAtomicAdd -d myAtomicAdd

    Сгенерированный код CUDA содержит myAtomicAdd_kernel1 ядро с вызовами atomicAdd() API CUDA.

    //
// File: myAtomicAdd.cu
//
...
static __global__ __launch_bounds__(1024, 1) void myAtomicAdd_kernel1(
    const float b_data[], const int b_size[3], int a_size[3],
    const int oldA_size[3], const int b_a_size, const int i, float oldA_data[],
    float a_data[])
{
  unsigned long loopEnd;
  unsigned long threadId;
...

  for (unsigned long idx{threadId}; idx <= loopEnd; idx += threadStride) {
    unsigned long tmpIndex;
    int b_i;
    int j;
    int k;
    k = static_cast<int>(idx % (static_cast<unsigned long>(b_a_size) + 1UL));
    tmpIndex = (idx - static_cast<unsigned long>(k)) /
               (static_cast<unsigned long>(b_a_size) + 1UL);
    j = static_cast<int>(tmpIndex % 30UL);
    tmpIndex = (tmpIndex - static_cast<unsigned long>(j)) / 30UL;
    b_i = static_cast<int>(tmpIndex);
    oldA_data[(b_i + oldA_size[0] * j) + oldA_size[0] * 30 * k] =
        atomicAdd(&a_data[(b_i + a_size[0] * j) + a_size[0] * 30 * k],
                  b_data[(b_i + b_size[0] * j) + b_size[0] * 30 * k]);
  }
}
...

void myAtomicAdd(float a_data[], int a_size[3], const float b_data[],
                 const int b_size[3], float oldA_data[], int oldA_size[3])
{
  dim3 block;
  dim3 grid;
...

    cudaMemcpy(gpu_a_data, a_data, a_size[0] * (30 * a_size[2]) * sizeof(float),
               cudaMemcpyHostToDevice);
    myAtomicAdd_kernel1<<<grid, block>>>(gpu_b_data, *gpu_b_size, *gpu_a_size,
                                         *gpu_oldA_size, b_a_size, i,
                                         gpu_oldA_data, gpu_a_data);
    oldA_data_dirtyOnGpu = true;
    cudaMemcpy(a_data, gpu_a_data, a_size[0] * (30 * a_size[2]) * sizeof(float),
               cudaMemcpyDeviceToHost);
  }
...

}

    Входные параметры

    свернуть все

    Операнды в виде скаляров, векторов, матриц или многомерных массивов. Входные параметры A и B должен удовлетворить следующим требованиям:

    • Имейте совпадающий тип данных.

    • Имейте тот же размер или имейте размеры, которые совместимы. Например, A M- N матрица и B скаляр или 1- N вектор-строка.

    • Требует, чтобы устройство CUDA с минимумом вычислило возможность 6,0, когда типом данных является double.

    Типы данных: double | single | int32 | uint32 | uint64

    Ограничения

    • Вход указателя на функцию к gpucoder.stencilKernel прагма не может содержать вызовы атомарных функций. Например,

      out1 = gpucoder.stencilKernel(@myAtomicAdd,A,[3 3],'same',B);

    Смотрите также

    Функции

    Темы

    Введенный в R2021b

    Документация GPU Coder

    Поддержка