// Copyright (c) 2017-2018, Lawrence Livermore National Security, LLC. // Produced at the Lawrence Livermore National Laboratory. LLNL-CODE-734707. // All Rights reserved. See files LICENSE and NOTICE for details. // // This file is part of CEED, a collection of benchmarks, miniapps, software // libraries and APIs for efficient high-order finite element and spectral // element discretizations for exascale applications. For more information and // source code availability see http://github.com/ceed. // // The CEED research is supported by the Exascale Computing Project 17-SC-20-SC, // a collaborative effort of two U.S. Department of Energy organizations (Office // of Science and the National Nuclear Security Administration) responsible for // the planning and preparation of a capable exascale ecosystem, including // software, applications, hardware, advanced system engineering and early // testbed platforms, in support of the nation's exascale computing imperative. #include #include #include #include #include #include "ceed-cuda-ref.h" #include "../cuda/ceed-cuda-compile.h" static const char *qReadWrite = QUOTE( template //------------------------------------------------------------------------------ // Read from quadrature points //------------------------------------------------------------------------------ inline __device__ void readQuads(const CeedInt quad, const CeedInt nquads, const CeedScalar* d_u, CeedScalar* r_u) { for(CeedInt comp = 0; comp < SIZE; ++comp) { r_u[comp] = d_u[quad + nquads * comp]; } } //------------------------------------------------------------------------------ // Write at quadrature points //------------------------------------------------------------------------------ template inline __device__ void writeQuads(const CeedInt quad, const CeedInt nquads, const CeedScalar* r_v, CeedScalar* d_v) { for(CeedInt comp = 0; comp < SIZE; ++comp) { d_v[quad + nquads * comp] = r_v[comp]; } } ); //------------------------------------------------------------------------------ // Build QFunction kernel //------------------------------------------------------------------------------ extern "C" int CeedCudaBuildQFunction(CeedQFunction qf) { CeedInt ierr; using std::ostringstream; using std::string; Ceed ceed; CeedQFunctionGetCeed(qf, &ceed); CeedQFunction_Cuda *data; ierr = CeedQFunctionGetData(qf, (void **)&data); CeedChkBackend(ierr); // QFunction is built if (data->qFunction) return CEED_ERROR_SUCCESS; if (!data->qFunctionSource) return CeedError(ceed, CEED_ERROR_BACKEND, "No QFunction source or CUfunction provided."); // QFunction kernel generation CeedInt numinputfields, numoutputfields, size; CeedQFunctionField *qfinputfields, *qfoutputfields; ierr = CeedQFunctionGetFields(qf, &numinputfields, &qfinputfields, &numoutputfields, &qfoutputfields); CeedChkBackend(ierr); // Build strings for final kernel string qFunction(data->qFunctionSource); string qReadWriteS(qReadWrite); ostringstream code; string qFunctionName(data->qFunctionName); string kernelName; kernelName = "CeedKernel_Cuda_ref_" + qFunctionName; // Defintions code << "\n#define CEED_QFUNCTION(name) inline __device__ int name\n"; code << "#define CEED_QFUNCTION_HELPER inline __device__\n"; code << "#define CeedPragmaSIMD\n"; code << "#define CEED_ERROR_SUCCESS 0\n"; code << "#define CEED_Q_VLA 1\n\n"; code << "typedef struct { const CeedScalar* inputs[16]; CeedScalar* outputs[16]; } Fields_Cuda;\n"; code << qReadWriteS; code << qFunction; code << "extern \"C\" __global__ void " << kernelName << "(void *ctx, CeedInt Q, Fields_Cuda fields) {\n"; // Inputs for (CeedInt i = 0; i < numinputfields; i++) { code << "// Input field "<(q, Q, fields.inputs["<(q, Q, r_qq"<module, 0); CeedChkBackend(ierr); ierr = CeedGetKernelCuda(ceed, data->module, kernelName.c_str(), &data->qFunction); CeedChkBackend(ierr); // Cleanup ierr = CeedFree(&data->qFunctionSource); CeedChkBackend(ierr); return CEED_ERROR_SUCCESS; } //------------------------------------------------------------------------------