Выполните простую атомарную операцию сложения при помощи gpucoder.atomicAdd функционируйте и сгенерируйте CUDA® код, который вызывает соответствующий atomicAdd() CUDA
В одном файле запишите функции точки входа myAtomicAdd это признает, что матрица вводит a и b.
Создать тип для матрицы удваивается для использования в генерации кода, используйте coder.newtype функция.
Чтобы сгенерировать библиотеку CUDA, используйте codegen функция.
Сгенерированный код 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);
}
...
}