15aed82e4SJeremy L Thompson // Copyright (c) 2017-2024, Lawrence Livermore National Security, LLC and other CEED contributors. 2a0154adeSJed Brown // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3a0154adeSJed Brown // 4a0154adeSJed Brown // SPDX-License-Identifier: BSD-2-Clause 5a0154adeSJed Brown // 6a0154adeSJed Brown // This file is part of CEED: http://github.com/ceed 7a0154adeSJed Brown 8b2165e7aSSebastian Grimberg /// @file 9b2165e7aSSebastian Grimberg /// Internal header for CUDA non-tensor product basis 10b2165e7aSSebastian Grimberg 11*c0b5abf0SJeremy L Thompson #include <ceed/types.h> 12a0154adeSJed Brown 13d075f50bSSebastian Grimberg #include "cuda-ref-basis-nontensor-templates.h" 14d075f50bSSebastian Grimberg 15a0154adeSJed Brown //------------------------------------------------------------------------------ 16a0154adeSJed Brown // Non-Tensor Basis Kernels 17a0154adeSJed Brown //------------------------------------------------------------------------------ 18a0154adeSJed Brown 19a0154adeSJed Brown //------------------------------------------------------------------------------ 20a0154adeSJed Brown // Interp 21a0154adeSJed Brown //------------------------------------------------------------------------------ 22d075f50bSSebastian Grimberg extern "C" __global__ void Interp(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, 23a0154adeSJed Brown CeedScalar *__restrict__ d_V) { 242b730f8bSJeremy L Thompson for (CeedInt elem = blockIdx.x * blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x * blockDim.z) { 25d075f50bSSebastian Grimberg Contract<BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem, BASIS_Q * num_elem, 26d075f50bSSebastian Grimberg BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); 27a0154adeSJed Brown } 28a0154adeSJed Brown } 29d075f50bSSebastian Grimberg 30d075f50bSSebastian Grimberg extern "C" __global__ void InterpTranspose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, 31d075f50bSSebastian Grimberg CeedScalar *__restrict__ d_V) { 32d075f50bSSebastian Grimberg for (CeedInt elem = blockIdx.x * blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x * blockDim.z) { 33d075f50bSSebastian Grimberg ContractTranspose<BASIS_NUM_COMP, BASIS_Q_COMP_INTERP, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q * num_elem, BASIS_P * num_elem, 34d075f50bSSebastian Grimberg BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); 35a0154adeSJed Brown } 36a0154adeSJed Brown } 37a0154adeSJed Brown 38a0154adeSJed Brown //------------------------------------------------------------------------------ 39d075f50bSSebastian Grimberg // Deriv 40a0154adeSJed Brown //------------------------------------------------------------------------------ 41d075f50bSSebastian Grimberg extern "C" __global__ void Deriv(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, 42a0154adeSJed Brown CeedScalar *__restrict__ d_V) { 432b730f8bSJeremy L Thompson for (CeedInt elem = blockIdx.x * blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x * blockDim.z) { 44d075f50bSSebastian Grimberg Contract<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_P, BASIS_Q, BASIS_P * num_elem, BASIS_Q * num_elem, 45d075f50bSSebastian Grimberg BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); 46d075f50bSSebastian Grimberg } 47a0154adeSJed Brown } 48a0154adeSJed Brown 49d075f50bSSebastian Grimberg extern "C" __global__ void DerivTranspose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, 50d075f50bSSebastian Grimberg CeedScalar *__restrict__ d_V) { 51d075f50bSSebastian Grimberg for (CeedInt elem = blockIdx.x * blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x * blockDim.z) { 52d075f50bSSebastian Grimberg ContractTranspose<BASIS_NUM_COMP, BASIS_Q_COMP_DERIV, BASIS_P, BASIS_Q>(elem, BASIS_Q, BASIS_P, BASIS_Q * num_elem, BASIS_P * num_elem, 53d075f50bSSebastian Grimberg BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); 54a0154adeSJed Brown } 55a0154adeSJed Brown } 56a0154adeSJed Brown 57a0154adeSJed Brown //------------------------------------------------------------------------------ 58a0154adeSJed Brown // Weight 59a0154adeSJed Brown //------------------------------------------------------------------------------ 602b730f8bSJeremy L Thompson extern "C" __global__ void Weight(const CeedInt num_elem, const CeedScalar *__restrict__ q_weight, CeedScalar *__restrict__ d_V) { 61a0154adeSJed Brown const CeedInt t_id = threadIdx.x; 62a0154adeSJed Brown // TODO load q_weight in shared memory if blockDim.z > 1? 63d075f50bSSebastian Grimberg 642b730f8bSJeremy L Thompson for (CeedInt elem = blockIdx.x * blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x * blockDim.z) { 65a0154adeSJed Brown d_V[elem * BASIS_Q + t_id] = q_weight[t_id]; 66a0154adeSJed Brown } 67a0154adeSJed Brown } 68