AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

use codexl marker interface to mark HIP function/begin end.

Browse Source 
- Creates markers in HIP group and they show up in CodeXL trace
- Marker text includes HIP functioin arguments
- (Add trace_helper to convert arguments to strings)
- Still need to add HIP_INIT_API for ~30 HIP functions.


[ROCm/hip commit: b6b3fea05b]
This commit is contained in:
Ben Sander
2016-03-23 01:17:53 -05:00
parent f9b2f0c840
commit 82cda5c087
5 changed files with 192 additions and 57 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/CMakeLists.txt
+15 -3
View File
@@ -19,11 +19,18 @@ MESSAGE ("HIP_PATH=" ${HIP_PATH})
if (${HIP_PLATFORM} STREQUAL "hcc")
MESSAGE ("HIP_PLATFORM=hcc")
set (HSA_PATH $ENV{HSA_PATH})
if (NOT DEFINED HSA_PATH)
set (HSA_PATH /opt/hsa)
endif()
set (CODEXL_PATH $ENV{CODEXL_PATH})
if (NOT DEFINED CODEXL_PATH)
set (CODEXL_PATH /opt/AMD/CodeXL)
endif()
set (CODEXL_SDK_ATAL_PATH ${CODEXL_PATH}/SDK/AMDTActivityLogger)
#---
# Add HSA library:
add_library(hsa-runtime64 SHARED IMPORTED)
@@ -37,7 +44,7 @@ if (${HIP_PLATFORM} STREQUAL "hcc")
# hip_hcc.o:
#add_library(hip_hcc STATIC ${HIP_PATH}/src/hip_hcc.cpp )
add_library(hip_hcc STATIC ${HIP_PATH}/src/hip_hcc.cpp ${HIP_PATH}/src/staging_buffer.cpp)
target_include_directories(hip_hcc PRIVATE ${HSA_PATH}/include)
target_include_directories(hip_hcc PRIVATE ${HSA_PATH}/include ${CODEXL_SDK_ATAL_PATH}/include)
elseif (${HIP_PLATFORM} STREQUAL "nvcc")
@@ -132,7 +139,8 @@ make_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp)
make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp)
make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp)
make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp)
make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp)
#TODO - re-enable. This uses the pointer add feature.
#make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp)
make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp)
make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp)
make_hip_executable (hipHostAlloc hipHostAlloc.cpp)
@@ -153,7 +161,8 @@ make_test(hipMemset --N 10013 --memsetval 0x5a ) # oddball size.
make_test(hipMemset --N 256M --memsetval 0xa6 ) # big copy
make_test(hipGridLaunch " " )
make_test(hipEnvVarDriver " " )
make_test(hipPointerAttrib " " )
#TODO -reenable
#make_test(hipPointerAttrib " " )
#make_test(hipMultiThreadStreams1 " " )
#make_test(hipMultiThreadStreams2 " " )
make_test(hipMemcpy_simple " " )
@@ -161,6 +170,9 @@ make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 )
make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 )
make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 )
# Debug synchronization, then enable.
#make_named_test(hipMemcpy_simple "hipMemcpyAsync-simple" --async)
make_test(hipHostAlloc " ")
make_test(hipMemcpyAsync " " )
make_test(hipHostGetFlags " ")
projects/hip/tests/src/hipMemcpy_simple.cpp
+47 -8
View File
@@ -22,6 +22,18 @@ THE SOFTWARE.
#include "hip_runtime.h"
#include "test_common.h"
bool p_async = false;
// ****************************************************************************
hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind)
{
if (p_async) {
return hipMemcpyAsync(dst, src, sizeBytes, kind, NULL);
} else {
return hipMemcpy(dst, src, sizeBytes, kind);
}
}
//---
// Test simple H2D copies and back.
@@ -40,12 +52,12 @@ void simpleTest1()
printf ("A_d=%p B_d=%p C_d=%p A_h=%p B_h=%p C_h=%p\n", A_d, B_d, C_d, A_h, B_d, C_h);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK ( memcopy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK ( memcopy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, N);
HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIPCHECK ( memcopy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIPCHECK (hipDeviceSynchronize());
@@ -86,9 +98,9 @@ void simpleTest2(size_t numElements, bool usePinnedHost)
A_h2[i] = 12345678.0 + i; // init output with something distincctive, to ensure we replace it.
}
HIPCHECK(hipMemcpy(A_d, A_h1, sizeElements, hipMemcpyHostToDevice));
HIPCHECK(memcopy(A_d, A_h1, sizeElements, hipMemcpyHostToDevice));
HIPCHECK(hipDeviceSynchronize());
HIPCHECK(hipMemcpy(A_h2, A_d, sizeElements, hipMemcpyDeviceToHost));
HIPCHECK(memcopy(A_h2, A_d, sizeElements, hipMemcpyDeviceToHost));
HIPCHECK(hipDeviceSynchronize());
for (size_t i=0; i<numElements; i++) {
@@ -104,9 +116,30 @@ void simpleTest2(size_t numElements, bool usePinnedHost)
free(A_h2);
}
}
//Parse arguments specific to this test.
void parseMyArguments(int argc, char *argv[])
{
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
// parse args for this test:
for (int i = 1; i < more_argc; i++) {
const char *arg = argv[i];
if (!strcmp(arg, "--async")) {
p_async = true;
} else {
failed("Bad argument '%s'", arg);
}
}
};
int main(int argc, char *argv[])
{
HipTest::parseStandardArguments(argc, argv, true);
parseMyArguments(argc, argv);
printf ("info: set device to %d, tests=%x\n", p_gpuDevice, p_tests);
HIPCHECK(hipSetDevice(p_gpuDevice));
@@ -120,16 +153,22 @@ int main(int argc, char *argv[])
}
if (p_tests & 0x2) {
printf ("\n\n=== tests&2 (copy pin-pong, pinned host)\n");
printf ("\n\n=== tests&2 (copy ping-pong, pinned host)\n");
simpleTest2<float>(N, true/*usePinnedHost*/);
simpleTest2<char>(N, true/*usePinnedHost*/);
}
if (p_tests & 0x4) {
printf ("\n\n=== tests&2 (copy pin-pong, unpinned host)\n");
printf ("\n\n=== tests&4 (copy ping-pong, unpinned host)\n");
simpleTest2<char>(N, false/*usePinnedHost*/);
simpleTest2<float>(N, false/*usePinnedHost*/);
}
hipDeviceSynchronize();
hipDeviceReset();
int v;
hipDriverGetVersion(&v);
passed();
};