AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

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

ソースを参照 
[ROCm/clr commit: b9544eabe8]
このコミットが含まれているのは:
Alex Voicu
2020-03-12 14:20:26 +00:00
親 1d009eaee4 666f217a6a
コミット 8acb4eaa33
7個のファイルの変更、60行の追加、18行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
projects/clr/hipamd/README.md
+2 -1
ファイルの表示
@@ -1,6 +1,7 @@
## What is this repository for? ###
HIP allows developers to convert CUDA code to portable C++. The same source code can be compiled to run on NVIDIA or AMD GPUs.
**HIP is a C++ Runtime API and Kernel Language that allows developers to create portable applications for AMD and NVIDIA GPUs from single source code.**
Key features include:
* HIP is very thin and has little or no performance impact over coding directly in CUDA or hcc "HC" mode.
projects/clr/hipamd/bin/hipcc
+1
ファイルの表示
@@ -133,6 +133,7 @@ if (defined $HIP_RUNTIME and $HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
} else {
$HIP_VDI_HOME = $HIP_PATH; # use HIP_PATH
}
$HIPCXXFLAGS .= "-D__HIP_VDI__";
}
if (defined $HIP_VDI_HOME) {
projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
+10 -1
ファイルの表示
@@ -33,6 +33,15 @@ THE SOFTWARE.
#include <hip/hip_vector_types.h>
#include <hip/hcc_detail/device_library_decls.h>
#include <hip/hcc_detail/llvm_intrinsics.h>
#if __HIP_CLANG_ONLY__
#if __HIP_VDI__
extern "C" __device__ int printf(const char *fmt, ...);
#else
static inline __device__ void printf(const char* format, All... all) {}
#endif
#endif
/*
Integer Intrinsics
*/
@@ -548,7 +557,7 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize)
return tmp1;
#else
static_assert(sizeof(long) == sizeof(int), "");
return static_cast<long>(__shfl_down(static_cast<int>(var), lane_delta, width));
return static_cast<long>(__shfl_xor(static_cast<int>(var), lane_mask, width));
#endif
}
__device__
projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h
+1
ファイルの表示
@@ -72,6 +72,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(
extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void);
extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void);
extern "C" __device__ void __ockl_global_atomic_add_f32(__attribute__((address_space(1))) float*, float);
// Introduce local address space
#define __local __attribute__((address_space(3)))
projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h
+8
ファイルの表示
@@ -73,6 +73,14 @@ float atomicAdd(float* address, float val)
return __uint_as_float(r);
}
__device__
inline
void atomicAddNoRet(float* address, float val)
{
__ockl_global_atomic_add_f32((__attribute__((address_space(1))) float*)address, val);
}
__device__
inline
double atomicAdd(double* address, double val)
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+7 -3
ファイルの表示
@@ -311,15 +311,19 @@ extern "C" __device__ void* __hip_free(void* ptr);
static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF)
// Declare printf only for the HCC compiler. hip-clang is handled in
// device_functions.h
#if __HCC_ACCELERATOR__
#if HC_FEATURE_PRINTF
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {
hc::printf(format, all...);
}
#elif defined(__HCC_ACCELERATOR__) || __HIP__
#else
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {}
#endif
#endif // HC_FEATURE_PRINTF
#endif // __HCC_ACCELERATOR__
#endif //__HCC_OR_HIP_CLANG__
projects/clr/hipamd/src/hip_memory.cpp
+31 -13
ファイルの表示
@@ -478,6 +478,10 @@ void* allocAndSharePtr(const char* msg, size_t sizeBytes, ihipCtx_t* ctx, bool s
hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned int flags) {
hipError_t hip_status = hipSuccess;
if (sizeBytes == 0) {
return hipSuccess;
}
if (HIP_SYNC_HOST_ALLOC) {
hipDeviceSynchronize();
}
@@ -485,10 +489,6 @@ hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned i
auto ctx = ihipGetTlsDefaultCtx();
if ((ctx == nullptr) || (ptr == nullptr)) {
hip_status = hipErrorInvalidValue;
}
else if (sizeBytes == 0) {
hip_status = hipSuccess;
// TODO - should size of 0 return err or be siliently ignored?
} else {
unsigned trueFlags = flags;
if (flags == hipHostMallocDefault) {
@@ -673,14 +673,15 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if (sizeBytes == 0) {
if (ptr) *ptr = NULL;
return ihipLogStatus(hipSuccess);
}
auto ctx = ihipGetTlsDefaultCtx();
// return NULL pointer when malloc size is 0
if ( nullptr == ctx || nullptr == ptr) {
hip_status = hipErrorInvalidValue;
}
else if (sizeBytes == 0) {
*ptr = NULL;
hip_status = hipSuccess;
} else {
auto device = ctx->getWriteableDevice();
*ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/,
@@ -700,14 +701,15 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
HIP_SET_DEVICE();
#if (__hcc_workweek__ >= 19115)
if (sizeBytes == 0) {
if (ptr) *ptr = NULL;
return ihipLogStatus(hipSuccess);
}
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
// return NULL pointer when malloc size is 0
if (sizeBytes == 0) {
*ptr = NULL;
hip_status = hipSuccess;
} else if ((ctx == nullptr) || (ptr == nullptr)) {
if ((ctx == nullptr) || (ptr == nullptr)) {
hip_status = hipErrorInvalidValue;
} else {
unsigned amFlags = 0;
@@ -736,6 +738,9 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
HIP_SET_DEVICE();
if (sizeBytes == 0) {
return ihipLogStatus(hipSuccess);
}
hipError_t hip_status = hipSuccess;
hip_status = hip_internal::ihipHostMalloc(tls, ptr, sizeBytes, flags);
return ihipLogStatus(hip_status);
@@ -744,6 +749,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) {
HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags);
HIP_SET_DEVICE();
if (size == 0) {
return ihipLogStatus(hipSuccess);
}
hipError_t hip_status = hipSuccess;
if(flags != hipMemAttachGlobal)
hip_status = hipErrorInvalidValue;
@@ -1224,6 +1232,7 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
if (count == 0) return ihipLogStatus(hipSuccess);
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
@@ -1246,6 +1255,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
if (count == 0) return ihipLogStatus(hipSuccess);
if (src == nullptr || dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
@@ -1269,6 +1279,7 @@ hipError_t hipMemcpyToSymbolAsync(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
if (count == 0) return ihipLogStatus(hipSuccess);
if (dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
@@ -1301,6 +1312,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, src);
if (count == 0) return ihipLogStatus(hipSuccess);
if (src == nullptr || dst == nullptr) {
return ihipLogStatus(hipErrorInvalidSymbol);
}
@@ -1592,6 +1604,7 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
srcXoffset = p->srcPos.x;
srcYoffset = p->srcPos.y;
srcZoffset = p->srcPos.z;
if (copyWidth == 0) return hipSuccess;
if (p->dstArray != nullptr) {
if ((p->dstArray->isDrv == true) ||( p->dstPtr.ptr!= nullptr)){
return hipErrorInvalidValue;
@@ -1933,6 +1946,7 @@ hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr)
// TODO - review and optimize
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
if (height == 0 || width == 0) return hipSuccess;
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1989,6 +2003,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
hipError_t ihipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
if (height == 0 || width == 0) return hipSuccess;
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
hipError_t e = hipSuccess;
int isLockedOrD2D = 0;
@@ -2043,6 +2058,7 @@ hipError_t ihip2dOffsetMemcpy(void* dst, size_t dpitch, const void* src, size_t
size_t height, size_t srcXOffsetInBytes, size_t srcYOffset,
size_t dstXOffsetInBytes, size_t dstYOffset,hipMemcpyKind kind,
hipStream_t stream, bool isAsync) {
if (height == 0 || width == 0) return hipSuccess;
if((spitch < width + srcXOffsetInBytes) || (srcYOffset >= height)){
return hipErrorInvalidValue;
} else if((dpitch < width + dstXOffsetInBytes) || (dstYOffset >= height)){
@@ -2061,6 +2077,7 @@ hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool
if (pCopy == nullptr) {
return hipErrorInvalidValue;
}
if (pCopy->Height == 0 || pCopy->WidthInBytes == 0) return hipSuccess;
void* dst; const void* src;
size_t spitch = pCopy->srcPitch;
size_t dpitch = pCopy->dstPitch;
@@ -2140,6 +2157,7 @@ hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t src,
hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream ){
HIP_INIT_SPECIAL_API(hipMemcpy2DFromArrayAsync, (TRACE_MCMD), dst, dpitch, src, wOffset, hOffset, width, height, kind, stream);
size_t byteSize;
if (height == 0 || width == 0) return ihipLogStatus(hipSuccess);
if(src) {
switch (src->desc.f) {
case hipChannelFormatKindSigned: