Initial commit for GPUOpen Launch

[ROCm/hip-tests commit: 4eafc92922]

projects/hip-tests/samples/0_Intro/bit_extract/Makefile

@@ -0,0 +1,24 @@
+#Dependencies : [MYHIP]/bin must be in user's path.
+
+HIP_PATH?=$(shell hipconfig --path)
+HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+# Show how to use PLATFORM to specify different options for each compiler:
+ifeq (${HIP_PLATFORM}, nvcc)
+	HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 
+endif
+ifeq (${HIP_PLATFORM}, hcc)
+	HIPCC_FLAGS =
+endif
+
+
+EXE=bit_extract
+
+$(EXE): bit_extract.cpp
+	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
+
+
+clean:
+	rm -f *.o $(EXE)
+	

projects/hip-tests/samples/0_Intro/bit_extract/README.md

@@ -0,0 +1,4 @@
+# bit_extract
+
+Show an application written directly in HIP which uses platform-specific check on __HIP_PLATFORM_HCC__ to enable use of 
+an instruction that only exists on the HCC platform.

projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp

@@ -0,0 +1,102 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <stdio.h>
+#include <iostream>
+#include <hip_runtime.h>
+
+
+#define CHECK(error) \
+    if (error != hipSuccess) { \
+      fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
+    exit(EXIT_FAILURE);\
+	}
+
+void __global__
+bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N)
+{
+    KERNELBEGIN;
+
+    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    size_t stride = hipBlockDim_x * hipGridDim_x ;
+
+    for (size_t i=offset; i<N; i+=stride) {
+#ifdef __HIP_PLATFORM_HCC__
+        C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4); 
+#else /* defined __HIP_PLATFORM_NVCC__ or other path */
+		C_d[i] = ((A_d[i] & 0xf00)  >> 8);
+#endif
+	}
+
+    KERNELEND;
+}
+
+
+int main(int argc, char *argv[])
+{
+	uint32_t *A_d, *C_d;
+	uint32_t *A_h, *C_h;
+	size_t N = 1000000;
+	size_t Nbytes = N * sizeof(uint32_t);
+
+    int deviceId;
+    CHECK (hipGetDevice(&deviceId));
+	hipDeviceProp_t props;
+	CHECK(hipDeviceGetProperties(&props, deviceId));
+	printf ("info: running on device #%d %s\n", deviceId, props.name);
+
+
+	printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+	A_h = (uint32_t*)malloc(Nbytes);
+	CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+	C_h = (uint32_t*)malloc(Nbytes);
+	CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+
+    for (size_t i=0; i<N; i++) 
+	{
+		A_h[i] = i;
+	}
+
+	printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+	CHECK(hipMalloc(&A_d, Nbytes));
+	CHECK(hipMalloc(&C_d, Nbytes));
+
+	printf ("info: copy Host2Device\n");
+    CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+
+	printf ("info: launch 'bit_extract_kernel' \n");
+	const unsigned blocks = 512;
+	const unsigned threadsPerBlock = 256;
+	hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0,   C_d, A_d, N);
+
+	printf ("info: copy Device2Host\n");
+    CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+
+	printf ("info: check result\n");
+    for (size_t i=0; i<N; i++)  {
+        unsigned Agold = ((A_h[i] & 0xf00) >> 8);
+		if (C_h[i] != Agold) {
+            fprintf (stderr, "mismatch detected.\n");
+            printf ("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
+			CHECK(hipErrorUnknown);
+		}
+	}
+}

projects/hip-tests/samples/0_Intro/square/Makefile

@@ -0,0 +1,18 @@
+HIP_PATH?=$(shell hipconfig --path)
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+all: square.hip.out
+
+square.cuda.out : square.cu
+	nvcc square.cu -o $@
+
+#hipify square.cu > square.cpp
+# Then review & finish port in square.cpp
+
+square.hip.out: square.hipref.cpp
+	$(HIPCC) square.hipref.cpp -o $@
+
+
+
+clean:
+	rm -f *.o *.out	

projects/hip-tests/samples/0_Intro/square/README.md

