AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I8ee24930b5b3c026f2882742c7effa7b9a73bc1f


[ROCm/hip commit: 8b756e343f]
This commit is contained in:
Jenkins
2019-08-15 05:11:01 -04:00
parent 28382e282d 3c8f84a5c3
commit 78cc0f14b3
14 changed files with 154 additions and 76 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+18 -18
View File
@@ -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` | |
projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+3 -3
View File
@@ -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` | | |
projects/hip/docs/markdown/hip_porting_guide.md
+9
View File
@@ -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(<min>,<max>)))```.
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
projects/hip/hip_prof_gen.py
+11 -6
View File
@@ -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' +
projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+6 -6
View File
@@ -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}},
@@ -178,9 +178,9 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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}},
projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+15 -15
View File
@@ -754,31 +754,31 @@ const std::map<llvm::StringRef, hipCounter> 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<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
projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+2 -2
View File
@@ -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}},
projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+2 -2
View File
@@ -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
projects/hip/include/hip/hcc_detail/hip_runtime.h
-18
View File
@@ -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 <hip/hcc_detail/hip_memory.h>
projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+1
View File
@@ -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;
projects/hip/src/hip_module.cpp
+1
View File
@@ -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"
projects/hip/tests/README.md
+5 -2
View File
@@ -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 <hipcc_specific_options> HCC_OPTIONS <hcc_specific_options> NVCC_OPTIONS <nvcc_specific_options> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
BUILD: %t %s HIPCC_OPTIONS <hipcc_specific_options> HCC_OPTIONS <hcc_specific_options> NVCC_OPTIONS <nvcc_specific_options> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> DEPENDS <dependencies>
```
%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: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> DEPENDS <dependencies>
```
%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: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
%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
projects/hip/tests/hipify-clang/unit_tests/samples/MallocManaged.cpp
+45
View File
@@ -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;
}
projects/hip/tests/hit/HIT.cmake
+36 -4
View File
@@ -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()