AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

Run clang-format on all source files

Kaynağa Gözat 
Change-Id: Ifb52ca306286b6b2d473821bed9db28e9f616d50
Bu işleme şunda yer alıyor:
Laurent Morichetti
2022-04-18 15:09:58 -07:00
işlemeyi yapan: Laurent Morichetti
ebeveyn 89f6880371
işleme 15ab5d9cda
40 değiştirilmiş dosya ile 1332 ekleme ve 1228 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
test/MatrixTranspose/MatrixTranspose.cpp
+49 -51
Dosyayı Görüntüle
@@ -37,57 +37,56 @@
#define THREADS_PER_BLOCK_Z 1
// Mark API
extern "C"
void roctracer_mark(const char* str);
extern "C" void roctracer_mark(const char* str);
// 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;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * width + y];
out[y * width + x] = in[x * width + y];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
for (unsigned int j = 0; j < width; j++) {
for (unsigned int i = 0; i < width; i++) {
output[i * width + j] = input[j * width + i];
}
for (unsigned int j = 0; j < width; j++) {
for (unsigned int i = 0; i < width; i++) {
output[i * width + j] = input[j * width + i];
}
}
}
int main() {
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
uint32_t iterations = 100;
while (iterations-- > 0) {
uint32_t iterations = 100;
while (iterations-- > 0) {
std::cout << "## Iteration (" << iterations << ") #################" << std::endl;
// Memory transfer from host to device
@@ -98,9 +97,9 @@ int main() {
int rangeId = roctxRangeStart("hipLaunchKernel range");
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);
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);
roctracer_mark("after HIP LaunchKernel");
roctxMark("after hipLaunchKernel");
@@ -109,8 +108,8 @@ int main() {
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
roctxRangePop(); // for "hipMemcpy"
roctxRangePop(); // for "hipLaunchKernel"
roctxRangePop(); // for "hipMemcpy"
roctxRangePop(); // for "hipLaunchKernel"
roctxRangeStop(rangeId);
// CPU MatrixTranspose computation
@@ -120,26 +119,25 @@ int main() {
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
errors++;
}
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
errors++;
}
}
if (errors != 0) {
printf("FAILED: %d errors\n", errors);
printf("FAILED: %d errors\n", errors);
} else {
printf("PASSED!\n");
printf("PASSED!\n");
}
}
}
// free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
return errors;
return errors;
}
test/MatrixTranspose_test/MatrixTranspose.cpp
+71 -82
Dosyayı Görüntüle
@@ -45,7 +45,7 @@ static char* message = NULL;
#endif
void SPRINT(const char* fmt, ...) {
if (msg_buf == NULL) {
msg_buf = (char*) calloc(msg_size, 1);
msg_buf = (char*)calloc(msg_size, 1);
message = msg_buf;
}
@@ -66,13 +66,18 @@ void SFLUSH() {
// hip header file
#include <hip/hip_runtime.h>
// Macro to call HIP API
#define HIP_CALL(call) do { call; } while(0)
#define HIP_CALL(call) \
do { \
call; \
} while (0)
#else
#define HIP_CALL(call) do {} while(0)
#define HIP_CALL(call) \
do { \
} while (0)
#endif
#ifndef ITERATIONS
# define ITERATIONS 101
#define ITERATIONS 101
#endif
#define WIDTH 1024
#define NUM (WIDTH * WIDTH)
@@ -83,20 +88,20 @@ void SFLUSH() {
#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;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * width + y];
out[y * width + x] = in[x * width + y];
}
#endif
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
for (unsigned int j = 0; j < width; j++) {
for (unsigned int i = 0; i < width; i++) {
output[i * width + j] = input[j * width + i];
}
for (unsigned int j = 0; j < width; j++) {
for (unsigned int i = 0; i < width; i++) {
output[i * width + j] = input[j * width + i];
}
}
}
int iterations = ITERATIONS;
@@ -105,28 +110,28 @@ void start_tracing();
void stop_tracing();
int main() {
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
int i;
int errors;
int i;
int errors;
init_tracing();
init_tracing();
#if HIP_TEST
int gpuCount = 1;
int gpuCount = 1;
#if MGPU_TEST
hipGetDeviceCount(&gpuCount);
printf("Number of GPUs: %d\n", gpuCount);
hipGetDeviceCount(&gpuCount);
printf("Number of GPUs: %d\n", gpuCount);
#endif
iterations *= gpuCount;
iterations *= gpuCount;
#endif
while (iterations-- > 0) {
while (iterations-- > 0) {
start_tracing();
#if HIP_TEST
@@ -145,7 +150,7 @@ int main() {
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
Matrix[i] = (float)i * 10.0f;
}
// allocate the memory on the device side
@@ -167,9 +172,10 @@ int main() {
roctxRangePush("hipLaunchKernel");
// Lauching kernel from host
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));
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");
@@ -179,10 +185,11 @@ int main() {
// Memory transfer from device to host
roctxRangePush("hipMemcpy");
HIP_CALL(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost));
HIP_CALL(
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost));
roctxRangePop(); // for "hipMemcpy"
roctxRangePop(); // for "hipLaunchKernel"
roctxRangePop(); // for "hipMemcpy"
roctxRangePop(); // for "hipLaunchKernel"
// correlation reagion end
roctracer_activity_pop_external_correlation_id(NULL);
@@ -194,15 +201,15 @@ int main() {
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
errors++;
}
if (abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
errors++;
}
}
if ((HIP_TEST != 0) && (errors != 0)) {
printf("FAILED: %d errors\n", errors);
printf("FAILED: %d errors\n", errors);
} else {
errors = 0;
printf("PASSED!\n");
errors = 0;
printf("PASSED!\n");
}
// free the resources on device side
@@ -218,11 +225,11 @@ int main() {
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
}
}
stop_tracing();
stop_tracing();
return errors;
return errors;
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -234,15 +241,15 @@ int main() {
#include <roctracer_hsa.h>
#include <roctracer_roctx.h>
#include <unistd.h>
#include <sys/syscall.h> /* For SYS_xxx definitions */
#include <unistd.h>
#include <sys/syscall.h> /* For SYS_xxx definitions */
// Macro to check ROC-tracer calls status
#define ROCTRACER_CALL(call) \
do { \
int err = call; \
if (err != 0) { \
fprintf(stderr, "%s\n", roctracer_error_string()); \
fprintf(stderr, "%s\n", roctracer_error_string()); \
abort(); \
} \
} while (0)
@@ -252,12 +259,7 @@ static inline uint32_t GetPid() { return syscall(__NR_getpid); }
// Runtime API callback function
void api_callback(
uint32_t domain,
uint32_t cid,
const void* callback_data,
void* arg)
{
void api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) {
(void)arg;
if (domain == ACTIVITY_DOMAIN_ROCTX) {
@@ -267,31 +269,25 @@ void api_callback(
}
const hip_api_data_t* data = (const hip_api_data_t*)(callback_data);
SPRINT("<%s id(%u)\tcorrelation_id(%lu) %s pid(%d) tid(%d)> ",
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
cid,
data->correlation_id,
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit", GetPid(), GetTid());
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0), cid, data->correlation_id,
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit", GetPid(), GetTid());
if (data->phase == ACTIVITY_API_PHASE_ENTER) {
switch (cid) {
case HIP_API_ID_hipMemcpy:
SPRINT("dst(%p) src(%p) size(0x%x) kind(%u)",
data->args.hipMemcpy.dst,
data->args.hipMemcpy.src,
(uint32_t)(data->args.hipMemcpy.sizeBytes),
(uint32_t)(data->args.hipMemcpy.kind));
SPRINT("dst(%p) src(%p) size(0x%x) kind(%u)", data->args.hipMemcpy.dst,
data->args.hipMemcpy.src, (uint32_t)(data->args.hipMemcpy.sizeBytes),
(uint32_t)(data->args.hipMemcpy.kind));
break;
case HIP_API_ID_hipMalloc:
SPRINT("ptr(%p) size(0x%x)",
data->args.hipMalloc.ptr,
(uint32_t)(data->args.hipMalloc.size));
SPRINT("ptr(%p) size(0x%x)", data->args.hipMalloc.ptr,
(uint32_t)(data->args.hipMalloc.size));
break;
case HIP_API_ID_hipFree:
SPRINT("ptr(%p)", data->args.hipFree.ptr);
break;
case HIP_API_ID_hipModuleLaunchKernel:
SPRINT("kernel(\"%s\") stream(%p)",
hipKernelNameRef(data->args.hipModuleLaunchKernel.f),
data->args.hipModuleLaunchKernel.stream);
SPRINT("kernel(\"%s\") stream(%p)", hipKernelNameRef(data->args.hipModuleLaunchKernel.f),
data->args.hipModuleLaunchKernel.stream);
break;
default:
break;
@@ -316,26 +312,17 @@ void activity_callback(const char* begin, const char* end, void* arg) {
SPRINT("\tActivity records:\n");
while (record < end_record) {
const char * name = roctracer_op_string(record->domain, record->op, record->kind);
SPRINT("\t%s\tcorrelation_id(%lu) time_ns(%lu:%lu)",
name,
record->correlation_id,
record->begin_ns,
record->end_ns);
const char* name = roctracer_op_string(record->domain, record->op, record->kind);
SPRINT("\t%s\tcorrelation_id(%lu) time_ns(%lu:%lu)", name, record->correlation_id,
record->begin_ns, record->end_ns);
if (record->domain == ACTIVITY_DOMAIN_HIP_API) {
SPRINT(" process_id(%u) thread_id(%u)",
record->process_id,
record->thread_id);
SPRINT(" process_id(%u) thread_id(%u)", record->process_id, record->thread_id);
} else if (record->domain == ACTIVITY_DOMAIN_HCC_OPS) {
SPRINT(" device_id(%d) queue_id(%lu)",
record->device_id,
record->queue_id);
SPRINT(" device_id(%d) queue_id(%lu)", record->device_id, record->queue_id);
if (record->op == HIP_OP_ID_COPY) SPRINT(" bytes(0x%zx)", record->bytes);
} else if (record->domain == ACTIVITY_DOMAIN_HSA_OPS) {
SPRINT(" se(%u) cycle(%lu) pc(%lx)",
record->pc_sample.se,
record->pc_sample.cycle,
record->pc_sample.pc);
SPRINT(" se(%u) cycle(%lu) pc(%lx)", record->pc_sample.se, record->pc_sample.cycle,
record->pc_sample.pc);
} else if (record->domain == ACTIVITY_DOMAIN_EXT_API) {
SPRINT(" external_id(%lu)", record->external_id);
} else {
@@ -377,8 +364,10 @@ void init_tracing() {
void start_tracing() {
printf("# START (%d) #############################\n", iterations);
// Start
if ((iterations & 1) == 1) roctracer_start();
else roctracer_stop();
if ((iterations & 1) == 1)
roctracer_start();
else
roctracer_stop();
}
// Stop tracing routine
test/app/codeobj_test.cpp
+10 -12
Dosyayı Görüntüle
@@ -43,19 +43,18 @@ void check_status(roctracer_status_t status) {
void codeobj_callback(uint32_t domain, uint32_t cid, const void* data, void* arg) {
const hsa_evt_data_t* evt_data = reinterpret_cast<const hsa_evt_data_t*>(data);
const char* uri = evt_data->codeobj.uri;
printf("codeobj_callback domain(%u) cid(%u): load_base(0x%lx) load_size(0x%lx) load_delta(0x%lx) uri(\"%s\")\n",
domain,
cid,
evt_data->codeobj.load_base,
evt_data->codeobj.load_size,
evt_data->codeobj.load_delta,
uri);
printf(
"codeobj_callback domain(%u) cid(%u): load_base(0x%lx) load_size(0x%lx) load_delta(0x%lx) "
"uri(\"%s\")\n",
domain, cid, evt_data->codeobj.load_base, evt_data->codeobj.load_size,
evt_data->codeobj.load_delta, uri);
free((void*)uri);
fflush(stdout);
}
void initialize() {
roctracer_status_t status = roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, codeobj_callback, NULL);
roctracer_status_t status = roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, codeobj_callback, NULL);
check_status(status);
}
@@ -79,9 +78,8 @@ extern "C" PUBLIC_API void OnUnloadTool() {
}
extern "C" CONSTRUCTOR_API void constructor() {
printf("constructor\n"); fflush(stdout);
printf("constructor\n");
fflush(stdout);
}
extern "C" DESTRUCTOR_API void destructor() {
OnUnloadTool();
}
extern "C" DESTRUCTOR_API void destructor() { OnUnloadTool(); }
test/app/hsaco_test.cpp
+44 -57
Dosyayı Görüntüle
@@ -28,25 +28,24 @@
#define CONSTRUCTOR_API __attribute__((constructor))
#define DESTRUCTOR_API __attribute__((destructor))
#define HSA_RT(call) \
do { \
const hsa_status_t status = call; \
if (status != HSA_STATUS_SUCCESS) { \
printf("error \"%s\"\n", #call); fflush(stdout); \
abort(); \
} \
} while(0)
#define HSA_RT(call) \
do { \
const hsa_status_t status = call; \
if (status != HSA_STATUS_SUCCESS) { \
printf("error \"%s\"\n", #call); \
fflush(stdout); \
abort(); \
} \
} while (0)
// HSA API intercepting primitives
decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
hsa_ven_amd_loader_1_01_pfn_t loader_api_table{};
hsa_status_t code_object_callback(
hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object,
void* arg)
{
printf("code_object_callback\n"); fflush(stdout);
hsa_status_t code_object_callback(hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object, void* arg) {
printf("code_object_callback\n");
fflush(stdout);
uint64_t load_base = 0;
uint64_t load_size = 0;
@@ -55,21 +54,13 @@ hsa_status_t code_object_callback(
char* uri_str = NULL;
HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
&load_base));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE, &load_base));
HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
&load_size));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, &load_size));
HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
&load_delta));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA, &load_delta));
HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
&uri_len));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH, &uri_len));
uri_str = (char*)calloc(uri_len + 1, sizeof(char));
if (!uri_str) {
@@ -78,63 +69,59 @@ hsa_status_t code_object_callback(
}
HSA_RT(loader_api_table.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI,
uri_str));
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, uri_str));
printf("load_base(0x%lx)\n", load_base); fflush(stdout);
printf("load_size(0x%lx)\n", load_size); fflush(stdout);
printf("load_delta(0x%lx)\n", load_delta); fflush(stdout);
printf("uri_len(%u)\n", uri_len); fflush(stdout);
printf("uri_str(\"%s\")\n", uri_str); fflush(stdout);
printf("load_base(0x%lx)\n", load_base);
fflush(stdout);
printf("load_size(0x%lx)\n", load_size);
fflush(stdout);
printf("load_delta(0x%lx)\n", load_delta);
fflush(stdout);
printf("uri_len(%u)\n", uri_len);
fflush(stdout);
printf("uri_str(\"%s\")\n", uri_str);
fflush(stdout);
free(uri_str);
return HSA_STATUS_SUCCESS;
}
hsa_status_t hsa_executable_freeze_interceptor(
hsa_executable_t executable,
const char *options)
{
hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable, const char* options) {
HSA_RT(loader_api_table.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
executable,
code_object_callback,
NULL));
HSA_RT(hsa_executable_freeze_fn(
executable,
options));
executable, code_object_callback, NULL));
HSA_RT(hsa_executable_freeze_fn(executable, options));
return HSA_STATUS_SUCCESS;
}
// HSA-runtime tool on-load method
extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table,
uint64_t runtime_version,
extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
uint64_t failed_tool_count,
const char* const* failed_tool_names)
{
printf("OnLoad: begin\n"); fflush(stdout);
const char* const* failed_tool_names) {
printf("OnLoad: begin\n");
fflush(stdout);
// intercepting hsa_executable_freeze API
hsa_executable_freeze_fn = table->core_->hsa_executable_freeze_fn;
table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
// Fetching AMD Loader HSA extension API
HSA_RT(hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER,
1,
sizeof(hsa_ven_amd_loader_1_01_pfn_t),
&loader_api_table));
printf("OnLoad: end\n"); fflush(stdout);
HSA_EXTENSION_AMD_LOADER, 1, sizeof(hsa_ven_amd_loader_1_01_pfn_t), &loader_api_table));
printf("OnLoad: end\n");
fflush(stdout);
return true;
}
extern "C" PUBLIC_API void OnUnload() {
printf("OnUnload\n"); fflush(stdout);
printf("OnUnload\n");
fflush(stdout);
}
extern "C" CONSTRUCTOR_API void constructor() {
printf("constructor\n"); fflush(stdout);
printf("constructor\n");
fflush(stdout);
}
extern "C" DESTRUCTOR_API void destructor() {
printf("destructor\n"); fflush(stdout);
printf("destructor\n");
fflush(stdout);
}
test/hsa/src/hsa_rsrc_factory.cpp
+99 -57
Dosyayı Görüntüle
@@ -67,13 +67,15 @@ static hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
if (HSA_AMD_SEGMENT_GLOBAL != segment) {
return HSA_STATUS_SUCCESS;
}
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
@@ -126,19 +128,22 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize
#ifdef ROCP_LD_AQLPROFILE
status = LoadAqlProfileLib(&aqlprofile_api_);
#else
status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, hsa_ven_amd_aqlprofile_VERSION_MAJOR, sizeof(aqlprofile_api_), &aqlprofile_api_);
status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE,
hsa_ven_amd_aqlprofile_VERSION_MAJOR,
sizeof(aqlprofile_api_), &aqlprofile_api_);
#endif
CHECK_STATUS("aqlprofile API table load failed", status);
// Get Loader API table
loader_api_ = {0};
status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(loader_api_), &loader_api_);
status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1,
sizeof(loader_api_), &loader_api_);
CHECK_STATUS("loader API table query failed", status);
// Instantiate HSA timer
timer_ = new HsaTimer(&hsa_api_);
CHECK_STATUS("HSA timer allocation failed",
(timer_ == NULL) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS);
(timer_ == NULL) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS);
// Time correlation
const uint32_t corr_iters = 1000;
@@ -146,7 +151,8 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize
CorrelateTime(HsaTimer::TIME_ID_CLOCK_MONOTONIC, corr_iters);
// System timeout
timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX) ? timeout_ns_ : timer_->ns_to_sysclock(timeout_ns_);
timeout_ =
(timeout_ns_ == HsaTimer::TIMESTAMP_MAX) ? timeout_ns_ : timer_->ns_to_sysclock(timeout_ns_);
}
// Destructor of the class
@@ -172,9 +178,12 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) {
hsa_api_.hsa_queue_create = table->core_->hsa_queue_create_fn;
hsa_api_.hsa_queue_destroy = table->core_->hsa_queue_destroy_fn;
hsa_api_.hsa_queue_load_write_index_relaxed = table->core_->hsa_queue_load_write_index_relaxed_fn;
hsa_api_.hsa_queue_store_write_index_relaxed = table->core_->hsa_queue_store_write_index_relaxed_fn;
hsa_api_.hsa_queue_load_read_index_relaxed = table->core_->hsa_queue_load_read_index_relaxed_fn;
hsa_api_.hsa_queue_load_write_index_relaxed =
table->core_->hsa_queue_load_write_index_relaxed_fn;
hsa_api_.hsa_queue_store_write_index_relaxed =
table->core_->hsa_queue_store_write_index_relaxed_fn;
hsa_api_.hsa_queue_load_read_index_relaxed =
table->core_->hsa_queue_load_read_index_relaxed_fn;
hsa_api_.hsa_signal_create = table->core_->hsa_signal_create_fn;
hsa_api_.hsa_signal_destroy = table->core_->hsa_signal_destroy_fn;
@@ -183,27 +192,34 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) {
hsa_api_.hsa_signal_wait_scacquire = table->core_->hsa_signal_wait_scacquire_fn;
hsa_api_.hsa_signal_store_screlease = table->core_->hsa_signal_store_screlease_fn;
hsa_api_.hsa_code_object_reader_create_from_file = table->core_->hsa_code_object_reader_create_from_file_fn;
hsa_api_.hsa_code_object_reader_create_from_file =
table->core_->hsa_code_object_reader_create_from_file_fn;
hsa_api_.hsa_executable_create_alt = table->core_->hsa_executable_create_alt_fn;
hsa_api_.hsa_executable_load_agent_code_object = table->core_->hsa_executable_load_agent_code_object_fn;
hsa_api_.hsa_executable_load_agent_code_object =
table->core_->hsa_executable_load_agent_code_object_fn;
hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn;
hsa_api_.hsa_executable_get_symbol = table->core_->hsa_executable_get_symbol_fn;
hsa_api_.hsa_executable_symbol_get_info = table->core_->hsa_executable_symbol_get_info_fn;
hsa_api_.hsa_executable_iterate_symbols = table->core_->hsa_executable_iterate_symbols_fn;
hsa_api_.hsa_system_get_info = table->core_->hsa_system_get_info_fn;
hsa_api_.hsa_system_get_major_extension_table = table->core_->hsa_system_get_major_extension_table_fn;
hsa_api_.hsa_system_get_major_extension_table =
table->core_->hsa_system_get_major_extension_table_fn;
hsa_api_.hsa_amd_agent_iterate_memory_pools = table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn;
hsa_api_.hsa_amd_agent_iterate_memory_pools =
table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn;
hsa_api_.hsa_amd_memory_pool_get_info = table->amd_ext_->hsa_amd_memory_pool_get_info_fn;
hsa_api_.hsa_amd_memory_pool_allocate = table->amd_ext_->hsa_amd_memory_pool_allocate_fn;
hsa_api_.hsa_amd_agents_allow_access = table->amd_ext_->hsa_amd_agents_allow_access_fn;
hsa_api_.hsa_amd_memory_async_copy = table->amd_ext_->hsa_amd_memory_async_copy_fn;
hsa_api_.hsa_amd_signal_async_handler = table->amd_ext_->hsa_amd_signal_async_handler_fn;
hsa_api_.hsa_amd_profiling_set_profiler_enabled = table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn;
hsa_api_.hsa_amd_profiling_get_async_copy_time = table->amd_ext_->hsa_amd_profiling_get_async_copy_time_fn;
hsa_api_.hsa_amd_profiling_get_dispatch_time = table->amd_ext_->hsa_amd_profiling_get_dispatch_time_fn;
hsa_api_.hsa_amd_profiling_set_profiler_enabled =
table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn;
hsa_api_.hsa_amd_profiling_get_async_copy_time =
table->amd_ext_->hsa_amd_profiling_get_async_copy_time_fn;
hsa_api_.hsa_amd_profiling_get_dispatch_time =
table->amd_ext_->hsa_amd_profiling_get_dispatch_time_fn;
} else {
hsa_api_.hsa_init = hsa_init;
hsa_api_.hsa_shut_down = hsa_shut_down;
@@ -298,10 +314,13 @@ const AgentInfo* HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent) {
agent_info->dev_type = HSA_DEVICE_TYPE_CPU;
agent_info->dev_index = cpu_list_.size();
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->cpu_pool);
status =
hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->cpu_pool);
if ((status == HSA_STATUS_INFO_BREAK) && (cpu_pool_ == NULL)) cpu_pool_ = &agent_info->cpu_pool;
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindKernArgPool, &agent_info->kern_arg_pool);
if ((status == HSA_STATUS_INFO_BREAK) && (kern_arg_pool_ == NULL)) kern_arg_pool_ = &agent_info->kern_arg_pool;
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindKernArgPool,
&agent_info->kern_arg_pool);
if ((status == HSA_STATUS_INFO_BREAK) && (kern_arg_pool_ == NULL))
kern_arg_pool_ = &agent_info->kern_arg_pool;
agent_info->gpu_pool = {};
cpu_list_.push_back(agent_info);
@@ -319,21 +338,26 @@ const AgentInfo* HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent) {
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &agent_info->max_queue_size);
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
agent_info->is_apu = (agent_info->profile == HSA_PROFILE_FULL) ? true : false;
hsa_api_.hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT),
&agent_info->cu_num);
hsa_api_.hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU),
&agent_info->waves_per_cu);
hsa_api_.hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU),
&agent_info->simds_per_cu);
hsa_api_.hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES),
&agent_info->se_num);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT),
&agent_info->cu_num);
hsa_api_.hsa_agent_get_info(agent,
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE),
&agent_info->shader_arrays_per_se);
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU),
&agent_info->waves_per_cu);
hsa_api_.hsa_agent_get_info(agent,
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU),
&agent_info->simds_per_cu);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES),
&agent_info->se_num);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE),
&agent_info->shader_arrays_per_se);
agent_info->cpu_pool = {};
agent_info->kern_arg_pool = {};
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->gpu_pool);
status =
hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool, &agent_info->gpu_pool);
CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(gpu pool)", status);
// GFX8 and GFX9 SGPR/VGPR block sizes
@@ -430,7 +454,7 @@ bool HsaRsrcFactory::CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts,
hsa_queue_t** queue) {
hsa_status_t status;
status = hsa_api_.hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, NULL,
UINT32_MAX, UINT32_MAX, queue);
UINT32_MAX, UINT32_MAX, queue);
return (status == HSA_STATUS_SUCCESS);
}
@@ -453,7 +477,8 @@ uint8_t* HsaRsrcFactory::AllocateLocalMemory(const AgentInfo* agent_info, size_t
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsa_api_.hsa_amd_memory_pool_allocate(agent_info->gpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
status = hsa_api_.hsa_amd_memory_pool_allocate(agent_info->gpu_pool, size, 0,
reinterpret_cast<void**>(&buffer));
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
return ptr;
}
@@ -468,7 +493,8 @@ uint8_t* HsaRsrcFactory::AllocateKernArgMemory(const AgentInfo* agent_info, size
uint8_t* buffer = NULL;
if (!cpu_agents_.empty()) {
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsa_api_.hsa_amd_memory_pool_allocate(*kern_arg_pool_, size, 0, reinterpret_cast<void**>(&buffer));
status = hsa_api_.hsa_amd_memory_pool_allocate(*kern_arg_pool_, size, 0,
reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the kernel arguments
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {agent_info->dev_id};
@@ -488,7 +514,8 @@ uint8_t* HsaRsrcFactory::AllocateSysMemory(const AgentInfo* agent_info, size_t s
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
if (!cpu_agents_.empty()) {
status = hsa_api_.hsa_amd_memory_pool_allocate(*cpu_pool_, size, 0, reinterpret_cast<void**>(&buffer));
status = hsa_api_.hsa_amd_memory_pool_allocate(*cpu_pool_, size, 0,
reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {agent_info->dev_id};
@@ -513,16 +540,18 @@ uint8_t* HsaRsrcFactory::AllocateCmdMemory(const AgentInfo* agent_info, size_t s
}
// Wait signal
hsa_signal_value_t HsaRsrcFactory::SignalWait(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const {
hsa_signal_value_t HsaRsrcFactory::SignalWait(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const {
const hsa_signal_value_t exp_value = signal_value - 1;
hsa_signal_value_t ret_value = signal_value;
while (1) {
ret_value =
hsa_api_.hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, signal_value, timeout_, HSA_WAIT_STATE_BLOCKED);
ret_value = hsa_api_.hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, signal_value,
timeout_, HSA_WAIT_STATE_BLOCKED);
if (ret_value == exp_value) break;
if (ret_value != signal_value) {
std::cerr << "Error: HsaRsrcFactory::SignalWait: signal_value(" << signal_value
<< "), ret_value(" << ret_value << ")" << std::endl << std::flush;
<< "), ret_value(" << ret_value << ")" << std::endl
<< std::flush;
abort();
}
}
@@ -530,7 +559,8 @@ hsa_signal_value_t HsaRsrcFactory::SignalWait(const hsa_signal_t& signal, const
}
// Wait signal with signal value restore
void HsaRsrcFactory::SignalWaitRestore(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const {
void HsaRsrcFactory::SignalWaitRestore(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const {
SignalWait(signal, signal_value);
hsa_api_.hsa_signal_store_relaxed(const_cast<hsa_signal_t&>(signal), signal_value);
}
@@ -594,13 +624,13 @@ bool HsaRsrcFactory::LoadAndFinalize(const AgentInfo* agent_info, const char* br
}
// Create executable.
status = hsa_api_.hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
NULL, executable);
status = hsa_api_.hsa_executable_create_alt(
HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, NULL, executable);
CHECK_STATUS("Error in creating executable object", status);
// Load code object.
status = hsa_api_.hsa_executable_load_agent_code_object(*executable, agent_info->dev_id, code_obj_rdr,
NULL, NULL);
status = hsa_api_.hsa_executable_load_agent_code_object(*executable, agent_info->dev_id,
code_obj_rdr, NULL, NULL);
CHECK_STATUS("Error in loading executable object", status);
// Freeze executable.
@@ -610,7 +640,7 @@ bool HsaRsrcFactory::LoadAndFinalize(const AgentInfo* agent_info, const char* br
// Get symbol handle.
hsa_executable_symbol_t kernelSymbol;
status = hsa_api_.hsa_executable_get_symbol(*executable, NULL, kernel_name, agent_info->dev_id, 0,
&kernelSymbol);
&kernelSymbol);
CHECK_STATUS("Error in looking up kernel symbol", status);
// Update output parameter
@@ -654,7 +684,8 @@ uint64_t HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet) {
}
uint32_t slot_idx = (uint32_t)(write_idx % queue->size);
uint32_t* queue_slot = reinterpret_cast<uint32_t*>((uintptr_t)(queue->base_address) + (slot_idx * slot_size_b));
uint32_t* queue_slot =
reinterpret_cast<uint32_t*>((uintptr_t)(queue->base_address) + (slot_idx * slot_size_b));
const uint32_t* slot_data = reinterpret_cast<const uint32_t*>(packet);
// Copy buffered commands into the queue slot.
@@ -704,18 +735,22 @@ void HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table) {
table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
}
hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data) {
hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec,
hsa_executable_symbol_t symbol, void* data) {
hsa_symbol_kind_t value = (hsa_symbol_kind_t)0;
hsa_status_t status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &value);
hsa_status_t status =
hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &value);
CHECK_STATUS("Error in getting symbol info", status);
if (value == HSA_SYMBOL_KIND_KERNEL) {
uint64_t addr = 0;
uint32_t len = 0;
status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &addr);
status = hsa_api_.hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &addr);
CHECK_STATUS("Error in getting kernel object", status);
status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len);
status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH,
&len);
CHECK_STATUS("Error in getting name len", status);
char *name = new char[len + 1];
char* name = new char[len + 1];
status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
CHECK_STATUS("Error in getting kernel name", status);
name[len] = 0;
@@ -728,12 +763,15 @@ hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_ex
return HSA_STATUS_SUCCESS;
}
hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options) {
hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable,
const char* options) {
std::lock_guard<mutex_t> lck(mutex_);
if (symbols_map_ == NULL) symbols_map_ = new symbols_map_t;
hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL);
hsa_status_t status =
hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL);
CHECK_STATUS("Error in iterating executable symbols", status);
return hsa_api_.hsa_executable_freeze(executable, options);;
return hsa_api_.hsa_executable_freeze(executable, options);
;
}
void HsaRsrcFactory::DumpHandles(FILE* file) {
@@ -741,10 +779,14 @@ void HsaRsrcFactory::DumpHandles(FILE* file) {
auto end = agent_map_.end();
for (auto it = beg; it != end; ++it) {
const AgentInfo* agent_info = it->second;
fprintf(file, "0x%lx agent %s\n", agent_info->dev_id.handle, (agent_info->dev_type == HSA_DEVICE_TYPE_CPU) ? "cpu" : "gpu");
if (agent_info->cpu_pool.handle != 0) fprintf(file, "0x%lx pool cpu\n", agent_info->cpu_pool.handle);
if (agent_info->kern_arg_pool.handle != 0) fprintf(file, "0x%lx pool cpu kernarg\n", agent_info->kern_arg_pool.handle);
if (agent_info->gpu_pool.handle != 0) fprintf(file, "0x%lx pool gpu\n", agent_info->gpu_pool.handle);
fprintf(file, "0x%lx agent %s\n", agent_info->dev_id.handle,
(agent_info->dev_type == HSA_DEVICE_TYPE_CPU) ? "cpu" : "gpu");
if (agent_info->cpu_pool.handle != 0)
fprintf(file, "0x%lx pool cpu\n", agent_info->cpu_pool.handle);
if (agent_info->kern_arg_pool.handle != 0)
fprintf(file, "0x%lx pool cpu kernarg\n", agent_info->kern_arg_pool.handle);
if (agent_info->gpu_pool.handle != 0)
fprintf(file, "0x%lx pool gpu\n", agent_info->gpu_pool.handle);
}
fflush(file);
}
test/hsa/src/hsa_rsrc_factory.h
+34 -30
Dosyayı Görüntüle
@@ -44,23 +44,25 @@
#define HSA_QUEUE_ALIGN_BYTES 64
#define HSA_PACKET_ALIGN_BYTES 64
#define CHECK_STATUS(msg, status) do { \
if ((status) != HSA_STATUS_SUCCESS) { \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
abort(); \
} \
} while (0)
#define CHECK_STATUS(msg, status) \
do { \
if ((status) != HSA_STATUS_SUCCESS) { \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
abort(); \
} \
} while (0)
#define CHECK_ITER_STATUS(msg, status) do { \
if ((status) != HSA_STATUS_INFO_BREAK) { \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
abort(); \
} \
} while (0)
#define CHECK_ITER_STATUS(msg, status) \
do { \
if ((status) != HSA_STATUS_INFO_BREAK) { \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
abort(); \
} \
} while (0)
static const size_t MEM_PAGE_BYTES = 0x1000;
static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
@@ -172,15 +174,12 @@ class HsaTimer {
static const timestamp_t TIMESTAMP_MAX = UINT64_MAX;
typedef long double freq_t;
enum time_id_t {
TIME_ID_CLOCK_REALTIME = 0,
TIME_ID_CLOCK_MONOTONIC = 1,
TIME_ID_NUMBER
};
enum time_id_t { TIME_ID_CLOCK_REALTIME = 0, TIME_ID_CLOCK_MONOTONIC = 1, TIME_ID_NUMBER };
HsaTimer(const hsa_pfn_t* hsa_api) : hsa_api_(hsa_api) {
timestamp_t sysclock_hz = 0;
hsa_status_t status = hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
hsa_status_t status =
hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY)", status);
sysclock_factor_ = (freq_t)1000000000 / (freq_t)sysclock_hz;
}
@@ -215,8 +214,8 @@ class HsaTimer {
// Return pair of correlated values of profiling timestamp and time with
// correlation error for a given time ID and number of iterations
void correlated_pair_ns(time_id_t time_id, uint32_t iters,
timestamp_t* timestamp_v, timestamp_t* time_v, timestamp_t* error_v) const {
void correlated_pair_ns(time_id_t time_id, uint32_t iters, timestamp_t* timestamp_v,
timestamp_t* time_v, timestamp_t* error_v) const {
clockid_t clock_id = 0;
switch (clock_id) {
case TIME_ID_CLOCK_REALTIME:
@@ -355,7 +354,8 @@ class HsaRsrcFactory {
uint8_t* AllocateCmdMemory(const AgentInfo* agent_info, size_t size);
// Wait signal
hsa_signal_value_t SignalWait(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const;
hsa_signal_value_t SignalWait(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const;
// Wait signal with signal value restore
void SignalWaitRestore(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const;
@@ -401,7 +401,9 @@ class HsaRsrcFactory {
const hsa_ven_amd_loader_1_00_pfn_t* LoaderApi() const { return &loader_api_; }
// Methods for system-clock/ns conversion and timestamp in 'ns'
timestamp_t SysclockToNs(const timestamp_t& sysclock) const { return timer_->sysclock_to_ns(sysclock); }
timestamp_t SysclockToNs(const timestamp_t& sysclock) const {
return timer_->sysclock_to_ns(sysclock);
}
timestamp_t NsToSysclock(const timestamp_t& time) const { return timer_->ns_to_sysclock(time); }
timestamp_t TimestampNs() const { return timer_->timestamp_ns(); }
@@ -480,8 +482,10 @@ class HsaRsrcFactory {
typedef std::map<uint64_t, const char*> symbols_map_t;
static symbols_map_t* symbols_map_;
static bool executable_tracking_on_;
static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options);
static hsa_status_t executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data);
static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable,
const char* options);
static hsa_status_t executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol,
void* data);
// HSA runtime API table
static hsa_pfn_t hsa_api_;
@@ -505,8 +509,8 @@ class HsaRsrcFactory {
timestamp_t time_error_[HsaTimer::TIME_ID_NUMBER];
// CPU/kern-arg memory pools
hsa_amd_memory_pool_t *cpu_pool_;
hsa_amd_memory_pool_t *kern_arg_pool_;
hsa_amd_memory_pool_t* cpu_pool_;
hsa_amd_memory_pool_t* kern_arg_pool_;
};
#endif // _HSA_RSRC_FACTORY_H_
test/hsa/test/ctrl/run_kernel.h
+3 -1
Dosyayı Görüntüle
@@ -24,7 +24,9 @@
#include "ctrl/test_hsa.h"
#include "util/test_assert.h"
template <class Kernel, class Test> bool RunKernel(int argc = 0, char* argv[] = NULL, const AgentInfo* agent_info = NULL, hsa_queue_t* queue = NULL, int count = 1) {
template <class Kernel, class Test>
bool RunKernel(int argc = 0, char* argv[] = NULL, const AgentInfo* agent_info = NULL,
hsa_queue_t* queue = NULL, int count = 1) {
bool ret_val = false;
if (getenv("ROC_TEST_TRACE") == NULL) std::clog.rdbuf(NULL);
test/hsa/test/ctrl/test_hsa.cpp
+6 -4
Dosyayı Görüntüle
@@ -63,7 +63,8 @@ bool TestHsa::Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
return false;
}
}
std::clog << "> Using agent[" << agent_info_->dev_index << "] : " << agent_info_->name << std::endl;
std::clog << "> Using agent[" << agent_info_->dev_index << "] : " << agent_info_->name
<< std::endl;
// Create an instance of Aql Queue
if (hsa_queue_ == NULL) {
@@ -116,8 +117,8 @@ bool TestHsa::Setup() {
size_t size_info = 0;
const hsa_status_t status = hsa_executable_symbol_get_info(
kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &size_info);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
size_info = kernarg_size;
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
size_info = kernarg_size;
const bool kernarg_missmatch = (kernarg_size > size_info);
if (kernarg_missmatch) {
std::cout << "kernarg_size = " << kernarg_size << ", size_info = " << size_info
@@ -209,7 +210,8 @@ bool TestHsa::Run() {
// Submit AQL packet to the queue
const uint64_t que_idx = hsa_rsrc_->Submit(hsa_queue_, &aql);
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl << std::flush;
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl
<< std::flush;
// Wait on the dispatch signal until the kernel is finished.
// Update wait condition to HSA_WAIT_STATE_ACTIVE for Polling
test/hsa/test/dummy_kernel/dummy_kernel.h
+4 -5
Dosyayı Görüntüle
@@ -33,10 +33,7 @@ class DummyKernel : public TestKernel {
enum { KERNARG_BUF_ID, LOCAL_BUF_ID };
// Constructor
DummyKernel() :
width_(64),
height_(64)
{
DummyKernel() : width_(64), height_(64) {
SetInDescr(KERNARG_BUF_ID, KERNARG_DES_ID, 0);
SetOutDescr(LOCAL_BUF_ID, LOCAL_DES_ID, 0);
}
@@ -57,7 +54,9 @@ class DummyKernel : public TestKernel {
// Reference CPU implementation
bool ReferenceImplementation(uint32_t* output, const uint32_t* input, const float* mask,
const uint32_t width, const uint32_t height,
const uint32_t maskWidth, const uint32_t maskHeight) { return true; }
const uint32_t maskWidth, const uint32_t maskHeight) {
return true;
}
// Width of the Input array
const uint32_t width_;
test/hsa/test/util/evt_stats.h
+13 -16
Dosyayı Görüntüle
@@ -28,9 +28,8 @@
#include <sstream>
#include <utility>
template <class evt_id_t, class evt_weight_t>
class EvtStatsT {
public:
template <class evt_id_t, class evt_weight_t> class EvtStatsT {
public:
typedef std::mutex mutex_t;
typedef uint64_t evt_count_t;
typedef double evt_avr_t;
@@ -51,7 +50,7 @@ class EvtStatsT {
inline void add_event(evt_id_t id, evt_weight_t weight) {
std::lock_guard<mutex_t> lck(mutex_);
//printf("EvtStats %p ::add_event %u %lu\n", this, id, weight); fflush(stdout);
// printf("EvtStats %p ::add_event %u %lu\n", this, id, weight); fflush(stdout);
evt_record_t& rec = map_[id];
const evt_count_t prev_count = rec.count;
@@ -65,7 +64,8 @@ class EvtStatsT {
void dump() {
std::lock_guard<mutex_t> lck(mutex_);
fprintf(stdout, "Dumping %s\n", path_); fflush(stdout);
fprintf(stdout, "Dumping %s\n", path_);
fflush(stdout);
typedef typename std::set<std::pair<evt_id_t, evt_record_t>, cmpfun> set_t;
set_t s_(map_.begin(), map_.end());
@@ -75,7 +75,8 @@ class EvtStatsT {
const evt_id_t id = e.first;
const char* label = get_label(id);
std::ostringstream oss;
oss << index << ",\"" << label << "\"," << e.second.count << "," << (uint64_t)(e.second.avr) << "," << (uint64_t)(e.second.count * e.second.avr);
oss << index << ",\"" << label << "\"," << e.second.count << "," << (uint64_t)(e.second.avr)
<< "," << (uint64_t)(e.second.count * e.second.avr);
fprintf(fdes_, "%s\n", oss.str().c_str());
index += 1;
}
@@ -88,24 +89,20 @@ class EvtStatsT {
const char* label = ret.first->second;
return label;
}
const char* get_label(const char* id) {
return id;
}
const char* get_label(const std::string& id) {
return id.c_str();
}
const char* get_label(const char* id) { return id; }
const char* get_label(const std::string& id) { return id.c_str(); }
void set_label(evt_id_t id, const char* label) {
//printf("EvtStats %p ::set_label %u %s\n", this, id, label); fflush(stdout);
// printf("EvtStats %p ::set_label %u %s\n", this, id, label); fflush(stdout);
labels_[id] = label;
}
EvtStatsT(FILE* f, const char* path) : fdes_(f), path_(path) {
//printf("EvtStats %p ::EvtStatsT()\n", this); fflush(stdout);
// printf("EvtStats %p ::EvtStatsT()\n", this); fflush(stdout);
fprintf(fdes_, "Index,Name,Count,Avr,Total\n");
}
private:
private:
mutex_t mutex_;
map_t map_;
labels_t labels_;
@@ -115,4 +112,4 @@ class EvtStatsT {
typedef EvtStatsT<uint32_t, uint64_t> EvtStats;
#endif // EVT_STATS_H_
#endif // EVT_STATS_H_
test/hsa/test/util/xml.h
+4 -2
Dosyayı Görüntüle
@@ -212,7 +212,8 @@ class Xml {
buf[size - 1] = '\0';
if (strncmp(buf, "#include \"", 10) == 0) {
for (ind = 0; (ind < size) && (buf[ind] != '\n'); ++ind) {}
for (ind = 0; (ind < size) && (buf[ind] != '\n'); ++ind) {
}
if (ind == size) {
fprintf(stderr, "XML PreProcess failed, line size limit %zu\n", kBufSize);
error = true;
@@ -222,7 +223,8 @@ class Xml {
size = ind;
lseek(fd_, pos + ind + 1, SEEK_SET);
for (ind = 10; (ind < size) && (buf[ind] != '"'); ++ind) {}
for (ind = 10; (ind < size) && (buf[ind] != '"'); ++ind) {
}
if (ind == size) {
error = true;
break;
test/tool/tracer_tool.cpp
+211 -191
Dosyayı Görüntüle
@@ -21,15 +21,15 @@
#include <sstream>
#include <string>
#include <cxxabi.h> /* names denangle */
#include <cxxabi.h> /* names denangle */
#include <dirent.h>
#include <pthread.h>
#include <stdarg.h>
#include <stdio.h>
#include <string.h>
#include <sys/syscall.h> /* SYS_xxx definitions */
#include <sys/syscall.h> /* SYS_xxx definitions */
#include <sys/types.h>
#include <unistd.h> /* usleep */
#include <unistd.h> /* usleep */
#include <roctracer_ext.h>
#include "src/util/exception.h"
@@ -70,10 +70,12 @@
} \
} while (0)
#define ONLOAD_TRACE(str) \
if (getenv("ROCP_ONLOAD_TRACE")) do { \
std::cout << "PID(" << GetPid() << "): TRACER_TOOL::" << __FUNCTION__ << " " << str << std::endl << std::flush; \
} while(0);
#define ONLOAD_TRACE(str) \
if (getenv("ROCP_ONLOAD_TRACE")) do { \
std::cout << "PID(" << GetPid() << "): TRACER_TOOL::" << __FUNCTION__ << " " << str \
<< std::endl \
<< std::flush; \
} while (0);
#define ONLOAD_TRACE_BEG() ONLOAD_TRACE("begin")
#define ONLOAD_TRACE_END() ONLOAD_TRACE("end")
@@ -88,7 +90,8 @@ inline static void DEBUG_TRACE(const char* fmt, ...) {
va_list valist;
va_start(valist, fmt);
vsnprintf(buf, size, fmt, valist);
printf("%u:%u %s", GetPid(), GetTid(), buf); fflush(stdout);
printf("%u:%u %s", GetPid(), GetTid(), buf);
fflush(stdout);
va_end(valist);
}
#else
@@ -107,18 +110,18 @@ bool trace_hip_activity = false;
bool trace_pcs = false;
// The below getter functions have been written intentionally to fix an issue
// with constructor ordering. Previously when hip_api_vec and hsa_api_vec
// with constructor ordering. Previously when hip_api_vec and hsa_api_vec
// were left as simple global variables, whenever the tool_load() function
// was called from "extern "C" CONSTRUCTOR_API void constructor()" of libtracer_tool.so
// the ordering of std::vector constructor becomes undefined. This meant that you could assign
// hip_api_vec and hsa_api_vec with a value in tool_load() and once the function returns, the std::vector
// default constructor would execute later, causing the values to be lost.
// hip_api_vec and hsa_api_vec with a value in tool_load() and once the function returns, the
// std::vector default constructor would execute later, causing the values to be lost.
static std::vector<std::string> &hsa_api_vec() {
static std::vector<std::string>& hsa_api_vec() {
static std::vector<std::string> hsa_api_vec;
return hsa_api_vec;
}
static std::vector<std::string> &hip_api_vec() {
static std::vector<std::string>& hip_api_vec() {
static std::vector<std::string> hip_api_vec;
return hip_api_vec;
}
@@ -167,7 +170,8 @@ void fatal(const std::string msg) {
static inline const char* cxx_demangle(const char* symbol) {
size_t funcnamesize;
int status;
const char* ret = (symbol != NULL) ? abi::__cxa_demangle(symbol, NULL, &funcnamesize, &status) : symbol;
const char* ret =
(symbol != NULL) ? abi::__cxa_demangle(symbol, NULL, &funcnamesize, &status) : symbol;
return (ret != NULL) ? ret : strdup(symbol);
}
@@ -208,7 +212,8 @@ void* control_thr_fun(void*) {
uint32_t control_flush_us = 0;
pthread_t flush_thread;
bool flush_thread_started = false;
std::mutex flush_thread_mutex;;
std::mutex flush_thread_mutex;
;
void* flush_thr_fun(void*) {
const uint32_t dist_sec = control_flush_us / 1000000;
@@ -218,7 +223,8 @@ void* flush_thr_fun(void*) {
sleep(dist_sec);
usleep(dist_us);
std::lock_guard<std::mutex> lock(flush_thread_mutex);
if (!flush_thread_started) while(1) sleep(1);
if (!flush_thread_started)
while (1) sleep(1);
ROCTRACER_CALL(roctracer_flush_activity());
roctracer::TraceBufferBase::FlushAll();
}
@@ -241,17 +247,13 @@ struct roctx_trace_entry_t {
};
void roctx_flush_cb(roctx_trace_entry_t* entry);
constexpr roctracer::TraceBuffer<roctx_trace_entry_t>::flush_prm_t roctx_flush_prm = {roctracer::DFLT_ENTRY_TYPE, roctx_flush_cb};
constexpr roctracer::TraceBuffer<roctx_trace_entry_t>::flush_prm_t roctx_flush_prm = {
roctracer::DFLT_ENTRY_TYPE, roctx_flush_cb};
roctracer::TraceBuffer<roctx_trace_entry_t>* roctx_trace_buffer = NULL;
// rocTX callback function
static inline void roctx_callback_fun(
uint32_t domain,
uint32_t cid,
uint32_t tid,
roctx_range_id_t rid,
const char* message)
{
static inline void roctx_callback_fun(uint32_t domain, uint32_t cid, uint32_t tid,
roctx_range_id_t rid, const char* message) {
#if ROCTX_CLOCK_TIME
const timestamp_t time = HsaTimer::clocktime_ns(HsaTimer::TIME_ID_CLOCK_MONOTONIC);
#else
@@ -267,12 +269,7 @@ static inline void roctx_callback_fun(
entry->valid.store(roctracer::TRACE_ENTRY_COMPL, std::memory_order_release);
}
void roctx_api_callback(
uint32_t domain,
uint32_t cid,
const void* callback_data,
void* arg)
{
void roctx_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) {
(void)arg;
const roctx_api_data_t* data = reinterpret_cast<const roctx_api_data_t*>(callback_data);
roctx_callback_fun(domain, cid, GetTid(), data->args.id, data->args.message);
@@ -280,27 +277,37 @@ void roctx_api_callback(
// rocTX Start/Stop callbacks
void roctx_range_start_callback(const roctx_range_data_t* data, void* arg) {
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, data->tid, 0, data->message);
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, data->tid, 0,
data->message);
}
void roctx_range_stop_callback(const roctx_range_data_t* data, void* arg) {
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePop, data->tid, 0, NULL);
}
void start_callback() { roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_start_callback, NULL); }
void stop_callback() { roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_stop_callback, NULL); }
void start_callback() {
roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_start_callback, NULL);
}
void stop_callback() {
roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_stop_callback, NULL);
}
// rocTX buffer flush function
void roctx_flush_cb(roctx_trace_entry_t* entry) {
#if ROCTX_CLOCK_TIME
timestamp_t timestamp = 0;
HsaRsrcFactory::Instance().GetTimestamp(HsaTimer::TIME_ID_CLOCK_MONOTONIC, entry->time, &timestamp);
HsaRsrcFactory::Instance().GetTimestamp(HsaTimer::TIME_ID_CLOCK_MONOTONIC, entry->time,
&timestamp);
#else
const timestamp_t timestamp = entry->time;
#endif
std::ostringstream os;
os << timestamp << " " << entry->pid << ":" << entry->tid << " " << entry->cid << ":" << entry->rid;
if (entry->message != NULL) os << ":\"" << entry->message << "\"";
else os << ":\"\"";
fprintf(roctx_file_handle, "%s\n", os.str().c_str()); fflush(roctx_file_handle);
os << timestamp << " " << entry->pid << ":" << entry->tid << " " << entry->cid << ":"
<< entry->rid;
if (entry->message != NULL)
os << ":\"" << entry->message << "\"";
else
os << ":\"\"";
fprintf(roctx_file_handle, "%s\n", os.str().c_str());
fflush(roctx_file_handle);
}
///////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -318,24 +325,20 @@ struct hsa_api_trace_entry_t {
};
void hsa_api_flush_cb(hsa_api_trace_entry_t* entry);
constexpr roctracer::TraceBuffer<hsa_api_trace_entry_t>::flush_prm_t hsa_flush_prm = {roctracer::DFLT_ENTRY_TYPE, hsa_api_flush_cb};
constexpr roctracer::TraceBuffer<hsa_api_trace_entry_t>::flush_prm_t hsa_flush_prm = {
roctracer::DFLT_ENTRY_TYPE, hsa_api_flush_cb};
roctracer::TraceBuffer<hsa_api_trace_entry_t>* hsa_api_trace_buffer = NULL;
// HSA API callback function
void hsa_api_callback(
uint32_t domain,
uint32_t cid,
const void* callback_data,
void* arg)
{
void hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) {
(void)arg;
const hsa_api_data_t* data = reinterpret_cast<const hsa_api_data_t*>(callback_data);
if (data->phase == ACTIVITY_API_PHASE_ENTER) {
hsa_begin_timestamp = timer->timestamp_fn_ns();
} else {
const timestamp_t end_timestamp = (cid == HSA_API_ID_hsa_shut_down) ? hsa_begin_timestamp : timer->timestamp_fn_ns();
const timestamp_t end_timestamp =
(cid == HSA_API_ID_hsa_shut_down) ? hsa_begin_timestamp : timer->timestamp_fn_ns();
hsa_api_trace_entry_t* entry = hsa_api_trace_buffer->GetEntry();
entry->cid = cid;
entry->begin = hsa_begin_timestamp;
@@ -349,17 +352,17 @@ void hsa_api_callback(
void hsa_api_flush_cb(hsa_api_trace_entry_t* entry) {
std::ostringstream os;
os << entry->begin << ":" << entry->end << " " << entry->pid << ":" << entry->tid << " " << hsa_api_data_pair_t(entry->cid, entry->data);
fprintf(hsa_api_file_handle, "%s\n", os.str().c_str()); fflush(hsa_api_file_handle);
os << entry->begin << ":" << entry->end << " " << entry->pid << ":" << entry->tid << " "
<< hsa_api_data_pair_t(entry->cid, entry->data);
fprintf(hsa_api_file_handle, "%s\n", os.str().c_str());
fflush(hsa_api_file_handle);
}
void hsa_activity_callback(
uint32_t op,
activity_record_t* record,
void* arg)
{
void hsa_activity_callback(uint32_t op, activity_record_t* record, void* arg) {
static uint64_t index = 0;
fprintf(hsa_async_copy_file_handle, "%lu:%lu async-copy:%lu:%u\n", record->begin_ns, record->end_ns, index, my_pid); fflush(hsa_async_copy_file_handle);
fprintf(hsa_async_copy_file_handle, "%lu:%lu async-copy:%lu:%u\n", record->begin_ns,
record->end_ns, index, my_pid);
fflush(hsa_async_copy_file_handle);
index++;
}
@@ -381,28 +384,21 @@ struct hip_api_trace_entry_t {
};
void hip_api_flush_cb(hip_api_trace_entry_t* entry);
constexpr roctracer::TraceBuffer<hip_api_trace_entry_t>::flush_prm_t hip_api_flush_prm = {roctracer::DFLT_ENTRY_TYPE, hip_api_flush_cb};
constexpr roctracer::TraceBuffer<hip_api_trace_entry_t>::flush_prm_t hip_api_flush_prm = {
roctracer::DFLT_ENTRY_TYPE, hip_api_flush_cb};
roctracer::TraceBuffer<hip_api_trace_entry_t>* hip_api_trace_buffer = NULL;
static inline bool is_hip_kernel_launch_api(const uint32_t& cid) {
bool ret =
(cid == HIP_API_ID_hipLaunchKernel) ||
(cid == HIP_API_ID_hipExtLaunchKernel) ||
(cid == HIP_API_ID_hipLaunchCooperativeKernel) ||
(cid == HIP_API_ID_hipLaunchCooperativeKernelMultiDevice) ||
(cid == HIP_API_ID_hipExtLaunchMultiKernelMultiDevice) ||
(cid == HIP_API_ID_hipModuleLaunchKernel) ||
(cid == HIP_API_ID_hipExtModuleLaunchKernel) ||
(cid == HIP_API_ID_hipHccModuleLaunchKernel);
bool ret = (cid == HIP_API_ID_hipLaunchKernel) || (cid == HIP_API_ID_hipExtLaunchKernel) ||
(cid == HIP_API_ID_hipLaunchCooperativeKernel) ||
(cid == HIP_API_ID_hipLaunchCooperativeKernelMultiDevice) ||
(cid == HIP_API_ID_hipExtLaunchMultiKernelMultiDevice) ||
(cid == HIP_API_ID_hipModuleLaunchKernel) || (cid == HIP_API_ID_hipExtModuleLaunchKernel) ||
(cid == HIP_API_ID_hipHccModuleLaunchKernel);
return ret;
}
void hip_api_callback(
uint32_t domain,
uint32_t cid,
const void* callback_data,
void* arg)
{
void hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) {
(void)arg;
const hip_api_data_t* data = reinterpret_cast<const hip_api_data_t*>(callback_data);
const timestamp_t timestamp = timer->timestamp_fn_ns();
@@ -428,16 +424,17 @@ void hip_api_callback(
if (cid == HIP_API_ID_hipMalloc) {
entry->ptr = *(data->args.hipMalloc.ptr);
} else if (is_hip_kernel_launch_api(cid)) {
switch(cid) {
switch (cid) {
case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice:
case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice:
{
const hipLaunchParams* listKernels = data->args.hipLaunchCooperativeKernelMultiDevice.launchParamsList;
case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: {
const hipLaunchParams* listKernels =
data->args.hipLaunchCooperativeKernelMultiDevice.launchParamsList;
std::string name_str = "";
for (int i = 0; i < data->args.hipLaunchCooperativeKernelMultiDevice.numDevices; ++i) {
const hipLaunchParams& lp = listKernels[i];
if (lp.func != NULL) {
const char* kernel_name = roctracer::HipLoader::Instance().KernelNameRefByPtr(lp.func, lp.stream);
const char* kernel_name =
roctracer::HipLoader::Instance().KernelNameRefByPtr(lp.func, lp.stream);
const int device_id = roctracer::HipLoader::Instance().GetStreamDeviceId(lp.stream);
name_str += std::string(kernel_name) + ":" + std::to_string(device_id) + ";";
}
@@ -446,22 +443,21 @@ void hip_api_callback(
break;
}
case HIP_API_ID_hipLaunchKernel:
case HIP_API_ID_hipLaunchCooperativeKernel:
{
case HIP_API_ID_hipLaunchCooperativeKernel: {
const void* f = data->args.hipLaunchKernel.function_address;
hipStream_t stream = data->args.hipLaunchKernel.stream;
if (f != NULL) entry->name = strdup(roctracer::HipLoader::Instance().KernelNameRefByPtr(f, stream));
if (f != NULL)
entry->name = strdup(roctracer::HipLoader::Instance().KernelNameRefByPtr(f, stream));
break;
}
case HIP_API_ID_hipExtLaunchKernel:
{
case HIP_API_ID_hipExtLaunchKernel: {
const void* f = data->args.hipExtLaunchKernel.function_address;
hipStream_t stream = data->args.hipExtLaunchKernel.stream;
if (f != NULL) entry->name = strdup(roctracer::HipLoader::Instance().KernelNameRefByPtr(f, stream));
if (f != NULL)
entry->name = strdup(roctracer::HipLoader::Instance().KernelNameRefByPtr(f, stream));
break;
}
default:
{
default: {
const hipFunction_t f = data->args.hipModuleLaunchKernel.f;
if (f != NULL) entry->name = strdup(roctracer::HipLoader::Instance().KernelNameRef(f));
}
@@ -471,17 +467,15 @@ void hip_api_callback(
entry->valid.store(roctracer::TRACE_ENTRY_COMPL, std::memory_order_release);
}
const char * name = roctracer_op_string(domain, cid, 0);
DEBUG_TRACE("hip_api_callback(\"%s\") phase(%d): cid(%u) data(%p) entry(%p) name(\"%s\") correlation_id(%lu) timestamp(%lu)\n",
name, data->phase, cid, data, entry, (entry) ? entry->name : NULL, data->correlation_id, timestamp);
const char* name = roctracer_op_string(domain, cid, 0);
DEBUG_TRACE(
"hip_api_callback(\"%s\") phase(%d): cid(%u) data(%p) entry(%p) name(\"%s\") "
"correlation_id(%lu) timestamp(%lu)\n",
name, data->phase, cid, data, entry, (entry) ? entry->name : NULL, data->correlation_id,
timestamp);
}
void mark_api_callback(
uint32_t domain,
uint32_t cid,
const void* callback_data,
void* arg)
{
void mark_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) {
(void)arg;
const char* name = reinterpret_cast<const char*>(callback_data);
@@ -513,22 +507,27 @@ void hip_api_flush_cb(hip_api_trace_entry_t* entry) {
std::ostringstream rec_ss;
std::ostringstream oss;
const char* str = (domain != ACTIVITY_DOMAIN_EXT_API) ? roctracer_op_string(domain, cid, 0) : strdup("MARK");
rec_ss << std::dec << begin_timestamp << ":" << end_timestamp << " " << entry->pid << ":" << entry->tid;
const char* str =
(domain != ACTIVITY_DOMAIN_EXT_API) ? roctracer_op_string(domain, cid, 0) : strdup("MARK");
rec_ss << std::dec << begin_timestamp << ":" << end_timestamp << " " << entry->pid << ":"
<< entry->tid;
oss << std::dec << rec_ss.str() << " " << str;
const char * name = roctracer_op_string(entry->domain, entry->cid, 0);
DEBUG_TRACE("hip_api_flush_cb(\"%s\"): domain(%u) cid(%u) entry(%p) name(\"%s\" correlation_id(%lu) beg(%lu) end(%lu))\n",
name, entry->domain, entry->cid, entry, entry->name, correlation_id, begin_timestamp, end_timestamp);
const char* name = roctracer_op_string(entry->domain, entry->cid, 0);
DEBUG_TRACE(
"hip_api_flush_cb(\"%s\"): domain(%u) cid(%u) entry(%p) name(\"%s\" correlation_id(%lu) "
"beg(%lu) end(%lu))\n",
name, entry->domain, entry->cid, entry, entry->name, correlation_id, begin_timestamp,
end_timestamp);
if (domain == ACTIVITY_DOMAIN_HIP_API) {
#if HIP_PROF_HIP_API_STRING
if (hip_api_stats != NULL) {
hip_api_stats->add_event(cid, end_timestamp - begin_timestamp);
if (is_hip_kernel_launch_api(cid)) {
hip_kernel_mutex.lock();
hip_kernel_mutex.lock();
(*hip_kernel_map)[correlation_id] = entry->name;
hip_kernel_mutex.unlock();
hip_kernel_mutex.unlock();
}
} else {
const char* str = hipApiString((hip_api_id_t)cid, data);
@@ -537,50 +536,36 @@ void hip_api_flush_cb(hip_api_trace_entry_t* entry) {
const char* kernel_name = cxx_demangle(entry->name);
rec_ss << " kernel=" << kernel_name;
}
rec_ss<< " :" << correlation_id;
rec_ss << " :" << correlation_id;
fprintf(hip_api_file_handle, "%s\n", rec_ss.str().c_str());
}
#else // !HIP_PROF_HIP_API_STRING
#else // !HIP_PROF_HIP_API_STRING
switch (cid) {
case HIP_API_ID_hipMemcpy:
fprintf(hip_api_file_handle, "%s(dst(%p) src(%p) size(0x%x) kind(%u))\n",
oss.str().c_str(),
data->args.hipMemcpy.dst,
data->args.hipMemcpy.src,
(uint32_t)(data->args.hipMemcpy.sizeBytes),
(uint32_t)(data->args.hipMemcpy.kind));
fprintf(hip_api_file_handle, "%s(dst(%p) src(%p) size(0x%x) kind(%u))\n", oss.str().c_str(),
data->args.hipMemcpy.dst, data->args.hipMemcpy.src,
(uint32_t)(data->args.hipMemcpy.sizeBytes), (uint32_t)(data->args.hipMemcpy.kind));
break;
case HIP_API_ID_hipMemcpyAsync:
fprintf(hip_api_file_handle, "%s(dst(%p) src(%p) size(0x%x) kind(%u) stream(%p))\n",
oss.str().c_str(),
data->args.hipMemcpyAsync.dst,
data->args.hipMemcpyAsync.src,
(uint32_t)(data->args.hipMemcpyAsync.sizeBytes),
(uint32_t)(data->args.hipMemcpyAsync.kind),
data->args.hipMemcpyAsync.stream);
oss.str().c_str(), data->args.hipMemcpyAsync.dst, data->args.hipMemcpyAsync.src,
(uint32_t)(data->args.hipMemcpyAsync.sizeBytes),
(uint32_t)(data->args.hipMemcpyAsync.kind), data->args.hipMemcpyAsync.stream);
break;
case HIP_API_ID_hipMalloc:
fprintf(hip_api_file_handle, "%s(ptr(%p) size(0x%x))\n",
oss.str().c_str(),
entry->ptr,
(uint32_t)(data->args.hipMalloc.size));
fprintf(hip_api_file_handle, "%s(ptr(%p) size(0x%x))\n", oss.str().c_str(), entry->ptr,
(uint32_t)(data->args.hipMalloc.size));
break;
case HIP_API_ID_hipFree:
fprintf(hip_api_file_handle, "%s(ptr(%p))\n",
oss.str().c_str(),
data->args.hipFree.ptr);
fprintf(hip_api_file_handle, "%s(ptr(%p))\n", oss.str().c_str(), data->args.hipFree.ptr);
break;
case HIP_API_ID_hipModuleLaunchKernel:
fprintf(hip_api_file_handle, "%s(kernel(%s) stream(%p))\n",
oss.str().c_str(),
cxx_demangle(entry->name),
data->args.hipModuleLaunchKernel.stream);
fprintf(hip_api_file_handle, "%s(kernel(%s) stream(%p))\n", oss.str().c_str(),
cxx_demangle(entry->name), data->args.hipModuleLaunchKernel.stream);
break;
case HIP_API_ID_hipExtModuleLaunchKernel:
fprintf(hip_api_file_handle, "%s(kernel(%s) stream(%p))\n",
oss.str().c_str(),
cxx_demangle(entry->name),
data->args.hipExtModuleLaunchKernel.hStream);
fprintf(hip_api_file_handle, "%s(kernel(%s) stream(%p))\n", oss.str().c_str(),
cxx_demangle(entry->name), data->args.hipExtModuleLaunchKernel.hStream);
break;
default:
fprintf(hip_api_file_handle, "%s()\n", oss.str().c_str());
@@ -605,23 +590,26 @@ struct hip_act_trace_entry_t {
};
void hip_act_flush_cb(hip_act_trace_entry_t* entry);
constexpr roctracer::TraceBuffer<hip_act_trace_entry_t>::flush_prm_t hip_act_flush_prm = {roctracer::DFLT_ENTRY_TYPE, hip_act_flush_cb};
constexpr roctracer::TraceBuffer<hip_act_trace_entry_t>::flush_prm_t hip_act_flush_prm = {
roctracer::DFLT_ENTRY_TYPE, hip_act_flush_cb};
roctracer::TraceBuffer<hip_act_trace_entry_t>* hip_act_trace_buffer = NULL;
// HIP ACT trace buffer flush callback
void hip_act_flush_cb(hip_act_trace_entry_t* entry) {
const uint32_t domain = ACTIVITY_DOMAIN_HCC_OPS;
const uint32_t op = 0;
const char * name = roctracer_op_string(domain, op, entry->kind);
const char* name = roctracer_op_string(domain, op, entry->kind);
if (name == NULL) {
printf("hip_act_flush_cb name is NULL\n"); fflush(stdout);
printf("hip_act_flush_cb name is NULL\n");
fflush(stdout);
abort();
}
if (strncmp("Kernel", name, 6) == 0) {
hip_kernel_mutex.lock();
if (hip_kernel_stats == NULL) {
printf("hip_act_flush_cb hip_kernel_stats is NULL\n"); fflush(stdout);
printf("hip_act_flush_cb hip_kernel_stats is NULL\n");
fflush(stdout);
abort();
}
name = (*hip_kernel_map)[entry->correlation_id];
@@ -640,11 +628,14 @@ void pool_activity_callback(const char* begin, const char* end, void* arg) {
const roctracer_record_t* end_record = reinterpret_cast<const roctracer_record_t*>(end);
while (record < end_record) {
const char * name = roctracer_op_string(record->domain, record->op, record->kind);
DEBUG_TRACE("pool_activity_callback(\"%s\"): domain(%u) op(%u) kind(%u) record(%p) correlation_id(%lu) beg(%lu) end(%lu)\n",
name, record->domain, record->op, record->kind, record, record->correlation_id, record->begin_ns, record->end_ns);
const char* name = roctracer_op_string(record->domain, record->op, record->kind);
DEBUG_TRACE(
"pool_activity_callback(\"%s\"): domain(%u) op(%u) kind(%u) record(%p) correlation_id(%lu) "
"beg(%lu) end(%lu)\n",
name, record->domain, record->op, record->kind, record, record->correlation_id,
record->begin_ns, record->end_ns);
switch(record->domain) {
switch (record->domain) {
case ACTIVITY_DOMAIN_HCC_OPS:
if (hip_memcpy_stats != NULL) {
hip_act_trace_entry_t* entry = hip_act_trace_buffer->GetEntry();
@@ -653,17 +644,16 @@ void pool_activity_callback(const char* begin, const char* end, void* arg) {
entry->correlation_id = record->correlation_id;
entry->valid.store(roctracer::TRACE_ENTRY_COMPL, std::memory_order_release);
} else {
fprintf(hcc_activity_file_handle, "%lu:%lu %d:%lu %s:%lu:%u\n",
record->begin_ns, record->end_ns,
record->device_id, record->queue_id,
name, record->correlation_id, my_pid);
fprintf(hcc_activity_file_handle, "%lu:%lu %d:%lu %s:%lu:%u\n", record->begin_ns,
record->end_ns, record->device_id, record->queue_id, name, record->correlation_id,
my_pid);
fflush(hcc_activity_file_handle);
}
break;
case ACTIVITY_DOMAIN_HSA_OPS:
if (record->op == HSA_OP_ID_RESERVED1) {
fprintf(pc_sample_file_handle, "%u %lu 0x%lx %s\n",
record->pc_sample.se, record->pc_sample.cycle, record->pc_sample.pc, name);
fprintf(pc_sample_file_handle, "%u %lu 0x%lx %s\n", record->pc_sample.se,
record->pc_sample.cycle, record->pc_sample.pc, name);
fflush(pc_sample_file_handle);
}
break;
@@ -682,11 +672,14 @@ std::string normalize_token(const std::string& token, bool not_empty, const std:
std::string error_str = "none";
if (first_pos != std::string::npos) {
const size_t last_pos = token.find_last_not_of(space_chars_set);
if (last_pos == std::string::npos) error_str = "token string error: \"" + token + "\"";
if (last_pos == std::string::npos)
error_str = "token string error: \"" + token + "\"";
else {
const size_t end_pos = last_pos + 1;
if (end_pos <= first_pos) error_str = "token string error: \"" + token + "\"";
else norm_len = end_pos - first_pos;
if (end_pos <= first_pos)
error_str = "token string error: \"" + token + "\"";
else
norm_len = end_pos - first_pos;
}
}
if (((first_pos != std::string::npos) && (norm_len == 0)) ||
@@ -696,7 +689,8 @@ std::string normalize_token(const std::string& token, bool not_empty, const std:
return (norm_len != 0) ? token.substr(first_pos, norm_len) : std::string("");
}
int get_xml_array(const xml::Xml::level_t* node, const std::string& field, const std::string& delim, std::vector<std::string>* vec, const char* label = NULL) {
int get_xml_array(const xml::Xml::level_t* node, const std::string& field, const std::string& delim,
std::vector<std::string>* vec, const char* label = NULL) {
int parse_iter = 0;
const auto& opts = node->opts;
auto it = opts.find(field);
@@ -706,7 +700,7 @@ int get_xml_array(const xml::Xml::level_t* node, const std::string& field, const
size_t pos1 = 0;
const size_t string_len = array_string.length();
while (pos1 < string_len) {
// set pos2 such that it also handles case of multiple delimiter options.
// set pos2 such that it also handles case of multiple delimiter options.
// For example- "hipLaunchKernel, hipExtModuleLaunchKernel, hipMemsetAsync"
// in this example delimiters are ' ' and also ','
const size_t pos2 = array_string.find_first_of(delim, pos1);
@@ -716,9 +710,9 @@ int get_xml_array(const xml::Xml::level_t* node, const std::string& field, const
const std::string norm_str = normalize_token(token, found, "get_xml_array");
if (norm_str.length() != 0) vec->push_back(norm_str);
if (!found) break;
// update pos2 such that it represents the first non-delimiter character
// in case multiple delimiters are specified in variable 'delim'
pos1 = array_string.find_first_not_of(delim, pos2);
// update pos2 such that it represents the first non-delimiter character
// in case multiple delimiters are specified in variable 'delim'
pos1 = array_string.find_first_not_of(delim, pos2);
++parse_iter;
}
}
@@ -742,7 +736,8 @@ FILE* open_output_file(const char* prefix, const char* name, const char** path =
}
if (path != NULL) *path = strdup(oss.str().c_str());
} else file_handle = stdout;
} else
file_handle = stdout;
return file_handle;
}
@@ -785,7 +780,7 @@ void tool_unload() {
flush_thread_started = false;
flush_thread_mutex.unlock();
PTHREAD_CALL(pthread_cancel(flush_thread));
void *res;
void* res;
PTHREAD_CALL(pthread_join(flush_thread, &res));
if (res != PTHREAD_CANCELED) FATAL("flush thread wasn't stopped correctly");
}
@@ -860,7 +855,8 @@ void tool_load() {
}
}
printf("ROCTracer (pid=%d): ", (int)GetPid()); fflush(stdout);
printf("ROCTracer (pid=%d): ", (int)GetPid());
fflush(stdout);
// XML input
const char* xml_name = getenv("ROCP_INPUT");
@@ -879,8 +875,10 @@ void tool_load() {
std::vector<std::string> api_vec;
for (const auto* node : entry->nodes) {
if (node->tag != "parameters") fatal("ROCTracer: trace node is not supported '" + name + ":" + node->tag + "'");
get_xml_array(node, "api", ", ", &api_vec); // delimiter options given as both spaces and commas (' ' and ',')
if (node->tag != "parameters")
fatal("ROCTracer: trace node is not supported '" + name + ":" + node->tag + "'");
get_xml_array(node, "api", ", ",
&api_vec); // delimiter options given as both spaces and commas (' ' and ',')
break;
}
@@ -917,14 +915,13 @@ void tool_load() {
roctx_file_handle = open_output_file(output_prefix, "roctx_trace.txt");
// initialize HSA tracing
roctracer_ext_properties_t properties {
start_callback,
stop_callback
};
roctracer_ext_properties_t properties{start_callback, stop_callback};
roctracer_set_properties(ACTIVITY_DOMAIN_EXT_API, &properties);
fprintf(stdout, " rocTX-trace()\n"); fflush(stdout);
ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, roctx_api_callback, NULL));
fprintf(stdout, " rocTX-trace()\n");
fflush(stdout);
ROCTRACER_CALL(
roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, roctx_api_callback, NULL));
}
const char* ctrl_str = getenv("ROCP_CTRL_RATE");
@@ -934,10 +931,13 @@ void tool_load() {
uint32_t ctrl_rate = 0;
if (sscanf(ctrl_str, "%d:%d:%d", &ctrl_delay, &ctrl_len, &ctrl_rate) != 3) {
EXC_RAISING(ROCTRACER_STATUS_ERROR, "Invalid ROCP_CTRL_RATE var(" << ctrl_str << "), expected ctrl_delay:ctrl_len:ctrl_rate");
EXC_RAISING(
ROCTRACER_STATUS_ERROR,
"Invalid ROCP_CTRL_RATE var(" << ctrl_str << "), expected ctrl_delay:ctrl_len:ctrl_rate");
}
if (ctrl_len > ctrl_rate) {
EXC_RAISING(ROCTRACER_STATUS_ERROR, "Control length value " << ctrl_len << " > rate value " << ctrl_rate);
EXC_RAISING(ROCTRACER_STATUS_ERROR,
"Control length value " << ctrl_len << " > rate value " << ctrl_rate);
}
control_dist_us = ctrl_rate - ctrl_len;
control_len_us = ctrl_len;
@@ -946,14 +946,21 @@ void tool_load() {
roctracer_stop();
if (ctrl_delay != UINT32_MAX) {
fprintf(stdout, "ROCTracer: trace control: delay(%uus), length(%uus), rate(%uus)\n", ctrl_delay, ctrl_len, ctrl_rate); fflush(stdout);
fprintf(stdout, "ROCTracer: trace control: delay(%uus), length(%uus), rate(%uus)\n",
ctrl_delay, ctrl_len, ctrl_rate);
fflush(stdout);
pthread_t thread;
pthread_attr_t attr;
int err = pthread_attr_init(&attr);
if (err) { errno = err; perror("pthread_attr_init"); abort(); }
if (err) {
errno = err;
perror("pthread_attr_init");
abort();
}
err = pthread_create(&thread, &attr, control_thr_fun, NULL);
} else {
fprintf(stdout, "ROCTracer: trace start disabled\n"); fflush(stdout);
fprintf(stdout, "ROCTracer: trace start disabled\n");
fflush(stdout);
}
}
@@ -965,10 +972,15 @@ void tool_load() {
abort();
}
fprintf(stdout, "ROCTracer: trace control flush rate(%uus)\n", control_flush_us); fflush(stdout);
fprintf(stdout, "ROCTracer: trace control flush rate(%uus)\n", control_flush_us);
fflush(stdout);
pthread_attr_t attr;
int err = pthread_attr_init(&attr);
if (err) { errno = err; perror("pthread_attr_init"); abort(); }
if (err) {
errno = err;
perror("pthread_attr_init");
abort();
}
std::lock_guard<std::mutex> lock(flush_thread_mutex);
PTHREAD_CALL(pthread_create(&flush_thread, &attr, flush_thr_fun, NULL));
flush_thread_started = true;
@@ -978,7 +990,8 @@ void tool_load() {
}
// HSA-runtime tool on-load method
extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
uint64_t failed_tool_count,
const char* const* failed_tool_names) {
ONLOAD_TRACE_BEG();
@@ -998,17 +1011,20 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
// initialize HSA tracing
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*)table);
fprintf(stdout, " HSA-trace("); fflush(stdout);
fprintf(stdout, " HSA-trace(");
fflush(stdout);
if (hsa_api_vec().size() != 0) {
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, NULL));
ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, NULL));
ROCTRACER_CALL(
roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, NULL));
printf(" %s", api);
}
} else {
ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, NULL));
ROCTRACER_CALL(
roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, NULL));
}
printf(")\n");
}
@@ -1018,21 +1034,20 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
hsa_async_copy_file_handle = open_output_file(output_prefix, "async_copy_trace.txt");
// initialize HSA tracing
roctracer::hsa_ops_properties_t ops_properties {
table,
reinterpret_cast<activity_async_callback_t>(hsa_activity_callback),
NULL,
output_prefix
};
roctracer::hsa_ops_properties_t ops_properties{
table, reinterpret_cast<activity_async_callback_t>(hsa_activity_callback), NULL,
output_prefix};
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties);
fprintf(stdout, " HSA-activity-trace()\n"); fflush(stdout);
fprintf(stdout, " HSA-activity-trace()\n");
fflush(stdout);
ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
}
// Enable HIP API callbacks/activity
if (trace_hip_api || trace_hip_activity) {
fprintf(stdout, " HIP-trace()\n"); fflush(stdout);
fprintf(stdout, " HIP-trace()\n");
fflush(stdout);
// roctracer properties
roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, (void*)mark_api_callback);
// Allocating tracing pool
@@ -1052,11 +1067,13 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
uint32_t cid = HIP_API_ID_NONE;
const char* api = hip_api_vec()[i].c_str();
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, api, &cid, NULL));
ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, cid, hip_api_callback, NULL));
ROCTRACER_CALL(
roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, cid, hip_api_callback, NULL));
printf(" %s", api);
}
} else {
ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, hip_api_callback, NULL));
ROCTRACER_CALL(
roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, hip_api_callback, NULL));
}
if (is_stats_opt) {
@@ -1087,7 +1104,8 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
// Enable PC sampling
if (trace_pcs) {
fprintf(stdout, " PCS-trace()\n"); fflush(stdout);
fprintf(stdout, " PCS-trace()\n");
fflush(stdout);
open_tracing_pool();
pc_sample_file_handle = open_output_file(output_prefix, "pcs_trace.txt");
ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_RESERVED1));
@@ -1103,17 +1121,19 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
}
// HSA-runtime on-unload method
extern "C" PUBLIC_API void OnUnload() {
ONLOAD_TRACE("");
}
extern "C" PUBLIC_API void OnUnload() { ONLOAD_TRACE(""); }
extern "C" CONSTRUCTOR_API void constructor() {
ONLOAD_TRACE_BEG();
roctracer::hip_support::HIP_depth_max = 0;
roctx_trace_buffer = new roctracer::TraceBuffer<roctx_trace_entry_t>("rocTX API", 0x200000, &roctx_flush_prm, 1);
hip_api_trace_buffer = new roctracer::TraceBuffer<hip_api_trace_entry_t>("HIP API", 0x200000, &hip_api_flush_prm, 1);
hip_act_trace_buffer = new roctracer::TraceBuffer<hip_act_trace_entry_t>("HIP ACT", 0x200000, &hip_act_flush_prm, 1, 1);
hsa_api_trace_buffer = new roctracer::TraceBuffer<hsa_api_trace_entry_t>("HSA API", 0x200000, &hsa_flush_prm, 1);
roctx_trace_buffer =
new roctracer::TraceBuffer<roctx_trace_entry_t>("rocTX API", 0x200000, &roctx_flush_prm, 1);
hip_api_trace_buffer =
new roctracer::TraceBuffer<hip_api_trace_entry_t>("HIP API", 0x200000, &hip_api_flush_prm, 1);
hip_act_trace_buffer = new roctracer::TraceBuffer<hip_act_trace_entry_t>(
"HIP ACT", 0x200000, &hip_act_flush_prm, 1, 1);
hsa_api_trace_buffer =
new roctracer::TraceBuffer<hsa_api_trace_entry_t>("HSA API", 0x200000, &hsa_flush_prm, 1);
roctracer_load();
tool_load();
ONLOAD_TRACE_END();