diff --git a/inc/ext/prof_protocol.h b/inc/ext/prof_protocol.h index 6d9cd62714..ee52e91082 100644 --- a/inc/ext/prof_protocol.h +++ b/inc/ext/prof_protocol.h @@ -60,7 +60,7 @@ typedef enum { typedef uint64_t activity_correlation_id_t; // Activity record type -struct activity_record_t { +typedef struct activity_record_s { uint32_t domain; // activity domain id activity_kind_t kind; // activity kind activity_op_t op; // activity op @@ -81,7 +81,7 @@ struct activity_record_t { }; }; size_t bytes; // data size bytes -}; +} activity_record_t; // Activity sync calback type typedef void* (*activity_sync_callback_t)(uint32_t cid, activity_record_t* record, const void* data, void* arg); diff --git a/inc/roctracer.h b/inc/roctracer.h index ce9b31add9..0fc6df0063 100644 --- a/inc/roctracer.h +++ b/inc/roctracer.h @@ -38,10 +38,13 @@ THE SOFTWARE. #include #include +#ifndef __cplusplus +#include +#endif #include "ext/prof_protocol.h" -#define ROCTRACER_VERSION_MAJOR 1 +#define ROCTRACER_VERSION_MAJOR 2 #define ROCTRACER_VERSION_MINOR 0 #ifdef __cplusplus @@ -89,7 +92,7 @@ roctracer_status_t roctracer_op_code( uint32_t domain, // tracing domain const char* str, // [in] op string uint32_t* op, // [out] op code - uint32_t* kind = NULL); // [out] op kind code + uint32_t* kind); // [out] op kind code if not NULL //////////////////////////////////////////////////////////////////////////////// // Callback API @@ -172,31 +175,59 @@ typedef void roctracer_pool_t; // Create tracer memory pool // The first invocation sets the default pool -roctracer_status_t roctracer_open_pool( +roctracer_status_t roctracer_open_pool_expl( const roctracer_properties_t* properties, // tracer pool properties - roctracer_pool_t** pool = NULL); // [out] returns tracer pool if not NULL, + roctracer_pool_t** pool); // [out] returns tracer pool if not NULL, // otherwise sets the default one if it is not set yet +roctracer_status_t roctracer_open_pool( + const roctracer_properties_t* properties) // tracer pool properties +{ + return roctracer_open_pool_expl(properties, NULL); +} // otherwise the error is generated // Close tracer memory pool -roctracer_status_t roctracer_close_pool( - roctracer_pool_t* pool = NULL); // [in] memory pool, NULL is a default one +roctracer_status_t roctracer_close_pool_expl( + roctracer_pool_t* pool); // [in] memory pool, NULL is a default one +roctracer_status_t roctracer_close_pool() +{ + return roctracer_close_pool_expl(NULL); +} // Return current default pool // Set new default pool if the argument is not NULL -roctracer_pool_t* roctracer_default_pool( - roctracer_pool_t* pool = NULL); // [in] new default pool if not NULL +roctracer_pool_t* roctracer_default_pool_expl( + roctracer_pool_t* pool); // [in] new default pool if not NULL +roctracer_pool_t* roctracer_default_pool() +{ + return roctracer_default_pool_expl(NULL); +} // Enable activity records logging -roctracer_status_t roctracer_enable_op_activity( +roctracer_status_t roctracer_enable_op_activity_expl( activity_domain_t domain, // tracing domain uint32_t op, // activity op ID - roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one -roctracer_status_t roctracer_enable_domain_activity( + roctracer_pool_t* pool); // memory pool, NULL is a default one +roctracer_status_t roctracer_enable_op_activity( activity_domain_t domain, // tracing domain - roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one -roctracer_status_t roctracer_enable_activity( - roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one + uint32_t op) // activity op ID +{ + return roctracer_enable_op_activity_expl(domain, op, NULL); +} +roctracer_status_t roctracer_enable_domain_activity_expl( + activity_domain_t domain, // tracing domain + roctracer_pool_t* pool); // memory pool, NULL is a default one +roctracer_status_t roctracer_enable_domain_activity( + activity_domain_t domain) // tracing domain +{ + return roctracer_enable_domain_activity_expl(domain, NULL); +} +roctracer_status_t roctracer_enable_activity_expl( + roctracer_pool_t* pool); // memory pool, NULL is a default one +roctracer_status_t roctracer_enable_activity() +{ + return roctracer_enable_activity_expl(NULL); +} // Disable activity records logging roctracer_status_t roctracer_disable_op_activity( @@ -207,8 +238,12 @@ roctracer_status_t roctracer_disable_domain_activity( roctracer_status_t roctracer_disable_activity(); // Flush available activity records -roctracer_status_t roctracer_flush_activity( - roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one +roctracer_status_t roctracer_flush_activity_expl( + roctracer_pool_t* pool); // memory pool, NULL is a default one +roctracer_status_t roctracer_flush_activity() +{ + return roctracer_flush_activity_expl(NULL); +} // Load/Unload methods // Set properties @@ -216,7 +251,7 @@ roctracer_status_t roctracer_set_properties( roctracer_domain_t domain, // tracing domain void* propertes); // tracing properties -struct HsaApiTable; +typedef struct HsaApiTable HsaApiTable; bool roctracer_load( HsaApiTable* table, uint64_t runtime_version, diff --git a/inc/roctracer_ext.h b/inc/roctracer_ext.h index c2f5c54542..2427336c7b 100644 --- a/inc/roctracer_ext.h +++ b/inc/roctracer_ext.h @@ -63,8 +63,8 @@ roctracer_status_t roctracer_activity_push_external_correlation_id(activity_corr // Notifies that the calling thread is leaving an external API region. // Pop an external correlation id for the calling thread. -// 'lastId' returns the last external correlation -roctracer_status_t roctracer_activity_pop_external_correlation_id(activity_correlation_id_t* last_id = NULL); +// 'lastId' returns the last external correlation if not NULL +roctracer_status_t roctracer_activity_pop_external_correlation_id(activity_correlation_id_t* last_id); #ifdef __cplusplus } // extern "C" block diff --git a/inc/roctracer_hcc.h b/inc/roctracer_hcc.h index 252b984d2a..0781460145 100644 --- a/inc/roctracer_hcc.h +++ b/inc/roctracer_hcc.h @@ -23,26 +23,22 @@ THE SOFTWARE. #ifndef INC_ROCTRACER_HCC_H_ #define INC_ROCTRACER_HCC_H_ -#if HIP_VDI -#define HIP_OP_ID_NUMBER 3 -#define HIP_OP_ID_COPY 1 +enum { + HIP_OP_ID_DISPATCH = 0, + HIP_OP_ID_COPY = 1, + HIP_OP_ID_BARRIER = 2, + HIP_OP_ID_NUMBER = 3 +}; + +#ifdef __cplusplus extern "C" { +#endif typedef void (hipInitAsyncActivityCallback_t)(void* id_callback, void* op_callback, void* arg); typedef bool (hipEnableAsyncActivityCallback_t)(unsigned op, bool enable); typedef const char* (hipGetOpName_t)(unsigned op); +#ifdef __cplusplus } -#else // !HIP_VDI -#if LOCAL_BUILD -#include -#else -#include #endif -#define HIP_OP_ID_NUMBER hc::HSA_OP_ID_NUMBER -#define HIP_OP_ID_COPY hc::HSA_OP_ID_COPY -typedef decltype(Kalmar::CLAMP::InitActivityCallback) hipInitAsyncActivityCallback_t; -typedef decltype(Kalmar::CLAMP::EnableActivityCallback) hipEnableAsyncActivityCallback_t; -typedef decltype(Kalmar::CLAMP::GetCmdName) hipGetOpName_t; -#endif // !HIP_VDI #include "roctracer.h" diff --git a/inc/roctracer_hip.h b/inc/roctracer_hip.h index d365dd9444..28e4868d59 100644 --- a/inc/roctracer_hip.h +++ b/inc/roctracer_hip.h @@ -33,7 +33,7 @@ extern "C" { #endif // __cplusplus // Traced calls ID enumeration -typedef hip_api_id_t roctracer_hip_api_cid_t; +typedef enum hip_api_id_t roctracer_hip_api_cid_t; #ifdef __cplusplus } // extern "C" block diff --git a/inc/roctracer_kfd.h b/inc/roctracer_kfd.h index 45113ce435..fcc1e3cd87 100644 --- a/inc/roctracer_kfd.h +++ b/inc/roctracer_kfd.h @@ -23,11 +23,11 @@ THE SOFTWARE. ///////////////////////////////////////////////////////////////////////////// #ifndef INC_ROCTRACER_KFD_H_ #define INC_ROCTRACER_KFD_H_ -#include - #include "roctracer.h" #include "hsakmt.h" +#ifdef __cplusplus #include "inc/kfd_ostream_ops.h" +#endif #include "inc/kfd_prof_str.h" #endif // INC_ROCTRACER_KFD_H_ diff --git a/inc/roctracer_roctx.h b/inc/roctracer_roctx.h index 329e974d61..accec45255 100644 --- a/inc/roctracer_roctx.h +++ b/inc/roctracer_roctx.h @@ -33,8 +33,6 @@ THE SOFTWARE. #ifndef INC_ROCTRACER_ROCTX_H_ #define INC_ROCTRACER_ROCTX_H_ -#include "cb_table.h" - // ROC-TX API ID enumeration enum roctx_api_id_t { ROCTX_API_ID_roctxMarkA = 0, @@ -45,7 +43,7 @@ enum roctx_api_id_t { }; // ROCTX callbacks data type -struct roctx_api_data_t { +typedef struct roctx_api_data_s { union { const char* message; struct { @@ -58,14 +56,15 @@ struct roctx_api_data_t { const char* message; } roctxRangePop; } args; -}; +} roctx_api_data_t; +#ifdef __cplusplus +#include "cb_table.h" namespace roctx { - // ROCTX callbacks table type typedef roctracer::CbTable cb_table_t; - } // namespace roctx +#endif #ifdef __cplusplus extern "C" { diff --git a/script/kfdap.py b/script/kfdap.py index c46bf60d78..06248d9f26 100755 --- a/script/kfdap.py +++ b/script/kfdap.py @@ -284,7 +284,6 @@ class API_DescrParser: self.content_h += '#include \n' self.content_h += '#include \"roctracer_kfd.h\"\n' self.content_h += '#include \"hsakmt.h\"\n' - self.content_h += '#include \"cb_table.h\"\n' self.content_h += '#define PUBLIC_API __attribute__((visibility(\"default\")))\n' @@ -293,6 +292,7 @@ class API_DescrParser: self.content_h += '\n' self.content_h += '#if PROF_API_IMPL\n' + self.content_h += '#include \"cb_table.h\"\n' self.content_h += 'namespace roctracer {\n' self.content_h += 'namespace kfd_support {\n' @@ -372,7 +372,7 @@ class API_DescrParser: # generate API args structure def gen_arg_struct(self, n, name, call, struct): if n == -1: - self.content_h += 'struct kfd_api_data_t {\n' + self.content_h += 'typedef struct kfd_api_data_s {\n' self.content_h += ' uint64_t correlation_id;\n' self.content_h += ' uint32_t phase;\n' if len(self.api_rettypes) != 0: @@ -394,7 +394,7 @@ class API_DescrParser: self.content_h += ' } ' + call + ';\n' else: self.content_h += ' } args;\n' - self.content_h += '};\n' + self.content_h += '} kfd_api_data_t;\n' # generate API callbacks def gen_callbacks(self, n, name, call, struct): @@ -476,6 +476,7 @@ class API_DescrParser: # generate stream operator def gen_out_stream(self, n, name, call, struct): if n == -1: + self.content_h += '#ifdef __cplusplus\n' self.content_h += 'typedef std::pair kfd_api_data_pair_t;\n' self.content_h += 'inline std::ostream& operator<< (std::ostream& out, const kfd_api_data_pair_t& data_pair) {\n' self.content_h += ' const uint32_t cid = data_pair.first;\n' @@ -509,6 +510,7 @@ class API_DescrParser: self.content_h += ' }\n' self.content_h += ' return out;\n' self.content_h += '}\n' + self.content_h += '#endif\n' self.content_cpp += 'inline std::ostream& operator<< (std::ostream& out, const HsaMemFlags& v) { out << "HsaMemFlags"; return out; }\n' # generate PUBLIC_API for all API fcts diff --git a/src/core/roctracer.cpp b/src/core/roctracer.cpp index 9199bad668..c006e6d88c 100644 --- a/src/core/roctracer.cpp +++ b/src/core/roctracer.cpp @@ -176,7 +176,7 @@ decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn; typedef decltype(roctracer_enable_op_callback)* roctracer_enable_op_callback_t; typedef decltype(roctracer_disable_op_callback)* roctracer_disable_op_callback_t; -typedef decltype(roctracer_enable_op_activity)* roctracer_enable_op_activity_t; +typedef decltype(roctracer_enable_op_activity_expl)* roctracer_enable_op_activity_t; typedef decltype(roctracer_disable_op_activity)* roctracer_disable_op_activity_t; struct cb_journal_data_t { @@ -774,7 +774,7 @@ PUBLIC_API roctracer_status_t roctracer_disable_callback() } // Return default pool and set new one if parameter pool is not NULL. -PUBLIC_API roctracer_pool_t* roctracer_default_pool(roctracer_pool_t* pool) { +PUBLIC_API roctracer_pool_t* roctracer_default_pool_expl(roctracer_pool_t* pool) { std::lock_guard lock(roctracer::memory_pool_mutex); roctracer_pool_t* p = reinterpret_cast(roctracer::memory_pool); if (pool != NULL) roctracer::memory_pool = reinterpret_cast(pool); @@ -782,7 +782,7 @@ PUBLIC_API roctracer_pool_t* roctracer_default_pool(roctracer_pool_t* pool) { } // Open memory pool -PUBLIC_API roctracer_status_t roctracer_open_pool( +PUBLIC_API roctracer_status_t roctracer_open_pool_expl( const roctracer_properties_t* properties, roctracer_pool_t** pool) { @@ -799,7 +799,7 @@ PUBLIC_API roctracer_status_t roctracer_open_pool( } // Close memory pool -PUBLIC_API roctracer_status_t roctracer_close_pool(roctracer_pool_t* pool) { +PUBLIC_API roctracer_status_t roctracer_close_pool_expl(roctracer_pool_t* pool) { API_METHOD_PREFIX std::lock_guard lock(roctracer::memory_pool_mutex); roctracer_pool_t* ptr = (pool == NULL) ? roctracer_default_pool() : pool; @@ -868,7 +868,7 @@ static void roctracer_enable_activity_impl( roctracer_enable_activity_fun((roctracer_domain_t)domain, op, pool); } -PUBLIC_API roctracer_status_t roctracer_enable_op_activity( +PUBLIC_API roctracer_status_t roctracer_enable_op_activity_expl( roctracer_domain_t domain, uint32_t op, roctracer_pool_t* pool) @@ -878,7 +878,7 @@ PUBLIC_API roctracer_status_t roctracer_enable_op_activity( API_METHOD_SUFFIX } -PUBLIC_API roctracer_status_t roctracer_enable_domain_activity( +PUBLIC_API roctracer_status_t roctracer_enable_domain_activity_expl( roctracer_domain_t domain, roctracer_pool_t* pool) { @@ -888,7 +888,7 @@ PUBLIC_API roctracer_status_t roctracer_enable_domain_activity( API_METHOD_SUFFIX } -PUBLIC_API roctracer_status_t roctracer_enable_activity( +PUBLIC_API roctracer_status_t roctracer_enable_activity_expl( roctracer_pool_t* pool) { API_METHOD_PREFIX @@ -970,7 +970,7 @@ PUBLIC_API roctracer_status_t roctracer_disable_activity() } // Flush available activity records -PUBLIC_API roctracer_status_t roctracer_flush_activity(roctracer_pool_t* pool) { +PUBLIC_API roctracer_status_t roctracer_flush_activity_expl(roctracer_pool_t* pool) { API_METHOD_PREFIX if (pool == NULL) pool = roctracer_default_pool(); roctracer::MemoryPool* memory_pool = reinterpret_cast(pool); diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 923384c444..ef0a8ea633 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -35,6 +35,8 @@ add_custom_target( mytest COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose/MatrixTranspose ${PROJECT_BINARY_DIR}/test" COMMAND HIP_VDI=${HIP_VDI} make -C "${TEST_DIR}/MatrixTranspose_test" COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_test" + COMMAND C_TEST=1 HIP_VDI=${HIP_VDI} make -C "${TEST_DIR}/MatrixTranspose_test" + COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose_ctest ${PROJECT_BINARY_DIR}/test/MatrixTranspose_ctest" ) ## Util sources diff --git a/test/MatrixTranspose_test/Makefile b/test/MatrixTranspose_test/Makefile index 2a767a5626..c59c497af1 100644 --- a/test/MatrixTranspose_test/Makefile +++ b/test/MatrixTranspose_test/Makefile @@ -14,24 +14,34 @@ HIPCC=$(HIP_PATH)/bin/hipcc TARGET=hcc -SOURCES = MatrixTranspose.cpp -OBJECTS = $(SOURCES:.cpp=.o) -EXECUTABLE=./MatrixTranspose + +FLAGS =-g -I$(ROOT_PATH) -I$(ROOT_PATH)/inc -I${HSA_KMT_INC_PATH} -DLOCAL_BUILD=1 -DHIP_VDI=${HIP_VDI} -DITERATIONS=$(ITERATIONS) +ifeq ($(C_TEST), 1) + COMP=gcc + SOURCES = MatrixTranspose.c + FLAGS += -DHIP_TEST=0 -D__HIP_PLATFORM_HCC__=1 -I/opt/rocm/hcc/include + EXECUTABLE=./MatrixTranspose_ctest +else + COMP=$(HIPCC) + FLAGS += -DHIP_TEST=1 + SOURCES = MatrixTranspose.cpp + EXECUTABLE=./MatrixTranspose +endif +OBJECTS = MatrixTranspose.o .PHONY: test - all: clean $(EXECUTABLE) -CXXFLAGS =-g -I$(ROOT_PATH) -I$(ROOT_PATH)/inc -I${HSA_KMT_INC_PATH} -DLOCAL_BUILD=1 -DHIP_VDI=${HIP_VDI} -DITERATIONS=$(ITERATIONS) -CXX=$(HIPCC) +$(OBJECTS): $(SOURCES) + $(COMP) $(FLAGS) -c -o $@ $< $(EXECUTABLE): $(OBJECTS) $(HIPCC) $(OBJECTS) -o $@ $(ROC_LIBS) test: $(EXECUTABLE) - $(EXECUTABLE) + LD_PRELOAD=$(LIB_PATH)/libkfdwrapper64.so $(EXECUTABLE) clean: rm -f $(EXECUTABLE) diff --git a/test/MatrixTranspose_test/MatrixTranspose.c b/test/MatrixTranspose_test/MatrixTranspose.c new file mode 120000 index 0000000000..14d96acbc8 --- /dev/null +++ b/test/MatrixTranspose_test/MatrixTranspose.c @@ -0,0 +1 @@ +MatrixTranspose.cpp \ No newline at end of file diff --git a/test/MatrixTranspose_test/MatrixTranspose.cpp b/test/MatrixTranspose_test/MatrixTranspose.cpp index 1a7b2cb97a..11ad71709a 100644 --- a/test/MatrixTranspose_test/MatrixTranspose.cpp +++ b/test/MatrixTranspose_test/MatrixTranspose.cpp @@ -20,32 +20,39 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include -// roctracer extension API -#include - -// hip header file -#include +#ifdef __cplusplus +#include +using namespace std; +#else +#include +#endif // roctx header file #include +// roctracer extension API +#include -// kfd header file -#include +#if HIP_TEST +// hip header file +#include +// Macro to call HIP API +#define HIP_CALL(call) do { call; } while(0) +#else +#define HIP_CALL(call) do {} while(0) +#endif #ifndef ITERATIONS # define ITERATIONS 101 #endif #define WIDTH 1024 - - #define NUM (WIDTH * WIDTH) - #define THREADS_PER_BLOCK_X 4 #define THREADS_PER_BLOCK_Y 4 #define THREADS_PER_BLOCK_Z 1 +#if HIP_TEST // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; @@ -53,6 +60,7 @@ __global__ void matrixTranspose(float* out, float* in, const int width) { out[y * width + x] = in[x * width + y]; } +#endif // CPU implementation of matrix transpose void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { @@ -76,10 +84,12 @@ int main() { float* gpuMatrix; float* gpuTransposeMatrix; +#if HIP_TEST hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + HIP_CALL(hipGetDeviceProperties(&devProp, 0)); - std::cout << "Device name " << devProp.name << std::endl; + printf("Device name %s\n", devProp.name); +#endif int i; int errors; @@ -99,8 +109,8 @@ int main() { } // allocate the memory on the device side - hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); - hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + HIP_CALL(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float))); + HIP_CALL(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float))); // correlation reagion32 roctracer_activity_push_external_correlation_id(31); @@ -108,7 +118,7 @@ int main() { roctracer_activity_push_external_correlation_id(32); // Memory transfer from host to device - hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + HIP_CALL(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); // correlation reagion33 roctracer_activity_push_external_correlation_id(33); @@ -117,9 +127,9 @@ int main() { roctxRangePush("hipLaunchKernel"); // Lauching kernel from host - hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, - gpuMatrix, WIDTH); + HIP_CALL(hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH)); roctxMark("after hipLaunchKernel"); @@ -129,39 +139,40 @@ int main() { // Memory transfer from device to host roctxRangePush("hipMemcpy"); - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + HIP_CALL(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); roctxRangePop(); // for "hipMemcpy" roctxRangePop(); // for "hipLaunchKernel" // correlation reagion end - roctracer_activity_pop_external_correlation_id(); + roctracer_activity_pop_external_correlation_id(NULL); // CPU MatrixTranspose computation - matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + HIP_CALL(matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH)); // verify the results errors = 0; double eps = 1.0E-6; for (i = 0; i < NUM; i++) { - if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + if (abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { errors++; } } - if (errors != 0) { + if ((HIP_TEST != 0) && (errors != 0)) { printf("FAILED: %d errors\n", errors); } else { + errors = 0; printf("PASSED!\n"); } // free the resources on device side - hipFree(gpuMatrix); - hipFree(gpuTransposeMatrix); + HIP_CALL(hipFree(gpuMatrix)); + HIP_CALL(hipFree(gpuTransposeMatrix)); // correlation reagion end - roctracer_activity_pop_external_correlation_id(); + roctracer_activity_pop_external_correlation_id(NULL); // correlation reagion end - roctracer_activity_pop_external_correlation_id(); + roctracer_activity_pop_external_correlation_id(NULL); // free the resources on host side free(Matrix); @@ -180,6 +191,7 @@ int main() { #if 1 #include #include +#include #include // Macro to check ROC-tracer calls status @@ -187,7 +199,7 @@ int main() { do { \ int err = call; \ if (err != 0) { \ - std::cerr << roctracer_error_string() << std::endl << std::flush; \ + fprintf(stderr, "%s\n", roctracer_error_string()); \ abort(); \ } \ } while (0) @@ -202,12 +214,12 @@ void api_callback( (void)arg; if (domain == ACTIVITY_DOMAIN_ROCTX) { - const roctx_api_data_t* data = reinterpret_cast(callback_data); + const roctx_api_data_t* data = (const roctx_api_data_t*)(callback_data); fprintf(stdout, "\n", data->args.message); return; } if (domain == ACTIVITY_DOMAIN_KFD_API) { - const kfd_api_data_t* data = reinterpret_cast(callback_data); + const kfd_api_data_t* data = (const kfd_api_data_t*)(callback_data); fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s> \n", roctracer_op_string(ACTIVITY_DOMAIN_KFD_API, cid, 0), cid, @@ -215,7 +227,7 @@ void api_callback( (data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit"); return; } - const hip_api_data_t* data = reinterpret_cast(callback_data); + const hip_api_data_t* data = (const hip_api_data_t*)(callback_data); fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s> ", roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0), cid, @@ -263,8 +275,8 @@ void api_callback( // Activity tracing callback // hipMalloc id(3) correlation_id(1): begin_ns(1525888652762640464) end_ns(1525888652762877067) void activity_callback(const char* begin, const char* end, void* arg) { - const roctracer_record_t* record = reinterpret_cast(begin); - const roctracer_record_t* end_record = reinterpret_cast(end); + const roctracer_record_t* record = (const roctracer_record_t*)(begin); + const roctracer_record_t* end_record = (const roctracer_record_t*)(end); fprintf(stdout, "\tActivity records:\n"); fflush(stdout); while (record < end_record) { const char * name = roctracer_op_string(record->domain, record->op, record->kind); @@ -274,7 +286,7 @@ void activity_callback(const char* begin, const char* end, void* arg) { record->begin_ns, record->end_ns ); - if (record->domain == ACTIVITY_DOMAIN_HIP_API or record->domain == ACTIVITY_DOMAIN_KFD_API) { + if ((record->domain == ACTIVITY_DOMAIN_HIP_API) || (record->domain == ACTIVITY_DOMAIN_KFD_API)) { fprintf(stdout, " process_id(%u) thread_id(%u)", record->process_id, record->thread_id @@ -301,11 +313,12 @@ void activity_callback(const char* begin, const char* end, void* arg) { // Init tracing routine void init_tracing() { - std::cout << "# INIT #############################" << std::endl << std::flush; + printf("# INIT #############################\n"); // roctracer properties roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, NULL); // Allocating tracing pool - roctracer_properties_t properties{}; + roctracer_properties_t properties; + memset(&properties, 0, sizeof(roctracer_properties_t)); properties.buffer_size = 0x1000; properties.buffer_callback_fun = activity_callback; ROCTRACER_CALL(roctracer_open_pool(&properties)); @@ -323,7 +336,7 @@ void init_tracing() { // Start tracing routine void start_tracing() { - std::cout << "# START (" << iterations << ") #############################" << std::endl << std::flush; + printf("# START (%d) #############################\n", iterations); // Start if ((iterations & 1) == 1) roctracer_start(); else roctracer_stop(); @@ -336,7 +349,7 @@ void stop_tracing() { ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS)); ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_KFD_API)); ROCTRACER_CALL(roctracer_flush_activity()); - std::cout << "# STOP #############################" << std::endl << std::flush; + printf("# STOP #############################\n"); } #else void init_tracing() {} diff --git a/test/run.sh b/test/run.sh index ccf646aed6..d634357516 100755 --- a/test/run.sh +++ b/test/run.sh @@ -65,6 +65,7 @@ eval_test() { # Standalone test # rocTrecer is used explicitely by test +eval_test "standalone C test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_ctest" eval_test "standalone HIP test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_test" # Tool test diff --git a/test/tool/tracer_tool.cpp b/test/tool/tracer_tool.cpp index 549debe97a..78751a7463 100644 --- a/test/tool/tracer_tool.cpp +++ b/test/tool/tracer_tool.cpp @@ -663,7 +663,7 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, for (unsigned i = 0; i < hsa_api_vec.size(); ++i) { uint32_t cid = HSA_API_ID_NUMBER; const char* api = hsa_api_vec[i].c_str(); - ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api, &cid)); + ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api, &cid, NULL)); ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, NULL)); printf(" %s", api); } @@ -753,7 +753,7 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, for (unsigned i = 0; i < kfd_api_vec.size(); ++i) { uint32_t cid = KFD_API_ID_NUMBER; const char* api = kfd_api_vec[i].c_str(); - ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_KFD_API, api, &cid)); + ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_KFD_API, api, &cid, NULL)); ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_KFD_API, cid, kfd_api_callback, NULL)); printf(" %s", api); }