api_cuda_joints.hpp | Ansys Developer Portal

#pragma once
 
#include <rocky20/api/device/api_device_joints.hpp>
#include <rocky20/api/device/api_cuda.hpp>
 
namespace rocky::cuda
{
 
template <typename Callable>
void dispatch_joints_kernel(
    IRockyModel &model,
    Callable h_joint_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));
        device::execute_joints_function<CudaBackend><<<CD_NBLOCKS, CD_NTHREADS>>>(
            device_data, i, model.get_current_time(),
            h_joint_f_ptr,
            udf_device_data
        );
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
template <typename Callable>
__global__ void execute_joint_forces_function_kernel(
    SDeviceModel *device_data, int device_id,
    Callable joint_f_ptr,
    void *udf_device_data
)
{
    device::compute_joint_forces<CudaBackend>(
        device_data, device_id,
        joint_f_ptr, udf_device_data
    );
}
 
template <typename Callable>
__global__ void execute_joint_breakage_function_kernel(
    SDeviceModel* device_data, int device_id,
    Callable joint_breakage_f_ptr,
    void* udf_device_data,
    typename CudaBackend::atomic_int_ptr nbreak_joints_a,
    cu_real time_curr
)
{
    device::compute_joint_breakage<CudaBackend>(
        device_data, device_id,
        joint_breakage_f_ptr, udf_device_data,
        nbreak_joints_a,
        time_curr
    );
}
 
template <typename Callable>
void dispatch_joint_forces_kernel(
    IRockyModel &model,
    Callable h_joint_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_joint_forces_function_kernel << <CD_NBLOCKS, CD_NTHREADS >> > (
            device_data, i,
            h_joint_f_ptr,
            udf_device_data
        );
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
template <typename Callable>
void dispatch_joint_breakage_kernel(
    IRockyModel& model,
    Callable h_joint_breakage_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_joint_breakage_function_kernel <<<CD_NBLOCKS, CD_NTHREADS >>> (
            device_data, i,
            h_joint_breakage_f_ptr,
            udf_device_data,
            &(model._impl->device_models[i].gen->nbreak_joints),
            model._impl->iteration.current_time
        );
    }
 
    //Sync all GPUs
    sync_devices(model);
}
 
} // namespace rocky::cuda
Rocky Solver SDK 2024 R2

api_cuda_joints

Connect with Ansys
