AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

normalizing C API

Przeglądaj źródła 
[ROCm/roctracer commit: 077f8ec6b6]
This commit is contained in:
Evgeny
2019-12-26 06:40:49 -06:00
rodzic 0f95193082
commit bca4e77df9
15 zmienionych plików z 162 dodań i 103 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/roctracer/inc/ext/prof_protocol.h
+2 -2
Wyświetl plik
@@ -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);
projects/roctracer/inc/roctracer.h
+52 -17
Wyświetl plik
@@ -38,10 +38,13 @@ THE SOFTWARE.
#include <stdint.h>
#include <stddef.h>
#ifndef __cplusplus
#include <stdbool.h>
#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,
projects/roctracer/inc/roctracer_ext.h
+2 -2
Wyświetl plik
@@ -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
projects/roctracer/inc/roctracer_hcc.h
+10 -14
Wyświetl plik
@@ -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 <hc_prof_runtime.h>
#else
#include <hcc/hc_prof_runtime.h>
#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"
projects/roctracer/inc/roctracer_hip.h
+1 -1
Wyświetl plik
@@ -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
projects/roctracer/inc/roctracer_kfd.h
+2 -2
Wyświetl plik
@@ -23,11 +23,11 @@ THE SOFTWARE.
/////////////////////////////////////////////////////////////////////////////
#ifndef INC_ROCTRACER_KFD_H_
#define INC_ROCTRACER_KFD_H_
#include <iostream>
#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_
projects/roctracer/inc/roctracer_roctx.h
+5 -6
Wyświetl plik
@@ -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<ROCTX_API_ID_NUMBER> cb_table_t;
} // namespace roctx
#endif
#ifdef __cplusplus
extern "C" {
projects/roctracer/script/kfdap.py
+5 -3
Wyświetl plik
@@ -284,7 +284,6 @@ class API_DescrParser:
self.content_h += '#include <string.h>\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<uint32_t, kfd_api_data_t> 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
projects/roctracer/src/core/roctracer.cpp
+8 -8
Wyświetl plik
@@ -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<roctracer::memory_pool_mutex_t> lock(roctracer::memory_pool_mutex);
roctracer_pool_t* p = reinterpret_cast<roctracer_pool_t*>(roctracer::memory_pool);
if (pool != NULL) roctracer::memory_pool = reinterpret_cast<roctracer::MemoryPool*>(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<roctracer::memory_pool_mutex_t> 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<roctracer::MemoryPool*>(pool);
projects/roctracer/test/CMakeLists.txt
+2
Wyświetl plik
@@ -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
projects/roctracer/test/MatrixTranspose_test/Makefile
+17 -7
Wyświetl plik
@@ -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)
projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.c
Dowiązanie symboliczne
+1
Wyświetl plik
@@ -0,0 +1 @@
MatrixTranspose.cpp
projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp
+52 -39
Wyświetl plik
@@ -20,32 +20,39 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <iostream>
#include <stdio.h>
// roctracer extension API
#include <inc/roctracer_ext.h>
// hip header file
#include <hip/hip_runtime.h>
#ifdef __cplusplus
#include <cstdlib>
using namespace std;
#else
#include <stdlib.h>
#endif
// roctx header file
#include <inc/roctx.h>
// roctracer extension API
#include <inc/roctracer_ext.h>
// kfd header file
#include <inc/roctracer_kfd.h>
#if HIP_TEST
// hip header file
#include <hip/hip_runtime.h>
// 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 <inc/roctracer_hip.h>
#include <inc/roctracer_hcc.h>
#include <inc/roctracer_kfd.h>
#include <inc/roctracer_roctx.h>
// 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<const roctx_api_data_t*>(callback_data);
const roctx_api_data_t* data = (const roctx_api_data_t*)(callback_data);
fprintf(stdout, "<rocTX \"%s\">\n", data->args.message);
return;
}
if (domain == ACTIVITY_DOMAIN_KFD_API) {
const kfd_api_data_t* data = reinterpret_cast<const kfd_api_data_t*>(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<const hip_api_data_t*>(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<const roctracer_record_t*>(begin);
const roctracer_record_t* end_record = reinterpret_cast<const roctracer_record_t*>(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() {}
projects/roctracer/test/run.sh
+1
Wyświetl plik
@@ -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
projects/roctracer/test/tool/tracer_tool.cpp
+2 -2
Wyświetl plik
@@ -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);
}