AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

[HIPIFY] Add cudaMallocManaged -> hipMallocManaged

Quellcode durchsuchen 
+ Add mapping for corresponding data types
+ Add a test
+ Update docs
Dieser Commit ist enthalten in:
Evgeny Mankov
2019-08-13 17:56:06 +03:00
Ursprung bb5fd89719
Commit d20ae3b50a
7 geänderte Dateien mit 59 neuen und 14 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+3 -3
Datei anzeigen
@@ -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` |
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+3 -3
Datei anzeigen
@@ -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` | | |
hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+2 -2
Datei anzeigen
@@ -164,8 +164,8 @@ const std::map<llvm::StringRef, hipCounter> 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}},
hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+2 -2
Datei anzeigen
@@ -959,9 +959,9 @@ const std::map<llvm::StringRef, hipCounter> 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
hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+2 -2
Datei anzeigen
@@ -271,8 +271,8 @@ const std::map<llvm::StringRef, hipCounter> 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}},
hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+2 -2
Datei anzeigen
@@ -1323,9 +1323,9 @@ const std::map<llvm::StringRef, hipCounter> 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
tests/hipify-clang/unit_tests/samples/MallocManaged.cpp
+45
Datei anzeigen
@@ -0,0 +1,45 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <math.h>
__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<<<dimGrid, dimBlock>>>(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;
}