@@ -0,0 +1,15 @@
+# square.md
+
+Simple test which shows how to use hipify to port CUDA code to HIP.  Covered in more detail in blog.
+
+1. Add hip/bin path to the PATH  :
+    export PATH=$PATH:[MYHIP]/bin
+
+2. hipify square.cu > square.cpp
+
+3. Manually edit square.cpp to add hipLaunchParms lp to kernel parms:
+    vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)
+
+    (see square.hipref.cpp for the correct output after running hipify and the above manual step)
+
+4. make

projects/hip-tests/samples/0_Intro/square/square.cu

@@ -0,0 +1,94 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <stdio.h>
+#include <cuda_runtime.h>
+
+#define CHECK(error) \
+    if (error != cudaSuccess) { \
+      fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \
+    exit(EXIT_FAILURE);\
+	}
+
+
+/* 
+ * Square each element in the array A and write to array C.
+ */
+template <typename T>
+__global__ void
+vector_square(T *C_d, const T *A_d, size_t N)
+{
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
+
+    for (size_t i=offset; i<N; i+=stride) {
+		C_d[i] = A_d[i] * A_d[i];
+	}
+}
+
+
+int main(int argc, char *argv[])
+{
+	float *A_d, *C_d;
+	float *A_h, *C_h;
+	size_t N = 1000000;
+	size_t Nbytes = N * sizeof(float);
+
+	cudaDeviceProp props;
+	CHECK(cudaGetDeviceProperties(&props, 0/*deviceID*/));
+	printf ("info: running on device %s\n", props.name);
+
+	printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+	A_h = (float*)malloc(Nbytes);
+	CHECK(A_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
+	C_h = (float*)malloc(Nbytes);
+	CHECK(C_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
+	// Fill with Phi + i
+    for (size_t i=0; i<N; i++) 
+	{
+		A_h[i] = 1.618f + i; 
+	}
+
+	printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+	CHECK(cudaMalloc(&A_d, Nbytes));
+	CHECK(cudaMalloc(&C_d, Nbytes));
+
+
+	printf ("info: copy Host2Device\n");
+    CHECK ( cudaMemcpy(A_d, A_h, Nbytes, cudaMemcpyHostToDevice));
+
+	const unsigned blocks = 512;
+	const unsigned threadsPerBlock = 256;
+
+	printf ("info: launch 'vector_square' kernel\n");
+	vector_square <<<blocks, threadsPerBlock>>> (C_d, A_d, N);
+
+	printf ("info: copy Device2Host\n");
+    CHECK ( cudaMemcpy(C_h, C_d, Nbytes, cudaMemcpyDeviceToHost));
+
+	printf ("info: check result\n");
+    for (size_t i=0; i<N; i++)  {
+		if (C_h[i] != A_h[i] * A_h[i]) {
+			CHECK(cudaErrorUnknown);
+		}
+	}
+	printf ("PASSED!\n");
+}

projects/hip-tests/samples/0_Intro/square/square.hipref.cpp

@@ -0,0 +1,94 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <stdio.h>
+#include <hip_runtime.h>
+
+#define CHECK(error) \
+    if (error != hipSuccess) { \
+      fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
+    exit(EXIT_FAILURE);\
+	}
+
+
+/* 
+ * Square each element in the array A and write to array C.
+ */
+template <typename T>
+__global__ void
+vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)
+{
+    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    size_t stride = hipBlockDim_x * hipGridDim_x ;
+
+    for (size_t i=offset; i<N; i+=stride) {
+		C_d[i] = A_d[i] * A_d[i];
+	}
+}
+
+
+int main(int argc, char *argv[])
+{
+	float *A_d, *C_d;
+	float *A_h, *C_h;
+	size_t N = 1000000;
+	size_t Nbytes = N * sizeof(float);
+
+	hipDeviceProp_t props;
+	CHECK(hipDeviceGetProperties(&props, 0/*deviceID*/));
+	printf ("info: running on device %s\n", props.name);
+
+	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++) 
+	{
+		A_h[i] = 1.618f + i; 
+	}
+
+	printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+	CHECK(hipMalloc(&A_d, Nbytes));
+	CHECK(hipMalloc(&C_d, Nbytes));
+
+
+	printf ("info: copy Host2Device\n");
+    CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+
+	const unsigned blocks = 512;
+	const unsigned threadsPerBlock = 256;
+
+	printf ("info: launch 'vector_square' kernel\n");
+	hipLaunchKernel(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
+
+	printf ("info: copy Device2Host\n");
+    CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+
+	printf ("info: check result\n");
+    for (size_t i=0; i<N; i++)  {
+		if (C_h[i] != A_h[i] * A_h[i]) {
+			CHECK(hipErrorUnknown);
+		}
+	}
+	printf ("PASSED!\n");
+}

projects/hip-tests/samples/1_Utils/hipInfo/Makefile

@@ -0,0 +1,16 @@
+HIP_PATH?=$(shell hipconfig -p)
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=hipInfo
+
+all: install
+
+$(EXE): hipInfo.cpp
+	$(HIPCC) hipInfo.cpp -o $@
+
+install: $(EXE)
+	cp $(EXE) $(HIP_PATH)/bin
+
+
+clean:
+	rm -f *.o $(EXE)

projects/hip-tests/samples/1_Utils/hipInfo/README.md

@@ -0,0 +1,6 @@
+# hipInfo
+
+Simple tool that prints properties for each device (from hipGetDeviceProperties), and compiler info.
+    Properties includes all of the architectural feature flags for each device.
+
+Also demonstrates how to use platform-specific compilation path (testing `__HIP_PLATFORM_NVCC__` or `__HIP_PLATFORM_HCC__`)

projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp

@@ -0,0 +1,154 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <iostream>
+#include <iomanip>
+#include "hip_runtime.h"
+
+#define KNRM  "\x1B[0m"
+#define KRED  "\x1B[31m"
+#define KGRN  "\x1B[32m"
+#define KYEL  "\x1B[33m"
+#define KBLU  "\x1B[34m"
+#define KMAG  "\x1B[35m"
+#define KCYN  "\x1B[36m"
+#define KWHT  "\x1B[37m"
+
+#define failed(...) \
+    printf ("%serror: ", KRED);\
+    printf (__VA_ARGS__);\
+    printf ("\n");\
+    printf ("error: TEST FAILED\n%s", KNRM );\
+    exit(EXIT_FAILURE);
+
+#define HIPCHECK(error) \
+    if (error != hipSuccess) { \
+      printf("%serror: '%s'(%d) at %s:%d%s\n", \
+          KRED,hipGetErrorString(error), error,\
+      __FILE__, __LINE__,KNRM); \
+        failed("API returned error code.");\
+    }
+
+void printCompilerInfo ()
+{
+#ifdef __HCC__
+    printf ("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__);
+#endif
+#ifdef __NVCC__
+    printf ("compiler: nvcc\n"); 
+#endif
+}
+
+
+double bytesToGB(size_t s) 
+{
+    return (double)s / (1024.0*1024.0*1024.0);
+}
+
+void printDeviceProp (int deviceId)
+{
+    using namespace std;
+
+    const int w1 = 30;
+
+    cout << left;
+
+    cout << setw(w1) << "--------------------------------------------------------------------------------" << endl;
+    cout << setw(w1) << "device#" << deviceId << endl;
+
+    hipDeviceProp_t props;
+    HIPCHECK(hipDeviceGetProperties(&props, deviceId));
+
+
+    cout << setw(w1) << "Name: " << props.name << endl;
+    cout << setw(w1) << "multiProcessorCount: " << props.multiProcessorCount << endl;
+    cout << setw(w1) << "clockRate: " << (float)props.clockRate / 1000.0 << " Mhz" << endl;
+    cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate  / 1000.0<< " Mhz" << endl;
+	cout << setw(w1) << "totalGlobalMem" << fixed << setprecision(2) << bytesToGB(props.totalGlobalMem) << " GB" << endl << defaultfloat;
+
+	cout << setw(w1) << "sharedMemPerBlock" << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl;
+	cout << setw(w1) << "regsPerBlock" << props.regsPerBlock  << endl; 
+	cout << setw(w1) << "warpSize" << props.warpSize  << endl;
+	cout << setw(w1) << "maxThreadsPerBlock" << props.maxThreadsPerBlock << endl;
+	cout << setw(w1) << "maxThreadsDim.x" << props.maxThreadsDim[0] << endl;
+	cout << setw(w1) << "maxThreadsDim.y" << props.maxThreadsDim[1] << endl;
+	cout << setw(w1) << "maxThreadsDim.z" << props.maxThreadsDim[2] << endl;
+
+	cout << setw(w1) << "maxGridSize.x" << props.maxGridSize[0] << endl;
+	cout << setw(w1) << "maxGridSize.y" << props.maxGridSize[1] << endl;
+	cout << setw(w1) << "maxGridSize.z" << props.maxGridSize[2] << endl;
+
+
+	cout << setw(w1) << "totalConstMem" << props.totalConstMem << endl;
+	cout << setw(w1) << "major" << props.major  << endl;
+	cout << setw(w1) << "minor" << props.minor << endl;
+	cout << setw(w1) << "l2CacheSize" << props.l2CacheSize << endl;
+	cout << setw(w1) << "maxThreadsPerMultiProcessor" << props.maxThreadsPerMultiProcessor << endl;
+	cout << setw(w1) << "computeMode" << props.computeMode << endl;
+
+    cout << setw(w1) << "arch.hasGlobalInt32Atomics" <<     props.arch.hasGlobalInt32Atomics    << endl;
+    cout << setw(w1) << "arch.hasGlobalFloatAtomicExch" <<  props.arch.hasGlobalFloatAtomicExch << endl;
+    cout << setw(w1) << "arch.hasSharedInt32Atomics" <<     props.arch.hasSharedInt32Atomics    << endl;
+    cout << setw(w1) << "arch.hasSharedFloatAtomicExch" <<  props.arch.hasSharedFloatAtomicExch << endl;
+    cout << setw(w1) << "arch.hasFloatAtomicAdd" <<         props.arch.hasFloatAtomicAdd        << endl;
+    cout << setw(w1) << "arch.hasGlobalInt64Atomics" <<     props.arch.hasGlobalInt64Atomics    << endl;
+    cout << setw(w1) << "arch.hasSharedInt64Atomics" <<     props.arch.hasSharedInt64Atomics    << endl;
+    cout << setw(w1) << "arch.hasDoubles" <<                props.arch.hasDoubles               << endl;
+    cout << setw(w1) << "arch.hasWarpVote" <<               props.arch.hasWarpVote              << endl;
+    cout << setw(w1) << "arch.hasWarpBallot" <<             props.arch.hasWarpBallot            << endl;
+    cout << setw(w1) << "arch.hasWarpShuffle" <<            props.arch.hasWarpShuffle           << endl;
+    cout << setw(w1) << "arch.hasFunnelShift" <<            props.arch.hasFunnelShift           << endl;
+    cout << setw(w1) << "arch.hasThreadFenceSystem" <<      props.arch.hasThreadFenceSystem     << endl;
+    cout << setw(w1) << "arch.hasSyncThreadsExt" <<         props.arch.hasSyncThreadsExt        << endl;
+    cout << setw(w1) << "arch.hasSurfaceFuncs" <<           props.arch.hasSurfaceFuncs          << endl;
+    cout << setw(w1) << "arch.has3dGrid" <<                 props.arch.has3dGrid                << endl;
+    cout << setw(w1) << "arch.hasDynamicParallelism" <<     props.arch.hasDynamicParallelism    << endl;
+
+    cout << endl;
+
+    size_t free, total;
+
+    hipMemGetInfo(&free, &total);
+
+    cout << fixed << setprecision(2);
+    cout << setw(w1) << "memInfo.total  " <<  bytesToGB(total) << " GB" << endl;
+    cout << setw(w1) << "memInfo.free   " <<  bytesToGB(free) << " GB (" << setprecision(0) << (float)free/total * 100.0 << "%)" << endl;
+    
+}
+
+int main(int argc, char *argv[])
+{
+    using namespace std;
+
+    cout << endl;
+
+    printCompilerInfo();
+
+    int deviceCnt;
+
+    HIPCHECK(hipGetDeviceCount(&deviceCnt));
+
+    for (int i=0; i< deviceCnt; i++) {
+        printDeviceProp(i);
+    }
+
+    std::cout << std::endl;
+}