api_cuda_contacts.hpp | Ansys Developer Portal

#pragma once
 
#include <rocky20/api/device/api_device_contacts.hpp>
#include <rocky20/api/device/api_cuda.hpp>
#include <rocky20/api/rocky_api_functions.h>
#include <rocky20/api/rocky_api.h>
 
namespace rocky::cuda
{
 
// =================================================================================================
// Contacts Loop
// =================================================================================================
 
template <typename Callable>
__global__ void execute_frictional_contacts_function(
    SDeviceModel *device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void *udf_device_data
)
{
    // Particles
    device::execute_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->ind_cnt_p_fr,
        device_data->n_contacts->particle_fr,
        true,
        particle_contact_f_ptr,
        udf_device_data);
 
    // Triangles
    device::execute_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->ind_cnt_t_fr,
        device_data->n_contacts->triangle_fr,
        false,
        triangle_contact_f_ptr,
        udf_device_data);
}
 
template <typename Callable>
__global__ void execute_finished_contacts_function(
    SDeviceModel* device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void* udf_device_data
)
{
    // particles -> particles
    device::execute_finished_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_p,
        true,
        particle_contact_f_ptr, udf_device_data
    );
 
    // particles -> triangles
    device::execute_finished_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_t,
        false,
        triangle_contact_f_ptr, udf_device_data
    );
}
 
template <typename Callable>
__global__ void execute_active_contacts_function(
    SDeviceModel* device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void* udf_device_data
)
{
    // particles -> particles
    device::execute_active_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_p,
        true,
        particle_contact_f_ptr,
        udf_device_data);
 
    // particles -> triangles
    device::execute_active_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_t,
        false,
        triangle_contact_f_ptr,
        udf_device_data);
}
 
template <typename Callable>
void dispatch_frictional_contacts_kernel(
    IRockyModel &model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void **udf_devices_data
)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        void *udf_device_data = udf_devices_data[i];
        SDeviceModel *device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i));
        rocky::cuda::execute_frictional_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
            device_data,
            model.get_current_time(),
            h_particle_contact_f_ptr,
            h_triangle_contact_f_ptr,
            udf_device_data);
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
// frictional and adhesive contacts just finished
template <typename Callable>
void dispatch_finished_contacts_kernel(
    IRockyModel& model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void** udf_devices_data
)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        void* udf_device_data = udf_devices_data[i];
        SDeviceModel* device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i));
        rocky::cuda::execute_finished_contacts_function <<<CD_NBLOCKS, CD_NTHREADS>>> (
            device_data, model.get_current_time(),
            h_particle_contact_f_ptr, h_triangle_contact_f_ptr,
            udf_device_data
        );
    }
 
    sync_devices(model);
}
 
// frictional and adhesive contacts still active
template <typename Callable>
void dispatch_active_contacts_kernel(
    IRockyModel& model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void** udf_devices_data)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        void* udf_device_data = udf_devices_data[i];
        SDeviceModel* device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i));
        rocky::cuda::execute_active_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
            device_data,
            model.get_current_time(),
            h_particle_contact_f_ptr,
            h_triangle_contact_f_ptr,
            udf_device_data);
    }
 
    sync_devices(model);
}
 
// Adhesive ----------------------------------------------------------------------------------------
 
template <typename Callable>
__global__ void execute_adhesive_contacts_function(
    SDeviceModel *device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void *udf_device_data
)
{
    // Particles
    if (device_data->n_contacts->particle_ad > 0)
    {
        device::execute_contacts_function<CudaBackend>(
            device_data,
            time_curr,
            device_data->ind_cnt_p_ad,
            device_data->n_contacts->particle_ad,
            true,
            particle_contact_f_ptr,
            udf_device_data);
    }
 
    // Triangles
    if (device_data->n_contacts->triangle_ad > 0)
    {
        device::execute_contacts_function<CudaBackend>(
            device_data,
            time_curr,
            device_data->ind_cnt_t_ad,
            device_data->n_contacts->triangle_ad,
            false,
            triangle_contact_f_ptr,
            udf_device_data);
    }
}
 
template <typename Callable>
void dispatch_adhesive_contacts_kernel(
    IRockyModel &model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void **udf_devices_data
)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    for (int i = 0; i<model.get_number_of_devices(); i++)
    {
        void *udf_device_data = udf_devices_data[i];
        SDeviceModel *device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i));
        rocky::cuda::execute_adhesive_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
            device_data,
            model.get_current_time(),
            h_particle_contact_f_ptr,
            h_triangle_contact_f_ptr,
            udf_device_data);
    }
 
    //Sync all GPUs
    sync_devices(model);
 
}
 
// =================================================================================================
// All Contacts Loop
// =================================================================================================
template <typename Callable>
__global__ void execute_all_contacts_function(
    SDeviceModel *device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void *udf_device_data
)
{
    // Particles
    device::execute_all_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_p,
        true,
        particle_contact_f_ptr,
        udf_device_data);
 
    // Triangles
    device::execute_all_contacts_function<CudaBackend>(
        device_data,
        time_curr,
        device_data->gen->ncont_t,
        false,
        triangle_contact_f_ptr,
        udf_device_data);
}
 
