AI/rocm-systems
2
0
포크 0
코드 이슈 풀 리퀘스트 액션 패키지 프로젝트 릴리즈 위키 활동

Merge branch 'amd-develop' into amd-master

소스 검색 
Change-Id: I8921e67e352e35e4c496e78a797fb309279ab7d0


[ROCm/hip-tests commit: b9d4c0be78]
이 커밋은 다음에 포함됨:
Maneesh Gupta
2017-03-14 15:57:38 +05:30
부모 874a182a6c ea06138990
커밋 c8321725f3
4개의 변경된 파일과 169개의 추가작업 그리고 65개의 파일을 삭제
패치 파일 다운로드 Diff 파일 다운로드 모든 파일 펼침 모든 파일 접기
projects/hip-tests/samples/0_Intro/square/square.hipref.cpp
+6 -4
파일 보기
@@ -32,7 +32,7 @@ THE SOFTWARE.
}\
}
/*
/*
* Square each element in the array A and write to array C.
*/
template <typename T>
@@ -58,16 +58,18 @@ int main(int argc, char *argv[])
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
printf ("info: running on device %s\n", props.name);
#ifdef __HIP_PLATFORM_HCC__
printf ("info: architecture on AMD GPU device is: %d\n",props.gcnArch);
#endif
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
// Fill with Phi + i
for (size_t i=0; i<N; i++)
for (size_t i=0; i<N; i++)
{
A_h[i] = 1.618f + i;
A_h[i] = 1.618f + i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+154 -59
파일 보기
@@ -6,9 +6,12 @@
#include "ResultDatabase.h"
enum MallocMode {MallocPinned, MallocUnpinned, MallocRegistered};
// Cmdline parms:
bool p_verbose = false;
bool p_pinned = true;
MallocMode p_malloc_mode = MallocPinned;
int p_numa_ctl = -1;
int p_iterations = 10;
int p_beatsperiteration=1;
int p_device = 0;
@@ -21,7 +24,7 @@ bool p_h2d = true;
bool p_d2h = true;
bool p_bidir = true;
//#define NO_CHECK
#define CHECK_HIP_ERROR() \
@@ -36,6 +39,14 @@ bool p_bidir = true;
}
std::string mallocModeString(int mallocMode) {
switch (mallocMode) {
case MallocPinned : return "pinned";
case MallocUnpinned: return "unpinned";
case MallocRegistered: return "registered";
default: return "mallocmode-UNKNOWN";
};
};
// ****************************************************************************
int sizeToBytes(int size) {
@@ -106,7 +117,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
// Create some host memory pattern
float *hostMem = NULL;
if (p_pinned)
if (p_malloc_mode == MallocPinned)
{
hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats);
while (hipGetLastError() != hipSuccess)
@@ -116,20 +127,33 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any pinned buffer\n";
std::cerr << "Error: Couldn't allocate any pinned buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats);
}
}
else
else if (p_malloc_mode == MallocUnpinned)
{
if (p_alignedhost) {
hostMem = (float*)aligned_alloc(p_alignedhost, numMaxFloats*sizeof(float));
} else {
hostMem = new float[numMaxFloats];
}
}
else if (p_malloc_mode == MallocRegistered)
{
if (p_numa_ctl == -1) {
hostMem = (float*)malloc(numMaxFloats*sizeof(float));
}
hipHostRegister(hostMem, numMaxFloats * sizeof(float), 0);
CHECK_HIP_ERROR();
}
else
{
assert(0);
}
for (int i = 0; i < numMaxFloats; i++)
@@ -146,7 +170,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any device buffer\n";
std::cerr << "Error: Couldn't allocate any device buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
@@ -199,8 +223,8 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
} else {
sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str());
}
resultDB.AddResult(std::string("H2D_Bandwidth") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("H2D_Time") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "ms", t);
resultDB.AddResult(std::string("H2D_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("H2D_Time") + mallocModeString(p_malloc_mode), sizeStr, "ms", t);
if (p_onesize) {
break;
@@ -212,6 +236,8 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
numMaxFloats = sizeToBytes(p_onesize) / sizeof(float);
}
#ifndef NO_CHECK
// Check. First reset the host memory, then copy-back result. Then compare against original ref value.
for (int i = 0; i < numMaxFloats; i++)
{
@@ -225,24 +251,36 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
printf ("error: H2D. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem[i]);
}
}
#endif
// Cleanup
hipFree((void*)device);
CHECK_HIP_ERROR();
if (p_pinned)
{
switch (p_malloc_mode) {
case MallocPinned:
hipHostFree((void*)hostMem);
CHECK_HIP_ERROR();
}
else
{
break;
case MallocUnpinned:
if (p_alignedhost) {
delete[] hostMem;
} else {
free(hostMem);
}
break;
case MallocRegistered:
hipHostUnregister(hostMem);
CHECK_HIP_ERROR();
free(hostMem);
break;
default:
assert(0);
}
hipEventDestroy(start);
hipEventDestroy(stop);
}
@@ -257,38 +295,56 @@ void RunBenchmark_D2H(ResultDatabase &resultDB)
// Create some host memory pattern
float *hostMem1;
float *hostMem2;
if (p_pinned)
if (p_malloc_mode == MallocPinned)
{
hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats);
hipError_t err1 = hipGetLastError();
hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats);
hipError_t err2 = hipGetLastError();
while (err1 != hipSuccess || err2 != hipSuccess)
{
// free the first buffer if only the second failed
if (err1 == hipSuccess)
hipHostFree((void*)hostMem1);
while (err1 != hipSuccess || err2 != hipSuccess)
{
// free the first buffer if only the second failed
if (err1 == hipSuccess)
hipHostFree((void*)hostMem1);
// drop the size and try again
if (p_verbose) std::cout << " - dropping size allocating pinned mem\n";
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any pinned buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats);
err1 = hipGetLastError();
hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats);
err2 = hipGetLastError();
}
}
else
// drop the size and try again
if (p_verbose) std::cout << " - dropping size allocating pinned mem\n";
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocate any pinned buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats);
err1 = hipGetLastError();
hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats);
err2 = hipGetLastError();
}
}
else if (p_malloc_mode == MallocUnpinned)
{
hostMem1 = new float[numMaxFloats];
hostMem2 = new float[numMaxFloats];
}
else if (p_malloc_mode == MallocRegistered)
{
if (p_numa_ctl == -1) {
hostMem1 = (float*)malloc(numMaxFloats*sizeof(float));
hostMem2 = (float*)malloc(numMaxFloats*sizeof(float));
}
hipHostRegister(hostMem1, numMaxFloats * sizeof(float), 0);
CHECK_HIP_ERROR();
hipHostRegister(hostMem2, numMaxFloats * sizeof(float), 0);
CHECK_HIP_ERROR();
}
else
{
assert(0);
}
for (int i=0; i<numMaxFloats; i++)
hostMem1[i] = i % 77;
@@ -301,7 +357,7 @@ void RunBenchmark_D2H(ResultDatabase &resultDB)
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any device buffer\n";
std::cerr << "Error: Couldn't allocate any device buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
@@ -358,8 +414,8 @@ void RunBenchmark_D2H(ResultDatabase &resultDB)
} else {
sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str());
}
resultDB.AddResult(std::string("D2H_Bandwidth") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("D2H_Time") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "ms", t);
resultDB.AddResult(std::string("D2H_Bandwidth") +"_" + mallocModeString(p_malloc_mode) , sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("D2H_Time") +"_" + mallocModeString(p_malloc_mode) , sizeStr, "ms", t);
if (p_onesize) {
break;
}
@@ -381,20 +437,31 @@ void RunBenchmark_D2H(ResultDatabase &resultDB)
// Cleanup
hipFree((void*)device);
CHECK_HIP_ERROR();
if (p_pinned)
{
switch (p_malloc_mode) {
case MallocPinned:
hipHostFree((void*)hostMem1);
CHECK_HIP_ERROR();
hipHostFree((void*)hostMem2);
CHECK_HIP_ERROR();
}
else
{
break;
case MallocUnpinned:
delete[] hostMem1;
delete[] hostMem2;
hipEventDestroy(start);
hipEventDestroy(stop);
break;
case MallocRegistered:
hipHostUnregister(hostMem1);
CHECK_HIP_ERROR();
free(hostMem1);
hipHostUnregister(hostMem2);
free(hostMem2);
break;
default:
assert(0);
}
hipEventDestroy(start);
hipEventDestroy(stop);
}
@@ -409,7 +476,7 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
// Create some host memory pattern
float *hostMem[2] = {NULL, NULL};
if (p_pinned)
if (p_malloc_mode == MallocPinned)
{
while (1)
{
@@ -424,18 +491,34 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any pinned buffer\n";
std::cerr << "Error: Couldn't allocate any pinned buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
}
}
}
else
else if (p_malloc_mode == MallocUnpinned)
{
hostMem[0] = new float[numMaxFloats];
hostMem[1] = new float[numMaxFloats];
}
else if (p_malloc_mode == MallocRegistered)
{
if (p_numa_ctl == -1) {
hostMem[0] = (float*)malloc(numMaxFloats*sizeof(float));
hostMem[1] = (float*)malloc(numMaxFloats*sizeof(float));
}
hipHostRegister(hostMem[0], numMaxFloats * sizeof(float), 0);
CHECK_HIP_ERROR();
hipHostRegister(hostMem[1], numMaxFloats * sizeof(float), 0);
CHECK_HIP_ERROR();
}
else
{
assert(0);
}
for (int i = 0; i < numMaxFloats; i++)
{
@@ -459,7 +542,7 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
--nSizes;
if (nSizes < 1)
{
std::cerr << "Error: Couldn't allocated any device buffer\n";
std::cerr << "Error: Couldn't allocate any device buffer\n";
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
@@ -512,8 +595,8 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
double speed = (double(sizeToBytes(thisSize)) / (1000*1000)) / t;
char sizeStr[256];
sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str());
resultDB.AddResult(std::string("Bidir_Bandwidth") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("Bidir_Time") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "ms", t);
resultDB.AddResult(std::string("Bidir_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed);
resultDB.AddResult(std::string("Bidir_Time") + "_" + mallocModeString(p_malloc_mode), sizeStr, "ms", t);
}
}
@@ -521,17 +604,27 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
hipFree((void*)deviceMem[0]);
hipFree((void*)deviceMem[1]);
CHECK_HIP_ERROR();
if (p_pinned)
{
switch (p_malloc_mode) {
case MallocPinned:
hipHostFree((void*)hostMem[0]);
hipHostFree((void*)hostMem[1]);
CHECK_HIP_ERROR();
}
else
{
break;
case MallocUnpinned:
delete[] hostMem[0];
delete[] hostMem[1];
}
break;
case MallocRegistered:
for (int i=0; i<2; i++) {
hipHostUnregister(hostMem[i]);
CHECK_HIP_ERROR();
free(hostMem[i]);
}
break;
default:
assert(0);
};
hipEventDestroy(start);
hipEventDestroy(stop);
hipStreamDestroy(stream[0]);
@@ -557,7 +650,7 @@ void printConfig() {
hipDeviceProp_t props;
hipGetDeviceProperties(&props, p_device);
printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz Pinned=%s\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0, p_pinned ? "YES" : "NO");
printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz MallocMode=%s\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0, mallocModeString(p_malloc_mode).c_str());
}
void help() {
@@ -601,7 +694,9 @@ int parseStandardArguments(int argc, char *argv[])
failed("Bad onesize argument");
}
} else if (!strcmp(arg, "--unpinned")) {
p_pinned = 0;
p_malloc_mode = MallocUnpinned;
} else if (!strcmp(arg, "--registered")) {
p_malloc_mode = MallocRegistered;
} else if (!strcmp(arg, "--h2d")) {
p_h2d = true;
p_d2h = false;
projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile
+5 -2
파일 보기
@@ -3,6 +3,10 @@ ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
ifeq (gfx701, $(findstring gfx701,$(HCC_AMDGPU_TARGET)))
$(error gfx701 is not a supported device for this sample)
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc
@@ -22,7 +26,7 @@ CXX=$(HIPCC)
$(EXECUTABLE): $(OBJECTS)
$(HIPCC) $(OBJECTS) -o $@
$(HIPCC) $(OBJECTS) -o $@
test: $(EXECUTABLE)
@@ -33,4 +37,3 @@ clean:
rm -f $(EXECUTABLE)
rm -f $(OBJECTS)
rm -f $(HIP_PATH)/src/*.o
projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile
+4
파일 보기
@@ -3,6 +3,10 @@ ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
ifeq (gfx701, $(findstring gfx701,$(HCC_AMDGPU_TARGET)))
$(error gfx701 is not a supported device for this sample)
endif
HIPCC=$(HIP_PATH)/bin/hipcc
TARGET=hcc