From 10f3cc1f27ec2991442cb4db71620aa2beff412a Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Tue, 10 Apr 2018 08:57:34 -0700 Subject: [PATCH 1/7] Create one hipModule_t per device [ROCm/hip commit: 91f996035c8cb58d38d414ba8232e9d578b835c2] --- projects/hip/src/hip_clang.cpp | 61 ++++++++++++++++++++-------------- 1 file changed, 36 insertions(+), 25 deletions(-) diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp index 5c4da0b39d..80b6111fc2 100644 --- a/projects/hip/src/hip_clang.cpp +++ b/projects/hip/src/hip_clang.cpp @@ -53,7 +53,7 @@ struct __CudaFatBinaryWrapper { }; -extern "C" std::unordered_map* +extern "C" std::vector* __hipRegisterFatBinary(const void* data) { HIP_INIT(); @@ -69,7 +69,7 @@ __hipRegisterFatBinary(const void* data) return nullptr; } - auto modules = new std::unordered_map{}; + auto modules = new std::vector{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( - reinterpret_cast(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( + reinterpret_cast(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> g_functions; extern "C" void __hipRegisterFunction( - std::unordered_map* modules, + std::vector* modules, const void* hostFunction, char* deviceFunction, const char* deviceName, @@ -111,18 +132,10 @@ extern "C" void __hipRegisterFunction( { std::vector 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* modules, char* hostVar, char* deviceVar, const char* deviceName, @@ -141,11 +154,9 @@ extern "C" void __hipRegisterVar( { } -extern "C" void __hipUnregisterFatBinary(std::unordered_map* modules) +extern "C" void __hipUnregisterFatBinary(std::vector* modules) { - for (auto&& it : *modules) { - delete it.second; - } + std::for_each(modules->begin(), modules->end(), [](hipModule_t module){ delete module; }); delete modules; } From 01e4cf1c68751a4b9a40158e2180e6d3a5fd66dd Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 11 Apr 2018 15:30:13 +0530 Subject: [PATCH 2/7] [cmake] Update target build dependencies Change-Id: Ib1c34135c0d04f88daa695e7a90272225d7e64b0 [ROCm/hip commit: b1be325745724ef9c9ba4c8621fc098242a1ccb9] --- projects/hip/CMakeLists.txt | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 4e4a295c35..e9a3451222 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -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) From 36c6e0019dbdea263ab15b52c23578d54dd72a54 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 11 Apr 2018 15:58:48 +0530 Subject: [PATCH 3/7] hipMemset refactoring [ROCm/hip commit: 412a35be20985c113b60c4993670375717f07e49] --- projects/hip/src/hip_memory.cpp | 181 +++++++++++--------------------- 1 file changed, 62 insertions(+), 119 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index ffa9f0cfdd..0644f8de99 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1542,6 +1542,63 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { sizeBytes, std::move(val)); } +typedef enum ihipMemsetCopyDataType { + ihipMemsetCopyDataTypeChar = 0, + ihipMemsetCopyDataTypeShort = 1, + ihipMemsetCopyDataTypeInt = 2 +}ihipMemsetCopyDataType; + +hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetCopyDataType copyDataType ) +{ + hipError_t e = hipSuccess; + + if (stream) { + if(copyDataType == ihipMemsetCopyDataTypeChar){ + 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 (stream, static_cast (dst), value32, sizeBytes/sizeof(uint32_t)); + } + catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } else { + // use a slow byte-per-workitem copy: + try { + ihipMemsetKernel (stream, static_cast (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 { + if(copyDataType == ihipMemsetCopyDataTypeInt) { // 4 Bytes value + try { + ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + } catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } else if(copyDataType == ihipMemsetCopyDataTypeShort) { + try { + value = value & 0xffff; + ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + } catch (std::exception &ex) { + e = hipErrorInvalidValue; + } + } + } + } 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(stream, static_cast(dst), value32, - sizeBytes / sizeof(uint32_t)); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } else { - // use a slow byte-per-workitem copy: - try { - ihipMemsetKernel(stream, static_cast(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, ihipMemsetCopyDataTypeChar); return ihipLogStatus(e); }; @@ -1590,42 +1619,12 @@ 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); - 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(stream, static_cast(dst), value32, - sizeBytes / sizeof(uint32_t)); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } else { - // use a slow byte-per-workitem copy: - try { - ihipMemsetKernel(stream, static_cast(dst), value, sizeBytes); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } - // TODO - is hipMemset supposed to be async? + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); 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 +1634,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(stream, static_cast(dst), value32, - sizeBytes / sizeof(uint32_t)); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } else { - // use a slow byte-per-workitem copy: - try { - ihipMemsetKernel(stream, static_cast(dst), value, sizeBytes); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } - // TODO - is hipMemset supposed to be async? + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); 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 +1652,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(stream, static_cast(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(stream, static_cast(dst), value, sizeBytes); - } catch (std::exception& ex) { - e = hipErrorInvalidValue; - } - } + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); 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); } From 6c4236dfb6a71ceb071c4a12d501ed9db0cd861c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 11 Apr 2018 19:01:53 +0530 Subject: [PATCH 4/7] Fix hipMemset stream resolution [ROCm/hip commit: 294bf50f6861d2ffec2857f4e4f0967decc2363c] --- projects/hip/src/hip_memory.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 0644f8de99..52a88e8b88 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1619,6 +1619,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { hipError_t e = hipSuccess; hipStream_t stream = hipStreamNull; + stream = ihipSyncAndResolveStream(stream); if (stream) { e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); stream->locked_wait(); From 3d9e3019d0d4e1f88c7a4730b3de8bcf0e1f2534 Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 11 Apr 2018 09:29:37 -0700 Subject: [PATCH 5/7] Add missing __device__ function attributes. [ROCm/hip commit: 211c18e67487ab714fb3bea6d5b7ac84192b9775] --- .../include/hip/hcc_detail/device_functions.h | 30 +++++++++---------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 99713ed4d7..28d874b27a 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -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); From abe14442a54f72c650b5b284f8e6a9118d8479b7 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 12 Apr 2018 09:29:22 +0530 Subject: [PATCH 6/7] Changed ihipMemsetCopyDataType to ihipMemsetDataType [ROCm/hip commit: 3d6eb758287289798e09a69f232cf839ae025220] --- projects/hip/src/hip_memory.cpp | 32 ++++++++++++++++---------------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 52a88e8b88..ff97869842 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1542,18 +1542,18 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { sizeBytes, std::move(val)); } -typedef enum ihipMemsetCopyDataType { - ihipMemsetCopyDataTypeChar = 0, - ihipMemsetCopyDataTypeShort = 1, - ihipMemsetCopyDataTypeInt = 2 +typedef enum ihipMemsetDataType { + ihipMemsetDataTypeChar = 0, + ihipMemsetDataTypeShort = 1, + ihipMemsetDataTypeInt = 2 }ihipMemsetCopyDataType; -hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetCopyDataType copyDataType ) +hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetDataType copyDataType ) { hipError_t e = hipSuccess; if (stream) { - if(copyDataType == ihipMemsetCopyDataTypeChar){ + if(copyDataType == ihipMemsetDataTypeChar){ if ((sizeBytes & 0x3) == 0) { // use a faster dword-per-workitem copy: try { @@ -1572,19 +1572,15 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea 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 { - if(copyDataType == ihipMemsetCopyDataTypeInt) { // 4 Bytes value + if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value try { ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); } catch (std::exception &ex) { e = hipErrorInvalidValue; } - } else if(copyDataType == ihipMemsetCopyDataTypeShort) { + } else if(copyDataType == ihipMemsetDataTypeShort) { try { value = value & 0xffff; ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); @@ -1593,6 +1589,10 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea } } } + if (HIP_API_BLOCKING) { + tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str()); + stream->locked_wait(); + } } else { e = hipErrorInvalidValue; } @@ -1608,7 +1608,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st stream = ihipSyncAndResolveStream(stream); - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); return ihipLogStatus(e); }; @@ -1621,7 +1621,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { hipStream_t stream = hipStreamNull; stream = ihipSyncAndResolveStream(stream); if (stream) { - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); stream->locked_wait(); } else { e = hipErrorInvalidValue; @@ -1638,7 +1638,7 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t stream = ihipSyncAndResolveStream(stream); if (stream) { size_t sizeBytes = pitch * height; - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); stream->locked_wait(); } else { e = hipErrorInvalidValue; @@ -1655,7 +1655,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes hipStream_t stream = hipStreamNull; stream = ihipSyncAndResolveStream(stream); if (stream) { - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetCopyDataTypeChar); + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); stream->locked_wait(); } else { e = hipErrorInvalidValue; From 89511823f0883042f17b77ec820f49a3a2f01d4f Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 12 Apr 2018 10:27:19 +0530 Subject: [PATCH 7/7] Correct missed ihipMemsetCopyDataType change [ROCm/hip commit: 16c89d101a34d5116f19a15b6f44ce8b57cc3758] --- projects/hip/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index ff97869842..e6d6ebd09e 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1546,7 +1546,7 @@ typedef enum ihipMemsetDataType { ihipMemsetDataTypeChar = 0, ihipMemsetDataTypeShort = 1, ihipMemsetDataTypeInt = 2 -}ihipMemsetCopyDataType; +}ihipMemsetDataType; hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetDataType copyDataType ) {