From d20ae3b50aa0e4c49a43bbdbb411840b9dbc49bf Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 13 Aug 2019 17:56:06 +0300 Subject: [PATCH] [HIPIFY] Add cudaMallocManaged -> hipMallocManaged + Add mapping for corresponding data types + Add a test + Update docs --- ...A_Driver_API_functions_supported_by_HIP.md | 6 +-- ..._Runtime_API_functions_supported_by_HIP.md | 6 +-- .../src/CUDA2HIP_Driver_API_functions.cpp | 4 +- .../src/CUDA2HIP_Driver_API_types.cpp | 4 +- .../src/CUDA2HIP_Runtime_API_functions.cpp | 4 +- .../src/CUDA2HIP_Runtime_API_types.cpp | 4 +- .../unit_tests/samples/MallocManaged.cpp | 45 +++++++++++++++++++ 7 files changed, 59 insertions(+), 14 deletions(-) create mode 100644 tests/hipify-clang/unit_tests/samples/MallocManaged.cpp diff --git a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 061105caf7..67bf19b07e 100644 --- a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -299,8 +299,8 @@ | 6 |*`CU_MEM_ADVISE_UNSET_ACCESSED_BY`* | | 8.0 | | enum |***`CUmemAttach_flags`*** | | | typedef |***`CUmemAttach_flags_enum`*** | | -| 0x1 |*`CU_MEM_ATTACH_GLOBAL`* | | -| 0x2 |*`CU_MEM_ATTACH_HOST`* | | +| 0x1 |*`CU_MEM_ATTACH_GLOBAL`* |*`hipMemAttachGlobal`* | +| 0x2 |*`CU_MEM_ATTACH_HOST`* |*`hipMemAttachHost`* | | 0x4 |*`CU_MEM_ATTACH_SINGLE`* | | | enum |***`CUmemorytype`*** |*`hipMemoryType`* | | typedef |***`CUmemorytype_enum`*** |*`hipMemoryType`* | @@ -858,7 +858,7 @@ | `cuIpcOpenMemHandle` | `hipIpcOpenMemHandle` | | `cuMemAlloc` | `hipMalloc` | | `cuMemAllocHost` | | -| `cuMemAllocManaged` | | +| `cuMemAllocManaged` | `hipMemAllocManaged` | | `cuMemAllocPitch` | | | `cuMemcpy` | | | `cuMemcpy2D` | `hipMemcpyParam2D` | diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index f35c451414..2b5ac181cc 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -151,7 +151,7 @@ | `cudaMalloc3DArray` | `hipMalloc3DArray` | | `cudaMallocArray` | `hipMallocArray` | | `cudaMallocHost` | `hipHostMalloc` | -| `cudaMallocManaged` | | +| `cudaMallocManaged` | `hipMallocManaged` | | `cudaMallocMipmappedArray` | | | `cudaMallocPitch` | | | `cudaMemGetInfo` | `hipMemGetInfo` | @@ -1024,8 +1024,8 @@ | define |`cudaHostRegisterMapped` | |`hipHostRegisterMapped` | | define |`cudaHostRegisterPortable` | |`hipHostRegisterPortable` | | define |`cudaIpcMemLazyEnablePeerAccess` | |`hipIpcMemLazyEnablePeerAccess` | 0 | -| define |`cudaMemAttachGlobal` | | | -| define |`cudaMemAttachHost` | | | +| define |`cudaMemAttachGlobal` | |`hipMemAttachGlobal` | +| define |`cudaMemAttachHost` | |`hipMemAttachHost` | | define |`cudaMemAttachSingle` | | | | define |`cudaOccupancyDefault` | | | | define |`cudaOccupancyDisableCachingOverride` | | | diff --git a/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 224aa06394..72534a96b6 100644 --- a/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -164,8 +164,8 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaHostAlloc {"cuMemAllocHost", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, {"cuMemAllocHost_v2", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - // no analogue - {"cuMemAllocManaged", {"hipMemAllocManaged", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + // cudaMallocManaged + {"cuMemAllocManaged", {"hipMemAllocManaged", "", CONV_MEMORY, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaMallocPitch due to different signatures {"cuMemAllocPitch", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 53dd4c0d80..7cfe2adc52 100644 --- a/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -959,9 +959,9 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUmemAttach_flags_enum", {"hipMemAttachFlags_t", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // CUmemAttach_flags enum values // cudaMemAttachGlobal - {"CU_MEM_ATTACH_GLOBAL", {"hipMemAttachGlobal", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + {"CU_MEM_ATTACH_GLOBAL", {"hipMemAttachGlobal", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x1 // cudaMemAttachHost - {"CU_MEM_ATTACH_HOST", {"hipMemAttachHost", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x2 + {"CU_MEM_ATTACH_HOST", {"hipMemAttachHost", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x2 // cudaMemAttachSingle {"CU_MEM_ATTACH_SINGLE", {"hipMemAttachSingle", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x4 diff --git a/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index bc31009285..f15220d0d1 100644 --- a/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp @@ -271,8 +271,8 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaMallocArray", {"hipMallocArray", "", CONV_MEMORY, API_RUNTIME}}, // cuMemHostAlloc {"cudaMallocHost", {"hipHostMalloc", "", CONV_MEMORY, API_RUNTIME}}, - // no analogue - {"cudaMallocManaged", {"hipMallocManaged", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, + // cuMemAllocManaged + {"cudaMallocManaged", {"hipMallocManaged", "", CONV_MEMORY, API_RUNTIME}}, // no analogue // NOTE: Not equal to cuMipmappedArrayCreate due to different signatures {"cudaMallocMipmappedArray", {"hipMallocMipmappedArray", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, diff --git a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 7e153a223f..8ebefb38cf 100644 --- a/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -1323,9 +1323,9 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CU_IPC_MEM_LAZY_ENABLE_PEER_ACCESS {"cudaIpcMemLazyEnablePeerAccess", {"hipIpcMemLazyEnablePeerAccess", "", CONV_DEFINE, API_RUNTIME}}, // 0x01 // CU_MEM_ATTACH_GLOBAL - {"cudaMemAttachGlobal", {"hipMemAttachGlobal", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 + {"cudaMemAttachGlobal", {"hipMemAttachGlobal", "", CONV_DEFINE, API_RUNTIME}}, // 0x01 // CU_MEM_ATTACH_HOST - {"cudaMemAttachHost", {"hipMemAttachHost", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x02 + {"cudaMemAttachHost", {"hipMemAttachHost", "", CONV_DEFINE, API_RUNTIME}}, // 0x02 // CU_MEM_ATTACH_SINGLE {"cudaMemAttachSingle", {"hipMemAttachSingle", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x04 // no analogue diff --git a/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp b/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp new file mode 100644 index 0000000000..d9122160b1 --- /dev/null +++ b/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp @@ -0,0 +1,45 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include +#include + +__global__ +void add(int n, float *x, float *y) +{ + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) + y[i] = x[i] + y[i]; +} + +int main(int argc, char *argv[]) +{ + int numElements = 10; + bool testResult = true; + float *A, *B; + // CHECK: hipMallocManaged(&A, numElements * sizeof(float)); + cudaMallocManaged(&A, numElements * sizeof(float)); + // CHECK: hipMallocManaged(&B, numElements * sizeof(float)); + cudaMallocManaged(&B, numElements * sizeof(float)); + for (int i = 0; i < numElements; i++) { + A[i] = 1.0f; + B[i] = 2.0f; + } + int blockSize = 256; + int numBlocks = (numElements + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + // CHECK: hipLaunchKernelGGL(add, dim3(dimGrid), dim3(dimBlock), 0, 0, numElements, A, B); + add<<>>(numElements, A, B); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + float maxError = 0.0f; + for (int i = 0; i < numElements; i++) + maxError = fmax(maxError, fabs(B[i]-3.0f)); + // CHECK: hipFree(A); + cudaFree(A); + // CHECK: hipFree(B); + cudaFree(B); + if(maxError == 0.0f) + return 0; + return -1; +}