template <typename Callable>
void dispatch_all_contacts_kernel(
    IRockyModel &model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void **udf_device_data
)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    for (int i = 0; i<model.get_number_of_devices(); i++)
    {
        CUDA_SET_DEVICE(model.get_device_id(i));
        rocky::cuda::execute_all_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
            model.get_device_data(i),
            model.get_current_time(),
            h_particle_contact_f_ptr,
            h_triangle_contact_f_ptr,
            udf_device_data[i]);
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
// =================================================================================================
// Contact forces
// =================================================================================================
template <typename CallableCollection>
__global__ void execute_frictional_particle_contact_forces_function(
    SDeviceModel* device_data,
    cu_real time_curr,
    const CallableCollection functions,
    void* udf_device_data)
{
    device::api_compute_particle_contact_frictional_forces<CudaBackend>(
        device_data, time_curr, functions, udf_device_data);
}
 
template <typename CallableCollection>
__global__ void execute_frictional_triangle_contact_forces_function(
    SDeviceModel* device_data,
    cu_real time_curr,
    const CallableCollection functions,
    void* udf_device_data)
{
    device::api_compute_triangle_contact_frictional_forces<CudaBackend>(
        device_data, time_curr, functions, udf_device_data);
}
 
 
template <typename CallableCollection>
void dispatch_frictional_contact_forces_kernel(
    IRockyModel& model, const CallableCollection& function_ptrs, void** udf_devices_data)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    auto& solver_model = *model._impl;
 
    solver_model.n_contacts.reset();
    std::vector<SNumberContacts> n_contacts(model.get_number_of_devices());
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        CUDA_SET_DEVICE(model.get_device_id(i));
        CUDA_COPY_D2H(n_contacts.data() + i, solver_model.GetDeviceModel(i).n_contacts, 1);
        solver_model.n_contacts.join(n_contacts[i]);
    }
 
    void* udf_device_data = nullptr;
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        if (udf_devices_data) udf_device_data = udf_devices_data[i];
        SDeviceModel* device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i))
        auto [blocks, threads] = CudaBackend::get_launch_params(n_contacts[i].particle_fr);
        rocky::cuda::execute_frictional_particle_contact_forces_function<<<blocks, threads>>>(
            device_data, model.get_current_time(), function_ptrs, udf_device_data);
    }
 
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        if (udf_devices_data) udf_device_data = udf_devices_data[i];
        SDeviceModel* device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i))
        auto [blocks, threads] = CudaBackend::get_launch_params(n_contacts[i].triangle_fr);
        rocky::cuda::execute_frictional_triangle_contact_forces_function <<<blocks, threads>>> (
            device_data, model.get_current_time(), function_ptrs, udf_device_data);
    }
 
    sync_devices(model);
}
 
template <typename Callable>
__global__ void execute_adhesive_contact_forces_function(
    SDeviceModel *device_data,
    cu_real time_curr,
    Callable particle_contact_f_ptr,
    Callable triangle_contact_f_ptr,
    void *udf_device_data
)
{
    // Particles
    device::api_compute_particle_contact_adhesive_forces<CudaBackend>(
        device_data, time_curr,
        particle_contact_f_ptr, udf_device_data
    );
 
    // Triangles
    device::api_compute_triangle_contact_adhesive_forces<CudaBackend>(
        device_data, time_curr,
        triangle_contact_f_ptr, udf_device_data
    );
}
 
template <typename Callable>
void dispatch_adhesive_contact_forces_kernel(
    IRockyModel &model,
    Callable h_particle_contact_f_ptr,
    Callable h_triangle_contact_f_ptr,
    void **udf_devices_data
)
{
    if (model.get_number_of_active_particles() <= 0)
        return;
 
    std::vector<SNumberContacts> n_contacts(model.get_number_of_devices());
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        CUDA_SET_DEVICE(model.get_device_id(i))
        CUDA_COPY_D2H(n_contacts.data() + i, model._impl->GetDeviceModel(i).n_contacts, 1)
    }
 
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        void *udf_device_data = udf_devices_data[i];
        SDeviceModel *device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i))
        auto [blocks, threads] = CudaBackend::get_launch_params(
            std::max(n_contacts[i].particle_ad, n_contacts[i].triangle_ad));
        rocky::cuda::execute_adhesive_contact_forces_function<<<blocks, threads>>>(
            device_data,
            model.get_current_time(),
            h_particle_contact_f_ptr,
            h_triangle_contact_f_ptr,
            udf_device_data);
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
// Boundary wear on particle-triangle contacts ----------------------------------------------------
 
template <typename Callable>
__global__ void execute_boundaries_wear_function_kernel(
    SDeviceModel* device_data,
    Callable triangle_wear_f_ptr,
    void* udf_device_data,
    cu_real time_curr)
{
    device::api_compute_boundaries_wear<CudaBackend>(
        device_data, time_curr, triangle_wear_f_ptr, udf_device_data);
}
 
template <typename Callable>
void dispatch_boundaries_wear_kernel(
    IRockyModel& model,
    Callable triangle_wear_f_ptr,
    void** udf_devices_data
)
{
    for (int i = 0; i < model.get_number_of_devices(); i++)
    {
        void* udf_device_data = udf_devices_data[i];
        SDeviceModel* device_data = model.get_device_data(i);
 
        CUDA_SET_DEVICE(model.get_device_id(i))
        rocky::cuda::execute_boundaries_wear_function_kernel<<<CD_NBLOCKS, CD_NTHREADS>>>(
            device_data,
            triangle_wear_f_ptr,
            udf_device_data,
            model._impl->iteration.current_time);
    }
 
    sync_devices(model);
}
 
} // namespace rocky::cuda
Rocky Solver SDK 2024 R2

api_cuda_contacts

Connect with Ansys
