AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: I68d26834a9e5e651e1be4e4219a7521f6da97415
Этот коммит содержится в:
Jenkins
2018-04-13 04:10:05 -05:00
родитель c516409a4b 2ab05a2069
Коммит 36173d50bb
4 изменённых файлов: 115 добавлений и 164 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/CMakeLists.txt
+1 -5
Просмотреть файл
@@ -201,19 +201,15 @@ if(HIP_PLATFORM STREQUAL "hcc")
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")
#find_package(LLVM HINTS ${HCC_HOME}/compiler/lib/cmake)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803 --amdgpu-target=gfx900")
if(COMPILE_HIP_ATP_MARKER)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
endif()
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc PRIVATE hc_am)
#target_link_libraries(hip_hcc PUBLIC LLVMAMDGPUUtils)
add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc_static PRIVATE hc_am)
add_dependencies(hip_hcc_static hip_hcc)
add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
add_dependencies(hip_device hip_hcc)
string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
foreach(TARGET hip_hcc hip_hcc_static hip_device)
@@ -309,7 +305,7 @@ add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR}
)
DEPENDS lpl ca)
# Packaging needs to wait for hipify-clang to build if it's enabled...
if (BUILD_HIPIFY_CLANG)
hipamd/include/hip/hcc_detail/device_functions.h
+15 -15
Просмотреть файл
@@ -104,28 +104,28 @@ __device__ double __fma_ru(double x, double y, double z);
__device__ double __fma_rz(double x, double y, double z);
// Single Precision Fast Math
extern __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32");
extern __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32");
extern __device__ __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32");
extern __device__ __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32");
__device__ float __hip_fast_exp10f(float);
__device__ float __hip_fast_expf(float);
__device__ float __hip_fast_frsqrt_rn(float);
extern __attribute__((const)) float __hip_fast_fsqrt_rd(float) __asm("llvm.sqrt.f32");
extern __device__ __attribute__((const)) float __hip_fast_fsqrt_rd(float) __asm("llvm.sqrt.f32");
__device__ float __hip_fast_fsqrt_rn(float);
__device__ float __hip_fast_fsqrt_ru(float);
__device__ float __hip_fast_fsqrt_rz(float);
__device__ float __hip_fast_log10f(float);
extern __attribute__((const)) float __hip_fast_log2f(float) __asm("llvm.log2.f32");
extern __device__ __attribute__((const)) float __hip_fast_log2f(float) __asm("llvm.log2.f32");
__device__ float __hip_fast_logf(float);
__device__ float __hip_fast_powf(float, float);
__device__ void __hip_fast_sincosf(float, float*, float*);
extern __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
extern __device__ __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
__device__ float __hip_fast_tanf(float);
extern __attribute__((const)) float __hip_fast_fmaf(float, float, float) __asm("llvm.fma.f32");
extern __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
extern __device__ __attribute__((const)) float __hip_fast_fmaf(float, float, float) __asm("llvm.fma.f32");
extern __device__ __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
extern __attribute__((const)) double __hip_fast_fma(double, double, double) __asm("llvm.fma.f64");
extern __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
extern __device__ __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
extern __device__ __attribute__((const)) double __hip_fast_fma(double, double, double) __asm("llvm.fma.f64");
extern __device__ __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
// Single Precision Fast Math
@@ -283,11 +283,11 @@ __device__ inline double __fma_ru(double x, double y, double z) { return __hip_f
__device__ inline double __fma_rz(double x, double y, double z) { return __hip_fast_fma(x, y, z); }
extern "C" unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);
extern "C" signed int __hip_hc_ir_mul24_int(signed int, signed int);
extern "C" signed int __hip_hc_ir_mulhi_int(signed int, signed int);
extern "C" unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int);
extern "C" unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int);
extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);
extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int);
extern "C" __device__ signed int __hip_hc_ir_mulhi_int(signed int, signed int);
extern "C" __device__ unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int);
extern "C" __device__ unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int);
// integer intrinsic function __poc __clz __ffs __brev
__device__ unsigned int __brev(unsigned int x);
hipamd/src/hip_clang.cpp
+36 -25
Просмотреть файл
@@ -53,7 +53,7 @@ struct __CudaFatBinaryWrapper {
};
extern "C" std::unordered_map<std::string, hipModule_t>*
extern "C" std::vector<hipModule_t>*
__hipRegisterFatBinary(const void* data)
{
HIP_INIT();
@@ -69,7 +69,7 @@ __hipRegisterFatBinary(const void* data)
return nullptr;
}
auto modules = new std::unordered_map<std::string, hipModule_t>{};
auto modules = new std::vector<hipModule_t>{g_deviceCnt};
if (!modules) {
return nullptr;
}
@@ -83,12 +83,33 @@ __hipRegisterFatBinary(const void* data)
if (triple.compare(AMDGCN_AMDHSA_TRIPLE))
continue;
hipModule_t module;
if (hipSuccess == hipModuleLoadData(&module, reinterpret_cast<const void*>(
reinterpret_cast<uintptr_t>(header) + desc->offset))) {
modules->emplace(std::string{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)},
module);
std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hsa_agent_t agent = g_allAgents[deviceId + 1];
char name[64] = {};
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
if (target.compare(name)) {
continue;
}
ihipModule_t* module = new ihipModule_t;
if (!module) {
continue;
}
hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr,
&module->executable);
std::string image{reinterpret_cast<const char*>(
reinterpret_cast<uintptr_t>(header) + desc->offset), desc->size};
module->executable = hip_impl::load_executable(image, module->executable, agent);
if (module->executable.handle) {
modules->at(deviceId) = module;
}
}
}
@@ -98,7 +119,7 @@ __hipRegisterFatBinary(const void* data)
std::map<const void*, std::vector<hipFunction_t>> g_functions;
extern "C" void __hipRegisterFunction(
std::unordered_map<std::string, hipModule_t>* modules,
std::vector<hipModule_t>* modules,
const void* hostFunction,
char* deviceFunction,
const char* deviceName,
@@ -111,18 +132,10 @@ extern "C" void __hipRegisterFunction(
{
std::vector<hipFunction_t> functions{g_deviceCnt};
for (auto&& it : *modules) {
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hipFunction_t function;
if (hipSuccess != hipModuleGetFunction(&function, it.second, deviceName)) {
continue;
}
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
char name[64] = {};
hsa_agent_get_info(g_allAgents[deviceId + 1], HSA_AGENT_INFO_NAME, name);
if (!it.first.compare(name)) {
functions[deviceId] = function;
}
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) {
functions[deviceId] = function;
}
}
@@ -130,7 +143,7 @@ extern "C" void __hipRegisterFunction(
}
extern "C" void __hipRegisterVar(
hipModule_t module,
std::vector<hipModule_t>* modules,
char* hostVar,
char* deviceVar,
const char* deviceName,
@@ -141,11 +154,9 @@ extern "C" void __hipRegisterVar(
{
}
extern "C" void __hipUnregisterFatBinary(std::unordered_map<std::string, hipModule_t>* modules)
extern "C" void __hipUnregisterFatBinary(std::vector<hipModule_t>* modules)
{
for (auto&& it : *modules) {
delete it.second;
}
std::for_each(modules->begin(), modules->end(), [](hipModule_t module){ delete module; });
delete modules;
}
hipamd/src/hip_memory.cpp
+63 -119
Просмотреть файл
@@ -1542,6 +1542,63 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
sizeBytes, std::move(val));
}
typedef enum ihipMemsetDataType {
ihipMemsetDataTypeChar = 0,
ihipMemsetDataTypeShort = 1,
ihipMemsetDataTypeInt = 2
}ihipMemsetDataType;
hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetDataType copyDataType )
{
hipError_t e = hipSuccess;
if (stream) {
if(copyDataType == ihipMemsetDataTypeChar){
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
value = value & 0xff;
uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value32, sizeBytes/sizeof(uint32_t));
}
catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
}
catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
}
} else {
if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value
try {
ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value, sizeBytes);
} catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
} else if(copyDataType == ihipMemsetDataTypeShort) {
try {
value = value & 0xffff;
ihipMemsetKernel<uint16_t> (stream, static_cast<uint16_t*> (dst), value, sizeBytes);
} catch (std::exception &ex) {
e = hipErrorInvalidValue;
}
}
}
if (HIP_API_BLOCKING) {
tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
stream->locked_wait();
}
} else {
e = hipErrorInvalidValue;
}
return 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) {
@@ -1551,35 +1608,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
stream = ihipSyncAndResolveStream(stream);
if (stream) {
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
value = value & 0xff;
uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value);
ihipMemsetKernel<uint32_t>(stream, static_cast<uint32_t*>(dst), value32,
sizeBytes / sizeof(uint32_t));
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
ihipMemsetKernel<char>(stream, static_cast<char*>(dst), value, sizeBytes);
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
}
if (HIP_API_BLOCKING) {
tprintf(DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n",
ToString(stream).c_str());
stream->locked_wait();
}
} else {
e = hipErrorInvalidValue;
}
e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
return ihipLogStatus(e);
};
@@ -1590,42 +1619,13 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
hipError_t e = hipSuccess;
hipStream_t stream = hipStreamNull;
// TODO - call an ihip memset so HIP_TRACE is correct.
stream = ihipSyncAndResolveStream(stream);
stream = ihipSyncAndResolveStream(stream);
if (stream) {
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
value = value & 0xff;
uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value);
ihipMemsetKernel<uint32_t>(stream, static_cast<uint32_t*>(dst), value32,
sizeBytes / sizeof(uint32_t));
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
ihipMemsetKernel<char>(stream, static_cast<char*>(dst), value, sizeBytes);
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
}
// TODO - is hipMemset supposed to be async?
e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
stream->locked_wait();
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset in %s.\n", __func__,
ToString(stream).c_str());
stream->locked_wait();
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed in %s.\n", __func__,
ToString(stream).c_str());
}
} else {
e = hipErrorInvalidValue;
}
}
return ihipLogStatus(e);
}
@@ -1635,39 +1635,11 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
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 = pitch * height;
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
value = value & 0xff;
uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value);
ihipMemsetKernel<uint32_t>(stream, static_cast<uint32_t*>(dst), value32,
sizeBytes / sizeof(uint32_t));
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
ihipMemsetKernel<char>(stream, static_cast<char*>(dst), value, sizeBytes);
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
}
// TODO - is hipMemset supposed to be async?
e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
stream->locked_wait();
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset in %s.\n", __func__,
ToString(stream).c_str());
stream->locked_wait();
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed in %s.\n", __func__,
ToString(stream).c_str());
}
} else {
e = hipErrorInvalidValue;
}
@@ -1681,41 +1653,13 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
hipError_t e = hipSuccess;
hipStream_t stream = hipStreamNull;
// TODO - call an ihip memset so HIP_TRACE is correct.
stream = ihipSyncAndResolveStream(stream);
if (stream) {
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
try {
uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value);
ihipMemsetKernel<uint32_t>(stream, static_cast<uint32_t*>(dst), value32,
sizeBytes / sizeof(uint32_t));
} catch (std::exception& ex) {
std::cout << ex.what() << std::endl;
e = hipErrorInvalidValue;
}
} else {
// use a slow byte-per-workitem copy:
try {
ihipMemsetKernel<char>(stream, static_cast<char*>(dst), value, sizeBytes);
} catch (std::exception& ex) {
e = hipErrorInvalidValue;
}
}
e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
stream->locked_wait();
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset in %s.\n", __func__,
ToString(stream).c_str());
stream->locked_wait();
tprintf(DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed in %s.\n", __func__,
ToString(stream).c_str());
}
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
}