AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_native_vector_types

Ver a proveniência
Este cometimento está contido em:
Alex Voicu
2018-07-15 11:44:48 +01:00
ascendente b3e6fcdf18 0c6fc28dba
cometimento fe289dbebe
19 ficheiros modificados com 158 adições e 60 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
bin/hipcc
+26 -3
Ver ficheiro
@@ -381,15 +381,35 @@ foreach $arg (@ARGV)
open my $out, ">", $new_file or die "$new_file: $!";
while (my $line = <$in>) {
chomp $line;
if ($line =~ m/\.a$/) {
if ($line =~ m/\.a$/ || $line =~ m/\.lo$/) {
my $libFile = $line;
my $path = abs_path($line);
my @objs = split ('\n', `cd $tmpdir; ar xv $path`);
## Check if all files in .a are object files.
my $allIsObj = 1;
my $realObjs = "";
foreach my $obj (@objs) {
chomp $obj;
$obj =~ s/^x - //;
$obj = "$tmpdir/$obj";
push (@inputs, $obj);
$new_arg = "$new_arg $obj";
my $fileType = `file $obj`;
my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
$allIsObj = ($allIsObj and $isObj);
if ($isObj) {
$realObjs = ($realObjs . " " . $obj);
} else {
push (@inputs, $obj);
$new_arg = "$new_arg $obj";
}
}
chomp $realObjs;
if ($allIsObj) {
print $out "$line\n";
} elsif ($realObjs) {
my($libBaseName, $libDir, $libExt) = fileparse($libFile);
$libBaseName = mktemp($libBaseName . "XXXX") . $libExt;
system("cd $tmpdir; ar c $libBaseName $realObjs");
print $out "$tmpdir/$libBaseName\n";
}
} else {
print $out "$line\n";
@@ -606,6 +626,9 @@ if ($needLDFLAGS and not $compileOnly) {
$CMD .= " $HIPLDFLAGS";
}
$CMD .= " $toolArgs";
if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") {
$CMD .= " -lgcc_s -lgcc";
}
if ($verbose & 0x1) {
print "hipcc-cmd: ", $CMD, "\n";
docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+6 -1
Ver ficheiro
@@ -231,6 +231,11 @@
| 0x02 |*`CU_MEMORYTYPE_DEVICE`* | |
| 0x03 |*`CU_MEMORYTYPE_ARRAY`* | |
| 0x04 |*`CU_MEMORYTYPE_UNIFIED`* | |
| enum |***`CUcomputemode`*** |***`hipComputeMode`*** |
| 0 |*`CU_COMPUTEMODE_DEFAULT`* |*`hipComputeModeDefault`* |
| 1 |*`CU_COMPUTEMODE_EXCLUSIVE`* |*`hipComputeModeExclusive`* |
| 2 |*`CU_COMPUTEMODE_PROHIBITED`* |*`hipComputeModeProhibited`* |
| 3 |*`CU_COMPUTEMODE_EXCLUSIVE_PROCESS`* |*`hipComputeModeExclusiveProcess`* |
| enum |***`CUoccupancy_flags`*** | |
| 0x00 |*`CU_OCCUPANCY_DEFAULT`* | |
| 0x01 |*`CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE`* | |
@@ -243,7 +248,7 @@
| 6 |*`CU_POINTER_ATTRIBUTE_SYNC_MEMOPS`* | |
| 7 |*`CU_POINTER_ATTRIBUTE_BUFFER_ID`* | |
| 8 |*`CU_POINTER_ATTRIBUTE_IS_MANAGED`* | |
| enum |***`CUmemorytype`*** | |
| enum |***`CUresourcetype`*** | |
| 0x00 |*`CU_RESOURCE_TYPE_ARRAY`* | |
| 0x01 |*`CU_RESOURCE_TYPE_MIPMAPPED_ARRAY`* | |
| 0x02 |*`CU_RESOURCE_TYPE_LINEAR`* | |
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+5 -5
Ver ficheiro
@@ -418,11 +418,11 @@
| 1 |*`cudaChannelFormatKindUnsigned`* |*`hipChannelFormatKindUnsigned`* |
| 2 |*`cudaChannelFormatKindFloat`* |*`hipChannelFormatKindFloat`* |
| 3 |*`cudaChannelFormatKindNone`* |*`hipChannelFormatKindNone`* |
| enum |***`cudaComputeMode`*** | |
| 0 |*`cudaComputeModeDefault`* | |
| 1 |*`cudaComputeModeExclusive`* | |
| 2 |*`cudaComputeModeProhibited`* | |
| 3 |*`cudaComputeModeExclusiveProcess`* | |
| enum |***`cudaComputeMode`*** |***`hipComputeMode`*** |
| 0 |*`cudaComputeModeDefault`* |*`hipComputeModeDefault`* |
| 1 |*`cudaComputeModeExclusive`* |*`hipComputeModeExclusive`* |
| 2 |*`cudaComputeModeProhibited`* |*`hipComputeModeProhibited`* |
| 3 |*`cudaComputeModeExclusiveProcess`* |*`hipComputeModeExclusiveProcess`* |
| enum |***`cudaDeviceAttr`*** |***`hipDeviceAttribute_t`*** |
| 1 |*`cudaDevAttrMaxThreadsPerBlock`* |*`hipDeviceAttributeMaxThreadsPerBlock`* |
| 2 |*`cudaDevAttrMaxBlockDimX`* |*`hipDeviceAttributeMaxBlockDimX`* |
docs/markdown/hip_deprecated_api_list.md
+17 -17
Ver ficheiro
@@ -1,22 +1,22 @@
# HIP Deprecated API List
# HIP Deprecated APIs
## HIP Context API
CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as separate entities. Contexts contain a single device, and a device can theoretically have multiple contexts. HIP initially added limited support for these API to facilitate easy porting from existing driver codes. These API are marked as deprecated now since there are better alternate interface (such as hipSetDevice or the stream API) to achieve the required functions.
###hipCtxCreate
###hipCtxDestroy
###hipCtxPopCurrent
###hipCtxPushCurrent
###hipCtxSetCurrent
###hipCtxGetCurrent
###hipCtxGetDevice
###hipCtxGetApiVersion
###hipCtxGetCacheConfig
###hipCtxSetCacheConfig
###hipCtxSetSharedMemConfig
###hipCtxGetSharedMemConfig
###hipCtxSynchronize
###hipCtxGetFlags
###hipCtxEnablePeerAccess
###hipCtxDisablePeerAccess
### hipCtxCreate
### hipCtxDestroy
### hipCtxPopCurrent
### hipCtxPushCurrent
### hipCtxSetCurrent
### hipCtxGetCurrent
### hipCtxGetDevice
### hipCtxGetApiVersion
### hipCtxGetCacheConfig
### hipCtxSetCacheConfig
### hipCtxSetSharedMemConfig
### hipCtxGetSharedMemConfig
### hipCtxSynchronize
### hipCtxGetFlags
### hipCtxEnablePeerAccess
### hipCtxDisablePeerAccess
hipify-clang/CMakeLists.txt
+5 -1
Ver ficheiro
@@ -51,6 +51,10 @@ if(WIN32)
target_link_libraries(hipify-clang version)
endif()
if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GREATER "7"))
target_link_libraries(hipify-clang clangToolingInclusions)
endif()
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
if(MSVC)
@@ -95,7 +99,7 @@ if (HIPIFY_CLANG_TESTS)
message(STATUS "Please install clang 4.0 or higher.")
elseif (CUDA_VERSION VERSION_EQUAL "9.0")
message(STATUS "Please install clang 6.0 or higher.")
elseif (CUDA_VERSION VERSION_EQUAL "9.1")
elseif ((CUDA_VERSION VERSION_EQUAL "9.1") OR (CUDA_VERSION VERSION_EQUAL "9.2"))
message(STATUS "Please install clang 7.0 or higher.")
endif()
endif()
hipify-clang/README.md
+16 -7
Ver ficheiro
@@ -31,7 +31,9 @@
| 4.0.1 | 8.0 |
| 5.0.0 | 8.0 |
| 5.0.1 | 8.0 |
| 5.0.2 | 8.0 |
| 6.0.0 | 9.0 |
| 6.0.1 | 9.0 |
In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
@@ -80,7 +82,7 @@ To run it:
cmake \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DCMAKE_BUILD_TYPE=Debug \
-DCMAKE_BUILD_TYPE=Release \
-Thost=x64 \
../llvm
@@ -114,9 +116,9 @@ To run it:
* Starting with LLVM 6.0.0 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option:
`-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py`,
`-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Release/bin/llvm-lit.py`,
where `f:/LLVM/6.0.0/build/Debug` is LLVM build directory.
where `f:/LLVM/6.0.0/build/Release` is LLVM build directory.
7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1.
8. `make test-hipify`
@@ -124,9 +126,15 @@ To run it:
### Windows
On Windows the following tested configuration is recommended:
On Windows the following configurations are tested:
LLVM 6.0.0 (exact), CUDA 9.0 (exact), cudnn-9.0 (exact), Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min).
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-9.0
LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0
Build system for the above configurations:
Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min).
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
@@ -134,13 +142,14 @@ Here is an example of building `hipify-clang` with testing support on `Windows 1
cmake
-G "Visual Studio 15 2017 Win64" \
-DHIPIFY_CLANG_TESTS=1 \
-DCMAKE_BUILD_TYPE=Debug \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=../dist \
-DCMAKE_PREFIX_PATH=f:/LLVM/6.0.0/dist \
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1 \
-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py \
-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Release/bin/llvm-lit.py \
-Thost=x64
..
```
A corresponding successful output:
hipify-clang/src/CUDA2HipMap.cpp
+10 -10
Ver ficheiro
@@ -24,7 +24,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
{"CUaddress_mode", {"hipAddress_mode", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUarray_cubemap_face", {"hipArray_cubemap_face", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUarray_format", {"hipArray_format", CONV_TYPE, API_DRIVER}},
{"CUcomputemode", {"hipComputemode", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaComputeMode)
{"CUcomputemode", {"hipComputeMode", CONV_TYPE, API_DRIVER}}, // API_RUNTIME ANALOGUE (cudaComputeMode)
{"CUmem_advise", {"hipMemAdvise", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaComputeMode)
{"CUmem_range_attribute", {"hipMemRangeAttribute", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaMemRangeAttribute)
{"CUctx_flags", {"hipCctx_flags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -236,7 +236,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
{"cudaDeviceAttr", {"hipDeviceAttribute_t", CONV_TYPE, API_RUNTIME}}, // API_DRIVER ANALOGUE (CUdevice_attribute)
{"cudaDeviceProp", {"hipDeviceProp_t", CONV_TYPE, API_RUNTIME}},
{"cudaDeviceP2PAttr", {"hipDeviceP2PAttribute", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // API_DRIVER ANALOGUE (CUdevice_P2PAttribute)
{"cudaComputeMode", {"hipComputeMode", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // API_DRIVER ANALOGUE (CUcomputemode)
{"cudaComputeMode", {"hipComputeMode", CONV_TYPE, API_RUNTIME}}, // API_DRIVER ANALOGUE (CUcomputemode)
{"cudaFuncCache", {"hipFuncCache_t", CONV_CACHE, API_RUNTIME}}, // API_Driver ANALOGUE (CUfunc_cache)
{"cudaFuncAttributes", {"hipFuncAttributes", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaSharedMemConfig", {"hipSharedMemConfig", CONV_TYPE, API_RUNTIME}},
@@ -628,10 +628,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
{"CU_AD_FORMAT_FLOAT", {"HIP_AD_FORMAT_FLOAT", CONV_TYPE, API_DRIVER}}, // 0x20
// CUcomputemode enum
{"CU_COMPUTEMODE_DEFAULT", {"hipComputeModeDefault", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0 // API_RUNTIME ANALOGUE (cudaComputeModeDefault = 0)
{"CU_COMPUTEMODE_EXCLUSIVE", {"hipComputeModeExclusive", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 1 // API_RUNTIME ANALOGUE (cudaComputeModeExclusive = 1)
{"CU_COMPUTEMODE_PROHIBITED", {"hipComputeModeProhibited", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 2 // API_RUNTIME ANALOGUE (cudaComputeModeProhibited = 2)
{"CU_COMPUTEMODE_EXCLUSIVE_PROCESS", {"hipComputeModeExclusiveProcess", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 3 // API_RUNTIME ANALOGUE (cudaComputeModeExclusiveProcess = 3)
{"CU_COMPUTEMODE_DEFAULT", {"hipComputeModeDefault", CONV_TYPE, API_DRIVER}}, // 0 // API_RUNTIME ANALOGUE (cudaComputeModeDefault = 0)
{"CU_COMPUTEMODE_EXCLUSIVE", {"hipComputeModeExclusive", CONV_TYPE, API_DRIVER}}, // 1 // API_RUNTIME ANALOGUE (cudaComputeModeExclusive = 1)
{"CU_COMPUTEMODE_PROHIBITED", {"hipComputeModeProhibited", CONV_TYPE, API_DRIVER}}, // 2 // API_RUNTIME ANALOGUE (cudaComputeModeProhibited = 2)
{"CU_COMPUTEMODE_EXCLUSIVE_PROCESS", {"hipComputeModeExclusiveProcess", CONV_TYPE, API_DRIVER}}, // 3 // API_RUNTIME ANALOGUE (cudaComputeModeExclusiveProcess = 3)
// Memory advise values
// {"CUmem_advise_enum", {"hipMemAdvise", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -1698,10 +1698,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
{"cudaDeviceGetP2PAttribute", {"hipDeviceGetP2PAttribute", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}}, // API_DRIVER ANALOGUE (cuDeviceGetP2PAttribute)
// enum cudaComputeMode
{"cudaComputeModeDefault", {"hipComputeModeDefault", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_DEFAULT = 0)
{"cudaComputeModeExclusive", {"hipComputeModeExclusive", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // 1 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_EXCLUSIVE = 1)
{"cudaComputeModeProhibited", {"hipComputeModeProhibited", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // 2 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_PROHIBITED = 2)
{"cudaComputeModeExclusiveProcess", {"hipComputeModeExclusiveProcess", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // 3 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_EXCLUSIVE_PROCESS = 3)
{"cudaComputeModeDefault", {"hipComputeModeDefault", CONV_TYPE, API_RUNTIME}}, // 0 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_DEFAULT = 0)
{"cudaComputeModeExclusive", {"hipComputeModeExclusive", CONV_TYPE, API_RUNTIME}}, // 1 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_EXCLUSIVE = 1)
{"cudaComputeModeProhibited", {"hipComputeModeProhibited", CONV_TYPE, API_RUNTIME}}, // 2 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_PROHIBITED = 2)
{"cudaComputeModeExclusiveProcess", {"hipComputeModeExclusiveProcess", CONV_TYPE, API_RUNTIME}}, // 3 // API_DRIVER ANALOGUE (CU_COMPUTEMODE_EXCLUSIVE_PROCESS = 3)
// Device Flags
{"cudaGetDeviceFlags", {"hipGetDeviceFlags", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}},
hipify-clang/src/HipifyAction.cpp
+5 -1
Ver ficheiro
@@ -428,7 +428,11 @@ public:
void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token,
StringRef file_name, bool is_angled, clang::CharSourceRange filename_range,
const clang::FileEntry* file, StringRef search_path, StringRef relative_path,
const clang::Module* imported) override {
const clang::Module* imported
#if LLVM_VERSION_MAJOR > 6
, clang::SrcMgr::CharacteristicKind FileType
#endif
) override {
hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported);
}
hipify-clang/src/LLVMCompat.h
+4
Ver ficheiro
@@ -23,6 +23,10 @@ namespace llcompat {
#define GET_NUM_ARGS() getNumArgs()
#endif
#if LLVM_VERSION_MAJOR < 7
#define LLVM_DEBUG(X) DEBUG(X)
#endif
void PrintStackTraceOnErrorSignal();
/**
hipify-clang/src/main.cpp
+1 -1
Ver ficheiro
@@ -132,7 +132,7 @@ int main(int argc, const char **argv) {
// Hipify _all_ the things!
if (Tool.runAndSave(&actionFactory)) {
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
}
// Either move the tmpfile to the output, or remove it.
include/hip/hcc_detail/hip_runtime_api.h
+7
Ver ficheiro
@@ -2621,6 +2621,13 @@ hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray
return ihipBindTextureToArrayImpl(dim, readMode, array, desc, &tex);
}
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
hipArray_const_t array,
const struct hipChannelFormatDesc* desc) {
return ihipBindTextureToArrayImpl(dim, readMode, array, *desc, tex);
}
// C API
hipError_t hipBindTextureToMipmappedArray(const textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
include/hip/hip_runtime_api.h
+6
Ver ficheiro
@@ -293,6 +293,12 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeIntegrated, ///< iGPU
} hipDeviceAttribute_t;
enum hipComputeMode {
hipComputeModeDefault = 0,
hipComputeModeExclusive = 1,
hipComputeModeProhibited = 2,
hipComputeModeExclusiveProcess = 3
};
/**
* @}
include/hip/nvcc_detail/hip_runtime_api.h
+22 -4
Ver ficheiro
@@ -163,6 +163,7 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
#define hipTextureType1D cudaTextureType1D
#define hipTextureType1DLayered cudaTextureType1DLayered
#define hipTextureType2D cudaTextureType2D
#define hipTextureType2DLayered cudaTextureType2DLayered
#define hipTextureType3D cudaTextureType3D
#define hipDeviceMapHost cudaDeviceMapHost
@@ -1168,20 +1169,20 @@ inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int
return hipCUDAErrorTohipError(cerror);
}
template <class T, int dim, enum cudaTextureReadMode readMode>
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
const void* devPtr, size_t size = UINT_MAX) {
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
}
template <class T, int dim, enum cudaTextureReadMode readMode>
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
const void* devPtr, const struct hipChannelFormatDesc& desc,
size_t size = UINT_MAX) {
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
}
template <class T, int dim, enum cudaTextureReadMode readMode>
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode>* tex) {
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
}
@@ -1198,7 +1199,14 @@ inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>&
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
}
template <class T, int dim, enum cudaTextureReadMode readMode>
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
hipArray_const_t array,
const struct hipChannelFormatDesc* desc) {
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc));
}
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
hipArray_const_t array) {
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
@@ -1239,6 +1247,16 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe
hipTextureObject_t textureObject) {
return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject));
}
inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref)
{
return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref));
}
inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)
{
return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array));
}
#endif //__CUDACC__
#endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H
packaging/hip-targets.cmake
+19 -1
Ver ficheiro
@@ -16,7 +16,7 @@ set(CMAKE_IMPORT_FILE_VERSION 1)
set(_targetsDefined)
set(_targetsNotDefined)
set(_expectedTargets)
foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device)
foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device hip::host hip::device)
list(APPEND _expectedTargets ${_expectedTarget})
if(NOT TARGET ${_expectedTarget})
list(APPEND _targetsNotDefined ${_expectedTarget})
@@ -65,6 +65,24 @@ set_target_properties(hip::hip_device PROPERTIES
INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
)
# Create imported target hip::host
add_library(hip::host INTERFACE IMPORTED)
set_target_properties(hip::host PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::hip_hcc"
)
# Create imported target hip::device
add_library(hip::device INTERFACE IMPORTED)
set_target_properties(hip::device PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am"
)
if(CMAKE_VERSION VERSION_LESS 3.0.0)
message(FATAL_ERROR "This file relies on consumers using CMake 3.0.0 or greater.")
endif()
# Load information for each installed configuration.
get_filename_component(_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH)
file(GLOB CONFIG_FILES "${_DIR}/hip-targets-*.cmake")
src/code_object_bundle.cpp
+2 -2
Ver ficheiro
@@ -38,7 +38,7 @@ std::string isa_name(std::string triple)
hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS};
if (is_old_rocr) {
auto tmp{triple.substr(triple.rfind('x') + 1)};
std::string tmp{triple.substr(triple.rfind('x') + 1)};
triple.replace(0, std::string::npos, "AMD:AMDGPU");
for (auto&& x : tmp) {
@@ -51,7 +51,7 @@ std::string isa_name(std::string triple)
}
hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) {
const auto isa{isa_name(std::move(triple))};
const std::string isa{isa_name(std::move(triple))};
if (isa.empty()) return hsa_isa_t({});
src/hip_hcc_internal.h
+1
Ver ficheiro
@@ -26,6 +26,7 @@ THE SOFTWARE.
#include <hc.hpp>
#include <hsa/hsa.h>
#include <unordered_map>
#include <stack>
#include "hsa/hsa_ext_amd.h"
#include "hip/hip_runtime.h"
src/hip_memory.cpp
+1 -4
Ver ficheiro
@@ -1680,12 +1680,9 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
actualDest = pinnedPtr;
}
}
#if 0
if((width == dpitch) && (width == spitch)) {
hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
} else
#endif
{
} else {
try {
if(!isLocked){
for (int i = 0; i < height; ++i)
src/hip_texture.cpp
+4 -2
Ver ficheiro
@@ -389,7 +389,8 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t
enum hipTextureFilterMode filterMode = tex->filterMode;
int normalizedCoords = tex->normalized;
hipTextureObject_t& textureObject = tex->textureObject;
*offset = 0;
if(offset != nullptr)
*offset = 0;
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
hc::accelerator acc = ctx->getDevice()->_acc;
@@ -459,7 +460,8 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size
enum hipTextureFilterMode filterMode = tex->filterMode;
int normalizedCoords = tex->normalized;
hipTextureObject_t& textureObject = tex->textureObject;
*offset = 0;
if(offset != nullptr)
*offset = 0;
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
hc::accelerator acc = ctx->getDevice()->_acc;
tests/src/deviceLib/hip_mbcnt.cpp
+1 -1
Ver ficheiro
@@ -88,7 +88,7 @@ int main() {
for (unsigned int i = 0; i < num_threads; i++) {
unsigned int this_lane_id = i % wave_size;
unsigned int this_mbcnt_lo = this_lane_id >= 32 ? 32 : this_lane_id;
unsigned int this_mbcnt_hi = this_lane_id < 32 ? 0 : (this_lane_id - 22);
unsigned int this_mbcnt_hi = this_lane_id < 32 ? 0 : (this_lane_id - 32);
if (host_mbcnt_lo[i] != this_mbcnt_lo)
mbcnt_lo_errors++;