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/.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 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/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-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/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/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 65c96bd7f0..34ce29bb2c 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); @@ -250,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 d98bed3eec..98a36a0865 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); @@ -579,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-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-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-vector.c b/interface/ceed-vector.c index 07e8a5e495..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; /// @} @@ -531,12 +531,42 @@ int CeedVectorNorm(CeedVector vec, CeedNormType norm_type, CeedScalar *norm) { return CEED_ERROR_SUCCESS; } +/** + @brief Compute x = alpha x + + @param[in,out] x vector for scaling + @param[in] alpha 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 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 ) 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) diff --git a/python/ceed_vector.py b/python/ceed_vector.py index eb1f3c2d32..55aae97335 100644 --- a/python/ceed_vector.py +++ b/python/ceed_vector.py @@ -341,6 +341,16 @@ def reciprocal(self): return self + # Compute self = alpha self + def scale(self, alpha): + """Compute self = alpha self.""" + + # libCEED call + err_code = lib.CeedVectorScale(self._pointer[0], alpha) + self._ceed._check_error(err_code) + + return self + # Compute self = alpha x + self def axpy(self, alpha, x): """Compute self = alpha x + self.""" diff --git a/python/tests/test-1-vector.py b/python/tests/test-1-vector.py index 9ba97d30ca..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 @@ -289,6 +288,24 @@ def test_122(ceed_resource, capsys): for i in range(len(b)): assert abs(b[i] - i * i) < 1e-14 +# ------------------------------------------------------------------------------- +# Test Scale +# ------------------------------------------------------------------------------- + + +def test_123(ceed_resource, capsys): + ceed = libceed.Ceed(ceed_resource) + + n = 10 + x = ceed.Vector(n) + + a = np.arange(10, 10 + n, dtype="float64") + x.set_array(a, cmode=libceed.COPY_VALUES) + + x.scale(-0.5) + with x.array() as b: + assert np.allclose(-.5 * a, b) + # ------------------------------------------------------------------------------- # Test modification of reshaped array # ------------------------------------------------------------------------------- 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) + } } // ----------------------------------------------------------------------------- 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/t123-vector.c b/tests/t123-vector.c new file mode 100644 index 0000000000..d1a38e57f7 --- /dev/null +++ b/tests/t123-vector.c @@ -0,0 +1,36 @@ +/// @file +/// Test scaling a vector +/// \test Test scaling of a vector +#include +#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; +} 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"