Skip to content

Commit 1ace732

Browse files
jeremyltjeremylt
committed
vec - make CopyStrided reflect SetStrided
1 parent 1bee084 commit 1ace732

File tree

7 files changed

+46
-38
lines changed

7 files changed

+46
-38
lines changed

‎backends/cuda-ref/ceed-cuda-ref-vector.c‎

Lines changed: 9 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -223,20 +223,20 @@ static int CeedVectorSetArray_Cuda(const CeedVector vec, const CeedMemType mem_t
223223
//------------------------------------------------------------------------------
224224
// Copy host array to value strided
225225
//------------------------------------------------------------------------------
226-
static int CeedHostCopyStrided_Cuda(CeedScalar *h_array, CeedSize start, CeedSize step, CeedSize length, CeedScalar *h_copy_array) {
227-
for (CeedSize i = start; i < length; i += step) h_copy_array[i] = h_array[i];
226+
static int CeedHostCopyStrided_Cuda(CeedScalar *h_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *h_copy_array) {
227+
for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
228228
return CEED_ERROR_SUCCESS;
229229
}
230230

231231
//------------------------------------------------------------------------------
232232
// Copy device array to value strided (impl in .cu file)
233233
//------------------------------------------------------------------------------
234-
int CeedDeviceCopyStrided_Cuda(CeedScalar *d_array, CeedSize start, CeedSize step, CeedSize length, CeedScalar *d_copy_array);
234+
int CeedDeviceCopyStrided_Cuda(CeedScalar *d_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *d_copy_array);
235235

236236
//------------------------------------------------------------------------------
237237
// Copy a vector to a value strided
238238
//------------------------------------------------------------------------------
239-
static int CeedVectorCopyStrided_Cuda(CeedVector vec, CeedSize start, CeedSize step, CeedVector vec_copy) {
239+
static int CeedVectorCopyStrided_Cuda(CeedVector vec, CeedSize start, CeedSize stop, CeedSize step, CeedVector vec_copy) {
240240
CeedSize length;
241241
CeedVector_Cuda *impl;
242242

@@ -248,6 +248,7 @@ static int CeedVectorCopyStrided_Cuda(CeedVector vec, CeedSize start, CeedSize s
248248
CeedCallBackend(CeedVectorGetLength(vec_copy, &length_copy));
249249
length = length_vec < length_copy ? length_vec : length_copy;
250250
}
251+
if (stop == -1) stop = length;
251252
// Set value for synced device/host array
252253
if (impl->d_array) {
253254
CeedScalar *copy_array;
@@ -260,21 +261,21 @@ static int CeedVectorCopyStrided_Cuda(CeedVector vec, CeedSize start, CeedSize s
260261
CeedCallBackend(CeedVectorGetCeed(vec, &ceed));
261262
CeedCallBackend(CeedGetCublasHandle_Cuda(ceed, &handle));
262263
#if defined(CEED_SCALAR_IS_FP32)
263-
CeedCallCublas(ceed, cublasScopy_64(handle, (int64_t)length, impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
264+
CeedCallCublas(ceed, cublasScopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
264265
#else /* CEED_SCALAR */
265-
CeedCallCublas(ceed, cublasDcopy_64(handle, (int64_t)length, impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
266+
CeedCallCublas(ceed, cublasDcopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
266267
#endif /* CEED_SCALAR */
267268
CeedCallBackend(CeedDestroy(&ceed));
268269
#else /* CUDA_VERSION */
269-
CeedCallBackend(CeedDeviceCopyStrided_Cuda(impl->d_array, start, step, length, copy_array));
270+
CeedCallBackend(CeedDeviceCopyStrided_Cuda(impl->d_array, start, stop, step, copy_array));
270271
#endif /* CUDA_VERSION */
271272
CeedCallBackend(CeedVectorRestoreArray(vec_copy, &copy_array));
272273
impl->h_array = NULL;
273274
} else if (impl->h_array) {
274275
CeedScalar *copy_array;
275276

276277
CeedCallBackend(CeedVectorGetArray(vec_copy, CEED_MEM_HOST, &copy_array));
277-
CeedCallBackend(CeedHostCopyStrided_Cuda(impl->h_array, start, step, length, copy_array));
278+
CeedCallBackend(CeedHostCopyStrided_Cuda(impl->h_array, start, stop, step, copy_array));
278279
CeedCallBackend(CeedVectorRestoreArray(vec_copy, &copy_array));
279280
impl->d_array = NULL;
280281
} else {

‎backends/cuda-ref/kernels/cuda-ref-vector.cu‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -11,24 +11,24 @@
1111
//------------------------------------------------------------------------------
1212
// Kernel for copy strided on device
1313
//------------------------------------------------------------------------------
14-
__global__ static void copyStridedK(CeedScalar *__restrict__ vec, CeedSize start, CeedSize step, CeedSize size, CeedScalar *__restrict__ vec_copy) {
14+
__global__ static void copyStridedK(CeedScalar *__restrict__ vec, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *__restrict__ vec_copy) {
1515
const CeedSize index = threadIdx.x + (CeedSize)blockDim.x * blockIdx.x;
1616

17-
if (index < size) {
18-
if ((index - start) % step == 0) vec_copy[index] = vec[index];
17+
if (index < stop - start) {
18+
if (index % step == 0) vec_copy[start + index] = vec[start + index];
1919
}
2020
}
2121

2222
//------------------------------------------------------------------------------
2323
// Copy strided on device memory
2424
//------------------------------------------------------------------------------
25-
extern "C" int CeedDeviceCopyStrided_Cuda(CeedScalar *d_array, CeedSize start, CeedSize step, CeedSize length, CeedScalar *d_copy_array) {
25+
extern "C" int CeedDeviceCopyStrided_Cuda(CeedScalar *d_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *d_copy_array) {
2626
const int block_size = 512;
27-
const CeedSize vec_size = length;
28-
int grid_size = vec_size / block_size;
27+
const CeedSize copy_size = stop - start;
28+
int grid_size = copy_size / block_size;
2929

30-
if (block_size * grid_size < vec_size) grid_size += 1;
31-
copyStridedK<<<grid_size, block_size>>>(d_array, start, step, length, d_copy_array);
30+
if (block_size * grid_size < copy_size) grid_size += 1;
31+
copyStridedK<<<grid_size, block_size>>>(d_array, start, stop, step, d_copy_array);
3232
return 0;
3333
}
3434

‎backends/hip-ref/ceed-hip-ref-vector.c‎

Lines changed: 9 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -223,20 +223,20 @@ static int CeedVectorSetArray_Hip(const CeedVector vec, const CeedMemType mem_ty
223223
//------------------------------------------------------------------------------
224224
// Copy host array to value strided
225225
//------------------------------------------------------------------------------
226-
static int CeedHostCopyStrided_Hip(CeedScalar *h_array, CeedSize start, CeedSize step, CeedSize length, CeedScalar *h_copy_array) {
227-
for (CeedSize i = start; i < length; i += step) h_copy_array[i] = h_array[i];
226+
static int CeedHostCopyStrided_Hip(CeedScalar *h_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *h_copy_array) {
227+
for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
228228
return CEED_ERROR_SUCCESS;
229229
}
230230

231231
//------------------------------------------------------------------------------
232232
// Copy device array to value strided (impl in .hip.cpp file)
233233
//------------------------------------------------------------------------------
234-
int CeedDeviceCopyStrided_Hip(CeedScalar *d_array, CeedSize start, CeedSize step, CeedSize length, CeedScalar *d_copy_array);
234+
int CeedDeviceCopyStrided_Hip(CeedScalar *d_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *d_copy_array);
235235

236236
//------------------------------------------------------------------------------
237237
// Copy a vector to a value strided
238238
//------------------------------------------------------------------------------
239-
static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize step, CeedVector vec_copy) {
239+
static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize stop, CeedSize step, CeedVector vec_copy) {
240240
CeedSize length;
241241
CeedVector_Hip *impl;
242242

@@ -248,6 +248,7 @@ static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize st
248248
CeedCallBackend(CeedVectorGetLength(vec_copy, &length_copy));
249249
length = length_vec < length_copy ? length_vec : length_copy;
250250
}
251+
if (stop == -1) stop = length;
251252
// Set value for synced device/host array
252253
if (impl->d_array) {
253254
CeedScalar *copy_array;
@@ -260,12 +261,12 @@ static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize st
260261
CeedCallBackend(CeedVectorGetCeed(vec, &ceed));
261262
CeedCallBackend(CeedGetHipblasHandle_Hip(ceed, &handle));
262263
#if defined(CEED_SCALAR_IS_FP32)
263-
CeedCallHipblas(ceed, hipblasScopy_64(handle, (int64_t)length, impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
264+
CeedCallHipblas(ceed, hipblasScopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
264265
#else /* CEED_SCALAR */
265-
CeedCallHipblas(ceed, hipblasDcopy_64(handle, (int64_t)length, impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
266+
CeedCallHipblas(ceed, hipblasDcopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int64_t)step, copy_array + start, (int64_t)step));
266267
#endif /* CEED_SCALAR */
267268
#else /* HIP_VERSION */
268-
CeedCallBackend(CeedDeviceCopyStrided_Hip(impl->d_array, start, step, length, copy_array));
269+
CeedCallBackend(CeedDeviceCopyStrided_Hip(impl->d_array, start, stop, step, copy_array));
269270
#endif /* HIP_VERSION */
270271
CeedCallBackend(CeedVectorRestoreArray(vec_copy, &copy_array));
271272
impl->h_array = NULL;
@@ -274,7 +275,7 @@ static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize st
274275
CeedScalar *copy_array;
275276

276277
CeedCallBackend(CeedVectorGetArray(vec_copy, CEED_MEM_HOST, &copy_array));
277-
CeedCallBackend(CeedHostCopyStrided_Hip(impl->h_array, start, step, length, copy_array));
278+
CeedCallBackend(CeedHostCopyStrided_Hip(impl->h_array, start, stop, step, copy_array));
278279
CeedCallBackend(CeedVectorRestoreArray(vec_copy, &copy_array));
279280
impl->d_array = NULL;
280281
} else {

‎include/ceed-impl.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -137,7 +137,7 @@ struct CeedVector_private {
137137
Ceed ceed;
138138
int (*HasValidArray)(CeedVector, bool *);
139139
int (*HasBorrowedArrayOfType)(CeedVector, CeedMemType, bool *);
140-
int (*CopyStrided)(CeedVector, CeedSize, CeedSize, CeedVector);
140+
int (*CopyStrided)(CeedVector, CeedSize, CeedSize, CeedSize, CeedVector);
141141
int (*SetArray)(CeedVector, CeedMemType, CeedCopyMode, CeedScalar *);
142142
int (*SetValue)(CeedVector, CeedScalar);
143143
int (*SetValueStrided)(CeedVector, CeedSize, CeedSize, CeedSize, CeedScalar);

‎include/ceed/ceed.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -181,7 +181,7 @@ CEED_EXTERN int CeedGetPreferredMemType(Ceed ceed, CeedMemType *type);
181181
CEED_EXTERN int CeedVectorCreate(Ceed ceed, CeedSize len, CeedVector *vec);
182182
CEED_EXTERN int CeedVectorReferenceCopy(CeedVector vec, CeedVector *vec_copy);
183183
CEED_EXTERN int CeedVectorCopy(CeedVector vec, CeedVector vec_copy);
184-
CEED_EXTERN int CeedVectorCopyStrided(CeedVector vec, CeedSize start, CeedInt step, CeedVector vec_copy);
184+
CEED_EXTERN int CeedVectorCopyStrided(CeedVector vec, CeedSize start, CeedSize stop, CeedSize step, CeedVector vec_copy);
185185
CEED_EXTERN int CeedVectorSetArray(CeedVector vec, CeedMemType mem_type, CeedCopyMode copy_mode, CeedScalar *array);
186186
CEED_EXTERN int CeedVectorSetValue(CeedVector vec, CeedScalar value);
187187
CEED_EXTERN int CeedVectorSetValueStrided(CeedVector vec, CeedSize start, CeedSize stop, CeedSize step, CeedScalar value);

‎interface/ceed-vector.c‎

Lines changed: 17 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -251,27 +251,21 @@ int CeedVectorCopy(CeedVector vec, CeedVector vec_copy) {
251251
@brief Copy a strided portion of `CeedVector` contents into a different `CeedVector`
252252
253253
@param[in] vec `CeedVector` to copy
254-
@param[in] start First index to copy
254+
@param[in] start First index to copy in the range `[start, stop)`
255+
@param[in] stop One past the last element to copy in the range, or `-1` for `length`
255256
@param[in] step Stride between indices to copy
256257
@param[in,out] vec_copy `CeedVector` to copy values to
257258
258259
@return An error code: 0 - success, otherwise - failure
259260
260261
@ref User
261262
**/
262-
int CeedVectorCopyStrided(CeedVector vec, CeedSize start, CeedInt step, CeedVector vec_copy) {
263+
int CeedVectorCopyStrided(CeedVector vec, CeedSize start, CeedSize stop, CeedSize step, CeedVector vec_copy) {
263264
CeedSize length;
264265
const CeedScalar *array = NULL;
265266
CeedScalar *array_copy = NULL;
266267

267-
// Backend version
268-
if (vec->CopyStrided && vec_copy->CopyStrided) {
269-
CeedCall(vec->CopyStrided(vec, start, step, vec_copy));
270-
vec_copy->state += 2;
271-
return CEED_ERROR_SUCCESS;
272-
}
273-
274-
// Get length
268+
// Check length
275269
{
276270
CeedSize length_vec, length_copy;
277271

@@ -280,11 +274,23 @@ int CeedVectorCopyStrided(CeedVector vec, CeedSize start, CeedInt step, CeedVect
280274
if (length_vec <= 0 || length_copy <= 0) return CEED_ERROR_SUCCESS;
281275
length = length_vec < length_copy ? length_vec : length_copy;
282276
}
277+
CeedCheck(stop >= -1 && stop <= length, CeedVectorReturnCeed(vec), CEED_ERROR_ACCESS,
278+
"Invalid value for stop %" CeedSize_FMT ", must be in the range [-1, length]", stop);
279+
CeedCheck(start >= 0 && start <= length && (start <= stop || stop == -1), CeedVectorReturnCeed(vec), CEED_ERROR_ACCESS,
280+
"Invalid value for start %" CeedSize_FMT ", must be in the range [0, stop]", start);
281+
282+
// Backend version
283+
if (vec->CopyStrided && vec_copy->CopyStrided) {
284+
CeedCall(vec->CopyStrided(vec, start, stop, step, vec_copy));
285+
vec_copy->state += 2;
286+
return CEED_ERROR_SUCCESS;
287+
}
283288

284289
// Copy
285290
CeedCall(CeedVectorGetArrayRead(vec, CEED_MEM_HOST, &array));
286291
CeedCall(CeedVectorGetArray(vec_copy, CEED_MEM_HOST, &array_copy));
287-
for (CeedSize i = start; i < length; i += step) array_copy[i] = array[i];
292+
if (stop == -1) stop = length;
293+
for (CeedSize i = start; i < stop; i += step) array_copy[i] = array[i];
288294

289295
// Cleanup
290296
CeedCall(CeedVectorRestoreArrayRead(vec, &array));

‎tests/t127-vector.c‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ int main(int argc, char **argv) {
3636

3737
// Copy strided
3838
CeedVectorSetValue(y, 0.0);
39-
CeedVectorCopyStrided(x, start, step, y);
39+
CeedVectorCopyStrided(x, start, -1, step, y);
4040
{
4141
const CeedScalar *read_array;
4242

0 commit comments

Comments
 (0)