From d5142afe5800fe1c0219926a2cba73fce6d0d477 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Wed, 7 Jul 2021 20:24:46 -0400 Subject: [PATCH] SWDEV-292273 - Add cmake tests Add test on CXX and Fortran build in cmake. Add test on hip::device interface linking in cmake. Change-Id: I3fe6dba05a7a140a9a19df107b7885df83d28042 --- .../18_cmake_hip_device/CMakeLists.txt | 32 +++++ .../2_Cookbook/18_cmake_hip_device/README.md | 16 +++ .../2_Cookbook/18_cmake_hip_device/square.cpp | 99 ++++++++++++++ .../2_Cookbook/19_cmake_lang/CMakeLists.txt | 14 ++ .../19_cmake_lang/MatrixTranspose.cpp | 122 ++++++++++++++++++ samples/2_Cookbook/19_cmake_lang/README.md | 20 +++ .../2_Cookbook/19_cmake_lang/TestFortran.F90 | 3 + 7 files changed, 306 insertions(+) create mode 100644 samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt create mode 100644 samples/2_Cookbook/18_cmake_hip_device/README.md create mode 100644 samples/2_Cookbook/18_cmake_hip_device/square.cpp create mode 100644 samples/2_Cookbook/19_cmake_lang/CMakeLists.txt create mode 100644 samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp create mode 100644 samples/2_Cookbook/19_cmake_lang/README.md create mode 100644 samples/2_Cookbook/19_cmake_lang/TestFortran.F90 diff --git a/samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt b/samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt new file mode 100644 index 0000000000..691c4682a1 --- /dev/null +++ b/samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt @@ -0,0 +1,32 @@ +# Copyright (C) 2020-2021 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. + +project(cmake_hip_device_test) + +cmake_minimum_required(VERSION 3.10.2) + +# Find hip +find_package(hip REQUIRED) + +# Create the excutable +add_executable(test_cpp square.cpp) + +# Link with HIP +target_link_libraries(test_cpp hip::device) \ No newline at end of file diff --git a/samples/2_Cookbook/18_cmake_hip_device/README.md b/samples/2_Cookbook/18_cmake_hip_device/README.md new file mode 100644 index 0000000000..d9e31d806e --- /dev/null +++ b/samples/2_Cookbook/18_cmake_hip_device/README.md @@ -0,0 +1,16 @@ +### This will test linking hip::device interface in cmake +I. Build +mkdir -p build; cd build +rm -rf *; CXX=`hipconfig -l`/clang++ cmake .. +make + +II. Test +$ ./test_cpp +info: running on device Vega 20 [Radeon Pro Vega 20] +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +info: copy Host2Device +info: launch 'vector_square' kernel +info: copy Device2Host +info: check result +PASSED! \ No newline at end of file diff --git a/samples/2_Cookbook/18_cmake_hip_device/square.cpp b/samples/2_Cookbook/18_cmake_hip_device/square.cpp new file mode 100644 index 0000000000..971488d994 --- /dev/null +++ b/samples/2_Cookbook/18_cmake_hip_device/square.cpp @@ -0,0 +1,99 @@ +#include "hip/hip_runtime.h" +/* +Copyright (c) 2015-2021 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(cmd) \ +{\ + hipError_t error = cmd;\ + 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(T *C_d, 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 + +// hip header file +#include "hip/hip_runtime.h" + + +#define WIDTH 1024 + + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +// 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; + + 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]; + } + } +} + +int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + int i; + int errors; + + 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; + } + + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + + // 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); + + // Memory transfer from device to host + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; +} diff --git a/samples/2_Cookbook/19_cmake_lang/README.md b/samples/2_Cookbook/19_cmake_lang/README.md new file mode 100644 index 0000000000..5ba5ef98e8 --- /dev/null +++ b/samples/2_Cookbook/19_cmake_lang/README.md @@ -0,0 +1,20 @@ +### This will test cmake lang support: CXX and Fortran +I. Prepare +1) You must install cmake version 3.18 or above to support LINK_LANGUAGE. + Otherwise, Fortran build will fail. + To download the latest cmake, see https://cmake.org/download/. +2) If there is no Fortran on your system, you must install it via, + sudo apt install gfortran + +II. Build +mkdir -p build; cd build +rm -rf *; CXX=`hipconfig -l`/clang++ FC=$(which gfortran) cmake .. +make + +III. Test +# ./test_fortran + Succeeded testing Fortran! + +# ./test_cpp +Device name Device 66a7 +PASSED! diff --git a/samples/2_Cookbook/19_cmake_lang/TestFortran.F90 b/samples/2_Cookbook/19_cmake_lang/TestFortran.F90 new file mode 100644 index 0000000000..14f29f7270 --- /dev/null +++ b/samples/2_Cookbook/19_cmake_lang/TestFortran.F90 @@ -0,0 +1,3 @@ + program fortran_test_program + print *, "Succeeded testing Fortran!" + end program \ No newline at end of file