From 0a8e1b62d794d48d682ee5b916318f14fda3b421 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 8 Aug 2024 17:01:44 +0100 Subject: [PATCH] [HIPIFY][HIP][6.2.0] Revise HIP APIs - Step 2 - Tests + Updated synthetic tests --- .../unit_tests/synthetic/driver_functions.cu | 30 ++++++++++++++++++- .../unit_tests/synthetic/runtime_functions.cu | 12 ++++++++ 2 files changed, 41 insertions(+), 1 deletion(-) diff --git a/tests/unit_tests/synthetic/driver_functions.cu b/tests/unit_tests/synthetic/driver_functions.cu index 347f2a1e..14c53578 100644 --- a/tests/unit_tests/synthetic/driver_functions.cu +++ b/tests/unit_tests/synthetic/driver_functions.cu @@ -993,13 +993,41 @@ int main() { // CHECK: result = hipTexRefGetArray(&array_, texref); result = cuTexRefGetArray(&array_, texref); - // CUDA:CUresult CUDAAPI cuMemcpyAtoA_v2(CUarray dstArray, size_t dstOffset, CUarray srcArray, size_t srcOffset, size_t ByteCount); + // CUDA: CUresult CUDAAPI cuMemcpyAtoA_v2(CUarray dstArray, size_t dstOffset, CUarray srcArray, size_t srcOffset, size_t ByteCount); // HIP: hipError_t hipMemcpyAtoA(hipArray_t dstArray, size_t dstOffset, hipArray_t srcArray, size_t srcOffset, size_t ByteCount); // CHECK: result = hipMemcpyAtoA(array_dst, offset_dst, array_, offset, bytes); // CHECK-NEXT: result = hipMemcpyAtoA(array_dst, offset_dst, array_, offset, bytes); result = cuMemcpyAtoA(array_dst, offset_dst, array_, offset, bytes); result = cuMemcpyAtoA_v2(array_dst, offset_dst, array_, offset, bytes); + // CUDA: CUresult CUDAAPI cuMemcpyAtoD_v2(CUdeviceptr dstDevice, CUarray srcArray, size_t srcOffset, size_t ByteCount); + // HIP: hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, hipArray_t srcArray, size_t srcOffset, size_t ByteCount); + // CHECK: result = hipMemcpyAtoD(deviceptr, array_, offset, bytes); + // CHECK-NEXT: result = hipMemcpyAtoD(deviceptr, array_, offset, bytes); + result = cuMemcpyAtoD(deviceptr, array_, offset, bytes); + result = cuMemcpyAtoD_v2(deviceptr, array_, offset, bytes); + + // CUDA: CUresult CUDAAPI cuMemcpyDtoA_v2(CUarray dstArray, size_t dstOffset, CUdeviceptr srcDevice, size_t ByteCount); + // HIP: hipError_t hipMemcpyDtoA(hipArray_t dstArray, size_t dstOffset, hipDeviceptr_t srcDevice, size_t ByteCount); + // CHECK: result = hipMemcpyDtoA(array_, offset, deviceptr, bytes); + // CHECK-NEXT: result = hipMemcpyDtoA(array_, offset, deviceptr, bytes); + result = cuMemcpyDtoA(array_, offset, deviceptr, bytes); + result = cuMemcpyDtoA_v2(array_, offset, deviceptr, bytes); + + // CUDA: CUresult CUDAAPI cuMemcpyAtoHAsync_v2(void *dstHost, CUarray srcArray, size_t srcOffset, size_t ByteCount, CUstream hStream); + // HIP: hipError_t hipMemcpyAtoHAsync(void* dstHost, hipArray_t srcArray, size_t srcOffset, size_t ByteCount, hipStream_t stream); + // CHECK: result = hipMemcpyAtoHAsync(dsthost, array_, offset, bytes, stream); + // CHECK-NEXT: result = hipMemcpyAtoHAsync(dsthost, array_, offset, bytes, stream); + result = cuMemcpyAtoHAsync(dsthost, array_, offset, bytes, stream); + result = cuMemcpyAtoHAsync_v2(dsthost, array_, offset, bytes, stream); + + // CUDA: CUresult CUDAAPI cuMemcpyHtoAAsync_v2(CUarray dstArray, size_t dstOffset, const void *srcHost, size_t ByteCount, CUstream hStream); + // HIP: hipError_t hipMemcpyHtoAAsync(hipArray_t dstArray, size_t dstOffset, const void* srcHost, size_t ByteCount, hipStream_t stream); + // CHECK: result = hipMemcpyHtoAAsync(array_, offset, dsthost, bytes, stream); + // CHECK-NEXT: result = hipMemcpyHtoAAsync(array_, offset, dsthost, bytes, stream); + result = cuMemcpyHtoAAsync(array_, offset, dsthost, bytes, stream); + result = cuMemcpyHtoAAsync_v2(array_, offset, dsthost, bytes, stream); + #if CUDA_VERSION >= 8000 // CHECK: hipMemRangeAttribute MemoryRangeAttribute; // CHECK-NEXT: hipMemoryAdvise MemoryAdvise; diff --git a/tests/unit_tests/synthetic/runtime_functions.cu b/tests/unit_tests/synthetic/runtime_functions.cu index 80f743fd..f9c92a39 100644 --- a/tests/unit_tests/synthetic/runtime_functions.cu +++ b/tests/unit_tests/synthetic/runtime_functions.cu @@ -25,7 +25,9 @@ int main() { size_t width = 0; size_t height = 0; size_t wOffset = 0; + size_t wOffset_src = 0; size_t hOffset = 0; + size_t hOffset_src = 0; size_t pitch = 0; size_t pitch_2 = 0; int device = 0; @@ -825,6 +827,16 @@ int main() { // CUDA: template static __inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(int* minGridSize, int* blockSize, T func, UnaryFunction blockSizeToDynamicSMemSize, int blockSizeLimit = 0, unsigned int flags = 0); // HIP: template static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(int* min_grid_size, int* block_size, T func, UnaryFunction block_size_to_dynamic_smem_size, int block_size_limit = 0, unsigned int flags = 0); + // CUDA: extern __host__ cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len); + // HIP: hipError_t hipSetValidDevices(int* device_arr, int len); + // CHECK: result = hipSetValidDevices(&device, intVal); + result = cudaSetValidDevices(&device, intVal); + + // CUDA: extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(cudaArray_t dst, size_t wOffsetDst, size_t hOffsetDst, cudaArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)); + // HIP: hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind); + // CHECK: result = hipMemcpy2DArrayToArray(Array_t, wOffset, hOffset, Array_const_t, wOffset_src, hOffset_src, width, height, MemcpyKind); + result = cudaMemcpy2DArrayToArray(Array_t, wOffset, hOffset, Array_const_t, wOffset_src, hOffset_src, width, height, MemcpyKind); + #if CUDA_VERSION >= 8000 // CHECK: hipDeviceP2PAttr DeviceP2PAttr; cudaDeviceP2PAttr DeviceP2PAttr;