From 603bb321ec2e8390b38ba82529d44e4ad22f4dcb Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 17 Nov 2016 01:09:12 -0600 Subject: [PATCH 01/31] Added i8 packed math intrinsics 1. Added add, sub, mul packed math i8 intrinsics 2. Removed c++ packed data structures included from HCC Change-Id: I1d109c5ce10c48b7cd3ea059478b88fc1de78499 TODO: Add better packed data structures support --- hipamd/include/hip/hcc_detail/hip_runtime.h | 4 +++ .../include/hip/hcc_detail/hip_vector_types.h | 22 ++++++++++-- hipamd/src/hip_hcc.cpp | 36 +++++++++++++++++++ 3 files changed, 59 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index b1edef18d7..b501f0b165 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -627,6 +627,10 @@ __device__ static inline void* free(void *ptr) return __hip_hc_free(ptr); } +extern "C" __device__ char4 __hip_hc_add8pk(char4, char4); +extern "C" __device__ char4 __hip_hc_sub8pk(char4, char4); +extern "C" __device__ char4 __hip_hc_mul8pk(char4, char4); + #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) #define HIP_KERNEL_NAME(...) __VA_ARGS__ diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 5c2e48026e..932e271527 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -32,7 +32,7 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif -#if __HCC__ +#if 0 #include using namespace hc::short_vector; @@ -137,8 +137,24 @@ struct uchar3 unsigned char x, y, z; }; -struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;); -struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;); +struct char4 +{ + union { + signed char x, y, z, w; + unsigned int val; + }; +}; + +struct uchar4 +{ + union { + unsigned char x, y, z, w; + unsigned int val; + }; +}; + +//struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;); +//struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;); struct __hip_align(short1, 2, signed short x;); struct __hip_align(ushort1, 2, unsigned short x;); diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index d066079fc0..3ebe9c3647 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -217,6 +217,42 @@ __device__ float __hip_ds_swizzlef(float src, int pattern) { __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) { return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); } + +#define MASK1 0x00ff00ff +#define MASK2 0xff00ff00 + +__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.val & MASK1; + unsigned one2 = in2.val & MASK1; + out.val = (one1 + one2) & MASK1; + one1 = in1.val & MASK2; + one2 = in2.val & MASK2; + out.val = out.val | ((one1 + one2) & MASK2); + return out; +} + +__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.val & MASK1; + unsigned one2 = in2.val & MASK1; + out.val = (one1 - one2) & MASK1; + one1 = in1.val & MASK2; + one2 = in2.val & MASK2; + out.val = out.val | ((one1 - one2) & MASK2); + return out; +} + +__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.val & MASK1; + unsigned one2 = in2.val & MASK1; + out.val = (one1 * one2) & MASK1; + one1 = in1.val & MASK2; + one2 = in2.val & MASK2; + out.val = out.val | ((one1 * one2) & MASK2); + return out; +} //================================================================================================= // Thread-local storage: //================================================================================================= From 94984470d4f8862e0c45ae9806511531d8e1833a Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 17 Nov 2016 11:55:29 -0600 Subject: [PATCH 02/31] make texture as seperate header as of now Change-Id: I3c65aa75f2f729eedd8c3292fa3cbc37709c1cfe --- hipamd/include/hip/hcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index b501f0b165..a0836dc06e 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -68,7 +68,7 @@ extern int HIP_TRACE_API; //typedef grid_launch_parm hipLaunchParm ; #define hipLaunchParm grid_launch_parm #ifdef __cplusplus -#include +//#include #include #endif #include From 84d0d10fad0b770fd885076e135d478c258bfda4 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 17 Nov 2016 11:57:53 -0600 Subject: [PATCH 03/31] added texture header to memory api source Change-Id: I1af6d60aca5a9a9ef1cadf8c304bea892acbe061 --- hipamd/src/hip_memory.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 672b9f2ee2..d3bbd9aeab 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "hip_hcc.h" #include "trace_helper.h" +#include "hip/hcc_detail/hip_texture.h" //------------------------------------------------------------------------------------------------- //------------------------------------------------------------------------------------------------- From 3b1f0e903ce7c997f1d9a9237153930f5b0d4350 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 17 Nov 2016 14:19:18 -0600 Subject: [PATCH 04/31] moved runtime macros to runtime_api.h Change-Id: Ib47e449328e8e6ec55d1b6ee19899de4b591ea8e --- hipamd/include/hip/hcc_detail/hip_runtime.h | 4 ---- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 4 ++++ 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index a0836dc06e..44a61e6924 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -58,10 +58,6 @@ THE SOFTWARE. #error (HCC must support GRID_LAUNCH_20) #endif -#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01) -#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*) 0x02) -#define HIP_LAUNCH_PARAM_END ((void*) 0x03) - extern int HIP_TRACE_API; //TODO-HCC-GL - change this to typedef. diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 82eba49771..09825e141e 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -39,6 +39,10 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif +#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01) +#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*) 0x02) +#define HIP_LAUNCH_PARAM_END ((void*) 0x03) + // Structure definitions: #ifdef __cplusplus extern "C" { From 1a67766dc5708938526087bb0db71a0c2b24f91b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 18 Nov 2016 12:20:47 +0530 Subject: [PATCH 05/31] Fix broken tests due to 9498447 Change-Id: I847c80f8462e1c955bdef957e6de2841a3a6ab29 --- hipamd/tests/src/test_common.h | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index 2b30b9d48f..1c4a6c7c23 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -23,6 +23,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" +#include "hip/hcc_detail/hip_texture.h" #define HC __attribute__((hc)) From 04049feaf438652eeb1e28af3353b5e509a70c8b Mon Sep 17 00:00:00 2001 From: scchan Date: Thu, 17 Nov 2016 18:44:42 -0500 Subject: [PATCH 06/31] Add extra linker flags to the shared library build Change-Id: I19e569d566fb5e25e343e364a3053a3f12659361 --- hipamd/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index ccd390fbe5..885897fb1b 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -191,6 +191,7 @@ if(HIP_PLATFORM STREQUAL "hcc") elseif(${HIP_LIB_TYPE} EQUAL 1) add_library(hip_hcc STATIC ${SOURCE_FILES}) else() + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic") add_library(hip_hcc SHARED ${SOURCE_FILES}) endif() From 2195e3c37dc1ca14053824a82849ecb396361e6f Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 18 Nov 2016 14:33:20 +0530 Subject: [PATCH 07/31] Refactor for building HIP as dynamic library Change-Id: I65a3d9d589c4fdbbdcf1611e5427224253be8260 --- hipamd/CMakeLists.txt | 23 ++- hipamd/bin/hipcc | 2 +- hipamd/include/hip/hcc_detail/hip_runtime.h | 34 +--- hipamd/packaging/hip_hcc.txt | 1 + hipamd/src/device_util.cpp | 182 ++++++++++++++++++++ hipamd/src/hip_hcc.cpp | 144 ---------------- hipamd/src/hip_memory.cpp | 4 - 7 files changed, 203 insertions(+), 187 deletions(-) diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index 885897fb1b..1b95145b96 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -173,33 +173,38 @@ if(HIP_PLATFORM STREQUAL "hcc") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}") set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}") - set(SOURCE_FILES src/device_util.cpp + set(SOURCE_FILES_SHARED src/hip_hcc.cpp src/hip_context.cpp src/hip_device.cpp src/hip_error.cpp src/hip_event.cpp - src/hip_ldg.cpp src/hip_memory.cpp src/hip_peer.cpp src/hip_stream.cpp - src/hip_fp16.cpp src/hip_module.cpp) + set(SOURCE_FILES_STATIC + src/device_util.cpp + src/hip_ldg.cpp + src/hip_fp16.cpp) + if(${HIP_LIB_TYPE} EQUAL 0) - add_library(hip_hcc OBJECT ${SOURCE_FILES}) + add_library(hip_hcc OBJECT ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC}) elseif(${HIP_LIB_TYPE} EQUAL 1) - add_library(hip_hcc STATIC ${SOURCE_FILES}) + add_library(hip_hcc STATIC ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC}) else() set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic") - add_library(hip_hcc SHARED ${SOURCE_FILES}) + add_library(hip_hcc SHARED ${SOURCE_FILES_SHARED}) + add_library(hip_device STATIC ${SOURCE_FILES_STATIC}) + add_dependencies(hip_device hip_hcc) endif() # Generate hcc_version.txt add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp) add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version) set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE) - set_source_files_properties(${SOURCE_FILES} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt) + set_source_files_properties(${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt) add_dependencies(hip_hcc check_hcc_version update_build_and_version_info) # Generate .hipInfo @@ -220,8 +225,10 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA if(HIP_PLATFORM STREQUAL "hcc") if(${HIP_LIB_TYPE} EQUAL 0) install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib) - else() + elseif(${HIP_LIB_TYPE} EQUAL 1) install(TARGETS hip_hcc DESTINATION lib) + else() + install(TARGETS hip_hcc hip_device DESTINATION lib) endif() install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index dee0894869..641f70b065 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -328,7 +328,7 @@ if ($needHipHcc) { } elsif ($HIP_LIB_TYPE eq 1) { substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ; } else { - substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc "; + substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc -lhip_device "; } } diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 44a61e6924..d7c8d3a675 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -588,40 +588,14 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask #define hipGridDim_y (hc_get_num_groups(1)) #define hipGridDim_z (hc_get_num_groups(2)) -// loop unrolling -__device__ static inline void* memcpy(void* dst, void* src, size_t size) -{ - uint8_t *dstPtr, *srcPtr; - dstPtr = (uint8_t*)dst; - srcPtr = (uint8_t*)src; - for(uint32_t i=0;i SIZE_OF_HEAP) + { + return (void*)nullptr; + } + uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z; + uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x; + + uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads; + uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads; + + uint32_t stride = size / SIZE_OF_PAGE; + uint32_t start = numHeapsPerWorkItem * currentWorkItem; + + uint32_t k=0; + + while(gpuFlags[k] > 0) + { + k++; + } + + for(uint32_t i=0;i g_lastShortTid(1); std::vector g_dbStartTriggers; std::vector g_dbStopTriggers; - - -/* - Implementation of malloc and free device functions. - - This is the best place to put them because the device - global variables need to be initialized at the start. -*/ - -#define NUM_PAGES_PER_THREAD 16 -#define SIZE_OF_PAGE 64 -#define NUM_THREADS_PER_CU 64 -#define NUM_CUS_PER_GPU 64 -#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU -#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE -#define SIZE_OF_HEAP SIZE_MALLOC - -size_t g_malloc_heap_size = SIZE_OF_HEAP; - -__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP]; -__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES]; - -__device__ void *__hip_hc_malloc(size_t size) -{ - char *heap = (char*)gpuHeap; - if(size > SIZE_OF_HEAP) - { - return (void*)nullptr; - } - uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z; - uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x; - - uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads; - uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads; - - uint32_t stride = size / SIZE_OF_PAGE; - uint32_t start = numHeapsPerWorkItem * currentWorkItem; - - uint32_t k=0; - - while(gpuFlags[k] > 0) - { - k++; - } - - for(uint32_t i=0;i Date: Tue, 15 Nov 2016 12:41:05 +0530 Subject: [PATCH 08/31] fix_format Change-Id: I34e265de434263a11654e5deba044c3f21e86578 --- .../samples/2_Cookbook/8_peer2peer/Makefile | 71 +++++++++---------- .../2_Cookbook/8_peer2peer/peer2peer.cpp | 12 ++-- 2 files changed, 41 insertions(+), 42 deletions(-) diff --git a/hipamd/samples/2_Cookbook/8_peer2peer/Makefile b/hipamd/samples/2_Cookbook/8_peer2peer/Makefile index a1dad7d1da..5cb7473921 100644 --- a/hipamd/samples/2_Cookbook/8_peer2peer/Makefile +++ b/hipamd/samples/2_Cookbook/8_peer2peer/Makefile @@ -1,36 +1,35 @@ -HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif - -HIPCC=$(HIP_PATH)/bin/hipcc - -TARGET=hcc - -SOURCES = peer2peer.cpp -OBJECTS = $(SOURCES:.cpp=.o) - -EXECUTABLE=./peer2peer - -.PHONY: test - - -all: $(EXECUTABLE) test - -CXXFLAGS =-g -CXX=$(HIPCC) - - -$(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ - - -test: $(EXECUTABLE) - $(EXECUTABLE) - - -clean: - rm -f $(EXECUTABLE) - rm -f $(OBJECTS) - rm -f $(HIP_PATH)/src/*.o - +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif + +HIPCC=$(HIP_PATH)/bin/hipcc + +TARGET=hcc + +SOURCES = peer2peer.cpp +OBJECTS = $(SOURCES:.cpp=.o) + +EXECUTABLE=./peer2peer + +.PHONY: test + + +all: $(EXECUTABLE) test + +CXXFLAGS =-g +CXX=$(HIPCC) + + +$(EXECUTABLE): $(OBJECTS) + $(HIPCC) $(OBJECTS) -o $@ + +test: $(EXECUTABLE) + $(EXECUTABLE) + + +clean: +rm -f $(EXECUTABLE) +rm -f $(OBJECTS) +rm -f $(HIP_PATH)/src/*.o + diff --git a/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp index 624de56cb0..990599e1cb 100644 --- a/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp +++ b/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp @@ -108,7 +108,7 @@ void disablePeer2Peer(int currentGpu, int peerGpu) HIPCHECK(hipSetDevice(currentGpu)); hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu); - + if(canAccessPeer==1){ HIPCHECK(hipDeviceDisablePeerAccess(peerGpu)); } @@ -155,7 +155,7 @@ __global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp, int main(){ checkPeer2PeerSupport(); - + int gpuCount; int currentGpu, peerGpu; @@ -191,10 +191,10 @@ int main(){ 0, 0, gpuTransposeMatrix[0], data[0], width); - HIPCHECK(hipSetDevice(peerGpu)); - TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float)); - hipMalloc((void**)&data[1], NUM * sizeof(float)); + HIPCHECK(hipSetDevice(peerGpu)); + TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float)); + hipMalloc((void**)&data[1], NUM * sizeof(float)); hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice); hipLaunchKernel(matrixTranspose_dynamic_shared, From 6692ee09d7f442d96275a48f888ae0316138f84d Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sat, 19 Nov 2016 21:36:28 -0600 Subject: [PATCH 09/31] added tests to check nvcc runtime api output Change-Id: Ifdd39b5d0a6a58d20a8e9745e59dd82d50a90e2f --- hipamd/include/hip/hcc_detail/hip_texture.h | 2 +- hipamd/include/hip/hip_texture.h | 13 +++++++++++++ hipamd/include/hip/nvcc_detail/hip_texture.h | 6 ++++++ hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp | 11 +++++++++++ hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp | 8 ++++++++ hipamd/tests/src/nvcc/Device/hipGetDevice.cpp | 8 ++++++++ hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp | 7 +++++++ .../src/nvcc/Device/hipGetDeviceProperties.cpp | 12 ++++++++++++ hipamd/tests/src/nvcc/Device/hipSetDevice.cpp | 10 ++++++++++ hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp | 8 ++++++++ hipamd/tests/src/test_common.h | 4 +++- 11 files changed, 87 insertions(+), 2 deletions(-) create mode 100644 hipamd/include/hip/hip_texture.h create mode 100644 hipamd/include/hip/nvcc_detail/hip_texture.h create mode 100644 hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipGetDevice.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipSetDevice.cpp create mode 100644 hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_texture.h b/hipamd/include/hip/hcc_detail/hip_texture.h index c83917b8d6..7bd14c7a40 100644 --- a/hipamd/include/hip/hcc_detail/hip_texture.h +++ b/hipamd/include/hip/hcc_detail/hip_texture.h @@ -32,7 +32,7 @@ THE SOFTWARE. #include -#include +#include //---- //Texture - TODO - likely need to move this to a separate file only included with kernel compilation. diff --git a/hipamd/include/hip/hip_texture.h b/hipamd/include/hip/hip_texture.h new file mode 100644 index 0000000000..b3d830a0df --- /dev/null +++ b/hipamd/include/hip/hip_texture.h @@ -0,0 +1,13 @@ +#ifndef HIP_TEXTURE_H +#define HIP_TEXTURE_H + +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) +#include +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif + + +#endif diff --git a/hipamd/include/hip/nvcc_detail/hip_texture.h b/hipamd/include/hip/nvcc_detail/hip_texture.h new file mode 100644 index 0000000000..cdd939bc1b --- /dev/null +++ b/hipamd/include/hip/nvcc_detail/hip_texture.h @@ -0,0 +1,6 @@ +#ifndef HIP_TEXTURE_H +#define HIP_TEXTURE_H + +#include + +#endif diff --git a/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp b/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp new file mode 100644 index 0000000000..d2fbfe1852 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp @@ -0,0 +1,11 @@ +#include +#include"test_common.h" + +int main() +{ + int dev; + hipDeviceProp_t prop; + HIP_PRINT_STATUS(hipChooseDevice(&dev, &prop)); + HIP_PRINT_STATUS(hipChooseDevice(0, &prop)); + HIP_PRINT_STATUS(hipChooseDevice(0, 0)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp new file mode 100644 index 0000000000..b17d5d6a07 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp @@ -0,0 +1,8 @@ +#include +#include "test_common.h" + +int main() +{ + hipLimit_t lim = hipLimitMallocHeapSize; + HIP_PRINT_STATUS(hipDeviceGetLimit(NULL, lim)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp new file mode 100644 index 0000000000..a009e337a6 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp @@ -0,0 +1,8 @@ +#include +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipGetDevice(NULL)); + HIP_PRINT_STATUS(hipGetDevice(0)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp new file mode 100644 index 0000000000..ec02e57ae3 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp @@ -0,0 +1,7 @@ +#include +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipGetDeviceCount(NULL)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp new file mode 100644 index 0000000000..944c4c1516 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp @@ -0,0 +1,12 @@ +#include +#include "test_common.h" + +int main() +{ + hipDeviceProp_t prop; + HIP_PRINT_STATUS(hipGetDeviceProperties(&prop, -1)); + int cnt; + hipGetDeviceCount(&cnt); + HIP_PRINT_STATUS(hipGetDeviceProperties(&prop, cnt+1)); + HIP_PRINT_STATUS(hipGetDeviceProperties(NULL, 0)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp new file mode 100644 index 0000000000..7a473ef042 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp @@ -0,0 +1,10 @@ +#include +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipSetDevice(-1)); + int count; + hipGetDeviceCount(&count); + HIP_PRINT_STATUS(hipSetDevice(count+1)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp new file mode 100644 index 0000000000..398041b84a --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp @@ -0,0 +1,8 @@ +#include +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipSetDeviceFlags(-1)); + HIP_PRINT_STATUS(hipSetDeviceFlags(11)); +} diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index 1c4a6c7c23..ee207db39d 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" -#include "hip/hcc_detail/hip_texture.h" +#include "hip/hip_texture.h" #define HC __attribute__((hc)) @@ -64,6 +64,8 @@ THE SOFTWARE. printf ("\n");\ printf ("warn: TEST WARNING\n%s", KNRM );\ +#define HIP_PRINT_STATUS(status) \ + std::cout< Date: Sat, 19 Nov 2016 23:02:56 -0600 Subject: [PATCH 10/31] added copy right to new header Change-Id: I16e1d02194551e4b20019bcb6850a3f84882ef18 --- hipamd/include/hip/hip_texture.h | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/hipamd/include/hip/hip_texture.h b/hipamd/include/hip/hip_texture.h index b3d830a0df..53a11d7d45 100644 --- a/hipamd/include/hip/hip_texture.h +++ b/hipamd/include/hip/hip_texture.h @@ -1,3 +1,27 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + + + #ifndef HIP_TEXTURE_H #define HIP_TEXTURE_H From e0aba8647f18a4bee232e712144a06d98d02a37c Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sun, 20 Nov 2016 11:53:16 -0600 Subject: [PATCH 11/31] added new test for getting attribute 1. Added copyright to all new tests 2. Added test for hipDeviceGetAttribute Change-Id: I7a070c5b8316ef6575b3f4c49bda2769aea2a7c4 --- .../tests/src/nvcc/Device/hipChooseDevice.cpp | 19 ++++++++++ .../src/nvcc/Device/hipDeviceGetAttribute.cpp | 38 +++++++++++++++++++ .../src/nvcc/Device/hipDeviceGetLimit.cpp | 19 ++++++++++ hipamd/tests/src/nvcc/Device/hipGetDevice.cpp | 19 ++++++++++ .../src/nvcc/Device/hipGetDeviceCount.cpp | 19 ++++++++++ .../nvcc/Device/hipGetDeviceProperties.cpp | 19 ++++++++++ hipamd/tests/src/nvcc/Device/hipSetDevice.cpp | 19 ++++++++++ .../src/nvcc/Device/hipSetDeviceFlags.cpp | 19 ++++++++++ 8 files changed, 171 insertions(+) create mode 100644 hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp diff --git a/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp b/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp index d2fbfe1852..3a64c6df33 100644 --- a/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp +++ b/hipamd/tests/src/nvcc/Device/hipChooseDevice.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include"test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp new file mode 100644 index 0000000000..f0e56bf274 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp @@ -0,0 +1,38 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include"test_common.h" + +int main() +{ + int val; + hipDeviceAttr_t attr = hipDevAttrMaxThreadsPerBlock; + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1)); + attr = 91; + + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1)); + HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1)); + +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp index b17d5d6a07..57b8594527 100644 --- a/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp index a009e337a6..0e9e5d96f2 100644 --- a/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp +++ b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp index ec02e57ae3..6565741b2f 100644 --- a/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp index 944c4c1516..325d9dd410 100644 --- a/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp index 7a473ef042..99c62a03be 100644 --- a/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp +++ b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" diff --git a/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp index 398041b84a..e933ab56fa 100644 --- a/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp +++ b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include #include "test_common.h" From b3c16ea7b51e0617e995f449e5312afd881fe2f1 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sun, 20 Nov 2016 12:18:08 -0600 Subject: [PATCH 12/31] Fixed hipDeviceGetCacheConfig on nvcc path 1. Changed test macro to emit line numbers 2. Added getcacheconfig api test for nvcc path 3. Fixed hipFuncCache_t data type TODO: With this commit, right now there are 2 func cache datatypes a. hipFuncCache_t for runtime API b. hipFuncCache for driver API Map these to a single data type Change-Id: Ia47c9f5d7c2633638051bf17b1103048a1ede973 --- .../include/hip/hcc_detail/hip_runtime_api.h | 4 +-- .../include/hip/nvcc_detail/hip_runtime_api.h | 5 ++++ hipamd/src/hip_device.cpp | 2 +- .../src/nvcc/Device/hipDeviceGetAttribute.cpp | 5 ++-- .../nvcc/Device/hipDeviceGetCacheConfig.cpp | 28 +++++++++++++++++++ hipamd/tests/src/test_common.h | 2 +- 6 files changed, 40 insertions(+), 6 deletions(-) create mode 100644 hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 09825e141e..65bec1ff1d 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -112,12 +112,12 @@ enum hipLimit_t /** * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored. */ -typedef enum hipFuncCache { +typedef enum hipFuncCache_t { hipFuncCachePreferNone, ///< no preference for shared memory or L1 (default) hipFuncCachePreferShared, ///< prefer larger shared memory and smaller L1 cache hipFuncCachePreferL1, ///< prefer larger L1 cache and smaller shared memory hipFuncCachePreferEqual, ///< prefer equal size L1 cache and shared memory -} hipFuncCache; +} hipFuncCache_t; /** diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 0d15dfcb01..c05d04ee29 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -76,6 +76,7 @@ typedef cudaStream_t hipStream_t; typedef cudaIpcEventHandle_t hipIpcEventHandle_t; typedef cudaIpcMemHandle_t hipIpcMemHandle_t; typedef cudaLimit hipLimit_t; +typedef cudaFuncCache hipFuncCache_t; typedef CUcontext hipCtx_t; typedef CUsharedconfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; @@ -323,6 +324,10 @@ inline static hipError_t hipDeviceSynchronize() { return hipCUDAErrorTohipError(cudaDeviceSynchronize()); } +inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *pCacheConfig) { + return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig)); +} + inline static const char* hipGetErrorString(hipError_t error){ return cudaGetErrorString(hipErrorToCudaError(error)); } diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 29ab0805b8..116f4ade20 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -71,7 +71,7 @@ hipError_t hipGetDeviceCount(int *count) return e; } -hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig) +hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { HIP_INIT_API(cacheConfig); diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp index f0e56bf274..a7e1814468 100644 --- a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp @@ -23,12 +23,13 @@ THE SOFTWARE. int main() { int val; - hipDeviceAttr_t attr = hipDevAttrMaxThreadsPerBlock; + hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock; ///< Maximum number of threads per block. + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1)); - attr = 91; + attr = hipDeviceAttribute_t(91); HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp new file mode 100644 index 0000000000..17f7ed7fbd --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp @@ -0,0 +1,28 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include"test_common.h" + +int main() +{ + hipFuncCache_t pCacheConfig; + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(&pCacheConfig)); + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(NULL)); +} diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index ee207db39d..90c130c56a 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -65,7 +65,7 @@ THE SOFTWARE. printf ("warn: TEST WARNING\n%s", KNRM );\ #define HIP_PRINT_STATUS(status) \ - std::cout< Date: Mon, 21 Nov 2016 08:56:30 -0600 Subject: [PATCH 13/31] fixed compilation bugs 1. Texture functions are now compiling fine 2. Fixed hipFuncCache to hipFuncCache_t Change-Id: I8f815887e4de43ee115bbaff249905b236541c39 --- hipamd/include/hip/hcc.h | 4 ++-- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 10 +++++----- hipamd/include/hip/hcc_detail/hip_texture.h | 8 ++++---- hipamd/include/hip/hip_texture.h | 6 ++++-- hipamd/src/hip_context.cpp | 4 ++-- hipamd/src/hip_device.cpp | 4 ++-- hipamd/tests/src/test_common.h | 1 + 7 files changed, 20 insertions(+), 17 deletions(-) diff --git a/hipamd/include/hip/hcc.h b/hipamd/include/hip/hcc.h index 1542d5b4f2..dba26aeab3 100644 --- a/hipamd/include/hip/hcc.h +++ b/hipamd/include/hip/hcc.h @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HCC_H -#define HCC_H +#ifndef HIP_HCC_H +#define HIP_HCC_H #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include "hip/hcc_detail/hcc_acc.h" diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 65bec1ff1d..7d3e1f78bd 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -321,7 +321,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ -hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ); +hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ); /** @@ -333,7 +333,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig ); * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ -hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig ); +hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t *cacheConfig ); /** * @brief Get Resource limits of current device @@ -357,7 +357,7 @@ hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit); * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache. This hint is ignored on those architectures. * */ -hipError_t hipFuncSetCacheConfig ( hipFuncCache config ); +hipError_t hipFuncSetCacheConfig ( hipFuncCache_t config ); /** * @brief Returns bank width of shared memory for current device @@ -1460,7 +1460,7 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion); * * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ); +hipError_t hipCtxGetCacheConfig ( hipFuncCache_t *cacheConfig ); /** * @brief Set L1/Shared cache partition. @@ -1473,7 +1473,7 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ); * * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ); +hipError_t hipCtxSetCacheConfig ( hipFuncCache_t cacheConfig ); /** * @brief Set Shared memory bank configuration. diff --git a/hipamd/include/hip/hcc_detail/hip_texture.h b/hipamd/include/hip/hcc_detail/hip_texture.h index 7bd14c7a40..fe13f11e49 100644 --- a/hipamd/include/hip/hcc_detail/hip_texture.h +++ b/hipamd/include/hip/hcc_detail/hip_texture.h @@ -22,8 +22,8 @@ THE SOFTWARE. //#pragma once -#ifndef HIP_TEXTURE_H -#define HIP_TEXTURE_H +#ifndef HIP_HCC_DETAIL_TEXTURE_H +#define HIP_HCC_DETAIL_TEXTURE_H /** * @file hcc_detail/hip_texture.h @@ -32,7 +32,7 @@ THE SOFTWARE. #include -#include +//#include //---- //Texture - TODO - likely need to move this to a separate file only included with kernel compilation. @@ -84,7 +84,7 @@ struct texture : public textureReference { }; #endif -typedef struct hipArray { +typedef struct { unsigned int width; unsigned int height; hipChannelFormatKind f; diff --git a/hipamd/include/hip/hip_texture.h b/hipamd/include/hip/hip_texture.h index 53a11d7d45..da30566500 100644 --- a/hipamd/include/hip/hip_texture.h +++ b/hipamd/include/hip/hip_texture.h @@ -22,13 +22,15 @@ THE SOFTWARE. -#ifndef HIP_TEXTURE_H -#define HIP_TEXTURE_H +#ifndef HIP_HIP_TEXTURE_H +#define HIP_HIP_TEXTURE_H #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include +#warning "Including hcc" #elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) #include +#warning "Including Nvcc" #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index a72ae9b58d..8a2451671c 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -217,7 +217,7 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) return ihipLogStatus(hipSuccess); } -hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) +hipError_t hipCtxGetCacheConfig ( hipFuncCache_t *cacheConfig ) { HIP_INIT_API(cacheConfig); @@ -226,7 +226,7 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) return ihipLogStatus(hipSuccess); } -hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) +hipError_t hipCtxSetCacheConfig ( hipFuncCache_t cacheConfig ) { HIP_INIT_API(cacheConfig); diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 116f4ade20..5bf7c5bf4e 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -80,7 +80,7 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) return ihipLogStatus(hipSuccess); } -hipError_t hipDeviceGetCacheConfig(hipFuncCache *cacheConfig) +hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig) { HIP_INIT_API(cacheConfig); @@ -105,7 +105,7 @@ hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit) } } -hipError_t hipFuncSetCacheConfig (hipFuncCache cacheConfig) +hipError_t hipFuncSetCacheConfig (hipFuncCache_t cacheConfig) { HIP_INIT_API(cacheConfig); diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index 90c130c56a..1250de4801 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -24,6 +24,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "hip/hip_texture.h" +#include "hip/hip_runtime_api.h" #define HC __attribute__((hc)) From fef766df881006c2937753764e4969e81fe3397f Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 21 Nov 2016 09:04:36 -0600 Subject: [PATCH 14/31] removed warnings in macros Change-Id: I992b11f6aee2bab09f46885a2d12234aa6814cc5 --- hipamd/include/hip/hip_texture.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/hipamd/include/hip/hip_texture.h b/hipamd/include/hip/hip_texture.h index da30566500..3e7802b457 100644 --- a/hipamd/include/hip/hip_texture.h +++ b/hipamd/include/hip/hip_texture.h @@ -27,10 +27,8 @@ THE SOFTWARE. #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include -#warning "Including hcc" #elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) #include -#warning "Including Nvcc" #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif From 2ded0ce30246f7b295fbad2b49b6e10d5bf0865c Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 21 Nov 2016 13:53:28 -0600 Subject: [PATCH 15/31] fixed texture header on nvcc Change-Id: Ibe19f94be5edf972b6b51dea263e1088b6c60c1d --- hipamd/include/hip/nvcc_detail/hip_texture.h | 2 +- .../Negative/Device/hipDeviceGetAttribute.cpp | 8 +++--- .../Device/hipDeviceGetSharedMemConfig.cpp | 28 +++++++++++++++++++ 3 files changed, 33 insertions(+), 5 deletions(-) create mode 100644 hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp diff --git a/hipamd/include/hip/nvcc_detail/hip_texture.h b/hipamd/include/hip/nvcc_detail/hip_texture.h index cdd939bc1b..388733e492 100644 --- a/hipamd/include/hip/nvcc_detail/hip_texture.h +++ b/hipamd/include/hip/nvcc_detail/hip_texture.h @@ -1,6 +1,6 @@ #ifndef HIP_TEXTURE_H #define HIP_TEXTURE_H -#include +#include #endif diff --git a/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp index b901f660ca..fbe26814e7 100644 --- a/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp +++ b/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp @@ -7,15 +7,15 @@ int main() int pi; int attr = 0; // hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock; - HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); + HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); attr = -1; - HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); + HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute); attr = 0; - HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); + HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); attr = -1; - HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); + HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute); } diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp new file mode 100644 index 0000000000..ce447aa828 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp @@ -0,0 +1,28 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include"test_common.h" + +int main() +{ + hipSharedMemConfig_t config; + HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(NULL)); + HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(&config)); +} From 912426716b17cdd643062d37d81ce2c343bf1363 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 21 Nov 2016 18:07:01 -0600 Subject: [PATCH 16/31] fixed error output for hipDeviceGetAttribute Change-Id: I1e343a4e4e20e1a550d419f701cc1e60e9d03af4 --- hipamd/src/hip_device.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 5bf7c5bf4e..51fd5e4f81 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -174,7 +174,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) hipError_t e = hipSuccess; - if(pi != nullptr) { + if(pi == nullptr) { + return ihipLogStatus(hipErrorInvalidValue); + } auto * hipDevice = ihipGetDevice(device); hipDeviceProp_t *prop = &hipDevice->_props; @@ -236,9 +238,6 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) } else { e = hipErrorInvalidDevice; } - }else{ - e = hipErrorInvalidDevice; - } return ihipLogStatus(e); } From afbd2788045b8f299f5d9c4230c692edc9be5671 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 22 Nov 2016 06:15:35 +0530 Subject: [PATCH 17/31] Removed hsaKmtReleaseSystemProperties call Change-Id: I7cb992cccf587c333f0ca0cb518409f3944bdb06 --- hipamd/src/hip_hcc.cpp | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index c7afd50171..e187a63887 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -774,18 +774,13 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) // Get Max Threads Per Multiprocessor - HsaSystemProperties props; - hsaKmtReleaseSystemProperties(); - if(HSAKMT_STATUS_SUCCESS == hsaKmtAcquireSystemProperties(&props)) { - HsaNodeProperties node_prop = {0}; - if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) { - uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD; - uint32_t simd_per_cu = node_prop.NumSIMDPerCU; - prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu; - } + HsaNodeProperties node_prop = {0}; + if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) { + uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD; + uint32_t simd_per_cu = node_prop.NumSIMDPerCU; + prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu; } - // Get memory properties err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop); DeviceErrorCheck(err); From 1a85762f53dd0de1728b56d37d9d4f4aef2a9ad7 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 22 Nov 2016 10:20:09 -0600 Subject: [PATCH 18/31] added fast math APIs 1. Added fast math apis for sin, cos, tan, sincos 2. Added test for trig math functions 3. Added logarithm fast math 4. Changed how hipGetDevice, hipDeviceGetCacheConfig emit errors Change-Id: Ie6ab594ddd5853cbe85e39a2f6d3479a807fa323 --- hipamd/include/hip/hcc_detail/hip_runtime.h | 8 +- hipamd/src/device_util.cpp | 12 +-- hipamd/src/hip_device.cpp | 8 +- hipamd/tests/src/deviceLib/hip_trig.cpp | 81 +++++++++++++++++++++ 4 files changed, 98 insertions(+), 11 deletions(-) create mode 100644 hipamd/tests/src/deviceLib/hip_trig.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index d7c8d3a675..079d681c39 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -478,7 +478,7 @@ __host__ __device__ int max(int arg1, int arg2); __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr(); //TODO - add a couple fast math operations here, the set here will grow : -__device__ float __cosf(float x); + __device__ float __expf(float x); __device__ float __frsqrt_rn(float x); __device__ float __fsqrt_rd(float x); @@ -486,11 +486,13 @@ __device__ float __fsqrt_rn(float x); __device__ float __fsqrt_ru(float x); __device__ float __fsqrt_rz(float x); __device__ float __log10f(float x); -__device__ float __log2f(float x); +//__device__ float __log2f(float x); __device__ float __logf(float x); __device__ float __powf(float base, float exponent); __device__ void __sincosf(float x, float *s, float *c) ; -__device__ float __sinf(float x); +extern __attribute__((const)) float __sinf(float) __asm("llvm.sin.f32"); +extern __attribute__((const)) float __cosf(float) __asm("llvm.cos.f32"); +extern __attribute__((const)) float __log2f(float) __asm("llvm.log2.f32"); __device__ float __tanf(float x); __device__ float __dsqrt_rd(double x); __device__ float __dsqrt_rn(double x); diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index deb0db0a34..fc966f3395 100644 --- a/hipamd/src/device_util.cpp +++ b/hipamd/src/device_util.cpp @@ -156,7 +156,7 @@ __device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) { one1 = in1.val & MASK2; one2 = in2.val & MASK2; out.val = out.val | ((one1 - one2) & MASK2); - return out; + return out; } __device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) { @@ -2045,7 +2045,7 @@ __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr() //TODO - add a couple fast math operations here, the set here will grow : -__device__ float __cosf(float x) {return hc::fast_math::cosf(x); }; +//__device__ float __cosf(float x) {return hc::fast_math::cosf(x); }; __device__ float __expf(float x) {return hc::fast_math::expf(x); }; __device__ float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); }; __device__ float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); }; @@ -2053,12 +2053,12 @@ __device__ float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); }; __device__ float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); }; __device__ float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); }; __device__ float __log10f(float x) {return hc::fast_math::log10f(x); }; -__device__ float __log2f(float x) {return hc::fast_math::log2f(x); }; +//__device__ float __log2f(float x) {return hc::fast_math::log2f(x); }; __device__ float __logf(float x) {return hc::fast_math::logf(x); }; __device__ float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); }; -__device__ void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); }; -__device__ float __sinf(float x) {return hc::fast_math::sinf(x); }; -__device__ float __tanf(float x) {return hc::fast_math::tanf(x); }; +__device__ void __sincosf(float x, float *s, float *c) { *s = __sinf(x); *c = __cosf(x); }; +//__device__ float __sinf(float x) {return hc::fast_math::sinf(x); }; +__device__ float __tanf(float x) {return __sinf(x)/__cosf(x); }; __device__ float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); }; __device__ float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); }; __device__ float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); }; diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 51fd5e4f81..371578ca2c 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -44,7 +44,7 @@ hipError_t hipGetDevice(int *deviceId) *deviceId = ctx->getDevice()->_deviceId; } }else{ - e = hipErrorInvalidDevice; + e = hipErrorInvalidValue; } return ihipLogStatus(e); @@ -66,7 +66,7 @@ hipError_t hipGetDeviceCount(int *count) e = ihipLogStatus(hipErrorNoDevice); } } else { - e = ihipLogStatus(hipErrorNoDevice); + e = ihipLogStatus(hipErrorInvalidValue); } return e; } @@ -84,6 +84,10 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig) { HIP_INIT_API(cacheConfig); + if(cacheConfig == nullptr) { + return ihipLogStatus(hipErrorInvalidValue); + } + *cacheConfig = hipFuncCachePreferNone; return ihipLogStatus(hipSuccess); diff --git a/hipamd/tests/src/deviceLib/hip_trig.cpp b/hipamd/tests/src/deviceLib/hip_trig.cpp new file mode 100644 index 0000000000..7f9b5d60b0 --- /dev/null +++ b/hipamd/tests/src/deviceLib/hip_trig.cpp @@ -0,0 +1,81 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s + * RUN: %t + * HIT_END + */ + +#include +#include +#include +#include"test_common.h" + +#define LEN 512 +#define SIZE LEN<<2 + +__global__ void kernel_trig(hipLaunchParm lp, float *In, float *sin_d, float *cos_d, float *tan_d, float *sin_pd, float *cos_pd){ + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + sin_d[tid] = __sinf(In[tid]); + cos_d[tid] = __cosf(In[tid]); + tan_d[tid] = __tanf(In[tid]); + __sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]); +} + +int main(){ + float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph; + float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd; + In = new float[LEN]; + sin_h = new float[LEN]; + cos_h = new float[LEN]; + tan_h = new float[LEN]; + sin_ph = new float[LEN]; + cos_ph = new float[LEN]; + for(int i=0;i Date: Tue, 22 Nov 2016 15:26:00 -0600 Subject: [PATCH 19/31] added fast math intrinsics to HIP 1. Added fast math intrinsics for single precision data types 2. Added test to check the intrinsics 3. Added HIP_PRECISE_MATH macro to enable precise math on fast math Change-Id: Iadacbb6182c31252c5e3252854372d1b80dfd27b --- hipamd/include/hip/hcc_detail/hip_runtime.h | 205 ++++++++++++++++++-- hipamd/src/device_util.cpp | 163 ++++++++++++++-- hipamd/tests/src/deviceLib/hipFloatMath.cpp | 61 ++++++ 3 files changed, 391 insertions(+), 38 deletions(-) create mode 100644 hipamd/tests/src/deviceLib/hipFloatMath.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 079d681c39..4b781b44ae 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -479,26 +479,193 @@ __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr(); //TODO - add a couple fast math operations here, the set here will grow : -__device__ float __expf(float x); -__device__ float __frsqrt_rn(float x); -__device__ float __fsqrt_rd(float x); -__device__ float __fsqrt_rn(float x); -__device__ float __fsqrt_ru(float x); -__device__ float __fsqrt_rz(float x); -__device__ float __log10f(float x); -//__device__ float __log2f(float x); -__device__ float __logf(float x); -__device__ float __powf(float base, float exponent); -__device__ void __sincosf(float x, float *s, float *c) ; -extern __attribute__((const)) float __sinf(float) __asm("llvm.sin.f32"); -extern __attribute__((const)) float __cosf(float) __asm("llvm.cos.f32"); -extern __attribute__((const)) float __log2f(float) __asm("llvm.log2.f32"); -__device__ float __tanf(float x); -__device__ float __dsqrt_rd(double x); -__device__ float __dsqrt_rn(double x); -__device__ float __dsqrt_ru(double x); -__device__ float __dsqrt_rz(double x); +// Single Precision Precise Math +__device__ float __hip_precise_cosf(float); +__device__ float __hip_precise_exp10f(float); +__device__ float __hip_precise_expf(float); +__device__ float __hip_precise_frsqrt_rn(float); +__device__ float __hip_precise_fsqrt_rd(float); +__device__ float __hip_precise_fsqrt_rn(float); +__device__ float __hip_precise_fsqrt_ru(float); +__device__ float __hip_precise_fsqrt_rz(float); +__device__ float __hip_precise_log10f(float); +__device__ float __hip_precise_log2f(float); +__device__ float __hip_precise_logf(float); +__device__ float __hip_precise_powf(float, float); +__device__ void __hip_precise_sincosf(float,float*,float*); +__device__ float __hip_precise_sinf(float); +__device__ float __hip_precise_tanf(float); +// Double Precision Precise Math +__device__ double __hip_precise_dsqrt_rd(double); +__device__ double __hip_precise_dsqrt_rn(double); +__device__ double __hip_precise_dsqrt_ru(double); +__device__ double __hip_precise_dsqrt_rz(double); + +// 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"); +__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"); +__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"); +__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"); +__device__ float __hip_fast_tanf(float); + +#ifdef HIP_PRECISE_MATH +// Single Precision Precise Math when enabled + +__device__ inline float __cosf(float x) { + return __hip_precise_cosf(x); +} + +__device__ inline float __exp10f(float x) { + return __hip_precise_exp10f(x); +} + +__device__ inline float __expf(float x) { + return __hip_precise_expf(x); +} + +__device__ inline float __frsqrt_rn(float x) { + return __hip_precise_frsqrt_rn(x); +} + +__device__ inline float __fsqrt_rd(float x) { + return __hip_precise_fsqrt_rd(x); +} + +__device__ inline float __fsqrt_rn(float x) { + return __hip_precise_fsqrt_rn(x); +} + +__device__ inline float __fsqrt_ru(float x) { + return __hip_precise_fsqrt_ru(x); +} + +__device__ inline float __fsqrt_rz(float x) { + return __hip_precise_fsqrt_rz(x); +} + +__device__ inline float __log10f(float x) { + return __hip_precise_log10f(x); +} + +__device__ inline float __log2f(float x) { + return __hip_precise_log2f(x); +} + +__device__ inline float __logf(float x) { + return __hip_precise_logf(x); +} + +__device__ inline float __powf(float base, float exponent) { + return __hip_precise_powf(base, exponent); +} + +__device__ inline void __sincosf(float x, float *s, float *c) { + return __hip_precise_sincosf(x, s, c); +} + +__device__ inline float __sinf(float x) { + return __hip_precise_sinf(x); +} + +__device__ inline float __tanf(float x) { + return __hip_precise_tanf(x); +} + +// Double Precision + +__device__ double __dsqrt_rd(double x) { + return __hip_precise_dsqrt_rd(x); +} + +__device__ double __dsqrt_rn(double x) { + return __hip_precise_dsqrt_rn(x); +} + +__device__ double __dsqrt_ru(double x) { + return __hip_precise_dsqrt_ru(x); +} + +__device__ double __dsqrt_rz(double x) { + return __hip_precise_dsqrt_rz(x); +} + +#else + +// Single Precision Fast Math +__device__ inline float __cosf(float x) { + return __hip_fast_cosf(x); +} + +__device__ inline float __exp10f(float x) { + return __hip_fast_exp10f(x); +} + +__device__ inline float __expf(float x) { + return __hip_fast_expf(x); +} + +__device__ inline float __frsqrt_rn(float x) { + return __hip_fast_frsqrt_rn(x); +} + +__device__ inline float __fsqrt_rd(float x) { + return __hip_fast_fsqrt_rd(x); +} + +__device__ inline float __fsqrt_rn(float x) { + return __hip_fast_fsqrt_rn(x); +} + +__device__ inline float __fsqrt_ru(float x) { + return __hip_fast_fsqrt_ru(x); +} + +__device__ inline float __fsqrt_rz(float x) { + return __hip_fast_fsqrt_rz(x); +} + +__device__ inline float __log10f(float x) { + return __hip_fast_log10f(x); +} + +__device__ inline float __log2f(float x) { + return __hip_fast_log2f(x); +} + +__device__ inline float __logf(float x) { + return __hip_fast_logf(x); +} + +__device__ inline float __powf(float base, float exponent) { + return __hip_fast_powf(base, exponent); +} + +__device__ inline void __sincosf(float x, float *s, float *c) { + return __hip_fast_sincosf(x, s, c); +} + +__device__ inline float __sinf(float x) { + return __hip_fast_sinf(x); +} + +__device__ inline float __tanf(float x) { + return __hip_fast_tanf(x); +} + + +#endif /** * CUDA 8 device function features diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index fc966f3395..203e7a7826 100644 --- a/hipamd/src/device_util.cpp +++ b/hipamd/src/device_util.cpp @@ -2043,26 +2043,151 @@ __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr() } +// Precise Math Functions +__device__ float __hip_precise_cosf(float x) { + return hc::precise_math::cosf(x); +} -//TODO - add a couple fast math operations here, the set here will grow : -//__device__ float __cosf(float x) {return hc::fast_math::cosf(x); }; -__device__ float __expf(float x) {return hc::fast_math::expf(x); }; -__device__ float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); }; -__device__ float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); }; -__device__ float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); }; -__device__ float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); }; -__device__ float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); }; -__device__ float __log10f(float x) {return hc::fast_math::log10f(x); }; -//__device__ float __log2f(float x) {return hc::fast_math::log2f(x); }; -__device__ float __logf(float x) {return hc::fast_math::logf(x); }; -__device__ float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); }; -__device__ void __sincosf(float x, float *s, float *c) { *s = __sinf(x); *c = __cosf(x); }; -//__device__ float __sinf(float x) {return hc::fast_math::sinf(x); }; -__device__ float __tanf(float x) {return __sinf(x)/__cosf(x); }; -__device__ float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); }; -__device__ float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); }; -__device__ float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); }; -__device__ float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); }; +__device__ float __hip_precise_exp10f(float x) { + return hc::precise_math::exp10f(x); +} + +__device__ float __hip_precise_expf(float x) { + return hc::precise_math::expf(x); +} + +__device__ float __hip_precise_frsqrt_rn(float x) { + return hc::precise_math::rsqrt(x); +} + +__device__ float __hip_precise_fsqrt_rd(float x) { + return hc::precise_math::sqrt(x); +} + +__device__ float __hip_precise_fsqrt_rn(float x) { + return hc::precise_math::sqrt(x); +} + +__device__ float __hip_precise_fsqrt_ru(float x) { + return hc::precise_math::sqrt(x); +} + +__device__ float __hip_precise_fsqrt_rz(float x) { + return hc::precise_math::sqrt(x); +} + +__device__ float __hip_precise_log10f(float x) { + return hc::precise_math::log10(x); +} + +__device__ float __hip_precise_log2f(float x) { + return hc::precise_math::log2(x); +} + +__device__ float __hip_precise_logf(float x) { + return hc::precise_math::logf(x); +} + +__device__ float __hip_precise_powf(float base, float exponent) { + return hc::precise_math::powf(base, exponent); +} + +__device__ void __hip_precise_sincosf(float x, float *s, float *c) { + hc::precise_math::sincosf(x, s, c); +} + +__device__ float __hip_precise_sinf(float x) { + return hc::precise_math::sinf(x); +} + +__device__ float __hip_precise_tanf(float x) { + return hc::precise_math::tanf(x); +} + +// Double Precision Math +__device__ double __hip_precise_dsqrt_rd(double x) { + return hc::precise_math::sqrt(x); +} + +__device__ double __hip_precise_dsqrt_rn(double x) { + return hc::precise_math::sqrt(x); +} + +__device__ double __hip_precise_dsqrt_ru(double x) { + return hc::precise_math::sqrt(x); +} + +__device__ double __hip_precise_dsqrt_rz(double x) { + return hc::precise_math::sqrt(x); +} + +#define LOG_BASE2_E_DIV_2 0.4426950408894701 +#define LOG_BASE2_5 2.321928094887362 +#define ONE_DIV_LOG_BASE2_E 0.69314718056 +#define ONE_DIV_LOG_BASE2_10 0.30102999566 + +// Fast Math Intrinsics +__device__ float __hip_fast_exp10f(float x) { + return __hip_fast_exp2f(x*LOG_BASE2_E_DIV_2); +} + +__device__ float __hip_fast_expf(float x) { + return __hip_fast_expf(x*LOG_BASE2_5); +} + +__device__ float __hip_fast_frsqrt_rn(float x) { + return 1 / __hip_fast_fsqrt_rd(x);; +} + +__device__ float __hip_fast_fsqrt_rn(float x) { + return __hip_fast_fsqrt_rd(x); +} + +__device__ float __hip_fast_fsqrt_ru(float x) { + return __hip_fast_fsqrt_rd(x); +} + +__device__ float __hip_fast_fsqrt_rz(float x) { + return __hip_fast_fsqrt_rd(x); +} + +__device__ float __hip_fast_log10f(float x) { + return ONE_DIV_LOG_BASE2_E * __hip_fast_log2f(x); +} + +__device__ float __hip_fast_logf(float x) { + return ONE_DIV_LOG_BASE2_10 * __hip_fast_log2f(x); +} + +__device__ float __hip_fast_powf(float base, float exponent) { + return hc::fast_math::powf(base, exponent); +} + +__device__ void __hip_fast_sincosf(float x, float *s, float *c) { + *s = __hip_fast_sinf(x); + *c = __hip_fast_cosf(x); +} + +__device__ float __hip_fast_tanf(float x) { + return hc::fast_math::tanf(x); +} + +// Double Precision Math +__device__ double __hip_fast_dsqrt_rd(double x) { + return hc::fast_math::sqrt(x); +} + +__device__ double __hip_fast_dsqrt_rn(double x) { + return hc::fast_math::sqrt(x); +} + +__device__ double __hip_fast_dsqrt_ru(double x) { + return hc::fast_math::sqrt(x); +} + +__device__ double __hip_fast_dsqrt_rz(double x) { + return hc::fast_math::sqrt(x); +} __HIP_DEVICE__ char1 make_char1(signed char x) { diff --git a/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/hipamd/tests/src/deviceLib/hipFloatMath.cpp new file mode 100644 index 0000000000..eb70eb6b0b --- /dev/null +++ b/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -0,0 +1,61 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +#define LEN 512 +#define SIZE LEN<<2 + + + +__global__ void floatMath(hipLaunchParm lp, float *In, float *Out) { + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + Out[tid] = __cosf(In[tid]); + Out[tid] = __exp10f(Out[tid]); + Out[tid] = __expf(Out[tid]); + Out[tid] = __frsqrt_rn(Out[tid]); + Out[tid] = __fsqrt_rd(Out[tid]); + Out[tid] = __fsqrt_rn(Out[tid]); + Out[tid] = __fsqrt_ru(Out[tid]); + Out[tid] = __fsqrt_rz(Out[tid]); + Out[tid] = __log10f(Out[tid]); + Out[tid] = __log2f(Out[tid]); + Out[tid] = __logf(Out[tid]); + Out[tid] = __powf(2.0f, Out[tid]); + __sincosf(Out[tid], &In[tid], &Out[tid]); + Out[tid] = __sinf(Out[tid]); + Out[tid] = __cosf(Out[tid]); + Out[tid] = __tanf(Out[tid]); +} + +int main(){ + float *Inh, *Outh, *Ind, *Outd; + hipMalloc((void**)&Ind, SIZE); + hipMalloc((void**)&Outd, SIZE); + hipLaunchKernel(floatMath, dim3(LEN,1,1), dim3(1,1,1), 0, 0, Ind, Outd); +} From 8a2685e6cdd253494c6b8aefb484c71fec330479 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 23 Nov 2016 18:37:06 +0530 Subject: [PATCH 20/31] Removed nested HIP calls from hip_device functions Change-Id: I18785b0ee27e32fb8950982fa5c3a64d1ae6a9b8 --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- hipamd/src/hip_device.cpp | 64 +++++++++++-------- 2 files changed, 39 insertions(+), 27 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 7d3e1f78bd..77853f02a2 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -413,7 +413,7 @@ hipError_t hipSetDeviceFlags ( unsigned flags); * * @returns #hipSuccess, #hipErrorInvalidValue */ -hipError_t hipChooseDevice(int *device, hipDeviceProp_t* prop); +hipError_t hipChooseDevice(int *device, const hipDeviceProp_t* prop); // end doxygen Device /** diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 371578ca2c..4f5729791e 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -51,10 +51,8 @@ hipError_t hipGetDevice(int *deviceId) } // TODO - does this initialize HIP runtime? -hipError_t hipGetDeviceCount(int *count) +hipError_t ihipGetDeviceCount(int *count) { - HIP_INIT_API(count); - hipError_t e = hipSuccess; if(count != nullptr) { @@ -71,6 +69,12 @@ hipError_t hipGetDeviceCount(int *count) return e; } +hipError_t hipGetDeviceCount(int *count) +{ + HIP_INIT_API(count); + return ihipGetDeviceCount(count); +} + hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { HIP_INIT_API(cacheConfig); @@ -172,10 +176,8 @@ hipError_t hipDeviceReset(void) return ihipLogStatus(hipSuccess); } -hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) +hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - HIP_INIT_API(pi, attr, device); - hipError_t e = hipSuccess; if(pi == nullptr) { @@ -245,10 +247,14 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) return ihipLogStatus(e); } -hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) +hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - HIP_INIT_API(props, device); + HIP_INIT_API(pi, attr, device); + return ihipDeviceGetAttribute(pi,attr,device); +} +hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device) +{ hipError_t e; if(props != nullptr){ @@ -267,6 +273,12 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) return ihipLogStatus(e); } +hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) +{ + HIP_INIT_API(props, device); + return ihipGetDeviceProperties(props, device); +} + hipError_t hipSetDeviceFlags( unsigned int flags) { HIP_INIT_API(flags); @@ -311,8 +323,8 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device HIP_INIT_API(major,minor, device); hipError_t e = hipSuccess; int deviceId= device->_deviceId; - e = hipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId); - e = hipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId); + e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId); + e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId); return ihipLogStatus(e); } @@ -331,7 +343,7 @@ hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device) HIP_INIT_API(pciBusId,len, device); hipError_t e = hipSuccess; int deviceId= device->_deviceId; - e = hipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId); + e = ihipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId); return ihipLogStatus(e); } @@ -349,10 +361,10 @@ hipError_t hipDeviceGetByPCIBusId (int* device, const int* pciBusId ) hipDeviceProp_t tempProp; int deviceCount; hipError_t e = hipErrorInvalidValue; - hipGetDeviceCount( &deviceCount ); + ihipGetDeviceCount( &deviceCount ); *device = 0; - for (int i=0; i< deviceCount; i++) { - hipGetDeviceProperties( &tempProp, i ); + for (int i = 0; i< deviceCount; i++) { + ihipGetDeviceProperties( &tempProp, i ); if(tempProp.pciBusID == *pciBusId) { *device =i; e = hipSuccess; @@ -367,19 +379,19 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) HIP_INIT_API(device,prop); hipDeviceProp_t tempProp; int deviceCount; - int inPropCount=0; - int matchedPropCount=0; + int inPropCount = 0; + int matchedPropCount = 0; hipError_t e = hipSuccess; - hipGetDeviceCount( &deviceCount ); + ihipGetDeviceCount( &deviceCount ); *device = 0; - for (int i=0; i< deviceCount; i++) { - hipGetDeviceProperties( &tempProp, i ); - if(prop->major !=0) { + for (int i = 0; i < deviceCount; i++) { + ihipGetDeviceProperties( &tempProp, i ); + if(prop->major != 0) { inPropCount++; if(tempProp.major >= prop->major) { matchedPropCount++; } - if(prop->minor !=0) { + if(prop->minor != 0) { inPropCount++; if(tempProp.minor >= prop->minor) { matchedPropCount++; @@ -398,31 +410,31 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) matchedPropCount++; } } - if(prop->maxThreadsPerBlock != 0) { + if(prop->maxThreadsPerBlock != 0) { inPropCount++; if(tempProp.maxThreadsPerBlock >= prop->maxThreadsPerBlock ) { matchedPropCount++; } } - if(prop->totalConstMem != 0) { + if(prop->totalConstMem != 0) { inPropCount++; if(tempProp.totalConstMem >= prop->totalConstMem ) { matchedPropCount++; } } - if(prop->multiProcessorCount != 0) { + if(prop->multiProcessorCount != 0) { inPropCount++; if(tempProp.multiProcessorCount >= prop->multiProcessorCount ) { matchedPropCount++; } } - if(prop->maxThreadsPerMultiProcessor != 0) { + if(prop->maxThreadsPerMultiProcessor != 0) { inPropCount++; if(tempProp.maxThreadsPerMultiProcessor >= prop->maxThreadsPerMultiProcessor ) { matchedPropCount++; } } - if(prop->memoryClockRate != 0) { + if(prop->memoryClockRate != 0) { inPropCount++; if(tempProp.memoryClockRate >= prop->memoryClockRate ) { matchedPropCount++; From 4bcb0fac228d02c535e52304048938a56d963b1b Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 23 Nov 2016 08:15:04 -0600 Subject: [PATCH 21/31] Improve profiler and debug documentation --- hipamd/docs/markdown/hip_profiling.md | 309 ++++++++++++++++++++++---- 1 file changed, 262 insertions(+), 47 deletions(-) diff --git a/hipamd/docs/markdown/hip_profiling.md b/hipamd/docs/markdown/hip_profiling.md index e4b88945e5..ee3a8d6656 100644 --- a/hipamd/docs/markdown/hip_profiling.md +++ b/hipamd/docs/markdown/hip_profiling.md @@ -1,31 +1,188 @@ -# Profiling HIP Code +# Profiling and Debugging HIP Code -HIP provides several capabilities to support debugging and profiling. Profiling information can be displayed to stderr or viewed in the CodeXl visualization tool. +This section describes the profiling and debugging capabilities that HIP provides. +Profiling information can viewed in the CodeXL visualization tool or printed directly to stderr as the application runs. +This document starts with some of the general capabilities of CodeXL and then describes some of the additional HIP marker and debug features. -### Usign CodeXL to profile a HIP Application -By defauly, CodeXL can trace all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls. -/opt/rocm/bin/rocm-profiler -o -A +## CodeXL Profiling -### Using CodeXL markers for HIP Functions -HIP can generate markers at function being/end which are displayed on the CodeXL timeline view. +### Collecting and Viewing Traces + +#### Using rocm-profiler timestamp profiling +rocm-profiler is a command-line tool for tracing any application that uses ROCr API, including HCC and HIP. +rocm-profiler's timeline trace will show the beginning and end for all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls. The trace results are saved into a file, which by convention uses the "atp" extension. Here is an example that shows how to run the command-line profiler: +```shell +$ /opt/rocm/bin/rocm-profiler -o -A -T +``` + +#### Using rocm-profiler performance counter collection: +rocm-profiler can record performance counter information to provide greater insight inside a kernel, such as the memory bandwidth, ALU busy percentage, and cache statistics. +Collecting the common set of useful counters requires passing the counter configuration files for two passes: +``` +$ /opt/rocm/bin/rocm-profiler -C -O --counterfile /opt/rocm/profiler/counterfiles/counters_HSA_Fiji_pass1 --counterfile /opt/rocm/profiler/counterfiles/counters_HSA_Fiji_pass2 +``` + + +#### Using CodeXL to view profiling results: +The trace can be loaded and viewed in the CodeXL visualization tool: + +- Open the CodeXL GUI, create an new project, and switch to "Profile Mode": + - $ CodeXL & + - [File->New Project, leave fields as is, just click "OK"] + - [Profile->Switch to Profile Mode] +- Load timestamp tracing results into a timeline view: + - Right click on the project + - Click "Import Session..." + - Select to $HOME/apitrace.atp (or appropriate .atp file if you used another file name) + +- Load the performance counter results + - Right click on the project + - Click "Import Session..." + - Select $HOME/Session1.csv (or appropriate .csv file if you used another file name) + + +#### More information on CodeXL +rocm-profiler --help will show additional options and usage guidelines. + +See this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) for more information on profiling ROCm apps (including HIP) with CodeXL. + +### HIP Markers +#### Profiling HIP APIs +HIP can generate markers at function beginning and end which are displayed on the CodeXL timeline view. HIP 1.0 compiles marker support by default, and you can enable it by setting the HIP_PROFILE_API environment variable and then running the rocm-profiler: ```shell # Use profile to generate timeline view: export HIP_PROFILE_API=1 -/opt/rocm/bin/rocm-profiler -o -A +$ /opt/rocm/bin/rocm-profiler -A -T Or -/opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -o -A +$ /opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -A -T ``` -#### Developer Builds +HIP_PROFILE_API supports two levels of information. +- HIP_PROFILE_API=1 : Short format. Print name of API but no arguments. For example: +`hipMemcpy` +- HIP_PROFILE_API=2 : Long format. Print name of API + values of all function arguments. For example: +`hipMemcpy (0x7f32154db010, 0x50446e000, 4000000, hipMemcpyDeviceToHost)` + +#### Adding markers to applications + +Markers can be used to define application-specific events that will be recorded in the ATP file and displayed in the CodeXL gui. +This can be particularly useful for visualizing how the higher-level phases of application behavior relate to the lower level HIP APIs, kernel launches, and data transfers. +For example, an instrumented machine learning framework could show the beginning and ending of each layer in the network. + +Markers have a specific begin and end time, and can be nested. Nested calls are displayed hierarchically in the CodeXL gui, with each level of the hierarchy occupying a different row. + +The HIP APis are defined in "hip_profile.h": +``` +#include + +HIP_BEGIN_MARKER(const char *markerName, const char *groupName); +HIP_END_MARKER(); + +HIP_BEGIN_MARKER("Setup", "MyAppGroup"); +// ... +// application code for setup +// ... +HIP_END_MARKER(); +``` + +For C++ codes, HIP also provides a scoped marker which records the start time when constructed and the end time when the scoped marker is destructed at the end of the scope. This provides a convenient, single-line mechanism to record an event that neatly corresponds to a region of code. + +```cxx +void FunctionFoo(...) +{ + HIP_SCOPED_MARKER("FunctionFoo", "MyAppGroup"); // Marker starts recording here. + + // ... + // Function implementation + // ... + + // Marker destroyed here and records end time stamp. +}; +``` + +The HIP marker API is only supported on ROCm platform. The marker macros are defined on CUDA platforms and will compile, but are silently ignored at runtime. + +This [HIP sample](samples/2_Cookbook/2_Profiler/) shows the profiler marker API used in a small application. + +More information on the marker API can be found in the profiler header file and PDF in a ROCM installation: +- /opt/rocm/profiler/CXLActivityLogger/include/CXLActivityLogger.h +- /opt/rocm/profiler/CXLActivityLogger/doc/CXLActivityLogger.pdf + +### Additional HIP Profiling Features +#### Demangling C++ Kernel Names +HIP includes the `hipdemangleatp` tool which can post-process an ATP file to "demangle" C++ names. +Mangled kernel names encode the C++ arguments and other information, and are guaranteed to be unique even for cases such as operator overloading. However, the mangled names can be quite verbose. For example: + +`ZZ39gemm_NoTransA_MICRO_NBK_M_N_K_TS16XMTS4RN2hc16accelerator_viewEPKflS3_lPfliiiiiiffEN3_EC__719__cxxamp_trampolineElililiiiiiiS3_iS3_S4_ff` + +`hipdemangleatp` will convert this into the more readable: +`gemm_NoTransA_MICRO_NBK_M_N_K_TS16XMTS4` + +The `hipdemangleatp` tool operates on the ATP file "in-place" and thus replaces the input file with the demangled version. + +``` +$ hipdemangleatp myfile.atp +``` + +The kernel name is also shown in some of the summary htlm files (Top10 kernels). These can be regenerated from the demangled ATP file by re-running rocm-profiler: +``` +$ rocm-profiler -T --atpfile myfile.atp +``` + +A future version of CodeXL may directly integrate demangle functionality. + + +#### Controlling when profiling starts and ends +hipProfilerStart() and hipProfilerEnd() can be inserted into an application to control which phases of the applications are profiled. +These APIs can be used to skip initialization code or to focus profiling on a desired region, and are particularly useful for large long-running applications. +See the API documentation for more information. These APIs work on both ROCm and CUDA paths. + +On ROCm, the following environment variables can be used to control when profiling occurs: + +``` +HIP_DB_START_API : Comma-separated list of tid.api_seq_num for when to start debug and profiling. +HIP_DB_STOP_API : Comma-separated list of tid.api_seq_num for when to stop debug and profiling. +``` + +HIP/ROCm assigns a monotonically increasing sequence number to the APIs called from each thread. The thread and API sequence number can be used in the above API to control when tracing starts and stops. These flags also control the HIP_DB messages (described below). + +When using these options, start the profiler with profiling disabled: +``` +# ROCm: +$ rocm-profiler --startdisabled ... + +# CUDA: +$ nvprof --profile-from-start-off ... +``` + +This feature is under development. + +#### Reducing timeline trace output file size +If the application is already recording the HIP APIs, the HSA APIs are somewhat redundant and the ATP file size can be substantially reduced by not recording these APIs. HIP includes a text file that lists all of the HSA APis and can assist in this filtering: + +``` +$ rocm-profiler -F hip/bin/hsa-api-filter-cxl.txt +``` + +This file can be copied and edited to provide more selective HSA event recording. + + +#### How to enable profiling at HIP build time +Recent pre-built packages of HIP are always built with profiling support enabled. For developer builds, you must enable marker support manually when compiling HIP. 1. Build HIP with ATP markers enabled HIP pre-built packages are enabled with ATP marker support by default. -To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step. +To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step. Build and install HIP. +```shell +$ mkdir build && cd build +$ cmake .. -DCOMPILE_HIP_ATP_MARKER +$ make install +``` 2. Install ROCm-Profiler Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well. @@ -36,20 +193,57 @@ Alternatively, you can build ROCm-Profiler using the instructions [here](https:/ Then follow the steps above to collect a marker-enabled trace. -### Using HIP_TRACE_API -You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided -by the HIP_DB switch. For example: -```shell -# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function. -HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp +## Tracing and Debug + +### Tracing HIP APIs +The HIP runtime can print the HIP function strings to stderr using HIP_TRACE_API environment variable. +The trace prints two messages for each API - one at the beginning of the API call (line starts with "<<") and one at the end of the API call (line ends with ">>"). +Here's an example for one API followed by a description for the sections of the trace: + +``` +<> ``` -#### Color -Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. -You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable. Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White. -None will disable use of color control codes and may be useful when saving the trace file or when a pure text trace is desired. +- `<> +info: running on device gfx803 +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +<> +<> +info: copy Host2Device +<> +info: launch 'vector_square' kernel +1.5 hipLaunchKernel 'HIP_KERNEL_NAME(vector_square)' gridDim:{512,1,1} groupDim:{256,1,1} sharedMem:+0 stream#0.0 +info: copy Device2Host +<> +info: check result +PASSED! +``` + + +#### Color +Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. +You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable. Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White. +None will disable use of color control codes for both the opening and closing and may be useful when saving the trace file or when a pure text trace is desired. -#### ### Using HIP_DB @@ -57,39 +251,60 @@ None will disable use of color control codes and may be useful when saving the t This flag is primarily targeted to assist HIP development team in the development of the HIP runtime, but in some situations may be useful to HIP application developers as well. The HIP debug information is designed to print important information during the execution of a HIP API. HIP provides different color-coded levels of debug informaton: - - api : Print the beginning and end of each HIP API, including the arguments and return codes. + - api : Print the beginning and end of each HIP API, including the arguments and return codes. This is equivalent to setting HIP_TRACE_API=1. - sync : Print multi-thread and other synchronization debug information. - copy : Print which engine is doing the copy, which copy flavor is selected, information on source and destination memory. - mem : Print information about memory allocation - which pointers are allocated, where they are allocated, peer mappings, and more. DB_MEM format is flags separated by '+' sign, or a hex code for the bitmask. Generally the + format is preferred. For example: -```shell -HIP_DB=api+copy+mem my-application -HIP_DB=0xF my-application ``` -HIP_DB=1 same as HIP_TRACE_API=1 +$ HIP_DB=api+copy+mem my-application +$ HIP_DB=0xF my-application +``` + +### Using ltrace +ltrace is a standard linux tool which provides a message to stderr on every dynamic library call. Since ROCr and the ROCt (the ROC thunk, which is the thin user-space interface to the ROC kernel driver) are both dynamic libraries, this provides an easy way to trace the activity in these libraries. Tracing can be a powerful way to quickly observe the flow of the application before diving into the details with a command-line debugger. +The trace can also show performance issues related to accidental calls to expensive API calls on the critical path. + +ltrace can be easily combined with the HIP_DB switches to visualize the runtime behavior of the entire ROCm software stack. Here's a sample command-line and output: + +``` +$ HIP_DB=api ltrace -C -e 'hsa*' + +... + +<hsa_signal_store_relaxed(0x1804000, 0, 0, 0x400000) = 0 +libmcwamp_hsa.so->hsa_signal_store_relaxed(0x1816000, 0, 0x7f777f85f2a0, 0x400000) = 0 +libmcwamp_hsa.so->hsa_amd_memory_lock(0x7f7776d3e010, 0x400000, 0x1213b70, 1 +libhsa-runtime64.so.1->hsaKmtRegisterMemoryToNodes(0x7f7776d3e010, 0x400000, 1, 0x1220c10) = 0 +libhsa-runtime64.so.1->hsaKmtMapMemoryToGPUNodes(0x7f7776d3e010, 0x400000, 0x7ffc32865400, 64) = 0 +<... hsa_amd_memory_lock resumed> ) = 0 +libmcwamp_hsa.so->hsa_signal_store_relaxed(0x1804000, 1, 0x7f777e95a770, 0x12205b0) = 0 +libmcwamp_hsa.so->hsa_amd_memory_async_copy(0x50411d010, 0x11e70d0, 0x503d1d000, 0x11e70d0) = 0 +libmcwamp_hsa.so->hsa_signal_wait_acquire(0x1804000, 2, 1, -1) = 0 +libmcwamp_hsa.so->hsa_amd_memory_unlock(0x7f7776d3e010, 0x1213c6c, 0x12c3c600000000, 0x1804000 +libhsa-runtime64.so.1->hsaKmtUnmapMemoryToGPU(0x7f7776d3e010, 0x7f7776d3e010, 0x12c3c600000000, 0x1804000) = 0 +libhsa-runtime64.so.1->hsaKmtDeregisterMemory(0x7f7776d3e010, 0x7f7776d3e010, 0x7f777f60f9e8, 0x1220580) = 0 +<... hsa_amd_memory_unlock resumed> ) = 0 + hip-api tid:1.17 hipMemcpy ret= 0 (hipSuccess)>> +``` + +Some key information from the trace above. + - The trace snippet shows the execution of a hipMemcpy API, bracketed by the first and last message in the trace output. The messages show the thread id and API sequence number (`1.17`). ltrace output intermixes messages from all threads, so the HIP debug information can be useful to determine which threads are executing. + - The code flows through HIP APIs into ROCr (HSA) APIs (hsa*) and into the thunk (hsaKmt*) calls. + - The HCC runtime is "libmcwamp_hsa.so" and the HSA/ROCr runtime is "libhsa-runtime64.so". + - In this particular case, the memory copy is for unpinned memory, and the selected copy algorithm is to pin the host memory "in-place" before performing the copy. The signaling APIs and calls to pin ("lock", "register") the memory are readily apparent in the trace output. +### Chicken bits +Chicken bits are environment variables which cause the HIP, HCC, or HSA driver to disable some feature or optimization. +These are not intended for production but can be useful diagnose synchronization problems in the application (or driver). +Some of the most useful chicken bits are described here: -Trace provides quick look at API. -Explain output of -Reference the cookbook example. -Command-line profile. -/// disable profiling at the start of the application you can start CodeXLGpuProfiler with the --startdisabled flag. - -Can use strace interleaved with HSA Debug calls . - -HIP_PROFILE_API=1 -HIP_PROFILE_API=2 : Will show the full API in the trace. This can be useful for lower-level debugging when you want to see all the parameters that are passed to a specific API. - -demangle atp - -Write how to collect performance counters. -- include how to compute bandwidth for copy and kernel activity. - -- How to disable HSA APIs. -- Do I need to use profiler with HSA enabled? Do I need to enable HSA profiling on the command line? - -Offline compile, how to visualize. +- HIP_LAUNCH_BLOCKING=1 : On ROCm, this flag waits on the host after each kernel launches and after each memory copy command. On CUDA, the waits are only enforced after each kernel launch. This is useful to isolate synchronization problems. Specifically, if the code works with this flag set, then it indicates the kernels and memory management code are correct, and any failures likely are causes by improper or missing synchronization. +- HSA_ENABLE_SDMA=0 : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine. This flag is useful to isolate issues with the hardware copy engines. +- HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts. Can be useful to diagnose interrupt storm issues in the driver. +- HSA_DISABLE_CACHE=1 : Disables the GPU L2 data cache. From 111b57ddd0fc126ff2c650e53c876277ce01b814 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 22 Nov 2016 18:31:25 -0600 Subject: [PATCH 22/31] Improve debug capabilities. Print TID mapping at init when HIP_TRACE_API=1. Print base host/dev info from tracker during copy. Change-Id: I84e26d7b801567e5a91baad36126fb590920ec87 --- hipamd/src/hip_hcc.cpp | 34 ++++++++++++++++++++++++---------- 1 file changed, 24 insertions(+), 10 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index e187a63887..971fa78423 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -214,7 +214,7 @@ ShortTid::ShortTid() : { _shortTid = g_lastShortTid.fetch_add(1); - if (HIP_DB & (1<getDeviceNum():-1, - dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, - src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem, + dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem, + src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem, sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); + tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes, + dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem); + tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes, + srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem); + #if USE_COPY_EXT_V2 crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); @@ -1831,13 +1841,17 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes ihipCtx_t *copyDevice; bool forceUnpinnedCopy; resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, ©Device, &forceUnpinnedCopy); - - - tprintf (DB_COPY, "copyASync copyEngine_dev:%d dst=%p(home_dev:%d, tracked:%d, isDevMem:%d) src=%p(home_dev:%d, tracked:%d, isDevMem:%d) sz=%zu dir=%s. forceUnpinnedCopy=%d \n", - copyDevice->getDeviceNum(), - dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, - src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem, + tprintf (DB_COPY, "copyASync copyDev:%d dst=%p (phys_dev:%d, isDevMem:%d) src=%p(phys_dev:%d, isDevMem:%d) sz=%zu dir=%s forceUnpinnedCopy=%d\n", + copyDevice ? copyDevice->getDeviceNum():-1, + dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem, + src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem, sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); + tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes, + dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem); + tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes, + srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem); // "tracked" really indicates if the pointer's virtual address is available in the GPU address space. // If both pointers are not tracked, we need to fall back to a sync copy. From 9db93a1b9619fff215dc8a1a9d01f7054e272be5 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 23 Nov 2016 08:14:42 -0600 Subject: [PATCH 23/31] Improve docs in some places Change-Id: If31e84fbf0c8595ca72edb842dce7ce47783579b --- hipamd/bin/hsa-api-filter-cxl.txt | 207 ++++++++++++++++++ .../include/hip/hcc_detail/hip_runtime_api.h | 2 + hipamd/src/hip_hcc.cpp | 2 +- 3 files changed, 210 insertions(+), 1 deletion(-) create mode 100644 hipamd/bin/hsa-api-filter-cxl.txt diff --git a/hipamd/bin/hsa-api-filter-cxl.txt b/hipamd/bin/hsa-api-filter-cxl.txt new file mode 100644 index 0000000000..ea3847a791 --- /dev/null +++ b/hipamd/bin/hsa-api-filter-cxl.txt @@ -0,0 +1,207 @@ +hsa_amd_coherency_get_type +hsa_amd_coherency_set_type +hsa_amd_profiling_get_dispatch_time +hsa_amd_profiling_get_async_copy_time +hsa_amd_image_get_info_max_dim +hsa_amd_queue_cu_set_mask +hsa_amd_agent_iterate_memory_pools +hsa_amd_memory_pool_free +hsa_amd_agent_memory_pool_get_info +hsa_amd_memory_migrate +hsa_amd_memory_lock +hsa_amd_memory_unlock +hsa_amd_interop_map_buffer +hsa_amd_interop_unmap_buffer +hsa_amd_image_create +hsa_ext_program_create +hsa_ext_program_destroy +hsa_ext_program_add_module +hsa_ext_program_iterate_modules +hsa_ext_program_get_info +hsa_ext_program_finalize +hsa_ext_image_data_get_info +hsa_ext_sampler_create +hsa_status_string +hsa_init +hsa_shut_down +hsa_system_get_info +hsa_extension_get_name +hsa_system_extension_supported +hsa_system_major_extension_supported +hsa_system_get_extension_table +hsa_system_get_major_extension_table +hsa_agent_get_info +hsa_iterate_agents +hsa_agent_get_info_size +hsa_agent_set_info +hsa_agent_get_exception_policies +hsa_cache_get_info +hsa_agent_iterate_caches +hsa_agent_extension_supported +hsa_agent_major_extension_supported +hsa_signal_create +hsa_signal_destroy +hsa_signal_load_scacquire +hsa_signal_load_relaxed +hsa_signal_load_acquire +hsa_signal_store_relaxed +hsa_signal_store_screlease +hsa_signal_store_release +hsa_signal_silent_store_relaxed +hsa_signal_silent_store_screlease +hsa_signal_exchange_scacq_screl +hsa_signal_exchange_acq_rel +hsa_signal_exchange_scacquire +hsa_signal_exchange_acquire +hsa_signal_exchange_relaxed +hsa_signal_exchange_screlease +hsa_signal_exchange_release +hsa_signal_cas_scacq_screl +hsa_signal_cas_acq_rel +hsa_signal_cas_scacquire +hsa_signal_cas_acquire +hsa_signal_cas_relaxed +hsa_signal_cas_screlease +hsa_signal_cas_release +hsa_signal_add_scacq_screl +hsa_signal_add_acq_rel +hsa_signal_add_scacquire +hsa_signal_add_acquire +hsa_signal_add_relaxed +hsa_signal_add_screlease +hsa_signal_add_release +hsa_signal_subtract_scacq_screl +hsa_signal_subtract_acq_rel +hsa_signal_subtract_scacquire +hsa_signal_subtract_acquire +hsa_signal_subtract_relaxed +hsa_signal_subtract_screlease +hsa_signal_subtract_release +hsa_signal_and_scacq_screl +hsa_signal_and_acq_rel +hsa_signal_and_scacquire +hsa_signal_and_acquire +hsa_signal_and_relaxed +hsa_signal_and_screlease +hsa_signal_and_release +hsa_signal_or_scacq_screl +hsa_signal_or_acq_rel +hsa_signal_or_scacquire +hsa_signal_or_acquire +hsa_signal_or_relaxed +hsa_signal_or_screlease +hsa_signal_or_release +hsa_signal_xor_scacq_screl +hsa_signal_xor_acq_rel +hsa_signal_xor_scacquire +hsa_signal_xor_acquire +hsa_signal_xor_relaxed +hsa_signal_xor_screlease +hsa_signal_xor_release +hsa_signal_wait_scacquire +hsa_signal_wait_relaxed +hsa_signal_wait_acquire +hsa_signal_group_create +hsa_signal_group_destroy +hsa_signal_group_wait_any_scacquire +hsa_signal_group_wait_any_relaxed +hsa_queue_create +hsa_soft_queue_create +hsa_queue_destroy +hsa_queue_inactivate +hsa_queue_load_read_index_acquire +hsa_queue_load_read_index_scacquire +hsa_queue_load_read_index_relaxed +hsa_queue_load_write_index_acquire +hsa_queue_load_write_index_scacquire +hsa_queue_load_write_index_relaxed +hsa_queue_store_write_index_relaxed +hsa_queue_store_write_index_release +hsa_queue_store_write_index_screlease +hsa_queue_cas_write_index_acq_rel +hsa_queue_cas_write_index_scacq_screl +hsa_queue_cas_write_index_acquire +hsa_queue_cas_write_index_scacquire +hsa_queue_cas_write_index_relaxed +hsa_queue_cas_write_index_release +hsa_queue_cas_write_index_screlease +hsa_queue_add_write_index_acq_rel +hsa_queue_add_write_index_scacq_screl +hsa_queue_add_write_index_acquire +hsa_queue_add_write_index_scacquire +hsa_queue_add_write_index_relaxed +hsa_queue_add_write_index_release +hsa_queue_add_write_index_screlease +hsa_queue_store_read_index_relaxed +hsa_queue_store_read_index_release +hsa_queue_store_read_index_screlease +hsa_region_get_info +hsa_agent_iterate_regions +hsa_memory_allocate +hsa_memory_free +hsa_memory_copy +hsa_memory_assign_agent +hsa_memory_register +hsa_memory_deregister +hsa_isa_from_name +hsa_agent_iterate_isas +hsa_isa_get_info +hsa_isa_get_info_alt +hsa_isa_get_exception_policies +hsa_isa_get_round_method +hsa_wavefront_get_info +hsa_isa_iterate_wavefronts +hsa_isa_compatible +hsa_code_object_reader_create_from_file +hsa_code_object_reader_create_from_memory +hsa_code_object_reader_destroy +hsa_executable_create +hsa_executable_create_alt +hsa_executable_destroy +hsa_executable_load_program_code_object +hsa_executable_load_agent_code_object +hsa_executable_freeze +hsa_executable_get_info +hsa_executable_global_variable_define +hsa_executable_agent_global_variable_define +hsa_executable_readonly_variable_define +hsa_executable_validate +hsa_executable_validate_alt +hsa_executable_get_symbol +hsa_executable_get_symbol_by_name +hsa_executable_symbol_get_info +hsa_executable_iterate_symbols +hsa_executable_iterate_agent_symbols +hsa_executable_iterate_program_symbols +hsa_code_object_serialize +hsa_code_object_deserialize +hsa_code_object_destroy +hsa_code_object_get_info +hsa_executable_load_code_object +hsa_code_object_get_symbol +hsa_code_object_get_symbol_from_name +hsa_code_symbol_get_info +hsa_code_object_iterate_symbols +hsa_ven_amd_loader_query_host_address +hsa_ven_amd_loader_query_segment_descriptors +hsa_amd_profiling_set_profiler_enabled +hsa_amd_profiling_async_copy_enable +hsa_amd_profiling_convert_tick_to_system_domain +hsa_amd_signal_async_handler +hsa_amd_async_function +hsa_amd_signal_wait_any +hsa_amd_memory_pool_get_info +hsa_amd_memory_pool_allocate +hsa_amd_memory_async_copy +hsa_amd_agents_allow_access +hsa_amd_memory_pool_can_migrate +hsa_amd_memory_fill + +hsa_ext_image_get_capability +hsa_ext_image_create +hsa_ext_image_destroy +hsa_ext_image_copy +hsa_ext_image_import +hsa_ext_image_export +hsa_ext_image_clear +hsa_ext_sampler_destroy diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 77853f02a2..ee703c4eec 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1761,6 +1761,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, // TODO - expand descriptions: /** * @brief Start recording of profiling information + * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStart API is under development. */ hipError_t hipProfilerStart(); @@ -1768,6 +1769,7 @@ hipError_t hipProfilerStart(); /** * @brief Stop recording of profiling information. + * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStop API is under development. */ hipError_t hipProfilerStop(); diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 971fa78423..79d117a0c4 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1255,7 +1255,7 @@ void ihipInit() READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes."); READ_ENV_S(release, HIP_TRACE_API_COLOR, 0, "Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White"); READ_ENV_I(release, HIP_PROFILE_API, 0, "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args."); - READ_ENV_S(release, HIP_DB_START_API, 0, "Comma-separted list of tid.api_seq_num for when to start debug and profiling."); + READ_ENV_S(release, HIP_DB_START_API, 0, "Comma-separated list of tid.api_seq_num for when to start debug and profiling."); READ_ENV_S(release, HIP_DB_STOP_API, 0, "Comma-separated list of tid.api_seq_num for when to stop debug and profiling."); READ_ENV_C(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the sequence are visible to HIP applications and they are enumerated in the order of sequence.", HIP_VISIBLE_DEVICES_callback ); From ddb1845ab5d458c5a42f797cd28d6b0eba062be8 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 23 Nov 2016 08:36:08 -0600 Subject: [PATCH 24/31] Add toc to hip_profiling.md Change-Id: I3ae100f12686d0398a0403b78ca571382acce135 --- hipamd/docs/markdown/hip_profiling.md | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/hipamd/docs/markdown/hip_profiling.md b/hipamd/docs/markdown/hip_profiling.md index ee3a8d6656..91acfafaf4 100644 --- a/hipamd/docs/markdown/hip_profiling.md +++ b/hipamd/docs/markdown/hip_profiling.md @@ -4,6 +4,27 @@ This section describes the profiling and debugging capabilities that HIP provide Profiling information can viewed in the CodeXL visualization tool or printed directly to stderr as the application runs. This document starts with some of the general capabilities of CodeXL and then describes some of the additional HIP marker and debug features. + * [CodeXL Profiling](#codexl-profiling) + * [Collecting and Viewing Traces](#collecting-and-viewing-traces) + * [Using rocm-profiler timestamp profiling](#using-rocm-profiler-timestamp-profiling) + * [Using rocm-profiler performance counter collection:](#using-rocm-profiler-performance-counter-collection) + * [Using CodeXL to view profiling results:](#using-codexl-to-view-profiling-results) + * [More information on CodeXL](#more-information-on-codexl) + * [HIP Markers](#hip-markers) + * [Profiling HIP APIs](#profiling-hip-apis) + * [Adding markers to applications](#adding-markers-to-applications) + * [Additional HIP Profiling Features](#additional-hip-profiling-features) + * [Demangling C Kernel Names](#demangling-c-kernel-names) + * [Controlling when profiling starts and ends](#controlling-when-profiling-starts-and-ends) + * [Reducing timeline trace output file size](#reducing-timeline-trace-output-file-size) + * [How to enable profiling at HIP build time](#how-to-enable-profiling-at-hip-build-time) + * [Tracing and Debug](#tracing-and-debug) + * [Tracing HIP APIs](#tracing-hip-apis) + * [Color](#color) + * [Using HIP_DB](#using-hip_db) + * [Using ltrace](#using-ltrace) + * [Chicken bits](#chicken-bits) + ## CodeXL Profiling ### Collecting and Viewing Traces From c2f6ecf2640474c6d00449e5f4195abc1172326d Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 23 Nov 2016 11:19:15 -0600 Subject: [PATCH 25/31] Added fast math flag 1. Use -DHIP_FAST_MATH to make precise math functions compiled to fast math 2. Added double fast math functions for sqrt 3. Changed hipcc to parse -use_fast_math (not working) 4. Added passed tag to hipFloatMath test Change-Id: I72884b2436b4efe61e9a9297346c1358fee38a2d --- hipamd/bin/hipcc | 5 + hipamd/include/hip/hcc_detail/hip_runtime.h | 119 +++++++---------- hipamd/tests/src/deviceLib/hipFloatMath.cpp | 1 + .../src/deviceLib/hipFloatMathPrecise.cpp | 122 ++++++++++++++++++ 4 files changed, 177 insertions(+), 70 deletions(-) create mode 100644 hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 641f70b065..09c4d813d0 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -274,6 +274,11 @@ foreach $arg (@ARGV) $buildDeps = 1; } + if($arg eq '-use_fast_math') { + print "In fast Math"; + $HIPCXXFLAGS .= " -DHIP_FAST_MATH "; + } + if ($arg =~ m/^-/) { # options start with - diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 4b781b44ae..45dbeff5a4 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -130,7 +130,6 @@ __device__ float atanhf(float x); __device__ float cbrtf(float x); __device__ float ceilf(float x); __device__ float copysignf(float x, float y); -__device__ float cosf(float x); __device__ float coshf(float x); __device__ float cyl_bessel_i0f(float x); __device__ float cyl_bessel_i1f(float x); @@ -142,9 +141,7 @@ __host__ float erfcxf(float x); __device__ float erff(float x); __device__ float erfinvf(float y); __host__ float erfinvf(float y); -__device__ float exp10f(float x); __device__ float exp2f(float x); -__device__ float expf(float x); __device__ float expm1f(float x); __device__ float fabsf(float x); __device__ float fdimf(float x, float y); @@ -167,11 +164,8 @@ __device__ float ldexpf(float x, int exp); __device__ float lgammaf(float x); __device__ long long int llrintf(float x); __device__ long long int llroundf(float x); -__device__ float log10f(float x); __device__ float log1pf(float x); -__device__ float log2f(float x); __device__ float logbf(float x); -__device__ float logf(float x); __device__ long int lrintf(float x); __device__ long int lroundf(float x); __device__ float modff(float x, float *iptr); @@ -187,7 +181,6 @@ __host__ float normcdff(float y); __device__ float normcdfinvf(float y); __host__ float normcdfinvf(float y); __device__ float normf(int dim, const float *a); -__device__ float powf(float x, float y); __device__ float rcbrtf(float x); __host__ float rcbrtf(float x); __device__ float remainderf(float x, float y); @@ -206,14 +199,11 @@ __device__ float rsqrtf(float x); __device__ float scalblnf(float x, long int n); __device__ float scalbnf(float x, int n); __host__ __device__ unsigned signbit(float a); -__device__ void sincosf(float x, float *sptr, float *cptr); __device__ void sincospif(float x, float *sptr, float *cptr); __host__ void sincospif(float x, float *sptr, float *cptr); -__device__ float sinf(float x); __device__ float sinhf(float x); __device__ float sinpif(float x); __device__ float sqrtf(float x); -__device__ float tanf(float x); __device__ float tanhf(float x); __device__ float tgammaf(float x); __device__ float truncf(float x); @@ -519,90 +509,65 @@ __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"); __device__ float __hip_fast_tanf(float); +extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64"); -#ifdef HIP_PRECISE_MATH +#ifdef HIP_FAST_MATH // Single Precision Precise Math when enabled -__device__ inline float __cosf(float x) { - return __hip_precise_cosf(x); +__device__ inline float cosf(float x) { + return __hip_fast_cosf(x); } -__device__ inline float __exp10f(float x) { - return __hip_precise_exp10f(x); +__device__ inline float exp10f(float x) { + return __hip_fast_exp10f(x); } -__device__ inline float __expf(float x) { - return __hip_precise_expf(x); +__device__ inline float expf(float x) { + return __hip_fast_expf(x); } -__device__ inline float __frsqrt_rn(float x) { - return __hip_precise_frsqrt_rn(x); +__device__ inline float log10f(float x) { + return __hip_fast_log10f(x); } -__device__ inline float __fsqrt_rd(float x) { - return __hip_precise_fsqrt_rd(x); +__device__ inline float log2f(float x) { + return __hip_fast_log2f(x); } -__device__ inline float __fsqrt_rn(float x) { - return __hip_precise_fsqrt_rn(x); +__device__ inline float logf(float x) { + return __hip_fast_logf(x); } -__device__ inline float __fsqrt_ru(float x) { - return __hip_precise_fsqrt_ru(x); +__device__ inline float powf(float base, float exponent) { + return __hip_fast_powf(base, exponent); } -__device__ inline float __fsqrt_rz(float x) { - return __hip_precise_fsqrt_rz(x); +__device__ inline void sincosf(float x, float *s, float *c) { + return __hip_fast_sincosf(x, s, c); } -__device__ inline float __log10f(float x) { - return __hip_precise_log10f(x); +__device__ inline float sinf(float x) { + return __hip_fast_sinf(x); } -__device__ inline float __log2f(float x) { - return __hip_precise_log2f(x); -} - -__device__ inline float __logf(float x) { - return __hip_precise_logf(x); -} - -__device__ inline float __powf(float base, float exponent) { - return __hip_precise_powf(base, exponent); -} - -__device__ inline void __sincosf(float x, float *s, float *c) { - return __hip_precise_sincosf(x, s, c); -} - -__device__ inline float __sinf(float x) { - return __hip_precise_sinf(x); -} - -__device__ inline float __tanf(float x) { - return __hip_precise_tanf(x); -} - -// Double Precision - -__device__ double __dsqrt_rd(double x) { - return __hip_precise_dsqrt_rd(x); -} - -__device__ double __dsqrt_rn(double x) { - return __hip_precise_dsqrt_rn(x); -} - -__device__ double __dsqrt_ru(double x) { - return __hip_precise_dsqrt_ru(x); -} - -__device__ double __dsqrt_rz(double x) { - return __hip_precise_dsqrt_rz(x); +__device__ inline float tanf(float x) { + return __hip_fast_tanf(x); } #else +__device__ float sinf(float); +__device__ float cosf(float); +__device__ float tanf(float); +__device__ void sincosf(float, float*, float*); +__device__ float logf(float); +__device__ float log2f(float); +__device__ float log10f(float); +__device__ float expf(float); +__device__ float exp10f(float); +__device__ float powf(float, float); + +#endif // Single Precision Fast Math __device__ inline float __cosf(float x) { return __hip_fast_cosf(x); @@ -664,8 +629,22 @@ __device__ inline float __tanf(float x) { return __hip_fast_tanf(x); } +__device__ inline double __dsqrt_rd(double x) { + return __hip_fast_dsqrt(x); +} + +__device__ inline double __dsqrt_rn(double x) { + return __hip_fast_dsqrt(x); +} + +__device__ inline double __dsqrt_ru(double x) { + return __hip_fast_dsqrt(x); +} + +__device__ inline double __dsqrt_rz(double x) { + return __hip_fast_dsqrt(x); +} -#endif /** * CUDA 8 device function features diff --git a/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/hipamd/tests/src/deviceLib/hipFloatMath.cpp index eb70eb6b0b..f137ca2602 100644 --- a/hipamd/tests/src/deviceLib/hipFloatMath.cpp +++ b/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -58,4 +58,5 @@ int main(){ hipMalloc((void**)&Ind, SIZE); hipMalloc((void**)&Outd, SIZE); hipLaunchKernel(floatMath, dim3(LEN,1,1), dim3(1,1,1), 0, 0, Ind, Outd); + passed(); } diff --git a/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp b/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp new file mode 100644 index 0000000000..4f6c2cd44a --- /dev/null +++ b/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp @@ -0,0 +1,122 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include "hip/hip_runtime.h" +#include "test_common.h" + +__global__ void FloatMathPrecise(hipLaunchParm lp) +{ + int iX; + float fX, fY; + + acosf(1.0f); + acoshf(1.0f); + asinf(0.0f); + asinhf(0.0f); + atan2f(0.0f, 1.0f); + atanf(0.0f); + atanhf(0.0f); + cbrtf(0.0f); + fX = ceilf(0.0f); + fX = copysignf(1.0f, -2.0f); + cosf(0.0f); + coshf(0.0f); + cospif(0.0f); + //cyl_bessel_i0f(0.0f); + //cyl_bessel_i1f(0.0f); + erfcf(0.0f); + erfcinvf(2.0f); + erfcxf(0.0f); + erff(0.0f); + erfinvf(1.0f); + exp10f(0.0f); + exp2f(0.0f); + expf(0.0f); + expm1f(0.0f); + fX = fabsf(1.0f); + fdimf(1.0f, 0.0f); + fdividef(0.0f, 1.0f); + fX = floorf(0.0f); + fmaf(1.0f, 2.0f, 3.0f); + fX = fmaxf(0.0f, 0.0f); + fX = fminf(0.0f, 0.0f); + fmodf(0.0f, 1.0f); + //frexpf(0.0f, &iX); + hypotf(1.0f, 0.0f); + ilogbf(1.0f); + isfinite(0.0f); + fX = isinf(0.0f); + fX = isnan(0.0f); + j0f(0.0f); + j1f(0.0f); + jnf(-1.0f, 1.0f); + ldexpf(0.0f, 0); + //lgammaf(1.0f); + llrintf(0.0f); + llroundf(0.0f); + log10f(1.0f); + log1pf(-1.0f); + log2f(1.0f); + logbf(1.0f); + logf(1.0f); + lrintf(0.0f); + lroundf(0.0f); + //modff(0.0f, &fX); + fX = nanf("1"); + fX = nearbyintf(0.0f); + //nextafterf(0.0f); + norm3df(1.0f, 0.0f, 0.0f); + norm4df(1.0f, 0.0f, 0.0f, 0.0f); + normcdff(0.0f); + normcdfinvf(1.0f); + fX = 1.0f; normf(1, &fX); + powf(1.0f, 0.0f); + rcbrtf(1.0f); + remainderf(2.0f, 1.0f); + //remquof(1.0f, 2.0f, &iX); + rhypotf(0.0f, 1.0f); + fY = rintf(1.0f); + rnorm3df(0.0f, 0.0f, 1.0f); + rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); + fX = 1.0f; rnormf(1, &fX); + fY = roundf(0.0f); + rsqrtf(1.0f); + scalblnf(0.0f, 1); + scalbnf(0.0f, 1); + signbit(1.0f); + sincosf(0.0f, &fX, &fY); + sincospif(0.0f, &fX, &fY); + sinf(0.0f); + sinhf(0.0f); + sinpif(0.0f); + sqrtf(0.0f); + tanf(0.0f); + tanhf(0.0f); + tgammaf(2.0f); + fY = truncf(0.0f); + y0f(1.0f); + y1f(1.0f); + ynf(1, 1.0f); +} + +int main() { + hipLaunchKernel(FloatMathPrecise, dim3(1,1,1), dim3(1,1,1), 0, 0); +} From a4b43a66105ffb2fa57d42e7b811ada68a802915 Mon Sep 17 00:00:00 2001 From: pensun Date: Wed, 23 Nov 2016 14:06:18 -0600 Subject: [PATCH 26/31] Add several missing APIs in hipify Change-Id: I58912871cb0b10128f221ef26a11b0d69fb7873c --- hipamd/bin/hipify | 59 ++++++++++++++++++++++++++--------------------- 1 file changed, 33 insertions(+), 26 deletions(-) diff --git a/hipamd/bin/hipify b/hipamd/bin/hipify index 5c56465040..a7fa75b0f4 100755 --- a/hipamd/bin/hipify +++ b/hipamd/bin/hipify @@ -20,19 +20,19 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. ## -#usage hipify [OPTIONS] INPUT_FILE +#usage hipify [OPTIONS] INPUT_FILE use Getopt::Long; my $warn_whitelist =""; GetOptions( "print-stats" => \$print_stats # print the command-line, like a header. , "count-conversions" => \$count_conversions # count conversions. - , "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions. - , "warn-whitelist=s"=> \$warn_whitelist + , "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions. + , "warn-whitelist=s"=> \$warn_whitelist , "no-translate-builtins" => \$no_translate_builtins # don't translate math functions. , "no-translate-textures" => \$no_translate_textures , "no-output" => \$no_output # don't write any translated output to stdout. - , "inplace" => \$inplace # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file. + , "inplace" => \$inplace # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file. # If .prehip file exists, use that as input to hip. , "n" => \$n # combination of print_stats + no-output. ); @@ -89,7 +89,7 @@ push (@warn_whitelist, split(',',$warn_whitelist)); #--- #Compute total of all individual counts: -sub totalStats +sub totalStats { my %count = %{ shift() }; @@ -113,11 +113,11 @@ sub printStats my $total = totalStats(\%counts); printf STDERR "%s %d CUDA->HIP refs( ", $label, $total; - + foreach $stat (@statNames) { printf STDERR "%s:%d ", $stat, $counts{$stat}; - } - + } + printf STDERR ") warn:%d LOC:%d", $warnings, $loc; } @@ -196,11 +196,11 @@ while (@ARGV) { # Note : \b is used in perl to indicate the start of a word - typically that is what we want in this case: # - + # count of transforms in this file, init to 0 here: my %ft; clearStats(\%ft, \@statNames); - my $countIncludes = 0; + my $countIncludes = 0; my $countKeywords = 0; # keywords like __global__, __shared__ - not converted by hipify but counted here. my $warnings = 0; my $warningsCublas = 0; @@ -218,11 +218,11 @@ while (@ARGV) { # __CUDACC__ is set by NVCC to indicate it is treating the input file as CUDA code (as opposed to host) # Typically we want any code treated as CUDA code to be treated as accelerator code by Kalmar too # __HIPCC__ will set KALMARCC - $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g; + $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g; # __CUDA_ARCH is often used to detect when a function or kernel is being compiled for the device. # Don't automaticall convert this - likely these will need special attention with HIP_ARCH_HAS_* macros - #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g; + #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g; @@ -294,6 +294,9 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; + $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g; + $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g; + $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; $ft{'mem'} += s/\bcudaMallocArray\b/hipMallocArray/g; $ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g; @@ -308,7 +311,7 @@ while (@ARGV) { $ft{'coord_func'} += s/\bblockIdx\.x\b/hipBlockIdx_x/g; $ft{'coord_func'} += s/\bblockIdx\.y\b/hipBlockIdx_y/g; $ft{'coord_func'} += s/\bblockIdx\.z\b/hipBlockIdx_z/g; - + $ft{'coord_func'} += s/\bblockDim\.x\b/hipBlockDim_x/g; $ft{'coord_func'} += s/\bblockDim\.y\b/hipBlockDim_y/g; $ft{'coord_func'} += s/\bblockDim\.z\b/hipBlockDim_z/g; @@ -322,7 +325,7 @@ while (@ARGV) { #-------- - # Events + # Events $ft{'event'} += s/\bcudaEvent_t\b/hipEvent_t/g; $ft{'event'} += s/\bcudaEventCreate\b/hipEventCreate/g; $ft{'event'} += s/\bcudaEventCreateWithFlags\b/hipEventCreateWithFlags/g; @@ -341,10 +344,10 @@ while (@ARGV) { $ft{'stream'} += s/\bcudaStreamSynchronize\b/hipStreamSynchronize/g; $ft{'stream'} += s/\bcudaStreamDefault\b/hipStreamDefault/g; $ft{'stream'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g; - + #-------- - # Other synchronization + # Other synchronization $ft{'dev'} += s/\bcudaDeviceSynchronize\b/hipDeviceSynchronize/g; $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g; # translate deprecated cudaThreadSynchronize $ft{'dev'} += s/\bcudaDeviceReset\b/hipDeviceReset/g; @@ -353,9 +356,10 @@ while (@ARGV) { $ft{'dev'} += s/\bcudaGetDevice\b/hipGetDevice/g; #-------- - # Device + # Device $ft{'dev'} += s/\bcudaDeviceProp\b/hipDeviceProp_t/g; $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipGetDeviceProperties/g; + $ft{'dev'} += s/\bcudaDeviceGetPCIBusId\b/hipDeviceGetPCIBusId/g; # Attribute $ft{'err'} += s/\bcudaDevAttrMaxThreadsPerBlock\b/hipDeviceAttributeMaxThreadsPerBlock/g; @@ -408,6 +412,9 @@ while (@ARGV) { $ft{'dev'} += s/\bcudaDeviceEnablePeerAccess\b/hipDeviceEnablePeerAccess/g; $ft{'mem'} += s/\bcudaMemcpyPeerAsync\b/hipMemcpyPeerAsync/g; $ft{'mem'} += s/\bcudaMemcpyPeer\b/hipMemcpyPeer/g; + $ft{'mem'} += s/\bcudaIpcOpenMemHandle\b/hipIpcOpenMemHandle/g; + $ft{'mem'} += s/\bcudaIpcCloseMemHandle\b/hipIpcCloseMemHandle/g; + $ft{'mem'} += s/\bcudaIpcGetMemHandle\b/hipIpcGetMemHandle/g; # Shared mem: @@ -441,7 +448,7 @@ while (@ARGV) { no warnings qw/uninitialized/; my $k = 0; - + # Match extern __shared__ type foo[]; syntax # Replace as HIP_DYNAMIC_SHARED() macro $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g; @@ -464,7 +471,7 @@ while (@ARGV) { #-------- # CUDA Launch Syntax # Note these only work if launch is on a single line. - + # Handle the <>> syntax: { # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. @@ -524,7 +531,7 @@ while (@ARGV) { unless ($quiet_warnings) { #print STDERR "Check WARNINGs\n"; # copy into array of lines, process line-by-line to show warnings: - if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) { + if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) { my @lines = split /\n/, $_; my $tmp = $_; # copies the whole file, could be a little smarter here... my $line_num = 0; @@ -541,7 +548,7 @@ while (@ARGV) { my $tag ; if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) { # flag any remaining code that look like cuda API calls, may want to add these to hipify - $tag = (defined $1) ? $1 : "Launch"; + $tag = (defined $1) ? $1 : "Launch"; } elsif (/(\bcublas[A-Z]\w+)/) { $warningsCublas++; $tag = $1; @@ -552,7 +559,7 @@ while (@ARGV) { if (defined $tag) { $warnings++; - $warningTags{$tag}++; + $warningTags{$tag}++; print STDERR " warning: $fileName:#$line_num : $_"; print STDERR "\n"; } @@ -599,7 +606,7 @@ while (@ARGV) { print $OUTFILE "$_"; } - $lineCount = $_ =~ tr/\n//; + $lineCount = $_ =~ tr/\n//; } @@ -653,7 +660,7 @@ if ($count_conversions) { -sub countSupportedSpecialFunctions +sub countSupportedSpecialFunctions { my $m = 0; @@ -671,9 +678,9 @@ sub countSupportedSpecialFunctions return $m; } -sub warnUnsupportedSpecialFunctions +sub warnUnsupportedSpecialFunctions { - my $line_num = shift; + my $line_num = shift; my $m = 0; From 69b43ec17c181cf7d46d2d4725c6ee4a28b649e6 Mon Sep 17 00:00:00 2001 From: pensun Date: Wed, 23 Nov 2016 14:36:30 -0600 Subject: [PATCH 27/31] Add some missing APIs on nv path and hipify Change-Id: Ic0f4740ab06bf70b1de61b39fedc7a6e7605cb61 --- hipamd/bin/hipify | 6 ++++++ hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 1 + 2 files changed, 7 insertions(+) diff --git a/hipamd/bin/hipify b/hipamd/bin/hipify index a7fa75b0f4..4d77fad3ed 100755 --- a/hipamd/bin/hipify +++ b/hipamd/bin/hipify @@ -230,6 +230,7 @@ while (@ARGV) { #Includes: $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime\.h[>"]/$1/; $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1/; + $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1/; #-------- @@ -249,6 +250,7 @@ while (@ARGV) { $ft{'err'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g; $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g; $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g; + $ft{'err'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g; # error APIs: $ft{'err'} += s/\bcudaGetLastError\b/hipGetLastError/g; @@ -294,6 +296,7 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; + $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g; $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g; $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g; $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; @@ -333,6 +336,7 @@ while (@ARGV) { $ft{'event'} += s/\bcudaEventRecord\b/hipEventRecord/g; $ft{'event'} += s/\bcudaEventElapsedTime\b/hipEventElapsedTime/g; $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g; + $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g; #-------- # Streams @@ -415,6 +419,8 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaIpcOpenMemHandle\b/hipIpcOpenMemHandle/g; $ft{'mem'} += s/\bcudaIpcCloseMemHandle\b/hipIpcCloseMemHandle/g; $ft{'mem'} += s/\bcudaIpcGetMemHandle\b/hipIpcGetMemHandle/g; + $ft{'mem'} += s/\bcudaIpcMemHandle_t\b/hipIpcMemHandle_t/g; + $ft{'mem'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g; # Shared mem: diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index c05d04ee29..625448094b 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -70,6 +70,7 @@ hipMemcpyHostToHost #define HIP_LAUNCH_PARAM_BUFFER_SIZE CU_LAUNCH_PARAM_BUFFER_SIZE #define HIP_LAUNCH_PARAM_END CU_LAUNCH_PARAM_END #define hipLimitMallocHeapSize cudaLimitMallocHeapSize +#define hipIpcMemLazyEnablePeerAccess cudaIpcMemLazyEnablePeerAccess typedef cudaEvent_t hipEvent_t; typedef cudaStream_t hipStream_t; From cc1f8a1011a0b09b0a7c7f052cdb7c3206e8ba99 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 23 Nov 2016 18:22:05 -0600 Subject: [PATCH 28/31] added fma for double and float 1. Added fma intrinsic support for double and float 2. Added test for fma Change-Id: I909fdbec34a3d12c03ba6eff3a39376a7128ee43 --- hipamd/include/hip/hcc_detail/hip_runtime.h | 35 +++++++++++++++++++ .../hipDoublePrecisionIntrinsics.cpp | 8 ++--- 2 files changed, 39 insertions(+), 4 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 45dbeff5a4..63cfb2ea3c 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -509,7 +509,10 @@ __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"); __device__ float __hip_fast_tanf(float); +extern __attribute__((const)) float __hip_fast_fmaf(float,float,float) __asm("llvm.fma.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"); #ifdef HIP_FAST_MATH // Single Precision Precise Math when enabled @@ -629,6 +632,22 @@ __device__ inline float __tanf(float x) { return __hip_fast_tanf(x); } +__device__ inline float __fmaf_rd(float x, float y, float z) { + return __hip_fast_fmaf(x, y, z); +} + +__device__ inline float __fmaf_rn(float x, float y, float z) { + return __hip_fast_fmaf(x, y, z); +} + +__device__ inline float __fmaf_ru(float x, float y, float z) { + return __hip_fast_fmaf(x, y, z); +} + +__device__ inline float __fmaf_rz(float x, float y, float z) { + return __hip_fast_fmaf(x, y, z); +} + __device__ inline double __dsqrt_rd(double x) { return __hip_fast_dsqrt(x); } @@ -645,6 +664,22 @@ __device__ inline double __dsqrt_rz(double x) { return __hip_fast_dsqrt(x); } +__device__ inline double __fma_rd(double x, double y, double z) { + return __hip_fast_fma(x, y, z); +} + +__device__ inline double __fma_rn(double x, double y, double z) { + return __hip_fast_fma(x, y, z); +} + +__device__ inline double __fma_ru(double x, double y, double z) { + return __hip_fast_fma(x, y, z); +} + +__device__ inline double __fma_rz(double x, double y, double z) { + return __hip_fast_fma(x, y, z); +} + /** * CUDA 8 device function features diff --git a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 0dab2d7106..0b4e0840a4 100644 --- a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -51,10 +51,10 @@ __device__ void double_precision_intrinsics() //__dsub_rn(2.0, 1.0); //__dsub_ru(2.0, 1.0); //__dsub_rz(2.0, 1.0); - //__fma_rd(1.0, 2.0, 3.0); - //__fma_rn(1.0, 2.0, 3.0); - //__fma_ru(1.0, 2.0, 3.0); - //__fma_rz(1.0, 2.0, 3.0); + __fma_rd(1.0, 2.0, 3.0); + __fma_rn(1.0, 2.0, 3.0); + __fma_ru(1.0, 2.0, 3.0); + __fma_rz(1.0, 2.0, 3.0); } __global__ void compileDoublePrecisionIntrinsics(hipLaunchParm lp, int ignored) From de89b25d52e38b3abc6e0b89636a677ac3b66326 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 23 Nov 2016 20:01:18 -0600 Subject: [PATCH 29/31] added support for rcp for float and double Change-Id: Ibeba3a9f64494fc0a176bcb4a854fb2f56567b55 --- hipamd/include/hip/hcc_detail/hip_runtime.h | 34 +++++++++++++++++++ .../hipDoublePrecisionIntrinsics.cpp | 8 ++--- 2 files changed, 38 insertions(+), 4 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 63cfb2ea3c..b1877ed0b3 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -510,9 +510,11 @@ __device__ void __hip_fast_sincosf(float,float*,float*); extern __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 __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"); #ifdef HIP_FAST_MATH // Single Precision Precise Math when enabled @@ -648,6 +650,22 @@ __device__ inline float __fmaf_rz(float x, float y, float z) { return __hip_fast_fmaf(x, y, z); } +__device__ inline float __frcp_rd(float x) { + return __hip_fast_frcp(x); +} + +__device__ inline float __frcp_rn(float x) { + return __hip_fast_frcp(x); +} + +__device__ inline float __frcp_ru(float x) { + return __hip_fast_frcp(x); +} + +__device__ inline float __frcp_rz(float x) { + return __hip_fast_frcp(x); +} + __device__ inline double __dsqrt_rd(double x) { return __hip_fast_dsqrt(x); } @@ -680,6 +698,22 @@ __device__ inline double __fma_rz(double x, double y, double z) { return __hip_fast_fma(x, y, z); } +__device__ inline double __drcp_rd(double x) { + return __hip_fast_drcp(x); +} + +__device__ inline double __drcp_rn(double x) { + return __hip_fast_drcp(x); +} + +__device__ inline double __drcp_ru(double x) { + return __hip_fast_drcp(x); +} + +__device__ inline double __drcp_rz(double x) { + return __hip_fast_drcp(x); +} + /** * CUDA 8 device function features diff --git a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 0b4e0840a4..330c4fc799 100644 --- a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -39,10 +39,10 @@ __device__ void double_precision_intrinsics() //__dmul_rn(1.0, 2.0); //__dmul_ru(1.0, 2.0); //__dmul_rz(1.0, 2.0); - //__drcp_rd(2.0); - //__drcp_rn(2.0); - //__drcp_ru(2.0); - //__drcp_rz(2.0); + __drcp_rd(2.0); + __drcp_rn(2.0); + __drcp_ru(2.0); + __drcp_rz(2.0); __dsqrt_rd(4.0); __dsqrt_rn(4.0); __dsqrt_ru(4.0); From a990806b6c0d4e2b436c950500903ce7e60a4b22 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 25 Nov 2016 14:05:06 -0600 Subject: [PATCH 30/31] Tweak profiler doc Change-Id: I7be7c44467510e38ae850e1e0a14209a3b4380f1 --- hipamd/docs/markdown/hip_profiling.md | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/hipamd/docs/markdown/hip_profiling.md b/hipamd/docs/markdown/hip_profiling.md index 91acfafaf4..db5d0fc425 100644 --- a/hipamd/docs/markdown/hip_profiling.md +++ b/hipamd/docs/markdown/hip_profiling.md @@ -52,12 +52,12 @@ The trace can be loaded and viewed in the CodeXL visualization tool: - [File->New Project, leave fields as is, just click "OK"] - [Profile->Switch to Profile Mode] - Load timestamp tracing results into a timeline view: - - Right click on the project + - Right click on the project in the CodeXL Explorer view - Click "Import Session..." - Select to $HOME/apitrace.atp (or appropriate .atp file if you used another file name) - Load the performance counter results - - Right click on the project + - Right click on the project in the CodeXL Explorer view - Click "Import Session..." - Select $HOME/Session1.csv (or appropriate .csv file if you used another file name) @@ -67,6 +67,8 @@ rocm-profiler --help will show additional options and usage guidelines. See this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) for more information on profiling ROCm apps (including HIP) with CodeXL. +The 2.2 version of Windows CodeXL does not correctly handle Linux line-endings. If you are collecting a trace on Linux and then viewing it with the 2.2 Windows CodeXL, first convert the line ending in the .atp file to Windows-style line endings. + ### HIP Markers #### Profiling HIP APIs HIP can generate markers at function beginning and end which are displayed on the CodeXL timeline view. @@ -230,10 +232,10 @@ Here's an example for one API followed by a description for the sections of the - `tid:1.6` indicates that this API call came from thread #1 and is the 6th API call in that thread. When the first API in a new thread is called, HIP will associates a short sequential ID with that thread. You can see the full thread ID (reported by C++) as 0x7f6183b097c0 in the example below. - `hipMemcpy` is the name of the API. - The first line then prints a comma-separated list of the arguments to the function. APIs which return values to the caller by writing to pointers will show the pointer addresses rather than the pointer contents. This behavior may change in the future. -- The second line shows the completio of the API, including the numeric return value (`ret= 0`) as well as an string representation for the error code (`hipSuccess`). If the returned error code is non-zero, then the csecond line message is shown in red (unless HIP_TRACE_API_COLOR is "none" - see below). +- The second line shows the completion of the API, including the numeric return value (`ret= 0`) as well as an string representation for the error code (`hipSuccess`). If the returned error code is non-zero, then the csecond line message is shown in red (unless HIP_TRACE_API_COLOR is "none" - see below). -Heres a specific example showing the output of the square program running on HIP: +Heres a specific example showing the output of the [square](samples/0_Intro/square) program running on HIP: ``` $ HIP_TRACE_API=1 ./square.hip.out From a504df955ec37073758e484829e4cb959ceeb82d Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 26 Nov 2016 08:55:51 -0600 Subject: [PATCH 31/31] Add more debug info --- .../samples/1_Utils/hipCommander/LICENSE.txt | 27 + hipamd/samples/1_Utils/hipCommander/Makefile | 33 + .../1_Utils/hipCommander/ResultDatabase.cpp | 527 ++++++++ .../1_Utils/hipCommander/ResultDatabase.h | 100 ++ hipamd/samples/1_Utils/hipCommander/TODO | 50 + hipamd/samples/1_Utils/hipCommander/c.cmd | 3 + .../samples/1_Utils/hipCommander/classic.cmd | 1 + .../1_Utils/hipCommander/hipCommander.cpp | 1096 +++++++++++++++++ hipamd/samples/1_Utils/hipCommander/l2.hcm | 3 + hipamd/samples/1_Utils/hipCommander/loop.hcm | 3 + hipamd/samples/1_Utils/hipCommander/loop2.hcm | 2 + .../1_Utils/hipCommander/nullkernel.hip.cpp | 7 + .../1_Utils/hipCommander/nullkernel.hsaco | Bin 0 -> 10265 bytes .../1_Utils/hipCommander/perf/latency2.hcm | 10 + .../hipCommander/perf/latency_hostsync.hcm | 8 + .../hipCommander/perf/latency_nosync.hcm | 5 + .../hipCommander/perf/latency_nullstream.hcm | 7 + .../perf/modulelaunch_latency.hcm | 5 + .../1_Utils/hipCommander/setstream.hcm | 3 + .../samples/1_Utils/hipCommander/testcase.cpp | 21 + hipamd/src/hip_hcc.cpp | 18 +- 21 files changed, 1922 insertions(+), 7 deletions(-) create mode 100644 hipamd/samples/1_Utils/hipCommander/LICENSE.txt create mode 100644 hipamd/samples/1_Utils/hipCommander/Makefile create mode 100644 hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp create mode 100644 hipamd/samples/1_Utils/hipCommander/ResultDatabase.h create mode 100644 hipamd/samples/1_Utils/hipCommander/TODO create mode 100644 hipamd/samples/1_Utils/hipCommander/c.cmd create mode 100644 hipamd/samples/1_Utils/hipCommander/classic.cmd create mode 100644 hipamd/samples/1_Utils/hipCommander/hipCommander.cpp create mode 100644 hipamd/samples/1_Utils/hipCommander/l2.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/loop.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/loop2.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp create mode 100755 hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco create mode 100644 hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/setstream.hcm create mode 100644 hipamd/samples/1_Utils/hipCommander/testcase.cpp diff --git a/hipamd/samples/1_Utils/hipCommander/LICENSE.txt b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt new file mode 100644 index 0000000000..5d0d603232 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt @@ -0,0 +1,27 @@ + +Copyright (c) 2011, UT-Battelle, LLC +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +* Redistributions of source code must retain the above copyright + notice, this list of conditions and the following disclaimer. +* Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimer in the + documentation and/or other materials provided with the distribution. +* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor + the names of its contributors may be used to endorse or promote products + derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + diff --git a/hipamd/samples/1_Utils/hipCommander/Makefile b/hipamd/samples/1_Utils/hipCommander/Makefile new file mode 100644 index 0000000000..e770c636a4 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/Makefile @@ -0,0 +1,33 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=hipCommander +OPT=-O3 +#CXXFLAGS = -O3 -g +CXXFLAGS = $(OPT) --std=c++11 + +HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) +ifeq (${HIP_PLATFORM}, hcc) + CXXFLAGS += " -stdlib=libc++" +endif + +CODE_OBJECTS=nullkernel.hsaco + +all: ${EXE} ${CODE_OBJECTS} + +$(EXE): hipCommander.cpp + $(HIPCC) $(CXXFLAGS) $^ -o $@ + +nullkernel.hsaco : nullkernel.hip.cpp + $(HIPCC) --genco nullkernel.hip -o nullkernel.hsaco + + +install: $(EXE) + cp $(EXE) $(HIP_PATH)/bin + + +clean: + rm -f *.o *.co $(EXE) diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp new file mode 100644 index 0000000000..2ec686f260 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp @@ -0,0 +1,527 @@ +#include "ResultDatabase.h" + +#include +#include +#include +#include + +using namespace std; + +bool ResultDatabase::Result::operator<(const Result &rhs) const +{ + if (test < rhs.test) + return true; + if (test > rhs.test) + return false; + if (atts < rhs.atts) + return true; + if (atts > rhs.atts) + return false; + return false; // less-operator returns false on equal +} + +double ResultDatabase::Result::GetMin() const +{ + double r = FLT_MAX; + for (int i=0; i= 100) + return value[n-1]; + + double index = ((n + 1.) * q / 100.) - 1; + + vector sorted = value; + sort(sorted.begin(), sorted.end()); + + if (n == 2) + return (sorted[0] * (1 - q/100.) + sorted[1] * (q/100.)); + + int index_lo = int(index); + double frac = index - index_lo; + if (frac == 0) + return sorted[index_lo]; + + double lo = sorted[index_lo]; + double hi = sorted[index_lo + 1]; + return lo + (hi-lo)*frac; +} + +double ResultDatabase::Result::GetMean() const +{ + double r = 0; + for (int i=0; i &values) +{ + for (int i=0; i= results.size()) + { + Result r; + r.test = test; + r.atts = atts; + r.unit = unit; + results.push_back(r); + } + + results[index].value.push_back(value); +} + +// **************************************************************************** +// Method: ResultDatabase::DumpDetailed +// +// Purpose: +// Writes the full results, including all trials. +// +// Arguments: +// out where to print +// +// Programmer: Jeremy Meredith +// Creation: August 14, 2009 +// +// Modifications: +// Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010 +// Renamed to DumpDetailed to make room for a DumpSummary. +// +// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010 +// Added note about (*) missing value tag. +// +// Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010 +// Changed note about missing values to be worded a little better. +// +// **************************************************************************** +void ResultDatabase::DumpDetailed(ostream &out) +{ + vector sorted(results); + sort(sorted.begin(), sorted.end()); + + const int testNameW = 24 ; + const int attW = 12; + const int fieldW = 11; + out << std::fixed << right << std::setprecision(4); + + int maxtrials = 1; + for (int i=0; i maxtrials) + maxtrials = sorted[i].value.size(); + } + + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << setw(testNameW) << "test\t" + << setw(attW) << "atts\t" + << setw(fieldW) + << "median\t" + << "mean\t" + << "stddev\t" + << "min\t" + << "max\t"; + for (int i=0; i sorted(results); + sort(sorted.begin(), sorted.end()); + + const int testNameW = 24 ; + const int attW = 12; + const int fieldW = 9; + out << std::fixed << right << std::setprecision(4); + + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << setw(testNameW) << "test\t" + << setw(attW) << "atts\t" + << setw(fieldW) + << "units\t" + << "median\t" + << "mean\t" + << "stddev\t" + << "min\t" + << "max\t"; + out << endl; + + for (int i=0; i sorted(results); + + sort(sorted.begin(), sorted.end()); + + //Check to see if the file is empty - if so, add the headers + emptyFile = this->IsFileEmpty(fileName); + + //Open file and append by default + ofstream out; + out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); + + //Add headers only for empty files + if(emptyFile) + { + // TODO: in big parallel runs, the "trials" are the procs + // and we really don't want to print them all out.... + out << "test, " + << "atts, " + << "units, " + << "median, " + << "mean, " + << "stddev, " + << "min, " + << "max, "; + out << endl; + } + + for (int i=0; i +ResultDatabase::GetResultsForTest(const string &test) +{ + // get only the given test results + vector retval; + for (int i=0; i & +ResultDatabase::GetResults() const +{ + return results; +} diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h new file mode 100644 index 0000000000..4b63a02a1f --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h @@ -0,0 +1,100 @@ +#ifndef RESULT_DATABASE_H +#define RESULT_DATABASE_H + +#include +#include +#include +#include +#include +using std::string; +using std::vector; +using std::ostream; +using std::ofstream; +using std::ifstream; + + +// **************************************************************************** +// Class: ResultDatabase +// +// Purpose: +// Track numerical results as they are generated. +// Print statistics of raw results. +// +// Programmer: Jeremy Meredith +// Creation: June 12, 2009 +// +// Modifications: +// Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010 +// Split timing reports into detailed and summary. E.g. for serial code, +// we might report all trial values, but skip them in parallel. +// +// Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010 +// Added check for missing value tag. +// +// Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010 +// Added percentile statistic. +// +// Jeremy Meredith, Fri Dec 3 16:30:31 EST 2010 +// Added a method to extract a subset of results based on test name. Also, +// the Result class is now public, so that clients can use them directly. +// Added a GetResults method as well, and made several functions const. +// +// **************************************************************************** +class ResultDatabase +{ + public: + // + // A performance result for a single SHOC benchmark run. + // + struct Result + { + string test; // e.g. "readback" + string atts; // e.g. "pagelocked 4k^2" + string unit; // e.g. "MB/sec" + vector value; // e.g. "837.14" + double GetMin() const; + double GetMax() const; + double GetMedian() const; + double GetPercentile(double q) const; + double GetMean() const; + double GetStdDev() const; + + bool operator<(const Result &rhs) const; + + bool HadAnyFLTMAXValues() const + { + for (int i=0; i= FLT_MAX) + return true; + } + return false; + } + }; + + protected: + vector results; + + public: + void AddResult(const string &test, + const string &atts, + const string &unit, + double value); + void AddResults(const string &test, + const string &atts, + const string &unit, + const vector &values); + vector GetResultsForTest(const string &test); + const vector &GetResults() const; + void ClearAllResults(); + void DumpDetailed(ostream&); + void DumpSummary(ostream&); + void DumpCsv(string fileName); + + private: + bool IsFileEmpty(string fileName); + +}; + + +#endif diff --git a/hipamd/samples/1_Utils/hipCommander/TODO b/hipamd/samples/1_Utils/hipCommander/TODO new file mode 100644 index 0000000000..4c835cfced --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/TODO @@ -0,0 +1,50 @@ +_ Add AQL kernel. +_ Fix &*kernel command so the kernel name/type is an argument not a new command. + +_ Add command to parse only. +_ Add regression to parse all the hcm files. + +_ Partition HCC, HIP, HSA, OpenCL commands into separate files. + + +_ Show time for back-to-back copies. +_ Add variables. + %loopcnt + + ./hipCommander %loopcnt=4 + +_ Add datasize command. + + +_ Add ( ) to parsing. +_ Add argument parsing and checking. + +_ Add verbose option to print each step of setup. + - print deliniater between setup and run. Add run start message. + + - print sizes of all buffers. + - print each command before running. + - show start/stop of timer routine. + +_ +_ Clear documentation on what each oepration does. +_ Add time instrumentation for each command. +_ Add pcie atomic. + + +_ Add tests for negative cases, ie endloop w/o opening loop. + + +README tips +--- +- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander. + + +Other ideas: +--- +[ ] Perf guide : stream creation very slow on HCC and should be avoided. + + +Scratch: + + diff --git a/hipamd/samples/1_Utils/hipCommander/c.cmd b/hipamd/samples/1_Utils/hipCommander/c.cmd new file mode 100644 index 0000000000..db11071203 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/c.cmd @@ -0,0 +1,3 @@ +loop,1000; H2D; NullKernel; D2H; endloop; +streamsync; +printTiming, 1000 diff --git a/hipamd/samples/1_Utils/hipCommander/classic.cmd b/hipamd/samples/1_Utils/hipCommander/classic.cmd new file mode 100644 index 0000000000..c149eec5f7 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/classic.cmd @@ -0,0 +1 @@ +H2D; NullKernel, D2H, streamsync diff --git a/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp new file mode 100644 index 0000000000..9c07d066ba --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp @@ -0,0 +1,1096 @@ +#include +#include +#include +#include +#include +#include +#include + +#include +#ifdef __HIP_PLATFORM_HCC__ +#include +#include +#include +#include +#endif + +#include + +#include "ResultDatabase.h" +#include "nullkernel.hip.cpp" + +bool g_printedTiming = false; + +// Cmdline parms: +int p_device = 0; +const char* p_command = "H2D; NullKernel; D2H"; +const char* p_file = nullptr; +unsigned p_verbose = 0x0; +unsigned p_db = 0x0; +unsigned p_blockingSync = 0x0; + +//--- +int p_iterations = 1; + +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" + + +#define failed(...) \ + printf ("error: ");\ + printf (__VA_ARGS__);\ + printf ("\n");\ + abort(); + + +#define HIPCHECK(error) \ +{\ + hipError_t localError = error; \ + if (localError != hipSuccess) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \ + KRED,hipGetErrorString(localError), localError,\ + #error,\ + __FILE__, __LINE__,KNRM); \ + failed("API returned error code.");\ + }\ +} +#define HIPASSERT(condition, msg) \ + if (! (condition) ) { \ + failed("%sassertion %s at %s:%d: %s%s\n", \ + KRED, #condition,\ + __FILE__, __LINE__,msg, KNRM); \ + } + + + + + + + +int parseInt(const char *str, int *output) +{ + char *next; + *output = strtol(str, &next, 0); + return !strlen(next); +} + + +void printConfig() { + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, p_device)); + + printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0); +} + + + + +void help() { + printf ("Usage: hipBusBandwidth [OPTIONS]\n"); + printf (" --file, -f : Read string of commands from file\n"); + printf (" --command, -c : String specifying commands to run.\n"); + printf (" --iterations, -i : Number of copy iterations to run.\n"); + printf (" --device, -d : Device ID to use (0..numDevices).\n"); + printf (" --verbose, -v : Verbose printing of status. Fore more info, combine with HIP_TRACE_API on ROCm\n"); +}; + + + +int parseStandardArguments(int argc, char *argv[]) +{ + for (int i = 1; i < argc; i++) { + const char *arg = argv[i]; + + if (!strcmp(arg, " ")) { + // skip NULL args. + } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { + if (++i >= argc || !parseInt(argv[i], &p_iterations)) { + failed("Bad --iterations argument"); + } + + } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) { + if (++i >= argc || !parseInt(argv[i], &p_device)) { + failed("Bad --device argument"); + } + + } else if (!strcmp(arg, "--file") || (!strcmp(arg, "-f"))) { + if (++i >= argc) { + failed("Bad --file argument"); + } else { + p_file = argv[i]; + } + + } else if (!strcmp(arg, "--commands") || (!strcmp(arg, "-c"))) { + if (++i >= argc) { + failed("Bad --commands argument"); + } else { + p_command = argv[i]; + } + + } else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) { + p_verbose = 1; + + } else if (!strcmp(arg, "--blockingSync") || (!strcmp(arg, "-B"))) { + p_blockingSync = 1; + + + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { + help(); + exit(EXIT_SUCCESS); + + } else { + failed("Bad argument '%s'", arg); + } + } + + return 0; +}; + +// Returns the current system time in microseconds +inline long long get_time() +{ + struct timeval tv; + gettimeofday(&tv, 0); + return (tv.tv_sec * 1000000) + tv.tv_usec; +} + + +class Command; + + +//================================================================================================= +// A stream of commands , specified as a string. +class CommandStream { +public: + // State that is inherited by sub-blocks: + struct CommandStreamState { + hipStream_t _currentStream; + std::vector _streams; + vector _subBlocks; + }; +public: + CommandStream(std::string commandStreamString, int iterations); + ~CommandStream(); + + hipStream_t currentStream() const { return _state._currentStream; }; + + void print(const std::string &indent="") const; + void printBrief(std::ostream &s=std::cout) const ; + void run(); + void recordTime(); + void printTiming(int iterations=0); + + CommandStream *currentCommandStream() { + return _parseInSubBlock ? _state._subBlocks.back() : this; + }; + + void enterSubBlock(CommandStream *commandStream) { + _parseInSubBlock = true; + _state._subBlocks.push_back(commandStream); + }; + + void exitSubBlock() { + _parseInSubBlock = false; + }; + + + void setParent(CommandStream *parentCmdStream) + { + _parentCommandStream = parentCmdStream; + _state = parentCmdStream->_state; + }; + CommandStream * getParent() { return _parentCommandStream; }; + + void setStream(int streamIndex); + + CommandStreamState &getState() { return _state; }; + +private: + static void tokenize(const std::string &s, char delim, std::vector &tokens); + void parse(const std::string fullCmd); + +protected: + CommandStreamState _state; +private: + + + // List of commands to run in this stream: + std::vector _commands; + + + + // Number of iterations to run the command loop + int _iterations; + + + + + // Us to run the the command-stream. Only valid after run is called. + long long _startTime; + double _elapsedUs; + + // Track nested loop of command streams: + CommandStream *_parentCommandStream; + + // Track if we are parsing commands in the subblock. + bool _parseInSubBlock; + +}; + + +//================================================================================================= +class Command { +public: + + // @p minArgs : Minimum arguments for command. -1 = don't check. + // @p maxArgs : Minimum arguments for command. 0 means min=max, ie exact #arguments expected. -1 = don't check max. + Command(CommandStream *cmdStream, const std::vector &args, int minArgs=0, int maxArgs=0) + : _commandStream(cmdStream), + _args(args) + { + int numArgs = args.size() - 1; + + if ((minArgs != -1 ) && (numArgs < minArgs)) { + // TODO - print full command here. + failed ("Not enough arguments for command %s. (Expected %d, got %d)", args[0].c_str(), minArgs, numArgs); + } + + // Check for an exact number of arguments: + if (maxArgs == 0) { + maxArgs = minArgs; + } + if ((maxArgs != -1 ) && (numArgs > maxArgs)) { + failed ("Too many arguments for command %s. (Expected %d, got %d)", args[0].c_str(), maxArgs, numArgs); + } + }; + + void printBrief(std::ostream &s=std::cout) const + { + s << _args[0]; + } + + virtual ~Command() {}; + + virtual void print(const std::string &indent = "") const { + std::cout << indent << "["; + std::for_each(_args.begin(), _args.end(), [] (const std::string &s) { + std::cout << s; + }); + std::cout << "]"; + }; + + + virtual void run() = 0; + +protected: + int readIntArg(int argIndex, const std::string &argName) + { + // TODO - catch references to non-existant arguments here. + int argVal; + try { + argVal = std::stoi(_args[argIndex]); + } catch (std::invalid_argument) { + failed ("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(), _args[argIndex].c_str()); + } + return argVal; + } +protected: + CommandStream *_commandStream; + std::vector _args; +}; + + +#define FILENAME "nullkernel.hsaco" +#define KERNEL_NAME "NullKernel" + + +#ifdef __HIP_PLATFORM_HCC__ +//================================================================================================= +// Use Aql to launch the NULL kernel. +class AqlKernelCommand : public Command +{ +public: + AqlKernelCommand(CommandStream *cmdStream, const std::vector args) : + Command(cmdStream, args) + { + hc::accelerator_view *av; + HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av)); + + hc::accelerator acc = av->get_accelerator(); + + hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region(); + + _hsaAgent = *(hsa_agent_t*) acc.get_hsa_agent(); + + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) + { + uint64_t elfSize = ElfSize(&buffer[0]); + + assert(fsize == elfSize); + + //TODO - replace module load code with explicit module load and unload. + + hipModule_t module; + HIPCHECK(hipModuleLoadData(&module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); + + } else { + failed("could not open code object '%s'\n", FILENAME); + } + }; + + ~AqlKernelCommand() {}; + + void run() override { +#define LEN 64 + uint32_t len = LEN; + uint32_t one = 1; + + float *Ad = NULL; + + size_t argSize = 36; + char argBuffer[argSize]; + *(uint32_t*) (&argBuffer[0]) = len; + *(uint32_t*) (&argBuffer[4]) = one; + *(uint32_t*) (&argBuffer[8]) = one; + *(uint32_t*) (&argBuffer[12]) = len; + *(uint32_t*) (&argBuffer[16]) = one; + *(uint32_t*) (&argBuffer[20]) = one; + *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + + + void *config[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, + HIP_LAUNCH_PARAM_END + }; + + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + }; + + +public: + hsa_queue_t _hsaQueue; + hsa_agent_t _hsaAgent; + + hipFunction_t _function; + +private: + static uint64_t ElfSize(const void *emi){ + const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; + const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); + + uint64_t max_offset = ehdr->e_shoff; + uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; + + for(uint16_t i=0;i < ehdr->e_shnum;++i){ + uint64_t cur_offset = static_cast(shdr[i].sh_offset); + if(max_offset < cur_offset){ + max_offset = cur_offset; + total_size = max_offset; + if(SHT_NOBITS != shdr[i].sh_type){ + total_size += static_cast(shdr[i].sh_size); + } + } + } + return total_size; + } +}; +#endif + +//================================================================================================= +// HCC optimizes away fully NULL kernel calls, so run one that is nearly null: +class ModuleKernelCommand : public Command +{ +public: + ModuleKernelCommand(CommandStream *cmdStream, const std::vector args) : + Command(cmdStream, args), + _stream (cmdStream->currentStream()) { + + hipModule_t module; + HIPCHECK(hipModuleLoad(&module, FILENAME)); + HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); + }; + ~ModuleKernelCommand() {}; + + void run() override { +#define LEN 64 + uint32_t len = LEN; + uint32_t one = 1; + + float *Ad = NULL; + + size_t argSize = 36; + char argBuffer[argSize]; + *(uint32_t*) (&argBuffer[0]) = len; + *(uint32_t*) (&argBuffer[4]) = one; + *(uint32_t*) (&argBuffer[8]) = one; + *(uint32_t*) (&argBuffer[12]) = len; + *(uint32_t*) (&argBuffer[16]) = one; + *(uint32_t*) (&argBuffer[20]) = one; + *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + + + void *config[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, + HIP_LAUNCH_PARAM_END + }; + + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + }; + + +public: + hipFunction_t _function; + hipStream_t _stream; +}; + + +class KernelCommand : public Command +{ +public: + enum Type {Null, VectorAdd}; + KernelCommand(CommandStream *cmdStream, const std::vector args, Type kind) : + Command(cmdStream, args), + _kind(kind), + _stream (cmdStream->currentStream()) { + }; + ~KernelCommand() {}; + + + void run() override { + static const int gridX = 64; + static const int groupX = 64; + + switch (_kind) { + case Null: + hipLaunchKernel(NullKernel, dim3(gridX/groupX), dim3(gridX), 0, _stream, nullptr); + break; + case VectorAdd: + assert(0); // TODO + break; + }; + } +private: + Type _kind; + hipStream_t _stream; +}; + + +#ifdef __HIP_PLATFORM_HCC__ +//================================================================================================= +class PfeCommand : public Command +{ +public: + PfeCommand(CommandStream *cmdStream, const std::vector args, hipStream_t stream = 0) : + Command(cmdStream, args) + { + HIPCHECK(hipHccGetAcceleratorView(stream, &_av)); + } + + ~PfeCommand() { } + + + void run() override { + static const int gridX = 64; + static const int groupX = 64; + auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX), + [=](hc::index<1>& idx) __HC__ { + }); + } +private: + hc::accelerator_view *_av; +}; +#endif + + + +//================================================================================================= +class CopyCommand : public Command +{ +enum MemType {PinnedHost, UnpinnedHost, Device} ; + +public: + CopyCommand(CommandStream *cmdStream, const std::vector &args, hipMemcpyKind kind, bool isAsync, bool isPinnedHost) ; + + ~CopyCommand() + { + if (_dst) { + dealloc(_dst, _dstType); + _dst = NULL; + }; + + if (_src) { + dealloc(_src, _srcType); + _src = NULL; + } + } + + + void run() override { + if (_isAsync) { + HIPCHECK(hipMemcpyAsync(_dst, _src, _sizeBytes, _kind, _stream)); + } else { + HIPCHECK(hipMemcpy(_dst, _src, _sizeBytes, _kind)); + } + }; + +private: + void * alloc(size_t size, MemType memType) { + void * p; + if (memType == Device) { + HIPCHECK(hipMalloc(&p, size)); + + } else if (memType == PinnedHost) { + HIPCHECK(hipHostMalloc(&p, size)); + + } else if (memType == UnpinnedHost) { + p = (char*)malloc(size); + HIPASSERT(p, "malloc failed"); + + } else { + HIPASSERT(0, "unsupported memType"); + } + + return p; + }; + + + void dealloc(void *p, MemType memType) { + if (memType == Device) { + HIPCHECK(hipFree(p)); + } else if (memType == PinnedHost) { + HIPCHECK(hipHostFree(p)); + } else if (memType == UnpinnedHost) { + free(p); + } else { + HIPASSERT(0, "unsupported memType"); + } + } + + +private: + bool _isAsync; + hipStream_t _stream; + hipMemcpyKind _kind; + + size_t _sizeBytes; + void *_dst; + MemType _dstType; + + void *_src; + MemType _srcType; +}; + + +//================================================================================================= +class DeviceSyncCommand : public Command +{ +public: + DeviceSyncCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args) {}; + + void run() override { + HIPCHECK(hipDeviceSynchronize()); + }; +}; + + +//================================================================================================= +class StreamSyncCommand : public Command +{ +public: + StreamSyncCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args), + _stream(cmdStream->currentStream()) + {}; + + const char *help() { + return "synchronizes the current stream"; + }; + + + + void run() override { + HIPCHECK(hipStreamSynchronize(_stream)); + }; + +private: + hipStream_t _stream; +}; + + +//================================================================================================= + +//================================================================================================= +class LoopCommand : public Command +{ +public: + LoopCommand(CommandStream *parentCmdStream, const std::vector &args) : + Command(parentCmdStream, args, 1) + { + int loopCnt; + try { + loopCnt = std::stoi(args[1]); + } catch (std::invalid_argument) { + failed ("bad LOOP_CNT=%s", args[1].c_str()); + } + + _commandStream = new CommandStream("", loopCnt); + _commandStream->setParent(parentCmdStream); + parentCmdStream->enterSubBlock(_commandStream); + + }; + + + void print(const std::string &indent = "") const override { + Command::print(); + _commandStream->print (indent + " "); + }; + + void run() override { + _commandStream->run(); + }; +}; + + +//================================================================================================= +class EndBlockCommand : public Command +{ +public: + EndBlockCommand(CommandStream *blockCmdStream, CommandStream *parentCmdStream, const std::vector &args) : + Command(parentCmdStream, args, 0, 1), + _blockCmdStream(blockCmdStream), + _printTiming(0) + { + int argCnt = args.size()-1; + if (argCnt >= 1 ) { + _printTiming = readIntArg(1, "PRINT_TIMING"); + } + + if (parentCmdStream == nullptr) { + failed ("%s without corresponding command to start block", args[0].c_str()); + } + parentCmdStream->exitSubBlock(); + }; + + void run() override { + if (_printTiming) { + _blockCmdStream->printTiming(); + } + + }; +private: + + CommandStream *_blockCmdStream; + + // print the stream when loop exits. + int _printTiming; +}; + + +//================================================================================================= +class SetStreamCommand : public Command +{ +public: + SetStreamCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args, 1) + { + int streamIndex = readIntArg(1, "STREAM_INDEX"); + + cmdStream->setStream(streamIndex); + + }; + + void run() override { + }; +}; + + +//================================================================================================= +class PrintTimingCommand : public Command +{ +public: + PrintTimingCommand(CommandStream *cmdStream, const std::vector &args) + : Command(cmdStream, args, 1) + { + _iterations = readIntArg(1, "ITERATIONS"); + }; + + void run() override { + _commandStream->printTiming(_iterations); + }; + +private: + int _iterations; +}; + + +//================================================================================================= +CopyCommand::CopyCommand(CommandStream *cmdStream, const std::vector &args, + hipMemcpyKind kind, bool isAsync, bool isPinnedHost) : + Command(cmdStream, args) , + _isAsync(isAsync), + _kind(kind), + _stream(cmdStream->currentStream()) + { + switch (kind) { + case hipMemcpyDeviceToHost: + _srcType = Device; + _dstType = isPinnedHost ? PinnedHost : UnpinnedHost; + break; + case hipMemcpyHostToDevice: + _srcType = isPinnedHost ? PinnedHost : UnpinnedHost; + _dstType = Device; + break; + default: + HIPASSERT(0, "Unknown hipMemcpyKind"); + }; + + _sizeBytes = 64; //TODO, support reading from arg. + + _dst = alloc(_sizeBytes, _dstType); + _src = alloc(_sizeBytes, _srcType); + }; + + +//================================================================================================= +//================================================================================================= +// Implementations: +//================================================================================================= + +//================================================================================================= +CommandStream::CommandStream(std::string commandStreamString, int iterations) + : _iterations(iterations), + _startTime(0), + _elapsedUs(0.0), + _parentCommandStream(nullptr), + _parseInSubBlock(false) +{ + std::vector tokens; + tokenize(commandStreamString, ';', tokens); + + + std::for_each(tokens.begin(), tokens.end(), [&] (const std::string s) { + this->parse(s); + }); + + setStream(0); +} + + +CommandStream::~CommandStream() +{ + std::for_each(_state._streams.begin(), _state._streams.end(), [&] (hipStream_t s) { + if (s) { + HIPCHECK(hipStreamDestroy(s)); + } + }); + + std::for_each(_commands.begin(), _commands.end(), [&] (Command *c) { + delete c; + }); + + +} + + +void CommandStream::setStream(int streamIndex) +{ + + if (streamIndex >= _state._streams.size()) { + _state._streams.resize(streamIndex+1); + } + + if (streamIndex && (_state._streams[streamIndex] == nullptr)) { + // Create new stream: + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + _state._streams[streamIndex] = stream; + _state._currentStream = stream; + } else { + // Use existing stream: + + _state._currentStream = _state._streams[streamIndex]; + } + +} + + +void CommandStream::tokenize(const std::string &s, char delim, std::vector &tokens) +{ + std::stringstream ss; + ss.str(s); + std::string item; + while (getline(ss, item, delim)) { + item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace. + tokens.push_back(item); + } +} + +void trim(std::string *s) +{ + // trim whitespace from begin and end: + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); + s->erase(s->find_last_not_of(t)+1); +} + +void ltrim(std::string *s) +{ + // trim whitespace from begin and end: + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); +} + +void CommandStream::parse(std::string fullCmd) +{ + //convert to lower-case: + std::transform(fullCmd.begin(), fullCmd.end(), fullCmd.begin(), ::tolower); + trim(&fullCmd); + + if (p_db) { + printf ("parse: <%s>\n", fullCmd.c_str()); + } + + + + std::string c; + std::vector args; + size_t leftParenZ = fullCmd.find_first_of('('); + if (leftParenZ == string::npos) { + c = fullCmd; + args.push_back(c); + } else { + c = fullCmd.substr(0, leftParenZ); + args.push_back(c); + size_t rightParenZ = fullCmd.find_first_of(')', leftParenZ); + std::string argStr = fullCmd.substr(leftParenZ+1, rightParenZ-leftParenZ-1); + //printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(), leftParenZ, rightParenZ); + tokenize(argStr, ',', args); + + } + + + + + + if ((args.size()==0) || (fullCmd.c_str()[0] == '#') ) { + if (p_db) { + printf (" skip comment\n"); + } + return; + } + + + + Command *cmd = NULL; + CommandStream *cmdStream = currentCommandStream(); + + if (c == "h2d") { + cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true/*isAsync*/, true/*isPinned*/); + //= h2d + //= Performs an async host-to-device copy of array A_h to A_d. + //= The size of these arrays may be set with the datasize command. + + } else if (c == "d2h") { + cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true/*isAsync*/, true/*isPinned*/); + //= d2h + //= Performs an async device-to-host copy of array A_d to A_h. + //= The size of these arrays may be set with the datasize command. + + } else if (c == "modulekernel") { + cmd = new ModuleKernelCommand(cmdStream, args); + + } else if (c == "nullkernel") { + cmd = new KernelCommand(cmdStream, args, KernelCommand::Null); + //= nullkernel + //= Dispatches a null kernel to the device. + + } else if (c == "vectoraddkernel") { + cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd); + +#ifdef __HIP_PLATFORM_HCC__ + } else if (c == "nullpfe") { + cmd = new PfeCommand(cmdStream, args); + + } else if (c == "aqlkernel") { + cmd = new AqlKernelCommand(cmdStream, args); +#endif + + } else if (c == "devicesync") { + cmd = new DeviceSyncCommand(cmdStream, args); + + } else if (c == "streamsync") { + //= streamsync + //= Execute hipStreamSynchronize. + //= This will cause the host thread to wait until the current stream + //= completes all pending operations. + cmd = new StreamSyncCommand(cmdStream, args); + + } else if (c == "setstream") { + //= setstream(STREAM_INDEX); + //= Set current stream used by subsequent commands. + //= STREAM_INDEX is index starting from 0...N. + //= This function will create new stream on first call to setstream or re-use previous + //= stream if setstream has already been called with STREAM_INDEX. + //= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will not create a new stream. + //= The default stream has special, conservative synchronization properties. + + cmd = new SetStreamCommand(cmdStream, args); + + } else if (c == "printtiming") { + cmd = new PrintTimingCommand(cmdStream, args); + + } else if (c == "loop") { + //= loop(LOOP_CNT) + //= Loop over next set of commands (until 'endloop' command) for LOOP_CNT iterations. + //= Loops can be nested. + + cmd = new LoopCommand(cmdStream, args); + + } else if (c == "endloop") { + //= endloop + //= End a looped sequence. Must be paired with a preceding loop command. + //= Command between the `loop` and `endloop` must be executed + + CommandStream * parentCmdStream = cmdStream->getParent() ; + cmd = new EndBlockCommand(cmdStream, parentCmdStream, args); + cmdStream = parentCmdStream; + + } else { + std::cerr << "error: Bad command '" << fullCmd << "\n"; + HIPASSERT(0, "bad command in command-stream"); + } + + if (cmd) { + cmdStream->_commands.push_back(cmd); + } +} + + + + +void CommandStream::print(const std::string &indent) const +{ + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + (*cmdI)->print(indent); + }; +} + + +void CommandStream::printBrief(std::ostream &s) const +{ + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + (*cmdI)->printBrief(s); + s << ";"; + }; +} + +void CommandStream::run() +{ + _startTime = get_time(); + for (int i=0; i<_iterations; i++) { + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + if (p_verbose) { + (*cmdI)->print(); + } + (*cmdI)->run(); + } + } + + // Record time, if not already stored. (an earlier printTime command will also store the time) + recordTime(); +}; + +void CommandStream::recordTime() +{ + if (_elapsedUs == 0.0) { + auto stopTime = get_time(); + _elapsedUs = stopTime - _startTime; + } +} + + +void CommandStream::printTiming(int iterations) +{ + + if ((_state._subBlocks.size() == 1) && (_commands.size()==1)) { + //printf ("print just the loop\n"); + _state._subBlocks.front()->printTiming(iterations); + } else { + g_printedTiming = true; + + recordTime(); + if (iterations == 0) { + iterations = _iterations; + } + std::cout << "command<"; printBrief(std::cout); + std::cout << ">," ; + printf (" iterations,%d, total_time,%6.3f, time/iteration,%6.3f\n", iterations, _elapsedUs, _elapsedUs/iterations); + } +}; + + + + + +//================================================================================================= +int main(int argc, char *argv[]) +{ + parseStandardArguments(argc, argv); + + printConfig(); + + CommandStream *cs; + + if (p_blockingSync) { +#ifdef __HIP_PLATFORM_HCC__ + printf ("setting BlockingSync for AMD\n"); + setenv("HIP_BLOCKING_SYNC", "1", 1); + +#endif +#ifdef __HIP_PLATFORM_NVCC__ + printf ("setting cudaDeviceBlockingSync\n"); + HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync)); +#endif + }; + + + if (p_file) { + // TODO - catch exception on file IO here: + std::ifstream file(p_file); + std::string str; + std::string file_contents; + while (std::getline(file, str)) + { + file_contents += str; + } + + cs = new CommandStream(file_contents, p_iterations); + + } else { + cs = new CommandStream(p_command, p_iterations); + } + + cs->print(); + printf ("------\n"); + + cs->run(); + if (!g_printedTiming) { + cs->printTiming(); + } + + delete cs; +} + + + +// TODO - add error checking for arguments. diff --git a/hipamd/samples/1_Utils/hipCommander/l2.hcm b/hipamd/samples/1_Utils/hipCommander/l2.hcm new file mode 100644 index 0000000000..b541bd6a66 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/l2.hcm @@ -0,0 +1,3 @@ +setstream,1; +NullKernel; streamsync; +loop,10000; H2D; NullKernel; streamsync; endloop,1; diff --git a/hipamd/samples/1_Utils/hipCommander/loop.hcm b/hipamd/samples/1_Utils/hipCommander/loop.hcm new file mode 100644 index 0000000000..db11071203 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/loop.hcm @@ -0,0 +1,3 @@ +loop,1000; H2D; NullKernel; D2H; endloop; +streamsync; +printTiming, 1000 diff --git a/hipamd/samples/1_Utils/hipCommander/loop2.hcm b/hipamd/samples/1_Utils/hipCommander/loop2.hcm new file mode 100644 index 0000000000..b8a14aa156 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/loop2.hcm @@ -0,0 +1,2 @@ +setstream,1; +loop,1000; NullKernel; syncstream; endloop,1, diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp new file mode 100644 index 0000000000..890e9bdc1e --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp @@ -0,0 +1,7 @@ +#include "hip/hip_runtime.h" + +extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){ + if (Ad) { + Ad[0] = 42; + } +} diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco b/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco new file mode 100755 index 0000000000000000000000000000000000000000..585b55cce587ed928acc1308d8a76361cdd87e57 GIT binary patch literal 10265 zcmeHNO>7%Q6n-9W){e7rDA1N4A(nuuP*E$pN!(QGAwP9XMB;{~MIu7lxa&>h;Gf9e zRH;N?J4e){oH(I02gIQ#BrYIvLE_l(-t3J3NKt=6)k-_E=gs?@ zo&DZ>E6;i^Cq_@QKmcS5;-JO>3<$-{8`6Sg{Sd;w9U`EjSi$drP&d;xF`1-cVMk)} zyI<=RN=e~4&yxIFXejt*dOYOWNhdcY$op{GQ5pX!{nq9&&*`z#++QqalPqA>_pO> zv5lj4Avc$G+)N>lv$^v{CoyiAWIf>)om9?p<U<1^Z(sJ!sobsZ=PAn7uXqa53X% z>{NCrRdOtMBst;sSeDV3n|GasL?SVM)|wcya*5=*Y;vO8G7LA9b<$QQpLP~TMl%MQ zB&^++pBuOB-IqT5^e=sOwi;c3;3ME8;3ME8;3ME8;3ME8;3ME8@T?;swu88Lqn#`6 z>ZIU4H?{=tzzx*3;oiT#y;#TW=X2ErilXf0mbNzW9*h{UP1@I9Vd^GLcNwm4z)o~! zX9pl04S#{CnCrkgBk;@OtiE=u99~;H2WuNqO(`#&gwLu%#VEL9i;PS%*?aek}+q+w3*kPhaH};D@b>%AgWfdococln}#6tZWBckhX-) zV2cvRz4jf}imN4Gd;t=y`u?=V(3h!dZ&dwi5%1+rWjR3aNRZydXyV*q21;A_5EJ7i zC1T!I!Y0kx(h{M4JfAzNhU4U~gdz}CRa347<9Aq5)x${RNaGdt(?!x%*v$a04d=PK z9aV7sfy%9QCETe&#B673IS|tNSpdvHE9+NTv>bt-7Xha3w%|LOS*WbeAx34mckQp8 zYeZMe;d`q9OB)E`6KgWvA2k1$A0Yn^Kx@PLpS}Foy!_X8&wmT^U#rXiE02=@FN^#~ z`pEf@zDtYKP|40gfr}70iFFF7GwWdtMK1RD)hE z{&)UkePaG=SYs-DwvEpOtm_rdN41FA!L)LPDdkYR*1xWXJ45jK_Sd{tmTv{CwV{*& z=I&y-1Jy*REn=$sw0?j;ws>{2<<&A1b;9>wbO{vLU|F*D~{|R>qy=PVHNl$xYG9Mf} zLBF|qBl)ZNrC*99SMWpu*-?`+l%Ei$H#jE$k`eNO>-gO@T zu1)JGAd~KxfuDHKG4n`Y`dF{ilRnaKPxtkhaeZ|3gE7ykbTR)nK&*5r=cdjhb&DRI zk)!mbJSrhJlPb+Xj5xOlG5j0Y3^n9MnicmN%x^lVpDJW zx9)MddLfRbpA-`>V=rYyzkSQ&WFFG{Fv48oFY}Myi0pR#xkH{my{W2iO>DiYVeHu{ zaQ*oPv52(-XPe~h+)RP*_BYyl$6XZoq)E + +static const int BLOCKSIZEX=32; +static const int BLOCKSIZEY=16; + +__global__ void fails(hipLaunchParm lp, float* pErrorI) +{ + if(pErrorI!=0) + { + pErrorI[0]=1; + } +} + +int main() +{ + dim3 blocks(1,1); + dim3 threads(BLOCKSIZEX,BLOCKSIZEY); + float error; + + hipLaunchKernel(HIP_KERNEL_NAME(fails), blocks, threads, 0, 0, &error); +} diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 79d117a0c4..36790271de 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1749,13 +1749,17 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, if (HIP_FORCE_P2P_HOST & 0x1) { *forceUnpinnedCopy = true; - tprintf (DB_COPY, "P2P. Copy engine (dev:%d) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", (*copyDevice)->getDeviceNum()); + tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", + (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); + } else { - tprintf (DB_COPY, "P2P. Copy engine (dev:%d) can see src and dst.\n", (*copyDevice)->getDeviceNum()); + tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst.\n", + (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); } } else { *forceUnpinnedCopy = true; - tprintf (DB_COPY, "P2P: copy engine(dev:%d) cannot see both host and device pointers - forcing copy with unpinned engine.\n", (*copyDevice)->getDeviceNum()); + tprintf (DB_COPY, "P2P: copy engine(dev:%d agent=0x%lx) cannot see both host and device pointers - forcing copy with unpinned engine.\n", + (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); } } @@ -1789,10 +1793,10 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem, src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem, sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); - tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n", dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem); - tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n", src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem); @@ -1846,10 +1850,10 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem, src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem, sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); - tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + tprintf (DB_COPY, " dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n", dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem); - tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n", + tprintf (DB_COPY, " src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n", src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem);