diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 7b66a484f2..67bf19b07e 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -190,19 +190,19 @@ | 0x01 |*`CU_FUNC_CACHE_PREFER_SHARED`* |*`hipFuncCachePreferShared`* | | 0x02 |*`CU_FUNC_CACHE_PREFER_L1`* |*`hipFuncCachePreferL1`* | | 0x03 |*`CU_FUNC_CACHE_PREFER_EQUAL`* |*`hipFuncCachePreferEqual`* | -| enum |***`CUfunction_attribute`*** | | -| typedef |***`CUfunction_attribute_enum`*** | | -| 0 |*`CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK`* | | -| 1 |*`CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES`* | | -| 2 |*`CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES`* | | -| 3 |*`CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES`* | | -| 4 |*`CU_FUNC_ATTRIBUTE_NUM_REGS`* | | -| 5 |*`CU_FUNC_ATTRIBUTE_PTX_VERSION`* | | -| 6 |*`CU_FUNC_ATTRIBUTE_BINARY_VERSION`* | | -| 7 |*`CU_FUNC_ATTRIBUTE_CACHE_MODE_CA`* | | -| 8 |*`CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES`* | | 9.0 | -| 9 |*`CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT`* | | 9.0 | -| 10 |*`CU_FUNC_ATTRIBUTE_MAX`* | | +| enum |***`CUfunction_attribute`*** |***`hipFunction_attribute`*** | +| typedef |***`CUfunction_attribute_enum`*** |***`hipFunction_attribute`*** | +| 0 |*`CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK`* |*`HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK`* | +| 1 |*`CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES`* |*`HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES`* | +| 2 |*`CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES`* |*`HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES`* | +| 3 |*`CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES`* |*`HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES`* | +| 4 |*`CU_FUNC_ATTRIBUTE_NUM_REGS`* |*`HIP_FUNC_ATTRIBUTE_NUM_REGS`* | +| 5 |*`CU_FUNC_ATTRIBUTE_PTX_VERSION`* |*`HIP_FUNC_ATTRIBUTE_PTX_VERSION`* | +| 6 |*`CU_FUNC_ATTRIBUTE_BINARY_VERSION`* |*`HIP_FUNC_ATTRIBUTE_BINARY_VERSION`* | +| 7 |*`CU_FUNC_ATTRIBUTE_CACHE_MODE_CA`* |*`HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA`* | +| 8 |*`CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES`* |*`HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES`* | 9.0 | +| 9 |*`CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT`* |*`HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT`* | 9.0 | +| 10 |*`CU_FUNC_ATTRIBUTE_MAX`* |*`HIP_FUNC_ATTRIBUTE_MAX`* | | enum |***`CUgraphicsMapResourceFlags`*** | | | typedef |***`CUgraphicsMapResourceFlags_enum`*** | | | 0x00 |*`CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE`* | | @@ -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,11 +858,11 @@ | `cuIpcOpenMemHandle` | `hipIpcOpenMemHandle` | | `cuMemAlloc` | `hipMalloc` | | `cuMemAllocHost` | | -| `cuMemAllocManaged` | | +| `cuMemAllocManaged` | `hipMemAllocManaged` | | `cuMemAllocPitch` | | | `cuMemcpy` | | | `cuMemcpy2D` | `hipMemcpyParam2D` | -| `cuMemcpy2DAsync` | | +| `cuMemcpy2DAsync` | `hipMemcpyParam2DAsync` | | `cuMemcpy2DUnaligned` | | | `cuMemcpy3D` | | | `cuMemcpy3DAsync` | | @@ -980,7 +980,7 @@ | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| -| `cuFuncGetAttribute` | | +| `cuFuncGetAttribute` | `hipFuncGetAttribute` | | `cuFuncSetAttribute` | | 9.0 | | `cuFuncSetCacheConfig` | `hipFuncSetCacheConfig` | | `cuFuncSetSharedMemConfig` | | diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index f35c451414..2b5ac181cc 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/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/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index a2b830ebf1..8f35ec776d 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -34,6 +34,7 @@ and provides practical suggestions on how to port CUDA code and work through com * [Choosing HIP File Extensions](#choosing-hip-file-extensions) - [Workarounds](#workarounds) * [warpSize](#warpsize) + * [Kernel launch with group size > 256](#kernel-launch-with-group-size--256) - [memcpyToSymbol](#memcpytosymbol) - [threadfence_system](#threadfence_system) * [Textures and Cache Control](#textures-and-cache-control) @@ -412,6 +413,14 @@ run hipcc when appropriate. ### warpSize Code should not assume a warp size of 32 or 64. See [Warp Cross-Lane Functions](hip_kernel_language.md#warp-cross-lane-functions) for information on how to write portable wave-aware code. +### Kernel launch with group size > 256 +Kernel code should use ``` __attribute__((amdgpu_flat_work_group_size(,)))```. + +For example: +``` +__global__ void dot(double *a,double *b,const int n) __attribute__((amdgpu_flat_work_group_size(1, 512))) +``` + ## memcpyToSymbol HIP support for hipMemcpyToSymbol is complete. This feature allows a kernel diff --git a/projects/hip/hip_prof_gen.py b/projects/hip/hip_prof_gen.py index e2876d87bc..6a1564cbb9 100755 --- a/projects/hip/hip_prof_gen.py +++ b/projects/hip/hip_prof_gen.py @@ -32,6 +32,11 @@ def fatal(msg): if errexit: sys.exit(1) ############################################################# +# Normalizing API name +def filtr_api_name(name): + name = re.sub(r'\s*$', r'', name); + return name + # Normalizing API arguments def filtr_api_args(args_str): args_str = re.sub(r'^\s*', r'', args_str); @@ -127,7 +132,7 @@ def parse_api(inp_file_p, out): if m: found = 0 if end_pattern.search(record): break - out[m.group(2)] = m.group(3) + out[filtr_api_name(m.group(2))] = m.group(3) else: continue hidden = 0 @@ -201,7 +206,7 @@ def parse_content(inp_file_p, api_map, out): # Checking if complete API matched if m: found = 2 - api_name = m.group(2); + api_name = filtr_api_name(m.group(2)); # Checking if API name is in the API map if api_name in api_map: # Getting API arguments @@ -353,10 +358,10 @@ def generate_prof_header(f, api_map, opts_map): if len(args) != 0: f.write(' struct {\n') for arg_tuple in args: - if arg_tuple[0] == "hipLimit_t": - f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') - else: - f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + if arg_tuple[0] == "hipLimit_t": + f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + else: + f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') f.write(' } ' + name + ';\n') f.write( ' } args;\n' + diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 0896f530e5..72534a96b6 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/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}}, @@ -178,9 +178,9 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuMemcpy2D", {"hipMemcpyParam2D", "", CONV_MEMORY, API_DRIVER}}, {"cuMemcpy2D_v2", {"hipMemcpyParam2D", "", CONV_MEMORY, API_DRIVER}}, // no analogue - // NOTE: Not equal to cudaMemcpy2DAsync due to different signatures - {"cuMemcpy2DAsync", {"hipMemcpy2DAsync_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuMemcpy2DAsync_v2", {"hipMemcpy2DAsync_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + // NOTE: Not equal to cudaMemcpy2DAsync/hipMemcpy2DAsync due to different signatures + {"cuMemcpy2DAsync", {"hipMemcpyParam2DAsync", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemcpy2DAsync_v2", {"hipMemcpyParam2DAsync", "", CONV_MEMORY, API_DRIVER}}, // no analogue {"cuMemcpy2DUnaligned", {"hipMemcpy2DUnaligned", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, {"cuMemcpy2DUnaligned_v2", {"hipMemcpy2DUnaligned", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, @@ -407,7 +407,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // 5.17.Execution Control // no analogue - {"cuFuncGetAttribute", {"hipFuncGetAttribute", "", CONV_EXECUTION, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuFuncGetAttribute", {"hipFuncGetAttribute", "", CONV_EXECUTION, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaFuncSetAttribute due to different signatures {"cuFuncSetAttribute", {"hipFuncSetAttribute", "", CONV_EXECUTION, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 717e86315d..7cfe2adc52 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -754,31 +754,31 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CU_FUNC_CACHE_PREFER_EQUAL", {"hipFuncCachePreferEqual", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x03 // cudaFuncAttribute - {"CUfunction_attribute", {"hipFuncAttribute", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUfunction_attribute_enum", {"hipFuncAttribute", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUfunction_attribute", {"hipFunction_attribute", "", CONV_TYPE, API_DRIVER}}, + {"CUfunction_attribute_enum", {"hipFunction_attribute", "", CONV_TYPE, API_DRIVER}}, // CUfunction_attribute enum values // no analogue - {"CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK", {"hipFuncAttributeMaxThreadsPerBlocks", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0 + {"CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK", {"HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0 // no analogue - {"CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES", {"hipFuncAttributeSharedSizeBytes", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 1 + {"CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES", {"HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1 // no analogue - {"CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES", {"hipFuncAttributeConstSizeBytes", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 2 + {"CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES", {"HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2 // no analogue - {"CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES", {"hipFuncAttributeLocalSizeBytes", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 3 + {"CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES", {"HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 3 // no analogue - {"CU_FUNC_ATTRIBUTE_NUM_REGS", {"hipFuncAttributeNumRegs", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 4 + {"CU_FUNC_ATTRIBUTE_NUM_REGS", {"HIP_FUNC_ATTRIBUTE_NUM_REGS", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 4 // no analogue - {"CU_FUNC_ATTRIBUTE_PTX_VERSION", {"hipFuncAttributePtxVersion", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 5 + {"CU_FUNC_ATTRIBUTE_PTX_VERSION", {"HIP_FUNC_ATTRIBUTE_PTX_VERSION", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 5 // no analogue - {"CU_FUNC_ATTRIBUTE_BINARY_VERSION", {"hipFuncAttributeBinaryVersion", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 6 + {"CU_FUNC_ATTRIBUTE_BINARY_VERSION", {"HIP_FUNC_ATTRIBUTE_BINARY_VERSION", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 6 // no analogue - {"CU_FUNC_ATTRIBUTE_CACHE_MODE_CA", {"hipFuncAttributeCacheModeCA", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 7 + {"CU_FUNC_ATTRIBUTE_CACHE_MODE_CA", {"HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 7 // cudaFuncAttributeMaxDynamicSharedMemorySize - {"CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES", {"hipFuncAttributeMaxDynamicSharedMemorySize", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 8 + {"CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES", {"HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8 // cudaFuncAttributePreferredSharedMemoryCarveout - {"CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT", {"hipFuncAttributePreferredSharedMemoryCarveout", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 9 + {"CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT", {"HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 9 // cudaFuncAttributeMax - {"CU_FUNC_ATTRIBUTE_MAX", {"hipFuncAttributeMax", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 10 + {"CU_FUNC_ATTRIBUTE_MAX", {"HIP_FUNC_ATTRIBUTE_MAX", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 10 //cudaGraphicsMapFlags {"CUgraphicsMapResourceFlags", {"hipGraphicsMapFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -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/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index bc31009285..f15220d0d1 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/projects/hip/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/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 7e153a223f..8ebefb38cf 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index c4ef1d829f..921d9951d8 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -462,24 +462,6 @@ hc_get_workitem_absolute_id(int dim) #undef __CUDA__ #pragma pop_macro("__CUDA__") -hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, - uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, - uint32_t localWorkSizeX, uint32_t localWorkSizeY, - uint32_t localWorkSizeZ, size_t sharedMemBytes, - hipStream_t hStream, void** kernelParams, void** extra, - hipEvent_t startEvent = nullptr, - hipEvent_t stopEvent = nullptr, - uint32_t flags = 0); - -hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, - uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, - uint32_t localWorkSizeX, uint32_t localWorkSizeY, - uint32_t localWorkSizeZ, size_t sharedMemBytes, - hipStream_t hStream, void** kernelParams, void** extra, - hipEvent_t startEvent = nullptr, - hipEvent_t stopEvent = nullptr) - __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); - #endif // defined(__clang__) && defined(__HIP__) #include diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index ec6ee855c7..e1ff25d9f2 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -160,6 +160,7 @@ typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; +typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; typedef cudaFuncAttributes hipFuncAttributes; typedef CUfunction_attribute hipFunction_attribute; diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 93751167ee..79e07161ac 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -25,6 +25,7 @@ THE SOFTWARE. #include "hip/hcc_detail/hsa_helpers.hpp" #include "hip/hcc_detail/program_state.hpp" #include "hip_hcc_internal.h" +#include "hip/hip_hcc.h" #include "program_state.inl" #include "trace_helper.h" diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md index 0c2c645085..e8f8116049 100644 --- a/projects/hip/tests/README.md +++ b/projects/hip/tests/README.md @@ -47,7 +47,7 @@ In the above, BUILD commands provide instructions on how to build the test case The supported syntax for the BUILD command is: ``` -BUILD: %t %s HIPCC_OPTIONS HCC_OPTIONS NVCC_OPTIONS EXCLUDE_HIP_PLATFORM +BUILD: %t %s HIPCC_OPTIONS HCC_OPTIONS NVCC_OPTIONS EXCLUDE_HIP_PLATFORM DEPENDS ``` %s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path). %t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified. @@ -55,13 +55,14 @@ HIPCC_OPTIONS: All options specified after this delimiter are passed to hipcc on HCC_OPTIONS: All options specified after this delimiter are passed to hipcc on HCC platform only. NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on NVCC platform only. EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. +DEPENDS: This can be used to specify dependencies that need to be built before building the current target. #### BUILD_CMD command The supported syntax for the BUILD_CMD command is: ``` -BUILD_CMD: EXCLUDE_HIP_PLATFORM +BUILD_CMD: EXCLUDE_HIP_PLATFORM DEPENDS ``` %s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path). %t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified. @@ -71,6 +72,8 @@ BUILD_CMD: EXCLUDE_HIP_PLATFORM %cxx: refers to system c compiler pointed to by /usr/bin/c++. %S: refers to path to current source file. %T: refers to path to current build target. +EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. +DEPENDS: This can be used to specify dependencies that need to be built before building the current target. #### TEST command diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp b/projects/hip/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp new file mode 100644 index 0000000000..d9122160b1 --- /dev/null +++ b/projects/hip/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; +} diff --git a/projects/hip/tests/hit/HIT.cmake b/projects/hip/tests/hit/HIT.cmake index d446742dd1..3d8607f0f5 100644 --- a/projects/hip/tests/hit/HIT.cmake +++ b/projects/hip/tests/hit/HIT.cmake @@ -3,7 +3,7 @@ find_package(HIP REQUIRED) #------------------------------------------------------------------------------- # Helper macro to parse BUILD instructions -macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _dir) +macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _depends _dir) set(${_target}) set(${_sources}) set(${_hipcc_options}) @@ -11,12 +11,14 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt set(${_nvcc_options}) set(${_link_options}) set(${_exclude_platforms}) + set(${_depends}) set(_target_found FALSE) set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) foreach(arg ${ARGN}) if(NOT _target_found) set(_target_found TRUE) @@ -27,30 +29,42 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt set(_nvcc_options_found FALSE) set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) elseif("x${arg}" STREQUAL "xHCC_OPTIONS") set(_hipcc_options_found FALSE) set(_hcc_options_found TRUE) set(_nvcc_options_found FALSE) set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) elseif("x${arg}" STREQUAL "xNVCC_OPTIONS") set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found TRUE) set(_link_options_found FALSE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) elseif("x${arg}" STREQUAL "xLINK_OPTIONS") set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) set(_link_options_found TRUE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM") set(_hipcc_options_found FALSE) set(_hcc_options_found FALSE) set(_nvcc_options_found FALSE) set(_link_options_found FALSE) set(_exclude_platforms_found TRUE) + set(_depends_found FALSE) + elseif("x${arg}" STREQUAL "xDEPENDS") + set(_hipcc_options_found FALSE) + set(_hcc_options_found FALSE) + set(_nvcc_options_found FALSE) + set(_link_options_found FALSE) + set(_exclude_platforms_found FALSE) + set(_depends_found TRUE) else() if(_hipcc_options_found) list(APPEND ${_hipcc_options} ${arg}) @@ -62,6 +76,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt list(APPEND ${_link_options} ${arg}) elseif(_exclude_platforms_found) set(${_exclude_platforms} ${arg}) + elseif(_depends_found) + list(APPEND ${_depends} ${arg}) else() list(APPEND ${_sources} "${_dir}/${arg}") endif() @@ -70,21 +86,29 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt endmacro() # Helper macro to parse CUSTOM BUILD instructions -macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms) +macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms _depends) set(${_target}) set(${_buildcmd} " ") set(${_exclude_platforms}) + set(${_depends}) set(_target_found FALSE) set(_exclude_platforms_found FALSE) + set(_depends_found FALSE) foreach(arg ${ARGN}) if(NOT _target_found) set(_target_found TRUE) set(${_target} ${arg}) elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM") set(_exclude_platforms_found TRUE) + set(_depends_found FALSE) + elseif("x${arg}" STREQUAL "xDEPENDS") + set(_exclude_platforms_found FALSE) + set(_depends_found TRUE) else() if(_exclude_platforms_found) set(${_exclude_platforms} ${arg}) + elseif(_depends_found) + list(APPEND ${_depends} ${arg}) else() list(APPEND ${_buildcmd} ${arg}) endif() @@ -179,7 +203,7 @@ macro(HIT_ADD_FILES _dir _label _parent) string(REGEX REPLACE "\n" ";" _contents "${_contents}") foreach(_cmd ${_contents}) string(REGEX REPLACE " " ";" _cmd "${_cmd}") - parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms ${_dir} ${_cmd}) + parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _depends ${_dir} ${_cmd}) string(REGEX REPLACE "/" "." target ${_label}/${_target}) insert_into_map("_exclude" "${target}" "${_exclude_platforms}") if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM}) @@ -190,6 +214,10 @@ macro(HIT_ADD_FILES _dir _label _parent) target_link_libraries(${target} PRIVATE ${_link_options}) set_target_properties(${target} PROPERTIES OUTPUT_NAME ${_target} RUNTIME_OUTPUT_DIRECTORY ${_label} LINK_DEPENDS "${HIP_LIB_FILES}") add_dependencies(${_parent} ${target}) + foreach(_dependency ${_depends}) + string(REGEX REPLACE "/" "." _dependency ${_label}/${_dependency}) + add_dependencies(${target} ${_dependency}) + endforeach() endif() endforeach() @@ -208,7 +236,7 @@ macro(HIT_ADD_FILES _dir _label _parent) string(REGEX REPLACE "%T" ${_label} _contents "${_contents}") foreach(_cmd ${_contents}) string(REGEX REPLACE " " ";" _cmd "${_cmd}") - parse_custombuild_command(_target _buildcmd _exclude_platforms ${_cmd}) + parse_custombuild_command(_target _buildcmd _exclude_platforms _depends ${_cmd}) string(REGEX REPLACE "/" "." target ${_label}/${_target}) insert_into_map("_exclude" "${target}" "${_exclude_platforms}") if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM}) @@ -219,6 +247,10 @@ macro(HIT_ADD_FILES _dir _label _parent) #add_custom_target(${target} COMMAND ${buildscript}) add_custom_target(${target} COMMAND sh -c "${_buildcmd}") add_dependencies(${_parent} ${target}) + foreach(_dependency ${_depends}) + string(REGEX REPLACE "/" "." _dependency ${_label}/${_dependency}) + add_dependencies(${target} ${_dependency}) + endforeach() endif() endforeach()