gpucoder.atomicCAS

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

    Описание

    пример

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

    Примеры

    свернуть все

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

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

    function a = myAtomicCAS(a,b,c)
    
    coder.gpu.kernelfun;
    for i =1:numel(a)
        [a(i),~] = gpucoder.atomicCAS(a(i), b, c);
    end
    
    end
    

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

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

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

    cfg = coder.gpuConfig('lib');
    cfg.GenerateReport = true;
    
    codegen -config cfg -args inputArgs myAtomicCAS -d myAtomicCAS
    

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

    //
    // File: myAtomicCAS.cu
    //
    ...
    
    static __global__ __launch_bounds__(1024, 1) void myAtomicCAS_kernel1(
        const uint32_T c, const uint32_T b, const int32_T i, uint32_T a_data[])
    {
      uint64_T loopEnd;
      uint64_T threadId;
    
    ...
    
      for (uint64_T idx{threadId}; idx <= loopEnd; idx += threadStride) {
        int32_T b_i;
        b_i = static_cast<int32_T>(idx);
        atomicCAS(&a_data[b_i], b, c);
      }
    }
    ...
    
    void myAtomicCAS(uint32_T a_data[], int32_T a_size[2], uint32_T b, uint32_T c)
    {
      dim3 block;
      dim3 grid;
    ...
    
      if (validLaunchParams) {
        cudaMemcpy(gpu_a_data, a_data, a_size[1] * sizeof(uint32_T),
                   cudaMemcpyHostToDevice);
        myAtomicCAS_kernel1<<<grid, block>>>(c, b, i, gpu_a_data);
        cudaMemcpy(a_data, gpu_a_data, a_size[1] * sizeof(uint32_T),
                   cudaMemcpyDeviceToHost);
    ...
    
    }
    

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

    свернуть все

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

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

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

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

    Введенный в R2021b
    Для просмотра документации необходимо авторизоваться на сайте