From d49392904a89a7fa117d42c06fadc42a32adfc44 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 26 Jan 2016 20:14:33 -0600 Subject: [PATCH] Initial commit for GPUOpen Launch [ROCm/hip-tests commit: 4eafc92922785546c796697f8040cbe9208879f4] --- .../samples/0_Intro/bit_extract/Makefile | 24 +++ .../samples/0_Intro/bit_extract/README.md | 4 + .../0_Intro/bit_extract/bit_extract.cpp | 102 ++++++++++++ .../hip-tests/samples/0_Intro/square/Makefile | 18 ++ .../samples/0_Intro/square/README.md | 15 ++ .../samples/0_Intro/square/square.cu | 94 +++++++++++ .../samples/0_Intro/square/square.hipref.cpp | 94 +++++++++++ .../samples/1_Utils/hipInfo/Makefile | 16 ++ .../samples/1_Utils/hipInfo/README.md | 6 + .../samples/1_Utils/hipInfo/hipInfo.cpp | 154 ++++++++++++++++++ 10 files changed, 527 insertions(+) create mode 100644 projects/hip-tests/samples/0_Intro/bit_extract/Makefile create mode 100644 projects/hip-tests/samples/0_Intro/bit_extract/README.md create mode 100644 projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp create mode 100644 projects/hip-tests/samples/0_Intro/square/Makefile create mode 100644 projects/hip-tests/samples/0_Intro/square/README.md create mode 100644 projects/hip-tests/samples/0_Intro/square/square.cu create mode 100644 projects/hip-tests/samples/0_Intro/square/square.hipref.cpp create mode 100644 projects/hip-tests/samples/1_Utils/hipInfo/Makefile create mode 100644 projects/hip-tests/samples/1_Utils/hipInfo/README.md create mode 100644 projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile new file mode 100644 index 0000000000..d0a2d4bc75 --- /dev/null +++ b/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) + diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/README.md b/projects/hip-tests/samples/0_Intro/bit_extract/README.md new file mode 100644 index 0000000000..025b7d3072 --- /dev/null +++ b/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. diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp new file mode 100644 index 0000000000..d32b65c15d --- /dev/null +++ b/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 +#include +#include + + +#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> 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> 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); + } + } +} diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile new file mode 100644 index 0000000000..9fb03d867e --- /dev/null +++ b/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 diff --git a/projects/hip-tests/samples/0_Intro/square/README.md b/projects/hip-tests/samples/0_Intro/square/README.md new file mode 100644 index 0000000000..56b558e5bc --- /dev/null +++ b/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 diff --git a/projects/hip-tests/samples/0_Intro/square/square.cu b/projects/hip-tests/samples/0_Intro/square/square.cu new file mode 100644 index 0000000000..996344ed49 --- /dev/null +++ b/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 +#include + +#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 +__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>> (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 +#include + +#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 +__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 +#include +#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; +}