From e0dd3b27d19566a0742e776b14ca8ce6d5794b11 Mon Sep 17 00:00:00 2001 From: jeremylt Date: Mon, 26 Apr 2021 15:02:41 -0600 Subject: [PATCH 01/11] vec - add CeedVectorScale --- backends/cuda/ceed-cuda-vector.c | 53 +++++++++++++++++++++++++ backends/cuda/kernels/cuda-vector.cu | 26 ++++++++++++ backends/hip/ceed-hip-vector.c | 53 +++++++++++++++++++++++++ backends/hip/kernels/hip-vector.hip.cpp | 25 ++++++++++++ include/ceed-impl.h | 1 + include/ceed/ceed.h | 1 + interface/ceed-vector.c | 30 ++++++++++++++ interface/ceed.c | 1 + python/ceed_vector.py | 10 +++++ python/tests/test-1-vector.py | 19 +++++++++ tests/t123-vector.c | 36 +++++++++++++++++ 11 files changed, 255 insertions(+) create mode 100644 tests/t123-vector.c diff --git a/backends/cuda/ceed-cuda-vector.c b/backends/cuda/ceed-cuda-vector.c index 823c8d21fb..b9e39dc1aa 100644 --- a/backends/cuda/ceed-cuda-vector.c +++ b/backends/cuda/ceed-cuda-vector.c @@ -450,6 +450,57 @@ static int CeedVectorReciprocal_Cuda(CeedVector vec) { return CEED_ERROR_SUCCESS; } +//------------------------------------------------------------------------------ +// Compute x = alpha x on the host +//------------------------------------------------------------------------------ +static int CeedHostScale_Cuda(CeedScalar *x_array, CeedScalar alpha, + CeedInt length) { + for (int i = 0; i < length; i++) + x_array[i] *= alpha; + return CEED_ERROR_SUCCESS; +} + +//------------------------------------------------------------------------------ +// Compute x = alpha x on device (impl in .cu file) +//------------------------------------------------------------------------------ +int CeedDeviceScale_Cuda(CeedScalar *x_array, CeedScalar alpha, + CeedInt length); + +//------------------------------------------------------------------------------ +// Compute x = alpha x +//------------------------------------------------------------------------------ +static int CeedVectorScale_Cuda(CeedVector x, CeedScalar alpha) { + int ierr; + Ceed ceed; + ierr = CeedVectorGetCeed(x, &ceed); CeedChkBackend(ierr); + CeedVector_Cuda *x_data; + ierr = CeedVectorGetData(x, &x_data); CeedChkBackend(ierr); + CeedInt length; + ierr = CeedVectorGetLength(x, &length); CeedChkBackend(ierr); + + // Set value for synced device/host array + switch(x_data->memState) { + case CEED_CUDA_HOST_SYNC: + ierr = CeedHostScale_Cuda(x_data->h_array, alpha, length); + CeedChkBackend(ierr); + break; + case CEED_CUDA_DEVICE_SYNC: + ierr = CeedDeviceScale_Cuda(x_data->d_array, alpha, length); + CeedChkBackend(ierr); + break; + case CEED_CUDA_BOTH_SYNC: + ierr = CeedDeviceScale_Cuda(x_data->d_array, alpha, length); + CeedChkBackend(ierr); + x_data->memState = CEED_CUDA_DEVICE_SYNC; + break; + // LCOV_EXCL_START + case CEED_CUDA_NONE_SYNC: + break; // Not possible, but included for completness + // LCOV_EXCL_STOP + } + return CEED_ERROR_SUCCESS; +} + //------------------------------------------------------------------------------ // Compute y = alpha x + y on the host //------------------------------------------------------------------------------ @@ -617,6 +668,8 @@ int CeedVectorCreate_Cuda(CeedInt n, CeedVector vec) { CeedVectorReciprocal_Cuda); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "AXPY", CeedVectorAXPY_Cuda); CeedChkBackend(ierr); + ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Scale", + CeedVectorScale_Cuda); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "PointwiseMult", CeedVectorPointwiseMult_Cuda); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Destroy", diff --git a/backends/cuda/kernels/cuda-vector.cu b/backends/cuda/kernels/cuda-vector.cu index 0a85cfdd41..c8aa916234 100644 --- a/backends/cuda/kernels/cuda-vector.cu +++ b/backends/cuda/kernels/cuda-vector.cu @@ -68,6 +68,32 @@ extern "C" int CeedDeviceReciprocal_Cuda(CeedScalar* d_array, CeedInt length) { return 0; } +//------------------------------------------------------------------------------ +// Kernel for scale +//------------------------------------------------------------------------------ +__global__ static void scaleValueK(CeedScalar * __restrict__ x, CeedScalar alpha, + CeedInt size) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= size) + return; + x[idx] *= alpha; +} + +//------------------------------------------------------------------------------ +// Compute x = alpha x on device +//------------------------------------------------------------------------------ +extern "C" int CeedDeviceScale_Cuda(CeedScalar *x_array, CeedScalar alpha, + CeedInt length) { + const int bsize = 512; + const int vecsize = length; + int gridsize = vecsize / bsize; + + if (bsize * gridsize < vecsize) + gridsize += 1; + scaleValueK<<>>(x_array, alpha, length); + return 0; +} + //------------------------------------------------------------------------------ // Kernel for axpy //------------------------------------------------------------------------------ diff --git a/backends/hip/ceed-hip-vector.c b/backends/hip/ceed-hip-vector.c index f830141ca8..faea62ece8 100644 --- a/backends/hip/ceed-hip-vector.c +++ b/backends/hip/ceed-hip-vector.c @@ -446,6 +446,57 @@ static int CeedVectorReciprocal_Hip(CeedVector vec) { return CEED_ERROR_SUCCESS; } +//------------------------------------------------------------------------------ +// Compute x = alpha x on the host +//------------------------------------------------------------------------------ +static int CeedHostScale_Hip(CeedScalar *x_array, CeedScalar alpha, + CeedInt length) { + for (int i = 0; i < length; i++) + x_array[i] *= alpha; + return CEED_ERROR_SUCCESS; +} + +//------------------------------------------------------------------------------ +// Compute x = alpha x on device (impl in .cu file) +//------------------------------------------------------------------------------ +int CeedDeviceScale_Hip(CeedScalar *x_array, CeedScalar alpha, + CeedInt length); + +//------------------------------------------------------------------------------ +// Compute x = alpha x +//------------------------------------------------------------------------------ +static int CeedVectorScale_Hip(CeedVector x, CeedScalar alpha) { + int ierr; + Ceed ceed; + ierr = CeedVectorGetCeed(x, &ceed); CeedChkBackend(ierr); + CeedVector_Hip *x_data; + ierr = CeedVectorGetData(x, &x_data); CeedChkBackend(ierr); + CeedInt length; + ierr = CeedVectorGetLength(x, &length); CeedChkBackend(ierr); + + // Set value for synced device/host array + switch(x_data->memState) { + case CEED_HIP_HOST_SYNC: + ierr = CeedHostScale_Hip(x_data->h_array, alpha, length); + CeedChkBackend(ierr); + break; + case CEED_HIP_DEVICE_SYNC: + ierr = CeedDeviceScale_Hip(x_data->d_array, alpha, length); + CeedChkBackend(ierr); + break; + case CEED_HIP_BOTH_SYNC: + ierr = CeedDeviceScale_Hip(x_data->d_array, alpha, length); + CeedChkBackend(ierr); + x_data->memState = CEED_HIP_DEVICE_SYNC; + break; + // LCOV_EXCL_START + case CEED_HIP_NONE_SYNC: + break; // Not possible, but included for completness + // LCOV_EXCL_STOP + } + return CEED_ERROR_SUCCESS; +} + //------------------------------------------------------------------------------ // Compute y = alpha x + y on the host //------------------------------------------------------------------------------ @@ -611,6 +662,8 @@ int CeedVectorCreate_Hip(CeedInt n, CeedVector vec) { CeedVectorNorm_Hip); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Reciprocal", CeedVectorReciprocal_Hip); CeedChkBackend(ierr); + ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Scale", + CeedVectorScale_Hip); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "AXPY", CeedVectorAXPY_Hip); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "Vector", vec, "PointwiseMult", diff --git a/backends/hip/kernels/hip-vector.hip.cpp b/backends/hip/kernels/hip-vector.hip.cpp index d3c8eee949..fe3c6cb38c 100644 --- a/backends/hip/kernels/hip-vector.hip.cpp +++ b/backends/hip/kernels/hip-vector.hip.cpp @@ -68,7 +68,32 @@ extern "C" int CeedDeviceReciprocal_Hip(CeedScalar* d_array, CeedInt length) { return 0; } +//------------------------------------------------------------------------------ +// Kernel for scale +//------------------------------------------------------------------------------ +__global__ static void scaleValueK(CeedScalar * __restrict__ x, CeedScalar alpha, + CeedInt size) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= size) + return; + x[idx] *= alpha; +} + +//------------------------------------------------------------------------------ +// Compute x = alpha x on device +//------------------------------------------------------------------------------ +extern "C" int CeedDeviceScale_Hip(CeedScalar *x_array, CeedScalar alpha, + CeedInt length) { + const int bsize = 512; + const int vecsize = length; + int gridsize = vecsize / bsize; + if (bsize * gridsize < vecsize) + gridsize += 1; + hipLaunchKernelGGL(scaleValueK, dim3(gridsize), dim3(bsize), 0, 0, x_array, alpha, + length); + return 0; +} //------------------------------------------------------------------------------ // Kernel for axpy diff --git a/include/ceed-impl.h b/include/ceed-impl.h index 65c96bd7f0..95a83d0764 100644 --- a/include/ceed-impl.h +++ b/include/ceed-impl.h @@ -138,6 +138,7 @@ struct CeedVector_private { int (*RestoreArray)(CeedVector); int (*RestoreArrayRead)(CeedVector); int (*Norm)(CeedVector, CeedNormType, CeedScalar *); + int (*Scale)(CeedVector, CeedScalar); int (*AXPY)(CeedVector, CeedScalar, CeedVector); int (*PointwiseMult)(CeedVector, CeedVector, CeedVector); int (*Reciprocal)(CeedVector); diff --git a/include/ceed/ceed.h b/include/ceed/ceed.h index d98bed3eec..22cce7150e 100644 --- a/include/ceed/ceed.h +++ b/include/ceed/ceed.h @@ -321,6 +321,7 @@ CEED_EXTERN int CeedVectorRestoreArrayRead(CeedVector vec, const CeedScalar **array); CEED_EXTERN int CeedVectorNorm(CeedVector vec, CeedNormType type, CeedScalar *norm); +CEED_EXTERN int CeedVectorScale(CeedVector x, CeedScalar alpha); CEED_EXTERN int CeedVectorAXPY(CeedVector y, CeedScalar alpha, CeedVector x); CEED_EXTERN int CeedVectorPointwiseMult(CeedVector w, CeedVector x, CeedVector y); CEED_EXTERN int CeedVectorReciprocal(CeedVector vec); diff --git a/interface/ceed-vector.c b/interface/ceed-vector.c index 07e8a5e495..b2f7edf976 100644 --- a/interface/ceed-vector.c +++ b/interface/ceed-vector.c @@ -531,6 +531,36 @@ int CeedVectorNorm(CeedVector vec, CeedNormType norm_type, CeedScalar *norm) { return CEED_ERROR_SUCCESS; } +/** + @brief Compute x = alpha x + + @param x[in,out] vector for scaling + @param alpha[in] scaling factor + + @return An error code: 0 - success, otherwise - failure + + @ref User +**/ +int CeedVectorScale(CeedVector x, CeedScalar alpha) { + int ierr; + CeedScalar *x_array; + CeedInt n_x; + + ierr = CeedVectorGetLength(x, &n_x); CeedChk(ierr); + + // Backend implementation + if (x->Scale) + return x->Scale(x, alpha); + + // Default implementation + ierr = CeedVectorGetArray(x, CEED_MEM_HOST, &x_array); CeedChk(ierr); + for (CeedInt i=0; i +#include + +int main(int argc, char **argv) { + Ceed ceed; + CeedVector x; + CeedInt n; + CeedScalar a[10]; + const CeedScalar *b; + + CeedInit(argv[1], &ceed); + + n = 10; + CeedVectorCreate(ceed, n, &x); + for (CeedInt i=0; i 1e-14) + // LCOV_EXCL_START + printf("Error in alpha x, computed: %f actual: %f\n", b[i], + -(10.0 + i)/2); + // LCOV_EXCL_STOP + CeedVectorRestoreArrayRead(x, &b); + + CeedVectorDestroy(&x); + CeedDestroy(&ceed); + return 0; +} From 126de23b4b614ff58f7bb7ce9f7a490dbe2d3a7b Mon Sep 17 00:00:00 2001 From: jeremylt Date: Tue, 27 Apr 2021 08:40:44 -0600 Subject: [PATCH 02/11] python - minor test update --- python/tests/test-1-vector.py | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/python/tests/test-1-vector.py b/python/tests/test-1-vector.py index d6e7ea282f..7f2b9b8286 100644 --- a/python/tests/test-1-vector.py +++ b/python/tests/test-1-vector.py @@ -248,8 +248,7 @@ def test_121(ceed_resource, capsys): y.axpy(-0.5, x) with y.array() as b: - for i in range(len(b)): - assert abs(b[i] - (10 + i) / 2) < 1e-14 + assert np.allclose(.5 * a, b) # ------------------------------------------------------------------------------- # Test pointwise multiplication @@ -305,8 +304,7 @@ def test_123(ceed_resource, capsys): x.scale(-0.5) with x.array() as b: - for i in range(len(b)): - assert abs(b[i] + (10 + i) / 2) < 1e-14 + assert np.allclose(-.5 * a, b) # ------------------------------------------------------------------------------- # Test modification of reshaped array From 74653b0dfd9ea5453a7e0d931542e0b08b318f8b Mon Sep 17 00:00:00 2001 From: jeremylt Date: Tue, 27 Apr 2021 09:14:04 -0600 Subject: [PATCH 03/11] rust - add vector convenience methods --- rust/src/vector.rs | 117 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 117 insertions(+) diff --git a/rust/src/vector.rs b/rust/src/vector.rs index 527c3d2d63..bc0350749f 100644 --- a/rust/src/vector.rs +++ b/rust/src/vector.rs @@ -349,6 +349,123 @@ impl<'a> Vector<'a> { self.ceed.check_error(ierr)?; Ok(res) } + + /// Compute x = alpha x for a CeedVector + /// + /// # arguments + /// + /// * `alpha` - scaling factor + /// + /// ``` + /// # use libceed::prelude::*; + /// # let ceed = libceed::Ceed::default_init(); + /// let mut vec = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// + /// vec = vec.scale(-1.0).unwrap(); + /// vec.view().iter().enumerate().for_each(|(i, &v)| { + /// assert_eq!(v, -(i as f64), "Value not set correctly"); + /// }); + /// ``` + #[allow(unused_mut)] + pub fn scale(mut self, alpha: f64) -> crate::Result { + let ierr = unsafe { bind_ceed::CeedVectorScale(self.ptr, alpha) }; + self.ceed.check_error(ierr)?; + Ok(self) + } + + /// Compute y = alpha x + y for a pair of CeedVectors + /// + /// # arguments + /// + /// * `alpha` - scaling factor + /// * `x` - second vector, must be different than self + /// + /// ``` + /// # use libceed::prelude::*; + /// # let ceed = libceed::Ceed::default_init(); + /// let x = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// let mut y = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// + /// y = y.axpy(-0.5, &x).unwrap(); + /// y.view().iter().enumerate().for_each(|(i, &v)| { + /// assert_eq!(v, (i as f64) / 2.0, "Value not set correctly"); + /// }); + /// ``` + #[allow(unused_mut)] + pub fn axpy(mut self, alpha: f64, x: &crate::Vector) -> crate::Result { + let ierr = unsafe { bind_ceed::CeedVectorAXPY(self.ptr, alpha, x.ptr) }; + self.ceed.check_error(ierr)?; + Ok(self) + } + + /// Compute the pointwise multiplication w = x .* y for CeedVectors + /// + /// # arguments + /// + /// * `x` - first vector for product + /// * `y` - second vector for product + /// + /// ``` + /// # use libceed::prelude::*; + /// # let ceed = libceed::Ceed::default_init(); + /// let mut w = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// let x = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// let y = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// + /// w = w.pointwise_mult(&x, &y).unwrap(); + /// w.view().iter().enumerate().for_each(|(i, &v)| { + /// assert_eq!(v, (i as f64).powf(2.0), "Value not set correctly"); + /// }); + /// ``` + #[allow(unused_mut)] + pub fn pointwise_mult(mut self, x: &crate::Vector, y: &crate::Vector) -> crate::Result { + let ierr = unsafe { bind_ceed::CeedVectorPointwiseMult(self.ptr, x.ptr, y.ptr) }; + self.ceed.check_error(ierr)?; + Ok(self) + } + + /// Compute the pointwise multiplication w = w .* x for CeedVectors + /// + /// # arguments + /// + /// * `x` - second vector for product + /// + /// ``` + /// # use libceed::prelude::*; + /// # let ceed = libceed::Ceed::default_init(); + /// let mut w = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// let x = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// + /// w = w.pointwise_scale(&x).unwrap(); + /// w.view().iter().enumerate().for_each(|(i, &v)| { + /// assert_eq!(v, (i as f64).powf(2.0), "Value not set correctly"); + /// }); + /// ``` + #[allow(unused_mut)] + pub fn pointwise_scale(mut self, x: &crate::Vector) -> crate::Result { + let ierr = unsafe { bind_ceed::CeedVectorPointwiseMult(self.ptr, self.ptr, x.ptr) }; + self.ceed.check_error(ierr)?; + Ok(self) + } + + /// Compute the pointwise multiplication w = w .* w for a CeedVector + /// + /// ``` + /// # use libceed::prelude::*; + /// # let ceed = libceed::Ceed::default_init(); + /// let mut w = ceed.vector_from_slice(&[0., 1., 2., 3., 4.]).unwrap(); + /// + /// w = w.pointwise_square().unwrap(); + /// w.view().iter().enumerate().for_each(|(i, &v)| { + /// assert_eq!(v, (i as f64).powf(2.0), "Value not set correctly"); + /// }); + /// ``` + #[allow(unused_mut)] + pub fn pointwise_square(mut self) -> crate::Result { + let ierr = unsafe { bind_ceed::CeedVectorPointwiseMult(self.ptr, self.ptr, self.ptr) }; + self.ceed.check_error(ierr)?; + Ok(self) + } } // ----------------------------------------------------------------------------- From 2ac872fe09e35217845f7cccbe5121b5cf6a6a3d Mon Sep 17 00:00:00 2001 From: jeremylt Date: Tue, 27 Apr 2021 09:32:27 -0600 Subject: [PATCH 04/11] actions - add rustfmt component --- .github/workflows/rust-test-with-style.yml | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/.github/workflows/rust-test-with-style.yml b/.github/workflows/rust-test-with-style.yml index 9cf84525fd..0bea5b4629 100644 --- a/.github/workflows/rust-test-with-style.yml +++ b/.github/workflows/rust-test-with-style.yml @@ -48,8 +48,10 @@ jobs: - name: Rust setup uses: actions-rs/toolchain@v1 with: - # Note: rustfmt not always included in nightly - toolchain: stable + # Note: rustfmt not always included in nightly, will attempt to downgrade until rustfmt found + toolchain: nightly + components: rustfmt - name: Rust style run: | - cargo fmt -- --check + cargo +nightly fmt --version + cargo +nightly fmt -- --check From 22785960d60f45e0c65c47523fd8ab76bf9537ed Mon Sep 17 00:00:00 2001 From: Will Pazner Date: Tue, 27 Apr 2021 12:40:11 -0700 Subject: [PATCH 05/11] [Julia] Update generated API --- julia/LibCEED.jl/src/generated/libceed_api.jl | 428 +++++++++++------- .../src/generated/libceed_common.jl | 25 +- 2 files changed, 288 insertions(+), 165 deletions(-) diff --git a/julia/LibCEED.jl/src/generated/libceed_api.jl b/julia/LibCEED.jl/src/generated/libceed_api.jl index a647d5a434..70f6bf660b 100644 --- a/julia/LibCEED.jl/src/generated/libceed_api.jl +++ b/julia/LibCEED.jl/src/generated/libceed_api.jl @@ -2,16 +2,24 @@ # Automatically generated using Clang.jl #! format: off +function CeedRegistryGetList(n, resources, array) + ccall((:CeedRegistryGetList, libceed), Cint, (Ptr{Csize_t}, Ptr{Ptr{Cstring}}, Ptr{Ptr{CeedInt}}), n, resources, array) +end + function CeedInit(resource, ceed) ccall((:CeedInit, libceed), Cint, (Cstring, Ptr{Ceed}), resource, ceed) end +function CeedReferenceCopy(ceed, ceed_copy) + ccall((:CeedReferenceCopy, libceed), Cint, (Ceed, Ptr{Ceed}), ceed, ceed_copy) +end + function CeedGetResource(ceed, resource) ccall((:CeedGetResource, libceed), Cint, (Ceed, Ptr{Cstring}), ceed, resource) end -function CeedIsDeterministic(ceed, isDeterministic) - ccall((:CeedIsDeterministic, libceed), Cint, (Ceed, Ptr{Bool}), ceed, isDeterministic) +function CeedIsDeterministic(ceed, is_deterministic) + ccall((:CeedIsDeterministic, libceed), Cint, (Ceed, Ptr{Bool}), ceed, is_deterministic) end function CeedView(ceed, stream) @@ -39,15 +47,19 @@ function CeedErrorExit(arg1, arg2, arg3, arg4, arg5, arg6, arg7) end function CeedSetErrorHandler(ceed, eh) - ccall((:CeedSetErrorHandler, libceed), Cint, (Ceed, Ptr{Cvoid}), ceed, eh) + ccall((:CeedSetErrorHandler, libceed), Cint, (Ceed, CeedErrorHandler), ceed, eh) end -function CeedGetErrorMessage(arg1, errmsg) - ccall((:CeedGetErrorMessage, libceed), Cint, (Ceed, Ptr{Cstring}), arg1, errmsg) +function CeedGetErrorMessage(arg1, err_msg) + ccall((:CeedGetErrorMessage, libceed), Cint, (Ceed, Ptr{Cstring}), arg1, err_msg) end -function CeedResetErrorMessage(arg1, errmsg) - ccall((:CeedResetErrorMessage, libceed), Cint, (Ceed, Ptr{Cstring}), arg1, errmsg) +function CeedResetErrorMessage(arg1, err_msg) + ccall((:CeedResetErrorMessage, libceed), Cint, (Ceed, Ptr{Cstring}), arg1, err_msg) +end + +function CeedGetVersion(major, minor, patch, release) + ccall((:CeedGetVersion, libceed), Cint, (Ptr{Cint}, Ptr{Cint}, Ptr{Cint}, Ptr{Bool}), major, minor, patch, release) end function CeedGetPreferredMemType(ceed, type) @@ -58,28 +70,32 @@ function CeedVectorCreate(ceed, len, vec) ccall((:CeedVectorCreate, libceed), Cint, (Ceed, CeedInt, Ptr{CeedVector}), ceed, len, vec) end -function CeedVectorSetArray(vec, mtype, cmode, array) - ccall((:CeedVectorSetArray, libceed), Cint, (CeedVector, CeedMemType, CeedCopyMode, Ptr{CeedScalar}), vec, mtype, cmode, array) +function CeedVectorReferenceCopy(vec, vec_copy) + ccall((:CeedVectorReferenceCopy, libceed), Cint, (CeedVector, Ptr{CeedVector}), vec, vec_copy) +end + +function CeedVectorSetArray(vec, mem_type, copy_mode, array) + ccall((:CeedVectorSetArray, libceed), Cint, (CeedVector, CeedMemType, CeedCopyMode, Ptr{CeedScalar}), vec, mem_type, copy_mode, array) end function CeedVectorSetValue(vec, value) ccall((:CeedVectorSetValue, libceed), Cint, (CeedVector, CeedScalar), vec, value) end -function CeedVectorSyncArray(vec, mtype) - ccall((:CeedVectorSyncArray, libceed), Cint, (CeedVector, CeedMemType), vec, mtype) +function CeedVectorSyncArray(vec, mem_type) + ccall((:CeedVectorSyncArray, libceed), Cint, (CeedVector, CeedMemType), vec, mem_type) end -function CeedVectorTakeArray(vec, mtype, array) - ccall((:CeedVectorTakeArray, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mtype, array) +function CeedVectorTakeArray(vec, mem_type, array) + ccall((:CeedVectorTakeArray, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mem_type, array) end -function CeedVectorGetArray(vec, mtype, array) - ccall((:CeedVectorGetArray, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mtype, array) +function CeedVectorGetArray(vec, mem_type, array) + ccall((:CeedVectorGetArray, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mem_type, array) end -function CeedVectorGetArrayRead(vec, mtype, array) - ccall((:CeedVectorGetArrayRead, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mtype, array) +function CeedVectorGetArrayRead(vec, mem_type, array) + ccall((:CeedVectorGetArrayRead, libceed), Cint, (CeedVector, CeedMemType, Ptr{Ptr{CeedScalar}}), vec, mem_type, array) end function CeedVectorRestoreArray(vec, array) @@ -94,12 +110,24 @@ function CeedVectorNorm(vec, type, norm) ccall((:CeedVectorNorm, libceed), Cint, (CeedVector, CeedNormType, Ptr{CeedScalar}), vec, type, norm) end +function CeedVectorScale(x, alpha) + ccall((:CeedVectorScale, libceed), Cint, (CeedVector, CeedScalar), x, alpha) +end + +function CeedVectorAXPY(y, alpha, x) + ccall((:CeedVectorAXPY, libceed), Cint, (CeedVector, CeedScalar, CeedVector), y, alpha, x) +end + +function CeedVectorPointwiseMult(w, x, y) + ccall((:CeedVectorPointwiseMult, libceed), Cint, (CeedVector, CeedVector, CeedVector), w, x, y) +end + function CeedVectorReciprocal(vec) ccall((:CeedVectorReciprocal, libceed), Cint, (CeedVector,), vec) end -function CeedVectorView(vec, fpfmt, stream) - ccall((:CeedVectorView, libceed), Cint, (CeedVector, Cstring, Ptr{FILE}), vec, fpfmt, stream) +function CeedVectorView(vec, fp_fmt, stream) + ccall((:CeedVectorView, libceed), Cint, (CeedVector, Cstring, Ptr{FILE}), vec, fp_fmt, stream) end function CeedVectorGetLength(vec, length) @@ -114,60 +142,64 @@ function CeedRequestWait(req) ccall((:CeedRequestWait, libceed), Cint, (Ptr{CeedRequest},), req) end -function CeedElemRestrictionCreate(ceed, nelem, elemsize, ncomp, compstride, lsize, mtype, cmode, offsets, rstr) - ccall((:CeedElemRestrictionCreate, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedMemType, CeedCopyMode, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, nelem, elemsize, ncomp, compstride, lsize, mtype, cmode, offsets, rstr) +function CeedElemRestrictionCreate(ceed, num_elem, elem_size, num_comp, comp_stride, l_size, mem_type, copy_mode, offsets, rstr) + ccall((:CeedElemRestrictionCreate, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedMemType, CeedCopyMode, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, num_elem, elem_size, num_comp, comp_stride, l_size, mem_type, copy_mode, offsets, rstr) +end + +function CeedElemRestrictionCreateStrided(ceed, num_elem, elem_size, num_comp, l_size, strides, rstr) + ccall((:CeedElemRestrictionCreateStrided, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, num_elem, elem_size, num_comp, l_size, strides, rstr) end -function CeedElemRestrictionCreateStrided(ceed, nelem, elemsize, ncomp, lsize, strides, rstr) - ccall((:CeedElemRestrictionCreateStrided, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, nelem, elemsize, ncomp, lsize, strides, rstr) +function CeedElemRestrictionCreateBlocked(ceed, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size, mem_type, copy_mode, offsets, rstr) + ccall((:CeedElemRestrictionCreateBlocked, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedMemType, CeedCopyMode, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, num_elem, elem_size, blk_size, num_comp, comp_stride, l_size, mem_type, copy_mode, offsets, rstr) end -function CeedElemRestrictionCreateBlocked(ceed, nelem, elemsize, blksize, ncomp, compstride, lsize, mtype, cmode, offsets, rstr) - ccall((:CeedElemRestrictionCreateBlocked, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, CeedMemType, CeedCopyMode, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, nelem, elemsize, blksize, ncomp, compstride, lsize, mtype, cmode, offsets, rstr) +function CeedElemRestrictionCreateBlockedStrided(ceed, num_elem, elem_size, blk_size, num_comp, l_size, strides, rstr) + ccall((:CeedElemRestrictionCreateBlockedStrided, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, num_elem, elem_size, blk_size, num_comp, l_size, strides, rstr) end -function CeedElemRestrictionCreateBlockedStrided(ceed, nelem, elemsize, blksize, ncomp, lsize, strides, rstr) - ccall((:CeedElemRestrictionCreateBlockedStrided, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedInt}, Ptr{CeedElemRestriction}), ceed, nelem, elemsize, blksize, ncomp, lsize, strides, rstr) +function CeedElemRestrictionReferenceCopy(rstr, rstr_copy) + ccall((:CeedElemRestrictionReferenceCopy, libceed), Cint, (CeedElemRestriction, Ptr{CeedElemRestriction}), rstr, rstr_copy) end function CeedElemRestrictionCreateVector(rstr, lvec, evec) ccall((:CeedElemRestrictionCreateVector, libceed), Cint, (CeedElemRestriction, Ptr{CeedVector}, Ptr{CeedVector}), rstr, lvec, evec) end -function CeedElemRestrictionApply(rstr, tmode, u, ru, request) - ccall((:CeedElemRestrictionApply, libceed), Cint, (CeedElemRestriction, CeedTransposeMode, CeedVector, CeedVector, Ptr{CeedRequest}), rstr, tmode, u, ru, request) +function CeedElemRestrictionApply(rstr, t_mode, u, ru, request) + ccall((:CeedElemRestrictionApply, libceed), Cint, (CeedElemRestriction, CeedTransposeMode, CeedVector, CeedVector, Ptr{CeedRequest}), rstr, t_mode, u, ru, request) end -function CeedElemRestrictionApplyBlock(rstr, block, tmode, u, ru, request) - ccall((:CeedElemRestrictionApplyBlock, libceed), Cint, (CeedElemRestriction, CeedInt, CeedTransposeMode, CeedVector, CeedVector, Ptr{CeedRequest}), rstr, block, tmode, u, ru, request) +function CeedElemRestrictionApplyBlock(rstr, block, t_mode, u, ru, request) + ccall((:CeedElemRestrictionApplyBlock, libceed), Cint, (CeedElemRestriction, CeedInt, CeedTransposeMode, CeedVector, CeedVector, Ptr{CeedRequest}), rstr, block, t_mode, u, ru, request) end -function CeedElemRestrictionGetCompStride(rstr, compstride) - ccall((:CeedElemRestrictionGetCompStride, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, compstride) +function CeedElemRestrictionGetCompStride(rstr, comp_stride) + ccall((:CeedElemRestrictionGetCompStride, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, comp_stride) end -function CeedElemRestrictionGetNumElements(rstr, numelem) - ccall((:CeedElemRestrictionGetNumElements, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, numelem) +function CeedElemRestrictionGetNumElements(rstr, num_elem) + ccall((:CeedElemRestrictionGetNumElements, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, num_elem) end -function CeedElemRestrictionGetElementSize(rstr, elemsize) - ccall((:CeedElemRestrictionGetElementSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, elemsize) +function CeedElemRestrictionGetElementSize(rstr, elem_size) + ccall((:CeedElemRestrictionGetElementSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, elem_size) end -function CeedElemRestrictionGetLVectorSize(rstr, lsize) - ccall((:CeedElemRestrictionGetLVectorSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, lsize) +function CeedElemRestrictionGetLVectorSize(rstr, l_size) + ccall((:CeedElemRestrictionGetLVectorSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, l_size) end -function CeedElemRestrictionGetNumComponents(rstr, numcomp) - ccall((:CeedElemRestrictionGetNumComponents, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, numcomp) +function CeedElemRestrictionGetNumComponents(rstr, num_comp) + ccall((:CeedElemRestrictionGetNumComponents, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, num_comp) end -function CeedElemRestrictionGetNumBlocks(rstr, numblk) - ccall((:CeedElemRestrictionGetNumBlocks, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, numblk) +function CeedElemRestrictionGetNumBlocks(rstr, num_blk) + ccall((:CeedElemRestrictionGetNumBlocks, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, num_blk) end -function CeedElemRestrictionGetBlockSize(rstr, blksize) - ccall((:CeedElemRestrictionGetBlockSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, blksize) +function CeedElemRestrictionGetBlockSize(rstr, blk_size) + ccall((:CeedElemRestrictionGetBlockSize, libceed), Cint, (CeedElemRestriction, Ptr{CeedInt}), rstr, blk_size) end function CeedElemRestrictionGetMultiplicity(rstr, mult) @@ -182,24 +214,28 @@ function CeedElemRestrictionDestroy(rstr) ccall((:CeedElemRestrictionDestroy, libceed), Cint, (Ptr{CeedElemRestriction},), rstr) end -function CeedBasisCreateTensorH1Lagrange(ceed, dim, ncomp, P, Q, qmode, basis) - ccall((:CeedBasisCreateTensorH1Lagrange, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedQuadMode, Ptr{CeedBasis}), ceed, dim, ncomp, P, Q, qmode, basis) +function CeedBasisCreateTensorH1Lagrange(ceed, dim, num_comp, P, Q, quad_mode, basis) + ccall((:CeedBasisCreateTensorH1Lagrange, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, CeedQuadMode, Ptr{CeedBasis}), ceed, dim, num_comp, P, Q, quad_mode, basis) end -function CeedBasisCreateTensorH1(ceed, dim, ncomp, P1d, Q1d, interp1d, grad1d, qref1d, qweight1d, basis) - ccall((:CeedBasisCreateTensorH1, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedBasis}), ceed, dim, ncomp, P1d, Q1d, interp1d, grad1d, qref1d, qweight1d, basis) +function CeedBasisCreateTensorH1(ceed, dim, num_comp, P_1d, Q_1d, interp_1d, grad_1d, q_ref_1d, q_weight_1d, basis) + ccall((:CeedBasisCreateTensorH1, libceed), Cint, (Ceed, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedBasis}), ceed, dim, num_comp, P_1d, Q_1d, interp_1d, grad_1d, q_ref_1d, q_weight_1d, basis) end -function CeedBasisCreateH1(ceed, topo, ncomp, nnodes, nqpts, interp, grad, qref, qweight, basis) - ccall((:CeedBasisCreateH1, libceed), Cint, (Ceed, CeedElemTopology, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedBasis}), ceed, topo, ncomp, nnodes, nqpts, interp, grad, qref, qweight, basis) +function CeedBasisCreateH1(ceed, topo, num_comp, num_nodes, nqpts, interp, grad, q_ref, q_weights, basis) + ccall((:CeedBasisCreateH1, libceed), Cint, (Ceed, CeedElemTopology, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedBasis}), ceed, topo, num_comp, num_nodes, nqpts, interp, grad, q_ref, q_weights, basis) +end + +function CeedBasisReferenceCopy(basis, basis_copy) + ccall((:CeedBasisReferenceCopy, libceed), Cint, (CeedBasis, Ptr{CeedBasis}), basis, basis_copy) end function CeedBasisView(basis, stream) ccall((:CeedBasisView, libceed), Cint, (CeedBasis, Ptr{FILE}), basis, stream) end -function CeedBasisApply(basis, nelem, tmode, emode, u, v) - ccall((:CeedBasisApply, libceed), Cint, (CeedBasis, CeedInt, CeedTransposeMode, CeedEvalMode, CeedVector, CeedVector), basis, nelem, tmode, emode, u, v) +function CeedBasisApply(basis, num_elem, t_mode, eval_mode, u, v) + ccall((:CeedBasisApply, libceed), Cint, (CeedBasis, CeedInt, CeedTransposeMode, CeedEvalMode, CeedVector, CeedVector), basis, num_elem, t_mode, eval_mode, u, v) end function CeedBasisGetDimension(basis, dim) @@ -210,60 +246,60 @@ function CeedBasisGetTopology(basis, topo) ccall((:CeedBasisGetTopology, libceed), Cint, (CeedBasis, Ptr{CeedElemTopology}), basis, topo) end -function CeedBasisGetNumComponents(basis, numcomp) - ccall((:CeedBasisGetNumComponents, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, numcomp) +function CeedBasisGetNumComponents(basis, num_comp) + ccall((:CeedBasisGetNumComponents, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, num_comp) end function CeedBasisGetNumNodes(basis, P) ccall((:CeedBasisGetNumNodes, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, P) end -function CeedBasisGetNumNodes1D(basis, P1d) - ccall((:CeedBasisGetNumNodes1D, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, P1d) +function CeedBasisGetNumNodes1D(basis, P_1d) + ccall((:CeedBasisGetNumNodes1D, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, P_1d) end function CeedBasisGetNumQuadraturePoints(basis, Q) ccall((:CeedBasisGetNumQuadraturePoints, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, Q) end -function CeedBasisGetNumQuadraturePoints1D(basis, Q1d) - ccall((:CeedBasisGetNumQuadraturePoints1D, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, Q1d) +function CeedBasisGetNumQuadraturePoints1D(basis, Q_1d) + ccall((:CeedBasisGetNumQuadraturePoints1D, libceed), Cint, (CeedBasis, Ptr{CeedInt}), basis, Q_1d) end -function CeedBasisGetQRef(basis, qref) - ccall((:CeedBasisGetQRef, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, qref) +function CeedBasisGetQRef(basis, q_ref) + ccall((:CeedBasisGetQRef, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, q_ref) end -function CeedBasisGetQWeights(basis, qweight) - ccall((:CeedBasisGetQWeights, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, qweight) +function CeedBasisGetQWeights(basis, q_weights) + ccall((:CeedBasisGetQWeights, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, q_weights) end function CeedBasisGetInterp(basis, interp) ccall((:CeedBasisGetInterp, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, interp) end -function CeedBasisGetInterp1D(basis, interp1d) - ccall((:CeedBasisGetInterp1D, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, interp1d) +function CeedBasisGetInterp1D(basis, interp_1d) + ccall((:CeedBasisGetInterp1D, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, interp_1d) end function CeedBasisGetGrad(basis, grad) ccall((:CeedBasisGetGrad, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, grad) end -function CeedBasisGetGrad1D(basis, grad1d) - ccall((:CeedBasisGetGrad1D, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, grad1d) +function CeedBasisGetGrad1D(basis, grad_1d) + ccall((:CeedBasisGetGrad1D, libceed), Cint, (CeedBasis, Ptr{Ptr{CeedScalar}}), basis, grad_1d) end function CeedBasisDestroy(basis) ccall((:CeedBasisDestroy, libceed), Cint, (Ptr{CeedBasis},), basis) end -function CeedGaussQuadrature(Q, qref1d, qweight1d) - ccall((:CeedGaussQuadrature, libceed), Cint, (CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), Q, qref1d, qweight1d) +function CeedGaussQuadrature(Q, q_ref_1d, q_weight_1d) + ccall((:CeedGaussQuadrature, libceed), Cint, (CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), Q, q_ref_1d, q_weight_1d) end -function CeedLobattoQuadrature(Q, qref1d, qweight1d) - ccall((:CeedLobattoQuadrature, libceed), Cint, (CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), Q, qref1d, qweight1d) +function CeedLobattoQuadrature(Q, q_ref_1d, q_weight_1d) + ccall((:CeedLobattoQuadrature, libceed), Cint, (CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), Q, q_ref_1d, q_weight_1d) end function CeedQRFactorization(ceed, mat, tau, m, n) @@ -274,28 +310,32 @@ function CeedSymmetricSchurDecomposition(ceed, mat, lambda, n) ccall((:CeedSymmetricSchurDecomposition, libceed), Cint, (Ceed, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedInt), ceed, mat, lambda, n) end -function CeedSimultaneousDiagonalization(ceed, matA, matB, x, lambda, n) - ccall((:CeedSimultaneousDiagonalization, libceed), Cint, (Ceed, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedInt), ceed, matA, matB, x, lambda, n) +function CeedSimultaneousDiagonalization(ceed, mat_A, mat_B, x, lambda, n) + ccall((:CeedSimultaneousDiagonalization, libceed), Cint, (Ceed, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedInt), ceed, mat_A, mat_B, x, lambda, n) end -function CeedQFunctionCreateInterior(ceed, vlength, f, source, qf) - ccall((:CeedQFunctionCreateInterior, libceed), Cint, (Ceed, CeedInt, CeedQFunctionUser, Cstring, Ptr{CeedQFunction}), ceed, vlength, f, source, qf) +function CeedQFunctionCreateInterior(ceed, vec_length, f, source, qf) + ccall((:CeedQFunctionCreateInterior, libceed), Cint, (Ceed, CeedInt, CeedQFunctionUser, Cstring, Ptr{CeedQFunction}), ceed, vec_length, f, source, qf) end function CeedQFunctionCreateInteriorByName(ceed, name, qf) ccall((:CeedQFunctionCreateInteriorByName, libceed), Cint, (Ceed, Cstring, Ptr{CeedQFunction}), ceed, name, qf) end -function CeedQFunctionCreateIdentity(ceed, size, inmode, outmode, qf) - ccall((:CeedQFunctionCreateIdentity, libceed), Cint, (Ceed, CeedInt, CeedEvalMode, CeedEvalMode, Ptr{CeedQFunction}), ceed, size, inmode, outmode, qf) +function CeedQFunctionCreateIdentity(ceed, size, in_mode, out_mode, qf) + ccall((:CeedQFunctionCreateIdentity, libceed), Cint, (Ceed, CeedInt, CeedEvalMode, CeedEvalMode, Ptr{CeedQFunction}), ceed, size, in_mode, out_mode, qf) +end + +function CeedQFunctionReferenceCopy(qf, qf_copy) + ccall((:CeedQFunctionReferenceCopy, libceed), Cint, (CeedQFunction, Ptr{CeedQFunction}), qf, qf_copy) end -function CeedQFunctionAddInput(qf, fieldname, size, emode) - ccall((:CeedQFunctionAddInput, libceed), Cint, (CeedQFunction, Cstring, CeedInt, CeedEvalMode), qf, fieldname, size, emode) +function CeedQFunctionAddInput(qf, field_name, size, eval_mode) + ccall((:CeedQFunctionAddInput, libceed), Cint, (CeedQFunction, Cstring, CeedInt, CeedEvalMode), qf, field_name, size, eval_mode) end -function CeedQFunctionAddOutput(qf, fieldname, size, emode) - ccall((:CeedQFunctionAddOutput, libceed), Cint, (CeedQFunction, Cstring, CeedInt, CeedEvalMode), qf, fieldname, size, emode) +function CeedQFunctionAddOutput(qf, field_name, size, eval_mode) + ccall((:CeedQFunctionAddOutput, libceed), Cint, (CeedQFunction, Cstring, CeedInt, CeedEvalMode), qf, field_name, size, eval_mode) end function CeedQFunctionSetContext(qf, ctx) @@ -318,12 +358,16 @@ function CeedQFunctionContextCreate(ceed, ctx) ccall((:CeedQFunctionContextCreate, libceed), Cint, (Ceed, Ptr{CeedQFunctionContext}), ceed, ctx) end -function CeedQFunctionContextSetData(ctx, mtype, cmode, size, data) - ccall((:CeedQFunctionContextSetData, libceed), Cint, (CeedQFunctionContext, CeedMemType, CeedCopyMode, Csize_t, Ptr{Cvoid}), ctx, mtype, cmode, size, data) +function CeedQFunctionContextReferenceCopy(ctx, ctx_copy) + ccall((:CeedQFunctionContextReferenceCopy, libceed), Cint, (CeedQFunctionContext, Ptr{CeedQFunctionContext}), ctx, ctx_copy) +end + +function CeedQFunctionContextSetData(ctx, mem_type, copy_mode, size, data) + ccall((:CeedQFunctionContextSetData, libceed), Cint, (CeedQFunctionContext, CeedMemType, CeedCopyMode, Csize_t, Ptr{Cvoid}), ctx, mem_type, copy_mode, size, data) end -function CeedQFunctionContextGetData(ctx, mtype, data) - ccall((:CeedQFunctionContextGetData, libceed), Cint, (CeedQFunctionContext, CeedMemType, Ptr{Cvoid}), ctx, mtype, data) +function CeedQFunctionContextGetData(ctx, mem_type, data) + ccall((:CeedQFunctionContextGetData, libceed), Cint, (CeedQFunctionContext, CeedMemType, Ptr{Cvoid}), ctx, mem_type, data) end function CeedQFunctionContextRestoreData(ctx, data) @@ -346,12 +390,16 @@ function CeedCompositeOperatorCreate(ceed, op) ccall((:CeedCompositeOperatorCreate, libceed), Cint, (Ceed, Ptr{CeedOperator}), ceed, op) end -function CeedOperatorSetField(op, fieldname, r, b, v) - ccall((:CeedOperatorSetField, libceed), Cint, (CeedOperator, Cstring, CeedElemRestriction, CeedBasis, CeedVector), op, fieldname, r, b, v) +function CeedOperatorReferenceCopy(op, op_copy) + ccall((:CeedOperatorReferenceCopy, libceed), Cint, (CeedOperator, Ptr{CeedOperator}), op, op_copy) end -function CeedCompositeOperatorAddSub(compositeop, subop) - ccall((:CeedCompositeOperatorAddSub, libceed), Cint, (CeedOperator, CeedOperator), compositeop, subop) +function CeedOperatorSetField(op, field_name, r, b, v) + ccall((:CeedOperatorSetField, libceed), Cint, (CeedOperator, Cstring, CeedElemRestriction, CeedBasis, CeedVector), op, field_name, r, b, v) +end + +function CeedCompositeOperatorAddSub(composite_op, sub_op) + ccall((:CeedCompositeOperatorAddSub, libceed), Cint, (CeedOperator, CeedOperator), composite_op, sub_op) end function CeedOperatorLinearAssembleQFunction(op, assembled, rstr, request) @@ -374,20 +422,32 @@ function CeedOperatorLinearAssembleAddPointBlockDiagonal(op, assembled, request) ccall((:CeedOperatorLinearAssembleAddPointBlockDiagonal, libceed), Cint, (CeedOperator, CeedVector, Ptr{CeedRequest}), op, assembled, request) end -function CeedOperatorMultigridLevelCreate(opFine, PMultFine, rstrCoarse, basisCoarse, opCoarse, opProlong, opRestrict) - ccall((:CeedOperatorMultigridLevelCreate, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), opFine, PMultFine, rstrCoarse, basisCoarse, opCoarse, opProlong, opRestrict) +function CeedOperatorLinearAssembleSymbolic(op, num_entries, rows, cols) + ccall((:CeedOperatorLinearAssembleSymbolic, libceed), Cint, (CeedOperator, Ptr{CeedInt}, Ptr{Ptr{CeedInt}}, Ptr{Ptr{CeedInt}}), op, num_entries, rows, cols) +end + +function CeedOperatorLinearAssemble(op, values) + ccall((:CeedOperatorLinearAssemble, libceed), Cint, (CeedOperator, CeedVector), op, values) +end + +function CeedOperatorMultigridLevelCreate(op_fine, p_mult_fine, rstr_coarse, basis_coarse, op_coarse, op_prolong, op_restrict) + ccall((:CeedOperatorMultigridLevelCreate, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), op_fine, p_mult_fine, rstr_coarse, basis_coarse, op_coarse, op_prolong, op_restrict) end -function CeedOperatorMultigridLevelCreateTensorH1(opFine, PMultFine, rstrCoarse, basisCoarse, interpCtoF, opCoarse, opProlong, opRestrict) - ccall((:CeedOperatorMultigridLevelCreateTensorH1, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedScalar}, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), opFine, PMultFine, rstrCoarse, basisCoarse, interpCtoF, opCoarse, opProlong, opRestrict) +function CeedOperatorMultigridLevelCreateTensorH1(op_fine, p_mult_fine, rstr_coarse, basis_coarse, interp_c_to_f, op_coarse, op_prolong, op_restrict) + ccall((:CeedOperatorMultigridLevelCreateTensorH1, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedScalar}, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), op_fine, p_mult_fine, rstr_coarse, basis_coarse, interp_c_to_f, op_coarse, op_prolong, op_restrict) end -function CeedOperatorMultigridLevelCreateH1(opFine, PMultFine, rstrCoarse, basisCoarse, interpCtoF, opCoarse, opProlong, opRestrict) - ccall((:CeedOperatorMultigridLevelCreateH1, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedScalar}, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), opFine, PMultFine, rstrCoarse, basisCoarse, interpCtoF, opCoarse, opProlong, opRestrict) +function CeedOperatorMultigridLevelCreateH1(op_fine, p_mult_fine, rstr_coarse, basis_coarse, interp_c_to_f, op_coarse, op_prolong, op_restrict) + ccall((:CeedOperatorMultigridLevelCreateH1, libceed), Cint, (CeedOperator, CeedVector, CeedElemRestriction, CeedBasis, Ptr{CeedScalar}, Ptr{CeedOperator}, Ptr{CeedOperator}, Ptr{CeedOperator}), op_fine, p_mult_fine, rstr_coarse, basis_coarse, interp_c_to_f, op_coarse, op_prolong, op_restrict) end -function CeedOperatorCreateFDMElementInverse(op, fdminv, request) - ccall((:CeedOperatorCreateFDMElementInverse, libceed), Cint, (CeedOperator, Ptr{CeedOperator}, Ptr{CeedRequest}), op, fdminv, request) +function CeedOperatorCreateFDMElementInverse(op, fdm_inv, request) + ccall((:CeedOperatorCreateFDMElementInverse, libceed), Cint, (CeedOperator, Ptr{CeedOperator}, Ptr{CeedRequest}), op, fdm_inv, request) +end + +function CeedOperatorSetNumQuadraturePoints(op, num_qpts) + ccall((:CeedOperatorSetNumQuadraturePoints, libceed), Cint, (CeedOperator, CeedInt), op, num_qpts) end function CeedOperatorView(op, stream) @@ -417,14 +477,22 @@ end function CeedIntMax(a, b) ccall((:CeedIntMax, libceed), CeedInt, (CeedInt, CeedInt), a, b) end -# Julia wrapper for header: ceed-cuda.h + +function CeedRegisterAll() + ccall((:CeedRegisterAll, libceed), Cint, ()) +end + +function CeedQFunctionRegisterAll() + ccall((:CeedQFunctionRegisterAll, libceed), Cint, ()) +end +# Julia wrapper for header: cuda.h # Automatically generated using Clang.jl function CeedQFunctionSetCUDAUserFunction(qf, f) ccall((:CeedQFunctionSetCUDAUserFunction, libceed), Cint, (CeedQFunction, Cint), qf, f) end -# Julia wrapper for header: ceed-backend.h +# Julia wrapper for header: backend.h # Automatically generated using Clang.jl @@ -448,8 +516,8 @@ function CeedRegister(prefix, init, priority) ccall((:CeedRegister, libceed), Cint, (Cstring, Ptr{Cvoid}, UInt32), prefix, init, priority) end -function CeedIsDebug(ceed, isDebug) - ccall((:CeedIsDebug, libceed), Cint, (Ceed, Ptr{Bool}), ceed, isDebug) +function CeedIsDebug(ceed, is_debug) + ccall((:CeedIsDebug, libceed), Cint, (Ceed, Ptr{Bool}), ceed, is_debug) end function CeedGetParent(ceed, parent) @@ -464,12 +532,12 @@ function CeedSetDelegate(ceed, delegate) ccall((:CeedSetDelegate, libceed), Cint, (Ceed, Ceed), ceed, delegate) end -function CeedGetObjectDelegate(ceed, delegate, objname) - ccall((:CeedGetObjectDelegate, libceed), Cint, (Ceed, Ptr{Ceed}, Cstring), ceed, delegate, objname) +function CeedGetObjectDelegate(ceed, delegate, obj_name) + ccall((:CeedGetObjectDelegate, libceed), Cint, (Ceed, Ptr{Ceed}, Cstring), ceed, delegate, obj_name) end -function CeedSetObjectDelegate(ceed, delegate, objname) - ccall((:CeedSetObjectDelegate, libceed), Cint, (Ceed, Ceed, Cstring), ceed, delegate, objname) +function CeedSetObjectDelegate(ceed, delegate, obj_name) + ccall((:CeedSetObjectDelegate, libceed), Cint, (Ceed, Ceed, Cstring), ceed, delegate, obj_name) end function CeedGetOperatorFallbackResource(ceed, resource) @@ -484,12 +552,12 @@ function CeedGetOperatorFallbackParentCeed(ceed, parent) ccall((:CeedGetOperatorFallbackParentCeed, libceed), Cint, (Ceed, Ptr{Ceed}), ceed, parent) end -function CeedSetDeterministic(ceed, isDeterministic) - ccall((:CeedSetDeterministic, libceed), Cint, (Ceed, Bool), ceed, isDeterministic) +function CeedSetDeterministic(ceed, is_deterministic) + ccall((:CeedSetDeterministic, libceed), Cint, (Ceed, Bool), ceed, is_deterministic) end -function CeedSetBackendFunction(ceed, type, object, fname, f) - ccall((:CeedSetBackendFunction, libceed), Cint, (Ceed, Cstring, Ptr{Cvoid}, Cstring, Ptr{Cvoid}), ceed, type, object, fname, f) +function CeedSetBackendFunction(ceed, type, object, func_name, f) + ccall((:CeedSetBackendFunction, libceed), Cint, (Ceed, Cstring, Ptr{Cvoid}, Cstring, Ptr{Cvoid}), ceed, type, object, func_name, f) end function CeedGetData(ceed, data) @@ -500,6 +568,10 @@ function CeedSetData(ceed, data) ccall((:CeedSetData, libceed), Cint, (Ceed, Ptr{Cvoid}), ceed, data) end +function CeedReference(ceed) + ccall((:CeedReference, libceed), Cint, (Ceed,), ceed) +end + function CeedVectorGetCeed(vec, ceed) ccall((:CeedVectorGetCeed, libceed), Cint, (CeedVector, Ptr{Ceed}), vec, ceed) end @@ -520,6 +592,10 @@ function CeedVectorSetData(vec, data) ccall((:CeedVectorSetData, libceed), Cint, (CeedVector, Ptr{Cvoid}), vec, data) end +function CeedVectorReference(vec) + ccall((:CeedVectorReference, libceed), Cint, (CeedVector,), vec) +end + function CeedElemRestrictionGetCeed(rstr, ceed) ccall((:CeedElemRestrictionGetCeed, libceed), Cint, (CeedElemRestriction, Ptr{Ceed}), rstr, ceed) end @@ -528,20 +604,20 @@ function CeedElemRestrictionGetStrides(rstr, strides) ccall((:CeedElemRestrictionGetStrides, libceed), Cint, (CeedElemRestriction, Ptr{NTuple{3, CeedInt}}), rstr, strides) end -function CeedElemRestrictionGetOffsets(rstr, mtype, offsets) - ccall((:CeedElemRestrictionGetOffsets, libceed), Cint, (CeedElemRestriction, CeedMemType, Ptr{Ptr{CeedInt}}), rstr, mtype, offsets) +function CeedElemRestrictionGetOffsets(rstr, mem_type, offsets) + ccall((:CeedElemRestrictionGetOffsets, libceed), Cint, (CeedElemRestriction, CeedMemType, Ptr{Ptr{CeedInt}}), rstr, mem_type, offsets) end function CeedElemRestrictionRestoreOffsets(rstr, offsets) ccall((:CeedElemRestrictionRestoreOffsets, libceed), Cint, (CeedElemRestriction, Ptr{Ptr{CeedInt}}), rstr, offsets) end -function CeedElemRestrictionIsStrided(rstr, isstrided) - ccall((:CeedElemRestrictionIsStrided, libceed), Cint, (CeedElemRestriction, Ptr{Bool}), rstr, isstrided) +function CeedElemRestrictionIsStrided(rstr, is_strided) + ccall((:CeedElemRestrictionIsStrided, libceed), Cint, (CeedElemRestriction, Ptr{Bool}), rstr, is_strided) end -function CeedElemRestrictionHasBackendStrides(rstr, hasbackendstrides) - ccall((:CeedElemRestrictionHasBackendStrides, libceed), Cint, (CeedElemRestriction, Ptr{Bool}), rstr, hasbackendstrides) +function CeedElemRestrictionHasBackendStrides(rstr, has_backend_strides) + ccall((:CeedElemRestrictionHasBackendStrides, libceed), Cint, (CeedElemRestriction, Ptr{Bool}), rstr, has_backend_strides) end function CeedElemRestrictionGetELayout(rstr, layout) @@ -560,20 +636,24 @@ function CeedElemRestrictionSetData(rstr, data) ccall((:CeedElemRestrictionSetData, libceed), Cint, (CeedElemRestriction, Ptr{Cvoid}), rstr, data) end -function CeedBasisGetCollocatedGrad(basis, colograd1d) - ccall((:CeedBasisGetCollocatedGrad, libceed), Cint, (CeedBasis, Ptr{CeedScalar}), basis, colograd1d) +function CeedElemRestrictionReference(rstr) + ccall((:CeedElemRestrictionReference, libceed), Cint, (CeedElemRestriction,), rstr) end -function CeedHouseholderApplyQ(A, Q, tau, tmode, m, n, k, row, col) - ccall((:CeedHouseholderApplyQ, libceed), Cint, (Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedTransposeMode, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt), A, Q, tau, tmode, m, n, k, row, col) +function CeedBasisGetCollocatedGrad(basis, colo_grad_1d) + ccall((:CeedBasisGetCollocatedGrad, libceed), Cint, (CeedBasis, Ptr{CeedScalar}), basis, colo_grad_1d) +end + +function CeedHouseholderApplyQ(A, Q, tau, t_mode, m, n, k, row, col) + ccall((:CeedHouseholderApplyQ, libceed), Cint, (Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedTransposeMode, CeedInt, CeedInt, CeedInt, CeedInt, CeedInt), A, Q, tau, t_mode, m, n, k, row, col) end function CeedBasisGetCeed(basis, ceed) ccall((:CeedBasisGetCeed, libceed), Cint, (CeedBasis, Ptr{Ceed}), basis, ceed) end -function CeedBasisIsTensor(basis, istensor) - ccall((:CeedBasisIsTensor, libceed), Cint, (CeedBasis, Ptr{Bool}), basis, istensor) +function CeedBasisIsTensor(basis, is_tensor) + ccall((:CeedBasisIsTensor, libceed), Cint, (CeedBasis, Ptr{Bool}), basis, is_tensor) end function CeedBasisGetData(basis, data) @@ -584,6 +664,10 @@ function CeedBasisSetData(basis, data) ccall((:CeedBasisSetData, libceed), Cint, (CeedBasis, Ptr{Cvoid}), basis, data) end +function CeedBasisReference(basis) + ccall((:CeedBasisReference, libceed), Cint, (CeedBasis,), basis) +end + function CeedBasisGetTopologyDimension(topo, dim) ccall((:CeedBasisGetTopologyDimension, libceed), Cint, (CeedElemTopology, Ptr{CeedInt}), topo, dim) end @@ -593,15 +677,15 @@ function CeedBasisGetTensorContract(basis, contract) end function CeedBasisSetTensorContract(basis, contract) - ccall((:CeedBasisSetTensorContract, libceed), Cint, (CeedBasis, Ptr{CeedTensorContract}), basis, contract) + ccall((:CeedBasisSetTensorContract, libceed), Cint, (CeedBasis, CeedTensorContract), basis, contract) end function CeedTensorContractCreate(ceed, basis, contract) ccall((:CeedTensorContractCreate, libceed), Cint, (Ceed, CeedBasis, Ptr{CeedTensorContract}), ceed, basis, contract) end -function CeedTensorContractApply(contract, A, B, C, J, t, tmode, Add, u, v) - ccall((:CeedTensorContractApply, libceed), Cint, (CeedTensorContract, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, CeedTransposeMode, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), contract, A, B, C, J, t, tmode, Add, u, v) +function CeedTensorContractApply(contract, A, B, C, J, t, t_mode, Add, u, v) + ccall((:CeedTensorContractApply, libceed), Cint, (CeedTensorContract, CeedInt, CeedInt, CeedInt, CeedInt, Ptr{CeedScalar}, CeedTransposeMode, CeedInt, Ptr{CeedScalar}, Ptr{CeedScalar}), contract, A, B, C, J, t, t_mode, Add, u, v) end function CeedTensorContractGetCeed(contract, ceed) @@ -616,6 +700,10 @@ function CeedTensorContractSetData(contract, data) ccall((:CeedTensorContractSetData, libceed), Cint, (CeedTensorContract, Ptr{Cvoid}), contract, data) end +function CeedTensorContractReference(contract) + ccall((:CeedTensorContractReference, libceed), Cint, (CeedTensorContract,), contract) +end + function CeedTensorContractDestroy(contract) ccall((:CeedTensorContractDestroy, libceed), Cint, (Ptr{CeedTensorContract},), contract) end @@ -632,12 +720,12 @@ function CeedQFunctionGetCeed(qf, ceed) ccall((:CeedQFunctionGetCeed, libceed), Cint, (CeedQFunction, Ptr{Ceed}), qf, ceed) end -function CeedQFunctionGetVectorLength(qf, vlength) - ccall((:CeedQFunctionGetVectorLength, libceed), Cint, (CeedQFunction, Ptr{CeedInt}), qf, vlength) +function CeedQFunctionGetVectorLength(qf, vec_length) + ccall((:CeedQFunctionGetVectorLength, libceed), Cint, (CeedQFunction, Ptr{CeedInt}), qf, vec_length) end -function CeedQFunctionGetNumArgs(qf, numinputfields, numoutputfields) - ccall((:CeedQFunctionGetNumArgs, libceed), Cint, (CeedQFunction, Ptr{CeedInt}, Ptr{CeedInt}), qf, numinputfields, numoutputfields) +function CeedQFunctionGetNumArgs(qf, num_input_fields, num_output_fields) + ccall((:CeedQFunctionGetNumArgs, libceed), Cint, (CeedQFunction, Ptr{CeedInt}, Ptr{CeedInt}), qf, num_input_fields, num_output_fields) end function CeedQFunctionGetSourcePath(qf, source) @@ -656,8 +744,8 @@ function CeedQFunctionGetInnerContext(qf, ctx) ccall((:CeedQFunctionGetInnerContext, libceed), Cint, (CeedQFunction, Ptr{CeedQFunctionContext}), qf, ctx) end -function CeedQFunctionIsIdentity(qf, isidentity) - ccall((:CeedQFunctionIsIdentity, libceed), Cint, (CeedQFunction, Ptr{Bool}), qf, isidentity) +function CeedQFunctionIsIdentity(qf, is_identity) + ccall((:CeedQFunctionIsIdentity, libceed), Cint, (CeedQFunction, Ptr{Bool}), qf, is_identity) end function CeedQFunctionGetData(qf, data) @@ -668,20 +756,24 @@ function CeedQFunctionSetData(qf, data) ccall((:CeedQFunctionSetData, libceed), Cint, (CeedQFunction, Ptr{Cvoid}), qf, data) end -function CeedQFunctionGetFields(qf, inputfields, outputfields) - ccall((:CeedQFunctionGetFields, libceed), Cint, (CeedQFunction, Ptr{Ptr{CeedQFunctionField}}, Ptr{Ptr{CeedQFunctionField}}), qf, inputfields, outputfields) +function CeedQFunctionReference(qf) + ccall((:CeedQFunctionReference, libceed), Cint, (CeedQFunction,), qf) end -function CeedQFunctionFieldGetName(qffield, fieldname) - ccall((:CeedQFunctionFieldGetName, libceed), Cint, (CeedQFunctionField, Ptr{Cstring}), qffield, fieldname) +function CeedQFunctionGetFields(qf, input_fields, output_fields) + ccall((:CeedQFunctionGetFields, libceed), Cint, (CeedQFunction, Ptr{Ptr{CeedQFunctionField}}, Ptr{Ptr{CeedQFunctionField}}), qf, input_fields, output_fields) end -function CeedQFunctionFieldGetSize(qffield, size) - ccall((:CeedQFunctionFieldGetSize, libceed), Cint, (CeedQFunctionField, Ptr{CeedInt}), qffield, size) +function CeedQFunctionFieldGetName(qf_field, field_name) + ccall((:CeedQFunctionFieldGetName, libceed), Cint, (CeedQFunctionField, Ptr{Cstring}), qf_field, field_name) end -function CeedQFunctionFieldGetEvalMode(qffield, emode) - ccall((:CeedQFunctionFieldGetEvalMode, libceed), Cint, (CeedQFunctionField, Ptr{CeedEvalMode}), qffield, emode) +function CeedQFunctionFieldGetSize(qf_field, size) + ccall((:CeedQFunctionFieldGetSize, libceed), Cint, (CeedQFunctionField, Ptr{CeedInt}), qf_field, size) +end + +function CeedQFunctionFieldGetEvalMode(qf_field, eval_mode) + ccall((:CeedQFunctionFieldGetEvalMode, libceed), Cint, (CeedQFunctionField, Ptr{CeedEvalMode}), qf_field, eval_mode) end function CeedQFunctionContextGetCeed(cxt, ceed) @@ -692,8 +784,8 @@ function CeedQFunctionContextGetState(ctx, state) ccall((:CeedQFunctionContextGetState, libceed), Cint, (CeedQFunctionContext, Ptr{UInt64}), ctx, state) end -function CeedQFunctionContextGetContextSize(ctx, ctxsize) - ccall((:CeedQFunctionContextGetContextSize, libceed), Cint, (CeedQFunctionContext, Ptr{Csize_t}), ctx, ctxsize) +function CeedQFunctionContextGetContextSize(ctx, ctx_size) + ccall((:CeedQFunctionContextGetContextSize, libceed), Cint, (CeedQFunctionContext, Ptr{Csize_t}), ctx, ctx_size) end function CeedQFunctionContextGetBackendData(ctx, data) @@ -704,40 +796,44 @@ function CeedQFunctionContextSetBackendData(ctx, data) ccall((:CeedQFunctionContextSetBackendData, libceed), Cint, (CeedQFunctionContext, Ptr{Cvoid}), ctx, data) end +function CeedQFunctionContextReference(ctx) + ccall((:CeedQFunctionContextReference, libceed), Cint, (CeedQFunctionContext,), ctx) +end + function CeedOperatorGetCeed(op, ceed) ccall((:CeedOperatorGetCeed, libceed), Cint, (CeedOperator, Ptr{Ceed}), op, ceed) end -function CeedOperatorGetNumElements(op, numelem) - ccall((:CeedOperatorGetNumElements, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, numelem) +function CeedOperatorGetNumElements(op, num_elem) + ccall((:CeedOperatorGetNumElements, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, num_elem) end -function CeedOperatorGetNumQuadraturePoints(op, numqpts) - ccall((:CeedOperatorGetNumQuadraturePoints, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, numqpts) +function CeedOperatorGetNumQuadraturePoints(op, num_qpts) + ccall((:CeedOperatorGetNumQuadraturePoints, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, num_qpts) end -function CeedOperatorGetNumArgs(op, numargs) - ccall((:CeedOperatorGetNumArgs, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, numargs) +function CeedOperatorGetNumArgs(op, num_args) + ccall((:CeedOperatorGetNumArgs, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, num_args) end -function CeedOperatorIsSetupDone(op, issetupdone) - ccall((:CeedOperatorIsSetupDone, libceed), Cint, (CeedOperator, Ptr{Bool}), op, issetupdone) +function CeedOperatorIsSetupDone(op, is_setup_done) + ccall((:CeedOperatorIsSetupDone, libceed), Cint, (CeedOperator, Ptr{Bool}), op, is_setup_done) end function CeedOperatorGetQFunction(op, qf) ccall((:CeedOperatorGetQFunction, libceed), Cint, (CeedOperator, Ptr{CeedQFunction}), op, qf) end -function CeedOperatorIsComposite(op, iscomposite) - ccall((:CeedOperatorIsComposite, libceed), Cint, (CeedOperator, Ptr{Bool}), op, iscomposite) +function CeedOperatorIsComposite(op, is_composite) + ccall((:CeedOperatorIsComposite, libceed), Cint, (CeedOperator, Ptr{Bool}), op, is_composite) end -function CeedOperatorGetNumSub(op, numsub) - ccall((:CeedOperatorGetNumSub, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, numsub) +function CeedOperatorGetNumSub(op, num_suboperators) + ccall((:CeedOperatorGetNumSub, libceed), Cint, (CeedOperator, Ptr{CeedInt}), op, num_suboperators) end -function CeedOperatorGetSubList(op, suboperators) - ccall((:CeedOperatorGetSubList, libceed), Cint, (CeedOperator, Ptr{Ptr{CeedOperator}}), op, suboperators) +function CeedOperatorGetSubList(op, sub_operators) + ccall((:CeedOperatorGetSubList, libceed), Cint, (CeedOperator, Ptr{Ptr{CeedOperator}}), op, sub_operators) end function CeedOperatorGetData(op, data) @@ -748,26 +844,30 @@ function CeedOperatorSetData(op, data) ccall((:CeedOperatorSetData, libceed), Cint, (CeedOperator, Ptr{Cvoid}), op, data) end +function CeedOperatorReference(op) + ccall((:CeedOperatorReference, libceed), Cint, (CeedOperator,), op) +end + function CeedOperatorSetSetupDone(op) ccall((:CeedOperatorSetSetupDone, libceed), Cint, (CeedOperator,), op) end -function CeedOperatorGetFields(op, inputfields, outputfields) - ccall((:CeedOperatorGetFields, libceed), Cint, (CeedOperator, Ptr{Ptr{CeedOperatorField}}, Ptr{Ptr{CeedOperatorField}}), op, inputfields, outputfields) +function CeedOperatorGetFields(op, input_fields, output_fields) + ccall((:CeedOperatorGetFields, libceed), Cint, (CeedOperator, Ptr{Ptr{CeedOperatorField}}, Ptr{Ptr{CeedOperatorField}}), op, input_fields, output_fields) end -function CeedOperatorFieldGetElemRestriction(opfield, rstr) - ccall((:CeedOperatorFieldGetElemRestriction, libceed), Cint, (CeedOperatorField, Ptr{CeedElemRestriction}), opfield, rstr) +function CeedOperatorFieldGetElemRestriction(op_field, rstr) + ccall((:CeedOperatorFieldGetElemRestriction, libceed), Cint, (CeedOperatorField, Ptr{CeedElemRestriction}), op_field, rstr) end -function CeedOperatorFieldGetBasis(opfield, basis) - ccall((:CeedOperatorFieldGetBasis, libceed), Cint, (CeedOperatorField, Ptr{CeedBasis}), opfield, basis) +function CeedOperatorFieldGetBasis(op_field, basis) + ccall((:CeedOperatorFieldGetBasis, libceed), Cint, (CeedOperatorField, Ptr{CeedBasis}), op_field, basis) end -function CeedOperatorFieldGetVector(opfield, vec) - ccall((:CeedOperatorFieldGetVector, libceed), Cint, (CeedOperatorField, Ptr{CeedVector}), opfield, vec) +function CeedOperatorFieldGetVector(op_field, vec) + ccall((:CeedOperatorFieldGetVector, libceed), Cint, (CeedOperatorField, Ptr{CeedVector}), op_field, vec) end -function CeedMatrixMultiply(ceed, matA, matB, matC, m, n, kk) - ccall((:CeedMatrixMultiply, libceed), Cint, (Ceed, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedInt, CeedInt, CeedInt), ceed, matA, matB, matC, m, n, kk) +function CeedMatrixMultiply(ceed, mat_A, mat_B, mat_C, m, n, kk) + ccall((:CeedMatrixMultiply, libceed), Cint, (Ceed, Ptr{CeedScalar}, Ptr{CeedScalar}, Ptr{CeedScalar}, CeedInt, CeedInt, CeedInt), ceed, mat_A, mat_B, mat_C, m, n, kk) end diff --git a/julia/LibCEED.jl/src/generated/libceed_common.jl b/julia/LibCEED.jl/src/generated/libceed_common.jl index aad84aa497..b9e68a619a 100644 --- a/julia/LibCEED.jl/src/generated/libceed_common.jl +++ b/julia/LibCEED.jl/src/generated/libceed_common.jl @@ -7,6 +7,13 @@ const FILE = Cvoid # Skipping MacroDefinition: CeedError ( ceed , ecode , ... ) ( CeedErrorImpl ( ( ceed ) , __FILE__ , __LINE__ , __func__ , ( ecode ) , __VA_ARGS__ ) ? : ( ecode ) ) +const CEED_VERSION_MAJOR = 0 +const CEED_VERSION_MINOR = 8 +const CEED_VERSION_PATCH = 0 +const CEED_VERSION_RELEASE = false + +# Skipping MacroDefinition: CEED_VERSION_GE ( major , minor , patch ) ( ! CEED_VERSION_RELEASE || ( CEED_VERSION_MAJOR > major || ( CEED_VERSION_MAJOR == major && ( CEED_VERSION_MINOR > minor || ( CEED_VERSION_MINOR == minor && CEED_VERSION_PATCH >= patch ) ) ) ) ) + const CeedInt = Int32 const CeedScalar = Cdouble const Ceed_private = Cvoid @@ -25,6 +32,19 @@ const CeedQFunctionContext_private = Cvoid const CeedQFunctionContext = Ptr{CeedQFunctionContext_private} const CeedOperator_private = Cvoid const CeedOperator = Ptr{CeedOperator_private} +const CeedErrorHandler = Ptr{Cvoid} + +@cenum CeedErrorType::Int32 begin + CEED_ERROR_SUCCESS = 0 + CEED_ERROR_MINOR = 1 + CEED_ERROR_DIMENSION = 2 + CEED_ERROR_INCOMPLETE = 3 + CEED_ERROR_INCOMPATIBLE = 4 + CEED_ERROR_ACCESS = 5 + CEED_ERROR_MAJOR = -1 + CEED_ERROR_BACKEND = -2 + CEED_ERROR_UNSUPPORTED = -3 +end @cenum CeedMemType::UInt32 begin CEED_MEM_HOST = 0 @@ -76,8 +96,10 @@ end const CeedQFunctionUser = Ptr{Cvoid} # Skipping MacroDefinition: CEED_INTERN CEED_EXTERN __attribute__ ( ( visibility ( "hidden" ) ) ) +# Skipping MacroDefinition: CEED_UNUSED __attribute__ ( ( unused ) ) const CEED_MAX_RESOURCE_LEN = 1024 +const CEED_MAX_BACKEND_PRIORITY = typemax(Cuint) const CEED_ALIGN = 64 const CEED_COMPOSITE_MAX = 16 const CEED_EPSILON = 1.0e-16 @@ -86,7 +108,8 @@ const CEED_DEBUG_COLOR = 0 # Skipping MacroDefinition: CeedDebug1 ( ceed , format , ... ) CeedDebugImpl ( ceed , format , ## __VA_ARGS__ ) # Skipping MacroDefinition: CeedDebug256 ( ceed , color , ... ) CeedDebugImpl256 ( ceed , color , ## __VA_ARGS__ ) # Skipping MacroDefinition: CeedDebug ( ... ) CeedDebug256 ( ceed , ( unsigned char ) CEED_DEBUG_COLOR , ## __VA_ARGS__ ) -# Skipping MacroDefinition: CeedChk ( ierr ) do { if ( ierr ) return ierr ; } while ( 0 ) +# Skipping MacroDefinition: CeedChk ( ierr ) do { int ierr_ = ierr ; if ( ierr_ ) return ierr_ ; } while ( 0 ) +# Skipping MacroDefinition: CeedChkBackend ( ierr ) do { int ierr_ = ierr ; if ( ierr_ ) { if ( ierr_ > CEED_ERROR_SUCCESS ) return CEED_ERROR_BACKEND ; else return ierr_ ; } } while ( 0 ) # Skipping MacroDefinition: CeedMalloc ( n , p ) CeedMallocArray ( ( n ) , sizeof ( * * ( p ) ) , p ) # Skipping MacroDefinition: CeedCalloc ( n , p ) CeedCallocArray ( ( n ) , sizeof ( * * ( p ) ) , p ) # Skipping MacroDefinition: CeedRealloc ( n , p ) CeedReallocArray ( ( n ) , sizeof ( * * ( p ) ) , p ) From d0293d3e9e5acb33f03f30a5189337eb6e116629 Mon Sep 17 00:00:00 2001 From: Will Pazner Date: Tue, 27 Apr 2021 12:40:46 -0700 Subject: [PATCH 06/11] [Julia] Access to version number at run time --- julia/LibCEED.jl/docs/src/Ceed.md | 2 ++ julia/LibCEED.jl/src/Ceed.jl | 28 ++++++++++++++++++++++++++++ julia/LibCEED.jl/src/LibCEED.jl | 2 ++ 3 files changed, 32 insertions(+) diff --git a/julia/LibCEED.jl/docs/src/Ceed.md b/julia/LibCEED.jl/docs/src/Ceed.md index 4111c9ce39..4dfbdfcd7a 100644 --- a/julia/LibCEED.jl/docs/src/Ceed.md +++ b/julia/LibCEED.jl/docs/src/Ceed.md @@ -6,4 +6,6 @@ getresource get_preferred_memtype isdeterministic iscuda +ceedversion +isrelease ``` diff --git a/julia/LibCEED.jl/src/Ceed.jl b/julia/LibCEED.jl/src/Ceed.jl index d919e75812..e7122bc5fa 100644 --- a/julia/LibCEED.jl/src/Ceed.jl +++ b/julia/LibCEED.jl/src/Ceed.jl @@ -37,6 +37,34 @@ function handle_ceed_error( throw(CeedError(fname, lineno, func, ecode, message)) end +""" + ceedversion() + +Returns a `VersionNumber` corresponding to the version of the libCEED library currently used. +""" +function ceedversion() + major = Ref{Cint}() + minor = Ref{Cint}() + patch = Ref{Cint}() + release = Ref{Bool}() + C.CeedGetVersion(major, minor, patch, release) + return VersionNumber(major[], minor[], patch[]) +end + +""" + isrelease() + +Returns true if the libCEED library is a release build, false otherwise. +""" +function isrelease() + major = Ref{Cint}() + minor = Ref{Cint}() + patch = Ref{Cint}() + release = Ref{Bool}() + C.CeedGetVersion(major, minor, patch, release) + return release[] +end + mutable struct Ceed ref::RefValue{C.Ceed} end diff --git a/julia/LibCEED.jl/src/LibCEED.jl b/julia/LibCEED.jl/src/LibCEED.jl index 6202d7c0e8..ebb3020163 100644 --- a/julia/LibCEED.jl/src/LibCEED.jl +++ b/julia/LibCEED.jl/src/LibCEED.jl @@ -71,6 +71,7 @@ export @interior_qf, assemble, assemble_add_diagonal!, assemble_diagonal!, + ceedversion, create_composite_operator, create_elem_restriction, create_elem_restriction_strided, @@ -111,6 +112,7 @@ export @interior_qf, getvoigt, iscuda, isdeterministic, + isrelease, lobatto_quadrature, norm, reciprocal!, From 557acef3f79ecffb65d04906a1f3c5c268aa9b4f Mon Sep 17 00:00:00 2001 From: Will Pazner Date: Tue, 27 Apr 2021 12:41:39 -0700 Subject: [PATCH 07/11] [Julia] Add scale!, axpy!, and pointwisemult! --- julia/LibCEED.jl/docs/src/CeedVector.md | 3 ++ julia/LibCEED.jl/src/CeedVector.jl | 38 +++++++++++++++++++++++-- julia/LibCEED.jl/src/LibCEED.jl | 3 ++ 3 files changed, 42 insertions(+), 2 deletions(-) diff --git a/julia/LibCEED.jl/docs/src/CeedVector.md b/julia/LibCEED.jl/docs/src/CeedVector.md index 75d94c6f1d..7deb01b1b9 100644 --- a/julia/LibCEED.jl/docs/src/CeedVector.md +++ b/julia/LibCEED.jl/docs/src/CeedVector.md @@ -15,4 +15,7 @@ witharray_read setarray! syncarray! takearray! +scale! +LinearAlgebra.axpy!(a::Real, x::CeedVector, y::CeedVector) +pointwisemult! ``` diff --git a/julia/LibCEED.jl/src/CeedVector.jl b/julia/LibCEED.jl/src/CeedVector.jl index 7d4dc0304a..47fd8300d7 100644 --- a/julia/LibCEED.jl/src/CeedVector.jl +++ b/julia/LibCEED.jl/src/CeedVector.jl @@ -1,4 +1,4 @@ -import LinearAlgebra: norm +import LinearAlgebra: norm, axpy! abstract type AbstractCeedVector end @@ -185,7 +185,7 @@ will be used. If the size is not specified, a flat vector will be assumed. # Examples Negate the contents of `CeedVector` `v`: ``` -@witharray v_arr=v v_arr *= -1.0 +@witharray v_arr=v v_arr .*= -1.0 ``` """ macro witharray(assignment, args...) @@ -306,3 +306,37 @@ function witharray_read(f, v::CeedVector, mtype::MemType=MEM_HOST) end return res end + +""" + scale!(v::CeedVector, a::Real) + +Overwrite `v` with `a*v` for scalar `a`. Returns `v`. +""" +function scale!(v::CeedVector, a::Real) + C.CeedVectorScale(v[], a) + return v +end + +""" + axpy!(a::Real, x::CeedVector, y::CeedVector) + +Overwrite `y` with `x*a + y`, where `a` is a scalar. Returns `y`. + +!!! warning "Different argument order" + In order to be consistent with `LinearAlgebra.axpy!`, the arguments are passed in order: `a`, + `x`, `y`. This is different than the order of arguments of the C function `CeedVectorAXPY`. +""" +function axpy!(a::Real, x::CeedVector, y::CeedVector) + C.CeedVectorAXPY(y[], a, x[]) + return y +end + +""" + pointwisemult!(w::CeedVector, x::CeedVector, y::CeedVector) + +Overwrite `w` with `x .* y`. Any subset of x, y, and w may be the same vector. Returns `w`. +""" +function pointwisemult!(w::CeedVector, x::CeedVector, y::CeedVector) + C.CeedVectorPointwiseMult(w[], x[], y[]) + return w +end diff --git a/julia/LibCEED.jl/src/LibCEED.jl b/julia/LibCEED.jl/src/LibCEED.jl index ebb3020163..90f8b5f8e6 100644 --- a/julia/LibCEED.jl/src/LibCEED.jl +++ b/julia/LibCEED.jl/src/LibCEED.jl @@ -71,6 +71,7 @@ export @interior_qf, assemble, assemble_add_diagonal!, assemble_diagonal!, + axpy!, ceedversion, create_composite_operator, create_elem_restriction, @@ -115,7 +116,9 @@ export @interior_qf, isrelease, lobatto_quadrature, norm, + pointwisemult!, reciprocal!, + scale!, set_context!, set_cufunction!, set_data!, From a697ff736c4bbf0dcf3b0c0690ba5a6b92dd6bdf Mon Sep 17 00:00:00 2001 From: Will Pazner Date: Tue, 27 Apr 2021 12:42:34 -0700 Subject: [PATCH 08/11] [Julia] Add tests for development features --- .github/workflows/julia-test-with-style.yml | 2 +- julia/LibCEED.jl/test/rundevtests.jl | 33 +++++++++++++++++++++ julia/LibCEED.jl/test/runtests.jl | 6 +++- 3 files changed, 39 insertions(+), 2 deletions(-) create mode 100644 julia/LibCEED.jl/test/rundevtests.jl diff --git a/.github/workflows/julia-test-with-style.yml b/.github/workflows/julia-test-with-style.yml index d413c1471a..749f68475c 100644 --- a/.github/workflows/julia-test-with-style.yml +++ b/.github/workflows/julia-test-with-style.yml @@ -25,6 +25,6 @@ jobs: make -j2 export JULIA_LIBCEED_LIB=$(find $PWD/lib -name "libceed.*") pushd julia/LibCEED.jl - julia --project -e 'import Pkg; Pkg.build(); Pkg.test("LibCEED"; coverage=true)' + julia --project -e 'import Pkg; Pkg.build(); Pkg.test("LibCEED"; coverage=true, test_args=["--run-dev-tests"])' unset JULIA_LIBCEED_LIB && julia --project -e 'import Pkg; Pkg.build(); Pkg.test("LibCEED")' julia --project=.style/ -e 'import Pkg; Pkg.instantiate()' && julia --project=.style/ .style/ceed_style.jl && git diff --exit-code src test examples diff --git a/julia/LibCEED.jl/test/rundevtests.jl b/julia/LibCEED.jl/test/rundevtests.jl new file mode 100644 index 0000000000..d863230ab1 --- /dev/null +++ b/julia/LibCEED.jl/test/rundevtests.jl @@ -0,0 +1,33 @@ +using Test, LibCEED, LinearAlgebra, StaticArrays + +@testset "LibCEED Development Tests" begin + @test ceedversion() isa VersionNumber + @test isrelease() == false + + @testset "CeedVector" begin + n = 10 + c = Ceed() + + v1 = rand(n) + v2 = rand(n) + v3 = rand(n) + + cv1 = CeedVector(c, v1) + cv2 = CeedVector(c, v2) + cv3 = CeedVector(c, v3) + + alpha = rand() + + scale!(cv1, alpha) + v1 .*= alpha + @test @witharray_read(a = cv1, a == v1) + + pointwisemult!(cv1, cv2, cv3) + v1 .= v2.*v3 + @test @witharray_read(a = cv1, a == v1) + + axpy!(alpha, cv2, cv1) + axpy!(alpha, v2, v1) + @test @witharray_read(a = cv1, a ≈ v1) + end +end diff --git a/julia/LibCEED.jl/test/runtests.jl b/julia/LibCEED.jl/test/runtests.jl index f1889e64f7..fb0a42838d 100644 --- a/julia/LibCEED.jl/test/runtests.jl +++ b/julia/LibCEED.jl/test/runtests.jl @@ -18,7 +18,11 @@ mutable struct CtxData x::Vector{Float64} end -@testset "LibCEED" begin +if "--run-dev-tests" in ARGS + include("rundevtests.jl") +end + +@testset "LibCEED Release Tests" begin @testset "Ceed" begin res = "/cpu/self/ref/serial" c = Ceed(res) From 278ca47885405bd548a48bfa436b167e787711c3 Mon Sep 17 00:00:00 2001 From: Will Pazner Date: Tue, 27 Apr 2021 14:28:35 -0700 Subject: [PATCH 09/11] [Julia] Use Real instead of CeedScalar where possible --- julia/LibCEED.jl/src/CeedVector.jl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/julia/LibCEED.jl/src/CeedVector.jl b/julia/LibCEED.jl/src/CeedVector.jl index 47fd8300d7..ac33c4d3e9 100644 --- a/julia/LibCEED.jl/src/CeedVector.jl +++ b/julia/LibCEED.jl/src/CeedVector.jl @@ -55,18 +55,18 @@ Base.size(v::CeedVector) = (length(Int, v),) Base.length(v::CeedVector) = length(Int, v) """ - setvalue!(v::CeedVector, val::CeedScalar) + setvalue!(v::CeedVector, val::Real) Set the [`CeedVector`](@ref) to a constant value. """ -setvalue!(v::CeedVector, val::CeedScalar) = C.CeedVectorSetValue(v[], val) +setvalue!(v::CeedVector, val::Real) = C.CeedVectorSetValue(v[], val) """ - setindex!(v::CeedVector, val::CeedScalar) + setindex!(v::CeedVector, val::Real) v[] = val Set the [`CeedVector`](@ref) to a constant value, synonymous to [`setvalue!`](@ref). """ -Base.setindex!(v::CeedVector, val::CeedScalar) = setvalue!(v, val) +Base.setindex!(v::CeedVector, val::Real) = setvalue!(v, val) """ norm(v::CeedVector, ntype::NormType) From 96b902e2a709bcaf161c65c2cb076b00357cddfa Mon Sep 17 00:00:00 2001 From: jeremylt Date: Wed, 28 Apr 2021 09:17:50 -0600 Subject: [PATCH 10/11] doc - small fixes --- interface/ceed-elemrestriction.c | 4 ++-- interface/ceed-vector.c | 18 +++++++++--------- interface/ceed.c | 3 ++- 3 files changed, 13 insertions(+), 12 deletions(-) diff --git a/interface/ceed-elemrestriction.c b/interface/ceed-elemrestriction.c index 572f48425e..2e6d931d35 100644 --- a/interface/ceed-elemrestriction.c +++ b/interface/ceed-elemrestriction.c @@ -170,8 +170,8 @@ int CeedElemRestrictionIsStrided(CeedElemRestriction rstr, bool *is_strided) { /** @brief Get the backend stride status of a CeedElemRestriction - @param rstr CeedElemRestriction - @param[out] status Variable to store stride status + @param rstr CeedElemRestriction + @param[out] has_backend_strides Variable to store stride status @return An error code: 0 - success, otherwise - failure diff --git a/interface/ceed-vector.c b/interface/ceed-vector.c index b2f7edf976..38fa6e6454 100644 --- a/interface/ceed-vector.c +++ b/interface/ceed-vector.c @@ -36,7 +36,7 @@ static struct CeedVector_private ceed_vector_none; /// CeedOperatorApply(). const CeedVector CEED_VECTOR_ACTIVE = &ceed_vector_active; -/// Indicate that no vector is applicable (i.e., for @ref CEED_EVAL_WEIGHTS). +/// Indicate that no vector is applicable (i.e., for @ref CEED_EVAL_WEIGHT). const CeedVector CEED_VECTOR_NONE = &ceed_vector_none; /// @} @@ -534,8 +534,8 @@ int CeedVectorNorm(CeedVector vec, CeedNormType norm_type, CeedScalar *norm) { /** @brief Compute x = alpha x - @param x[in,out] vector for scaling - @param alpha[in] scaling factor + @param[in,out] x vector for scaling + @param[in] alpha scaling factor @return An error code: 0 - success, otherwise - failure @@ -564,9 +564,9 @@ int CeedVectorScale(CeedVector x, CeedScalar alpha) { /** @brief Compute y = alpha x + y - @param y[in,out] target vector for sum - @param alpha[in] scaling factor - @param x[in] second vector, must be different than y + @param[in,out] y target vector for sum + @param[in] alpha scaling factor + @param[in] x second vector, must be different than y @return An error code: 0 - success, otherwise - failure @@ -621,9 +621,9 @@ int CeedVectorAXPY(CeedVector y, CeedScalar alpha, CeedVector x) { @brief Compute the pointwise multiplication w = x .* y. Any subset of x, y, and w may be the same vector. - @param w[out] target vector for the product - @param x[in] first vector for product - @param y[in] second vector for the product + @param[out] w target vector for the product + @param[in] x first vector for product + @param[in] y second vector for the product @return An error code: 0 - success, otherwise - failure diff --git a/interface/ceed.c b/interface/ceed.c index ea61dd6e96..de724c5e6e 100644 --- a/interface/ceed.c +++ b/interface/ceed.c @@ -497,7 +497,8 @@ int CeedGetOperatorFallbackParentCeed(Ceed ceed, Ceed *parent) { /** @brief Flag Ceed context as deterministic - @param ceed Ceed to flag as deterministic + @param ceed Ceed to flag as deterministic + @param[out] is_deterministic Deterministic status to set @return An error code: 0 - success, otherwise - failure From 891038deaa55d3f0c57599b0bbe43569c9f61de7 Mon Sep 17 00:00:00 2001 From: jeremylt Date: Wed, 28 Apr 2021 09:57:55 -0600 Subject: [PATCH 11/11] context - add QFunctionContextTakeData to silence docs warning about missing ref --- backends/cuda/ceed-cuda-qfunctioncontext.c | 57 ++++++++++++++++++-- backends/hip/ceed-hip-qfunctioncontext.c | 53 +++++++++++++++++- backends/occa/ceed-occa-qfunctioncontext.cpp | 43 +++++++++++++++ backends/occa/ceed-occa-qfunctioncontext.hpp | 5 ++ backends/ref/ceed-ref-qfunctioncontext.c | 28 ++++++++++ include/ceed-impl.h | 1 + include/ceed/ceed.h | 5 +- interface/ceed-qfunctioncontext.c | 37 +++++++++++++ interface/ceed.c | 1 + tests/junit.py | 2 + tests/t404-qfunction.c | 35 ++++++++++++ tests/tap.sh | 51 ++++++++++-------- 12 files changed, 290 insertions(+), 28 deletions(-) create mode 100644 tests/t404-qfunction.c diff --git a/backends/cuda/ceed-cuda-qfunctioncontext.c b/backends/cuda/ceed-cuda-qfunctioncontext.c index 6c210d4878..61f130ef2f 100644 --- a/backends/cuda/ceed-cuda-qfunctioncontext.c +++ b/backends/cuda/ceed-cuda-qfunctioncontext.c @@ -136,9 +136,7 @@ static int CeedQFunctionContextSetDataDevice_Cuda(const CeedQFunctionContext // freeing any previously allocated array if applicable //------------------------------------------------------------------------------ static int CeedQFunctionContextSetData_Cuda(const CeedQFunctionContext ctx, - const CeedMemType mtype, - const CeedCopyMode cmode, - CeedScalar *data) { + const CeedMemType mtype, const CeedCopyMode cmode, CeedScalar *data) { int ierr; Ceed ceed; ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); @@ -153,7 +151,56 @@ static int CeedQFunctionContextSetData_Cuda(const CeedQFunctionContext ctx, } //------------------------------------------------------------------------------ -// Get array +// Take data +//------------------------------------------------------------------------------ +static int CeedQFunctionContextTakeData_Cuda(const CeedQFunctionContext ctx, + const CeedMemType mtype, CeedScalar *data) { + int ierr; + Ceed ceed; + ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); + CeedQFunctionContext_Cuda *impl; + ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); + if(impl->h_data == NULL && impl->d_data == NULL) + // LCOV_EXCL_START + return CeedError(ceed, CEED_ERROR_BACKEND, "No context data set"); + // LCOV_EXCL_STOP + + // Sync array to requested memtype and update pointer + switch (mtype) { + case CEED_MEM_HOST: + if (impl->h_data == NULL) { + ierr = CeedMalloc(bytes(ctx), &impl->h_data_allocated); + CeedChkBackend(ierr); + impl->h_data = impl->h_data_allocated; + } + if (impl->memState == CEED_CUDA_DEVICE_SYNC) { + ierr = CeedQFunctionContextSyncD2H_Cuda(ctx); CeedChkBackend(ierr); + } + impl->memState = CEED_CUDA_HOST_SYNC; + *(void **)data = impl->h_data; + impl->h_data = NULL; + impl->h_data_allocated = NULL; + break; + case CEED_MEM_DEVICE: + if (impl->d_data == NULL) { + ierr = cudaMalloc((void **)&impl->d_data_allocated, bytes(ctx)); + CeedChk_Cu(ceed, ierr); + impl->d_data = impl->d_data_allocated; + } + if (impl->memState == CEED_CUDA_HOST_SYNC) { + ierr = CeedQFunctionContextSyncH2D_Cuda(ctx); CeedChkBackend(ierr); + } + impl->memState = CEED_CUDA_DEVICE_SYNC; + *(void **)data = impl->d_data; + impl->d_data = NULL; + impl->d_data_allocated = NULL; + break; + } + return CEED_ERROR_SUCCESS; +} + +//------------------------------------------------------------------------------ +// Get data //------------------------------------------------------------------------------ static int CeedQFunctionContextGetData_Cuda(const CeedQFunctionContext ctx, const CeedMemType mtype, CeedScalar *data) { @@ -232,6 +279,8 @@ int CeedQFunctionContextCreate_Cuda(CeedQFunctionContext ctx) { ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "SetData", CeedQFunctionContextSetData_Cuda); CeedChkBackend(ierr); + ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "TakeData", + CeedQFunctionContextTakeData_Cuda); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "GetData", CeedQFunctionContextGetData_Cuda); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "RestoreData", diff --git a/backends/hip/ceed-hip-qfunctioncontext.c b/backends/hip/ceed-hip-qfunctioncontext.c index f1a7fbf935..3dbc133203 100644 --- a/backends/hip/ceed-hip-qfunctioncontext.c +++ b/backends/hip/ceed-hip-qfunctioncontext.c @@ -149,7 +149,56 @@ static int CeedQFunctionContextSetData_Hip(const CeedQFunctionContext ctx, } //------------------------------------------------------------------------------ -// Get array +// Take data +//------------------------------------------------------------------------------ +static int CeedQFunctionContextTakeData_Hip(const CeedQFunctionContext ctx, + const CeedMemType mtype, CeedScalar *data) { + int ierr; + Ceed ceed; + ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); + CeedQFunctionContext_Hip *impl; + ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); + if(impl->h_data == NULL && impl->d_data == NULL) + // LCOV_EXCL_START + return CeedError(ceed, CEED_ERROR_BACKEND, "No context data set"); + // LCOV_EXCL_STOP + + // Sync array to requested memtype and update pointer + switch (mtype) { + case CEED_MEM_HOST: + if (impl->h_data == NULL) { + ierr = CeedMalloc(bytes(ctx), &impl->h_data_allocated); + CeedChkBackend(ierr); + impl->h_data = impl->h_data_allocated; + } + if (impl->memState == CEED_HIP_DEVICE_SYNC) { + ierr = CeedQFunctionContextSyncD2H_Hip(ctx); CeedChkBackend(ierr); + } + impl->memState = CEED_HIP_HOST_SYNC; + *(void **)data = impl->h_data; + impl->h_data = NULL; + impl->h_data_allocated = NULL; + break; + case CEED_MEM_DEVICE: + if (impl->d_data == NULL) { + ierr = hipMalloc((void **)&impl->d_data_allocated, bytes(ctx)); + CeedChk_Hip(ceed, ierr); + impl->d_data = impl->d_data_allocated; + } + if (impl->memState == CEED_HIP_HOST_SYNC) { + ierr = CeedQFunctionContextSyncH2D_Hip(ctx); CeedChkBackend(ierr); + } + impl->memState = CEED_HIP_DEVICE_SYNC; + *(void **)data = impl->d_data; + impl->d_data = NULL; + impl->d_data_allocated = NULL; + break; + } + return CEED_ERROR_SUCCESS; +} + +//------------------------------------------------------------------------------ +// Get data //------------------------------------------------------------------------------ static int CeedQFunctionContextGetData_Hip(const CeedQFunctionContext ctx, const CeedMemType mtype, CeedScalar *data) { @@ -227,6 +276,8 @@ int CeedQFunctionContextCreate_Hip(CeedQFunctionContext ctx) { ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "SetData", CeedQFunctionContextSetData_Hip); CeedChkBackend(ierr); + ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "TakeData", + CeedQFunctionContextTakeData_Hip); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "GetData", CeedQFunctionContextGetData_Hip); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "RestoreData", diff --git a/backends/occa/ceed-occa-qfunctioncontext.cpp b/backends/occa/ceed-occa-qfunctioncontext.cpp index 0218938853..03169db543 100644 --- a/backends/occa/ceed-occa-qfunctioncontext.cpp +++ b/backends/occa/ceed-occa-qfunctioncontext.cpp @@ -150,9 +150,42 @@ namespace ceed { return ceedError("Invalid CeedMemType passed"); } + int QFunctionContext::takeData(CeedMemType mtype, + void *data) { + if (hostBuffer == NULL && memory == ::occa::null) + return ceedError("No context data set"); + switch (mtype) { + case CEED_MEM_HOST: + setCurrentHostCtxBufferIfNeeded(); + if (syncState == SyncState::device) { + setCurrentCtxMemoryIfNeeded(); + currentMemory.copyTo(currentHostBuffer); + } + syncState = SyncState::host; + *(void **)data = currentHostBuffer; + hostBuffer = NULL; + currentHostBuffer = NULL; + return CEED_ERROR_SUCCESS; + case CEED_MEM_DEVICE: + setCurrentCtxMemoryIfNeeded(); + if (syncState == SyncState::host) { + setCurrentHostCtxBufferIfNeeded(); + currentMemory.copyFrom(currentHostBuffer); + } + syncState = SyncState::device; + *(void **)data = memoryToData(currentMemory); + memory = ::occa::null; + currentMemory = ::occa::null; + return CEED_ERROR_SUCCESS; + } + return ceedError("Invalid CeedMemType passed"); + } + int QFunctionContext::getData(CeedMemType mtype, void *data) { // The passed `data` might be modified before restoring + if (hostBuffer == NULL && memory == ::occa::null) + return ceedError("No context data set"); switch (mtype) { case CEED_MEM_HOST: setCurrentHostCtxBufferIfNeeded(); @@ -203,6 +236,7 @@ namespace ceed { ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChk(ierr); CeedOccaRegisterFunction(ctx, "SetData", QFunctionContext::ceedSetData); + CeedOccaRegisterFunction(ctx, "TakeData", QFunctionContext::ceedTakeData); CeedOccaRegisterFunction(ctx, "GetData", QFunctionContext::ceedGetData); CeedOccaRegisterFunction(ctx, "RestoreData", QFunctionContext::ceedRestoreData); CeedOccaRegisterFunction(ctx, "Destroy", QFunctionContext::ceedDestroy); @@ -222,6 +256,15 @@ namespace ceed { return ctx_->setData(mtype, cmode, data); } + int QFunctionContext::ceedTakeData(CeedQFunctionContext ctx, CeedMemType mtype, + void *data) { + QFunctionContext *ctx_ = QFunctionContext::from(ctx); + if (!ctx_) { + return staticCeedError("Invalid CeedQFunctionContext passed"); + } + return ctx_->takeData(mtype, data); + } + int QFunctionContext::ceedGetData(CeedQFunctionContext ctx, CeedMemType mtype, void *data) { QFunctionContext *ctx_ = QFunctionContext::from(ctx); diff --git a/backends/occa/ceed-occa-qfunctioncontext.hpp b/backends/occa/ceed-occa-qfunctioncontext.hpp index 113ef63255..9185b8e9ba 100644 --- a/backends/occa/ceed-occa-qfunctioncontext.hpp +++ b/backends/occa/ceed-occa-qfunctioncontext.hpp @@ -72,6 +72,8 @@ namespace ceed { int useDataPointer(CeedMemType mtype, void *data); + int takeData(CeedMemType mtype, void *data); + int getData(CeedMemType mtype, void *data); int restoreData(); @@ -87,6 +89,9 @@ namespace ceed { static int ceedSetData(CeedQFunctionContext ctx, CeedMemType mtype, CeedCopyMode cmode, void *data); + static int ceedTakeData(CeedQFunctionContext ctx, CeedMemType mtype, + void *data); + static int ceedGetData(CeedQFunctionContext ctx, CeedMemType mtype, void *data); diff --git a/backends/ref/ceed-ref-qfunctioncontext.c b/backends/ref/ceed-ref-qfunctioncontext.c index 92659522e5..e5818925b9 100644 --- a/backends/ref/ceed-ref-qfunctioncontext.c +++ b/backends/ref/ceed-ref-qfunctioncontext.c @@ -55,6 +55,32 @@ static int CeedQFunctionContextSetData_Ref(CeedQFunctionContext ctx, return CEED_ERROR_SUCCESS; } +//------------------------------------------------------------------------------ +// QFunctionContext Take Data +//------------------------------------------------------------------------------ +static int CeedQFunctionContextTakeData_Ref(CeedQFunctionContext ctx, + CeedMemType mem_type, CeedScalar *data) { + int ierr; + CeedQFunctionContext_Ref *impl; + ierr = CeedQFunctionContextGetBackendData(ctx, (void *)&impl); + CeedChkBackend(ierr); + Ceed ceed; + ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); + + if (mem_type != CEED_MEM_HOST) + // LCOV_EXCL_START + return CeedError(ceed, CEED_ERROR_BACKEND, "Can only provide to HOST memory"); + // LCOV_EXCL_STOP + if (!impl->data) + // LCOV_EXCL_START + return CeedError(ceed, CEED_ERROR_BACKEND, "No context data set"); + // LCOV_EXCL_STOP + *(void **)data = impl->data; + impl->data = NULL; + impl->data_allocated = NULL; + return CEED_ERROR_SUCCESS; +} + //------------------------------------------------------------------------------ // QFunctionContext Get Data //------------------------------------------------------------------------------ @@ -110,6 +136,8 @@ int CeedQFunctionContextCreate_Ref(CeedQFunctionContext ctx) { ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "SetData", CeedQFunctionContextSetData_Ref); CeedChkBackend(ierr); + ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "TakeData", + CeedQFunctionContextTakeData_Ref); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "GetData", CeedQFunctionContextGetData_Ref); CeedChkBackend(ierr); ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "RestoreData", diff --git a/include/ceed-impl.h b/include/ceed-impl.h index 95a83d0764..34ce29bb2c 100644 --- a/include/ceed-impl.h +++ b/include/ceed-impl.h @@ -251,6 +251,7 @@ struct CeedQFunctionContext_private { Ceed ceed; int ref_count; int (*SetData)(CeedQFunctionContext, CeedMemType, CeedCopyMode, void *); + int (*TakeData)(CeedQFunctionContext, CeedMemType, void *); int (*GetData)(CeedQFunctionContext, CeedMemType, void *); int (*RestoreData)(CeedQFunctionContext); int (*Destroy)(CeedQFunctionContext); diff --git a/include/ceed/ceed.h b/include/ceed/ceed.h index 22cce7150e..98a36a0865 100644 --- a/include/ceed/ceed.h +++ b/include/ceed/ceed.h @@ -580,9 +580,10 @@ CEED_EXTERN int CeedQFunctionContextReferenceCopy(CeedQFunctionContext ctx, CeedQFunctionContext *ctx_copy); CEED_EXTERN int CeedQFunctionContextSetData(CeedQFunctionContext ctx, CeedMemType mem_type, CeedCopyMode copy_mode, size_t size, void *data); +CEED_EXTERN int CeedQFunctionContextTakeData(CeedQFunctionContext ctx, + CeedMemType mem_type, void *data); CEED_EXTERN int CeedQFunctionContextGetData(CeedQFunctionContext ctx, - CeedMemType mem_type, - void *data); + CeedMemType mem_type, void *data); CEED_EXTERN int CeedQFunctionContextRestoreData(CeedQFunctionContext ctx, void *data); CEED_EXTERN int CeedQFunctionContextView(CeedQFunctionContext ctx, diff --git a/interface/ceed-qfunctioncontext.c b/interface/ceed-qfunctioncontext.c index e6f2648dea..4fb3dcda2c 100644 --- a/interface/ceed-qfunctioncontext.c +++ b/interface/ceed-qfunctioncontext.c @@ -197,6 +197,7 @@ int CeedQFunctionContextReferenceCopy(CeedQFunctionContext ctx, @param ctx CeedQFunctionContext @param mem_type Memory type of the data being passed @param copy_mode Copy mode for the data + @param size Size of data, in bytes @param data Data to be used @return An error code: 0 - success, otherwise - failure @@ -227,6 +228,42 @@ int CeedQFunctionContextSetData(CeedQFunctionContext ctx, CeedMemType mem_type, return CEED_ERROR_SUCCESS; } +/** + @brief Take ownership of the data in a CeedQFunctionContext via the specified memory type. + The caller is responsible for managing and freeing the memory. + + @param ctx CeedQFunctionContext to access + @param mem_type Memory type on which to access the data. If the backend + uses a different memory type, this will perform a copy. + @param[out] data Data on memory type mem_type + + @return An error code: 0 - success, otherwise - failure + + @ref User +**/ +int CeedQFunctionContextTakeData(CeedQFunctionContext ctx, CeedMemType mem_type, + void *data) { + int ierr; + + if (!ctx->TakeData) + // LCOV_EXCL_START + return CeedError(ctx->ceed, CEED_ERROR_UNSUPPORTED, + "Backend does not support TakeData"); + // LCOV_EXCL_STOP + + if (ctx->state % 2 == 1) + // LCOV_EXCL_START + return CeedError(ctx->ceed, 1, + "Cannot grant CeedQFunctionContext data access, the " + "access lock is already in use"); + // LCOV_EXCL_STOP + + void *temp_data = NULL; + ierr = ctx->TakeData(ctx, mem_type, &temp_data); CeedChk(ierr); + if (data) (*(void **)data) = temp_data; + return CEED_ERROR_SUCCESS; +} + /** @brief Get read/write access to a CeedQFunctionContext via the specified memory type. Restore access with @ref CeedQFunctionContextRestoreData(). diff --git a/interface/ceed.c b/interface/ceed.c index de724c5e6e..37f83a0cd4 100644 --- a/interface/ceed.c +++ b/interface/ceed.c @@ -789,6 +789,7 @@ int CeedInit(const char *resource, Ceed *ceed) { CEED_FTABLE_ENTRY(CeedQFunction, SetHIPUserFunction), CEED_FTABLE_ENTRY(CeedQFunction, Destroy), CEED_FTABLE_ENTRY(CeedQFunctionContext, SetData), + CEED_FTABLE_ENTRY(CeedQFunctionContext, TakeData), CEED_FTABLE_ENTRY(CeedQFunctionContext, GetData), CEED_FTABLE_ENTRY(CeedQFunctionContext, RestoreData), CEED_FTABLE_ENTRY(CeedQFunctionContext, Destroy), diff --git a/tests/junit.py b/tests/junit.py index 650398baa9..4ce5b35389 100755 --- a/tests/junit.py +++ b/tests/junit.py @@ -116,6 +116,8 @@ def run(test, backends): check_required_failure(case, proc.stderr, 'Cannot destroy CeedElemRestriction, a process has read access to the offset data') if test[:4] in 't303'.split(): check_required_failure(case, proc.stderr, 'Length of input/output vectors incompatible with basis dimensions') + if test[:4] in 't404'.split(): + check_required_failure(case, proc.stderr, 'No context data set') if not case.is_skipped() and not case.status: if proc.stderr: diff --git a/tests/t404-qfunction.c b/tests/t404-qfunction.c new file mode 100644 index 0000000000..4458f3e64f --- /dev/null +++ b/tests/t404-qfunction.c @@ -0,0 +1,35 @@ +/// @file +/// Test creation, setting, and taking data for QFunctionContext +/// \test Test creation, setting, and taking data for QFunctionContext +#include +#include + +int main(int argc, char **argv) { + Ceed ceed; + CeedQFunctionContext ctx; + CeedScalar ctxData[5] = {1, 2, 3, 4, 5}, *ctxDataCopy; + + CeedInit(argv[1], &ceed); + + CeedQFunctionContextCreate(ceed, &ctx); + CeedQFunctionContextSetData(ctx, CEED_MEM_HOST, CEED_USE_POINTER, + sizeof(ctxData), &ctxData); + + CeedQFunctionContextGetData(ctx, CEED_MEM_HOST, &ctxDataCopy); + ctxDataCopy[4] = 6; + CeedQFunctionContextRestoreData(ctx, &ctxDataCopy); + if (fabs(ctxData[4] - 6) > 1.e-14) + // LCOV_EXCL_START + printf("error modifying data: %f != 6.0\n", ctxData[4]); + // LCOV_EXCL_STOP + + // Verify that taking the data revokes access + CeedQFunctionContextTakeData(ctx, CEED_MEM_HOST, &ctxDataCopy); + CeedQFunctionContextGetData(ctx, CEED_MEM_HOST, &ctxDataCopy); + + // LCOV_EXCL_START + CeedQFunctionContextDestroy(&ctx); + CeedDestroy(&ceed); + return 0; + // LCOV_EXCL_STOP +} diff --git a/tests/tap.sh b/tests/tap.sh index 8dae6ce4c4..739fa144dd 100755 --- a/tests/tap.sh +++ b/tests/tap.sh @@ -115,12 +115,12 @@ for ((i=0;i<${#backends[@]};++i)); do continue fi - # grep to pass test t215 on error - if grep -F -q -e 'access' ${output}.err \ - && [[ "$1" = "t215"* ]] ; then - printf "ok $i0 PASS - expected failure $1 $backend\n" - printf "ok $i1 PASS - expected failure $1 $backend stdout\n" - printf "ok $i2 PASS - expected failure $1 $backend stderr\n" + # grep to skip test if Device memory is not supported + if grep -F -q -e 'Can only provide to HOST memory' \ + ${output}.err ; then + printf "ok $i0 # SKIP - not supported $1 $backend\n" + printf "ok $i1 # SKIP - not supported $1 $backend stdout\n" + printf "ok $i2 # SKIP - not supported $1 $backend stderr\n" continue fi @@ -151,36 +151,45 @@ for ((i=0;i<${#backends[@]};++i)); do continue fi - # grep to pass test t303 on error - if grep -F -q -e 'vectors incompatible' ${output}.err \ - && [[ "$1" = "t303"* ]] ; then + # grep to pass test t215 on error + if grep -F -q -e 'access' ${output}.err \ + && [[ "$1" = "t215"* ]] ; then printf "ok $i0 PASS - expected failure $1 $backend\n" printf "ok $i1 PASS - expected failure $1 $backend stdout\n" printf "ok $i2 PASS - expected failure $1 $backend stderr\n" continue fi - # grep to skip test if Device memory is not supported - if grep -F -q -e 'Can only provide to HOST memory' \ - ${output}.err ; then - printf "ok $i0 # SKIP - not supported $1 $backend\n" - printf "ok $i1 # SKIP - not supported $1 $backend stdout\n" - printf "ok $i2 # SKIP - not supported $1 $backend stderr\n" + # grep to pass test t303 on error + if grep -F -q -e 'vectors incompatible' ${output}.err \ + && [[ "$1" = "t303"* ]] ; then + printf "ok $i0 PASS - expected failure $1 $backend\n" + printf "ok $i1 PASS - expected failure $1 $backend stdout\n" + printf "ok $i2 PASS - expected failure $1 $backend stderr\n" continue fi - # grep to skip t506 for MAGMA, range of basis kernels limited for now - if [[ "$backend" = *magma* ]] \ - && [[ "$1" = t506* ]] ; then + # grep to skip t318 for cuda/ref and MAGMA, Q is too large for these backends + if [[ "$backend" = *magma* || "$backend" = *cuda/ref ]] \ + && [[ "$1" = t318* ]] ; then printf "ok $i0 # SKIP - backend basis kernel not available $1 $backend\n" printf "ok $i1 # SKIP - backend basis kernel not available $1 $backend stdout\n" printf "ok $i2 # SKIP - backend basis kernel not available $1 $backend stderr\n" continue fi - # grep to skip t318 for cuda/ref and MAGMA, Q is too large for these backends - if [[ "$backend" = *magma* || "$backend" = *cuda/ref ]] \ - && [[ "$1" = t318* ]] ; then + # grep to pass test t404 on error + if grep -F -q -e 'No context data set' ${output}.err \ + && [[ "$1" = "t404"* ]] ; then + printf "ok $i0 PASS - expected failure $1 $backend\n" + printf "ok $i1 PASS - expected failure $1 $backend stdout\n" + printf "ok $i2 PASS - expected failure $1 $backend stderr\n" + continue + fi + + # grep to skip t506 for MAGMA, range of basis kernels limited for now + if [[ "$backend" = *magma* ]] \ + && [[ "$1" = t506* ]] ; then printf "ok $i0 # SKIP - backend basis kernel not available $1 $backend\n" printf "ok $i1 # SKIP - backend basis kernel not available $1 $backend stdout\n" printf "ok $i2 # SKIP - backend basis kernel not available $1 $backend stderr\n"