[HIPIFY][HIP][6.2.0] Revise HIP APIs - Step 2 - Tests

+ Updated synthetic tests
emankov · Aug 8, 2024 · 0a8e1b6 · 0a8e1b6
1 parent 0ecf3bc
commit 0a8e1b6
Show file tree

Hide file tree

Showing 2 changed files with 41 additions and 1 deletion.
diff --git 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
@@ -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<typename UnaryFunction, class T> 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<typename UnaryFunction, class T> 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;