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

Re-sync with upstream. Add integer abs.

Browse Source
This commit is contained in:
Alex Voicu
2018-05-31 16:38:00 +01:00
parent 59db16fd36 449ddba0aa
commit ab4b2a650b
30 changed files with 1573 additions and 346 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/device_functions.cpp
+20 -27
View File
@@ -278,21 +278,6 @@ struct uchar2Holder {
};
} __attribute__((aligned(8)));
struct intHolder {
union {
signed int si[2];
signed int long sl;
};
} __attribute__((aligned(8)));
struct uintHolder {
union {
signed int ui[2];
signed int long ul;
};
} __attribute__((aligned(8)));
__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) {
struct uchar2Holder cHoldVal;
struct ucharHolder cHoldKey;
@@ -308,21 +293,29 @@ __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int
}
__device__ long long __mul64hi(long long int x, long long int y) {
struct intHolder iHold1;
struct intHolder iHold2;
iHold1.sl = x;
iHold2.sl = y;
iHold1.sl = iHold1.si[1] * iHold2.si[1];
return iHold1.sl;
ulong x0 = (ulong)x & 0xffffffffUL;
long x1 = x >> 32;
ulong y0 = (ulong)y & 0xffffffffUL;
long y1 = y >> 32;
ulong z0 = x0*y0;
long t = x1*y0 + (z0 >> 32);
long z1 = t & 0xffffffffL;
long z2 = t >> 32;
z1 = x0*y1 + z1;
return x1*y1 + z2 + (z1 >> 32);
}
__device__ unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) {
struct uintHolder uHold1;
struct uintHolder uHold2;
uHold1.ul = x;
uHold2.ul = y;
uHold1.ul = uHold1.ui[1] * uHold2.ui[1];
return uHold1.ul;
ulong x0 = x & 0xffffffffUL;
ulong x1 = x >> 32;
ulong y0 = y & 0xffffffffUL;
ulong y1 = y >> 32;
ulong z0 = x0*y0;
ulong t = x1*y0 + (z0 >> 32);
ulong z1 = t & 0xffffffffUL;
ulong z2 = t >> 32;
z1 = x0*y1 + z1;
return x1*y1 + z2 + (z1 >> 32);
}
/*
src/functional_grid_launch.inl
+2 -2
View File
@@ -107,8 +107,8 @@ namespace hip_impl
it0->second.cbegin(),
it0->second.cend(),
[=](const pair<hsa_agent_t, Kernel_descriptor>& x) {
return x.first.handle == agent.handle;
});
return x.first == agent;
});
if (it1 == it0->second.cend()) {
throw runtime_error{
src/hip_clang.cpp
+3 -3
View File
@@ -27,10 +27,10 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
constexpr unsigned __cudaFatMAGIC2 = 0x466243b1;
constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
#define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__"
#define AMDGCN_AMDHSA_TRIPLE "openmp-amdgcn--amdhsa"
#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
struct __ClangOffloadBundleDesc {
uint64_t offset;
@@ -59,7 +59,7 @@ __hipRegisterFatBinary(const void* data)
HIP_INIT();
const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) {
if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
return nullptr;
}
src/hip_memory.cpp
+172 -64
View File
@@ -1132,15 +1132,19 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);
hipError_t e = hipSuccess;
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return ihipLogStatus(e);
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
if(dst==NULL || src==NULL)
{
e=hipErrorInvalidValue;
return e;
return ihipLogStatus(e);
}
try {
stream->locked_copySync(dst, src, sizeBytes, kind);
@@ -1208,7 +1212,6 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
return ihipLogStatus(e);
}
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes);
@@ -1227,7 +1230,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
return ihipLogStatus(e);
}
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream);
@@ -1235,7 +1237,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
}
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream);
@@ -1258,65 +1259,6 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
}
// 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 (width > dpitch || width > spitch) return hipErrorUnknown;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
for (int i = 0; i < height; ++i) {
stream->locked_copySync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind);
}
} catch (ihipException& ex) {
e = ex._code;
}
return e;
}
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
}
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
return ihipLogStatus(e);
}
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
if (width > dpitch || width > spitch) return ihipLogStatus(hipErrorUnknown);
hipError_t e = hipSuccess;
try {
for (int i = 0; i < height; ++i) {
e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind, stream);
}
} catch (ihipException& ex) {
e = ex._code;
}
return ihipLogStatus(e);
}
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind);
@@ -1547,6 +1489,24 @@ inline const T& clamp_integer(const T& x, const T& lower, const T& upper) {
return std::min(upper, std::max(x, lower));
}
template <typename T>
__global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) {
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
size_t idy = blockIdx.y * blockDim.y + threadIdx.y;
size_t floorWidth = (width/sizeof(T));
T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
if((idx < floorWidth) && (idy < height)){
dstPtr[idx] = srcPtr[idx];
} else if((idx < width) && (idy < height)){
size_t bytesToCopy = width - (floorWidth * sizeof(T));
dstPtr += floorWidth;
srcPtr += floorWidth;
__builtin_memcpy(reinterpret_cast<uint8_t*>(dstPtr), reinterpret_cast<const uint8_t*>(srcPtr),bytesToCopy);
}
}
} // namespace
template <typename T>
@@ -1559,6 +1519,16 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
sizeBytes, std::move(val));
}
template <typename T>
void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) {
size_t threadsPerBlock_x = 64;
size_t threadsPerBlock_y = 4;
uint32_t grid_dim_x = clamp_integer<size_t>( (width+(threadsPerBlock_x*sizeof(T)-1)) / (threadsPerBlock_x*sizeof(T)), 1, UINT32_MAX);
uint32_t grid_dim_y = clamp_integer<size_t>( (height+(threadsPerBlock_y-1)) / threadsPerBlock_y, 1, UINT32_MAX);
hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock_x,threadsPerBlock_y), 0u, stream, dst, src,
width, height, destPitch, srcPitch);
}
typedef enum ihipMemsetDataType {
ihipMemsetDataTypeChar = 0,
ihipMemsetDataTypeShort = 1,
@@ -1616,6 +1586,108 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea
return e;
};
int isLockedPointer(const void *ptr)
{
hsa_amd_pointer_info_t info;
int isLocked = 0;
info.size = sizeof(info);
hsa_status_t hsa_status = hsa_amd_pointer_info(const_cast<void*>(ptr), &info, nullptr, nullptr, nullptr);
if(hsa_status != HSA_STATUS_SUCCESS) {
return -1;
}
if((info.type == HSA_EXT_POINTER_TYPE_HSA) || (info.type == HSA_EXT_POINTER_TYPE_LOCKED)) {
isLocked = 1;
}
return isLocked;
}
// 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 (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
int isLocked = 0;
if(kind == hipMemcpyHostToDevice) {
isLocked = isLockedPointer(src);
} else if(kind == hipMemcpyDeviceToHost) {
isLocked = isLockedPointer(dst);
}
hc::completion_future marker;
hipError_t e = hipSuccess;
if((width == dpitch) && (width == spitch)) {
stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false);
} else {
try {
if(isLocked) {
for (int i = 0; i < height; ++i)
stream->locked_copySync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind);
} else {
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch);
stream->locked_wait();
}
} catch (ihipException& ex) {
e = ex._code;
}
}
return e;
}
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
}
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
hipError_t e = hipSuccess;
int isLocked = 0;
if(kind == hipMemcpyHostToDevice) {
isLocked = isLockedPointer(src);
} else if(kind == hipMemcpyDeviceToHost) {
isLocked = isLockedPointer(dst);
}
if((width == dpitch) && (width == spitch)) {
hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
} else {
try {
if(!isLocked){
for (int i = 0; i < height; ++i)
e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind, stream);
} else{
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch);
}
} catch (ihipException& ex) {
e = ex._code;
}
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
}
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
return ihipLogStatus(e);
}
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
@@ -1698,6 +1770,42 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
return ihipLogStatus(e);
}
hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
hipError_t e = hipSuccess;
hipStream_t stream = hipStreamNull;
// TODO - call an ihip memset so HIP_TRACE is correct.
stream = ihipSyncAndResolveStream(stream);
if (stream) {
size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar);
stream->locked_wait();
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
}
hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent);
hipError_t e = hipSuccess;
// TODO - call an ihip memset so HIP_TRACE is correct.
stream = ihipSyncAndResolveStream(stream);
if (stream) {
size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar);
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
}
hipError_t hipMemGetInfo(size_t* free, size_t* total) {
HIP_INIT_API(free, total);
src/hip_module.cpp
+72 -11
View File
@@ -27,10 +27,11 @@ THE SOFTWARE.
#include "hsa_helpers.hpp"
#include "trace_helper.h"
#include <hsa/amd_hsa_kernel_code.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/amd_hsa_kernel_code.h>
#include <algorithm>
#include <cassert>
#include <cstdint>
#include <cstdio>
@@ -71,9 +72,8 @@ struct ihipKernArgInfo {
map<string, ihipKernArgInfo> kernelArguments;
struct ihipModuleSymbol_t {
uint64_t _object; // The kernel object.
uint32_t _groupSegmentSize;
uint32_t _privateSegmentSize;
uint64_t _object{}; // The kernel object.
amd_kernel_code_t const* _header{};
string _name; // TODO - review for performance cost. Name is just used for debug.
};
@@ -179,8 +179,10 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
aql.grid_size_x = globalWorkSizeX;
aql.grid_size_y = globalWorkSizeY;
aql.grid_size_z = globalWorkSizeZ;
aql.group_segment_size = f->_groupSegmentSize + sharedMemBytes;
aql.private_segment_size = f->_privateSegmentSize;
aql.group_segment_size =
f->_header->workgroup_group_segment_byte_size + sharedMemBytes;
aql.private_segment_size =
f->_header->workitem_private_segment_byte_size;
aql.kernel_object = f->_object;
aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
aql.header =
@@ -444,10 +446,10 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
if (kernel.handle == 0u) return hipErrorNotFound;
(*func)->_object = kernel_object(kernel);
(*func)->_groupSegmentSize = group_size(kernel);
(*func)->_privateSegmentSize = private_size(kernel);
(*func)->_name = name;
// TODO: refactor the whole ihipThisThat, which is a mess and yields the
// below, due to hipFunction_t being a pointer to ihipModuleSymbol_t.
func[0][0] = *static_cast<hipFunction_t>(
Kernel_descriptor{kernel_object(kernel), name});
return hipSuccess;
}
@@ -471,6 +473,65 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h
return ihipLogStatus(r);
}
namespace
{
inline
hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header)
{
hipFuncAttributes r{};
hipDeviceProp_t prop{};
hipGetDeviceProperties(
&prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId);
// TODO: at the moment there is no way to query the count of registers
// available per CU, therefore we hardcode it to 64 KiRegisters.
prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024;
r.localSizeBytes = header.workitem_private_segment_byte_size;
r.sharedSizeBytes = header.workgroup_group_segment_byte_size;
r.maxDynamicSharedSizeBytes =
prop.sharedMemPerBlock - r.sharedSizeBytes;
r.numRegs = header.workitem_vgpr_count;
r.maxThreadsPerBlock = r.numRegs ?
std::min(prop.maxThreadsPerBlock, prop.regsPerBlock / r.numRegs) :
prop.maxThreadsPerBlock;
r.binaryVersion =
header.amd_machine_version_major * 10 +
header.amd_machine_version_minor;
r.ptxVersion = prop.major * 10 + prop.minor; // HIP currently presents itself as PTX 3.0.
return r;
}
}
hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
{
if (!attr) return hipErrorInvalidValue;
if (!func) return hipErrorInvalidDeviceFunction;
const auto it0 = functions().find(reinterpret_cast<uintptr_t>(func));
if (it0 == functions().cend()) return hipErrorInvalidDeviceFunction;
auto agent = this_agent();
const auto it1 = find_if(
it0->second.cbegin(),
it0->second.cend(),
[=](const pair<hsa_agent_t, Kernel_descriptor>& x) {
return x.first == agent;
});
if (it1 == it0->second.cend()) return hipErrorInvalidDeviceFunction;
const auto header = static_cast<hipFunction_t>(it1->second)->_header;
if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."};
*attr = make_function_attributes(*header);
return hipSuccess;
}
hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
if (!module) return hipErrorInvalidValue;
@@ -487,7 +548,7 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
(*module)->executable = hip_impl::load_executable(
tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent());
return (*module)->executable.handle ? hipSuccess : hipErrorUnknown;
}
src/program_state.cpp
+1 -2
View File
@@ -382,8 +382,7 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
for (auto&& kernel_symbol : it->second) {
r[function.first].emplace_back(
agent(kernel_symbol),
Kernel_descriptor{kernel_object(kernel_symbol), group_size(kernel_symbol),
private_size(kernel_symbol), it->first});
Kernel_descriptor{kernel_object(kernel_symbol), it->first});
}
}
}