diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/README.md b/projects/hip-tests/samples/2_Cookbook/15_static_library/README.md new file mode 100644 index 0000000000..3a347d2cf7 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/README.md @@ -0,0 +1,175 @@ +# Emitting Static Library + +This sample shows how to generate a static library for a simple HIP application. We will evaluate two types of static libraries: the first type exports host functions in a static library generated with --emit-static-lib and is compatible with host linkers, and second type exports device functions in a static library made with system ar. + +Please refer to the hip_programming_guide for limitations. + +## Static libraries with host functions + +### Source files +The static library source files may contain host functions and kernel `__global__` and `__device__` functions. Here is an example (please refer to the directory host_functions). + +hipOptLibrary.cpp: +``` +#define HIP_ASSERT(status) assert(status == hipSuccess) +#define LEN 512 + +__global__ void copy(uint32_t* A, uint32_t* B) { + size_t tid = threadIdx.x + blockIdx.x * blockDim.x; + B[tid] = A[tid]; +} + +void run_test1() { + uint32_t *A_h, *B_h, *A_d, *B_d; + size_t valbytes = LEN * sizeof(uint32_t); + + A_h = (uint32_t*)malloc(valbytes); + B_h = (uint32_t*)malloc(valbytes); + for (uint32_t i = 0; i < LEN; i++) { + A_h[i] = i; + B_h[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&A_d, valbytes)); + HIP_ASSERT(hipMalloc((void**)&B_d, valbytes)); + + HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(copy, dim3(LEN/64), dim3(64), 0, 0, A_d, B_d); + HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + assert(A_h[i] == B_h[i]); + } + + HIP_ASSERT(hipFree(A_d)); + HIP_ASSERT(hipFree(B_d)); + free(A_h); + free(B_h); + std::cout << "Test Passed!\n"; +} +``` + +The above source file can be compiled into a static library, libHipOptLibrary.a, using the --emit-static-lib flag, like so: +``` +hipcc hipOptLibrary.cpp --emit-static-lib -fPIC -o libHipOptLibrary.a +``` + +### Main source files +The main() program source file may link with the above static library using either hipcc or a host compiler (such as g++). A simple source file that calls the host function inside libHipOptLibrary.a: + +hipMain1.cpp: +``` +extern void run_test1(); + +int main(){ + run_test1(); +} +``` + +To link to the static library: + +Using hipcc: +``` +hipcc hipMain1.cpp -L. -lHipOptLibrary -o test_emit_static_hipcc_linker.out +``` +Using g++: +``` +g++ hipMain1.cpp -L. -lHipOptLibrary -L/opt/rocm/hip/lib -lamdhip64 -o test_emit_static_host_linker.out +``` + +## Static libraries with device functions + +### Source files +The static library source files which contain only `__device__` functions need to be created using ar. Here is an example (please refer to the directory device_functions). + +hipDevice.cpp: +``` +#include + +__device__ int square_me(int A) { + return A*A; +} +``` + +The above source file may be compiled into a static library, libHipDevice.a, by first compiling into a relocatable object, and then placed in an archive using ar: +``` +hipcc hipDevice.cpp -c -fgpu-rdc -fPIC -o hipDevice.o +ar rcsD libHipDevice.a hipDevice.o +``` + +### Main source files +The main() program source file can link with the static library using hipcc. A simple source file that calls the device function inside libHipDevice.a: + +hipMain2.cpp: +``` +#include +#include +#include + +#define HIP_ASSERT(status) assert(status == hipSuccess) +#define LEN 512 + +extern __device__ int square_me(int); + +__global__ void square_and_save(int* A, int* B) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + B[tid] = square_me(A[tid]); +} + +void run_test2() { + int *A_h, *B_h, *A_d, *B_d; + A_h = new int[LEN]; + B_h = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A_h[i] = i; + B_h[i] = 0; + } + size_t valbytes = LEN*sizeof(int); + + HIP_ASSERT(hipMalloc((void**)&A_d, valbytes)); + HIP_ASSERT(hipMalloc((void**)&B_d, valbytes)); + + HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(square_and_save, dim3(LEN/64), dim3(64), + 0, 0, A_d, B_d); + HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A_h[i]*A_h[i] == B_h[i]); + } + + HIP_ASSERT(hipFree(A_d)); + HIP_ASSERT(hipFree(B_d)); + free(A_h); + free(B_h); + std::cout << "Test Passed!\n"; +} + +int main(){ + // Run test that generates static lib with ar + run_test2(); +} +``` + +To link to the static library: +``` +hipcc libHipDevice.a hipMain2.cpp -fgpu-rdc -o test_device_static_hipcc.out +``` + +## How to build and run this sample: +Use the make command to build the static libraries, link with it, and execute it. +- Change directory to either host or device functions folder. +- To build the static library and link the main executable, use `make all`. +- To execute, run the generated executable `./test_*.out`. + +Alternatively, use these CMake commands. +``` +cd device_functions +mkdir -p build +cd build +cmake .. +make +./test_*.out +``` + +## For More Infomation, please refer to the HIP FAQ. diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt new file mode 100644 index 0000000000..f73ba6766e --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt @@ -0,0 +1,36 @@ +project(static_lib) + +cmake_minimum_required(VERSION 3.10) + +# Search for rocm in common locations +list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm) + +# Find hip +find_package(hip REQUIRED) + +# Set compiler and linker +set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_BUILD_TYPE Release) + +# Turn static library generation ON +option(BUILD_SHARED_LIBS "Build as a shared library" OFF) + +set(CPP_SOURCES hipDevice.cpp) + +# Generate static lib libHipDevice.a +add_library(HipDevice STATIC ${CPP_SOURCES}) + +target_compile_options(HipDevice PRIVATE -fgpu-rdc) +target_link_libraries(HipDevice PRIVATE -fgpu-rdc) +target_include_directories(HipDevice PRIVATE /opt/rocm/hsa/include) + +# Create test executable that uses libHipDevice.a +set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain2.cpp) + +add_executable(test_device_static ${TEST_SOURCES}) +add_dependencies(test_device_static HipDevice) +target_compile_options(test_device_static PRIVATE -fgpu-rdc) +target_link_libraries(test_device_static PRIVATE HipDevice) +target_link_libraries(test_device_static PRIVATE -fgpu-rdc hip::host) + diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/Makefile b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/Makefile new file mode 100644 index 0000000000..26c9ed122f --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/Makefile @@ -0,0 +1,31 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif + +HIPCC=$(HIP_PATH)/bin/hipcc + +.PHONY: test + +all: $(RDC_EXE) test + +STATIC_LIB_SRC=hipDevice.cpp +STATIC_LIB=./libHipDevice.a +STATIC_MAIN_SRC=hipMain2.cpp +RDC_EXE=./test_device_static.out + +$(STATIC_LIB): + $(HIPCC) $(STATIC_LIB_SRC) -c -fgpu-rdc -fPIC -o hipDevice.o + ar rcsD $@ hipDevice.o + +# Compiles hipMain2 with hipcc and links with libHipDevice.a which contains device function. +$(RDC_EXE): $(STATIC_LIB) + $(HIPCC) $(STATIC_LIB) $(STATIC_MAIN_SRC) -fgpu-rdc -o $@ + +test: $(RDC_EXE) + $(RDC_EXE) + +clean: + rm -f $(RDC_EXE) + rm -f $(STATIC_LIB) + rm -f *.o diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipDevice.cpp b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipDevice.cpp new file mode 100644 index 0000000000..bafbc41373 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipDevice.cpp @@ -0,0 +1,28 @@ +/* + * Copyright (c) 2020-present 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 + +__device__ int square_me(int A) { + return A*A; +} + diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp new file mode 100644 index 0000000000..b041becf22 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp @@ -0,0 +1,69 @@ +/* + * Copyright (c) 2020-present 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 HIP_ASSERT(status) assert(status == hipSuccess) +#define LEN 512 + +extern __device__ int square_me(int); + +__global__ void square_and_save(int* A, int* B) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + B[tid] = square_me(A[tid]); +} + +void run_test2() { + int *A_h, *B_h, *A_d, *B_d; + A_h = new int[LEN]; + B_h = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A_h[i] = i; + B_h[i] = 0; + } + size_t valbytes = LEN*sizeof(int); + + HIP_ASSERT(hipMalloc((void**)&A_d, valbytes)); + HIP_ASSERT(hipMalloc((void**)&B_d, valbytes)); + + HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(square_and_save, dim3(LEN/64), dim3(64), + 0, 0, A_d, B_d); + HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A_h[i]*A_h[i] == B_h[i]); + } + + HIP_ASSERT(hipFree(A_d)); + HIP_ASSERT(hipFree(B_d)); + free(A_h); + free(B_h); + std::cout << "Test Passed!\n"; +} + +int main(){ + // Run test that generates static lib with ar + run_test2(); +} diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt new file mode 100644 index 0000000000..bbb738f0f1 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt @@ -0,0 +1,39 @@ +project(static_lib) + +cmake_minimum_required(VERSION 3.10) + +# Search for rocm in common locations +list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm) + +# Find hip +find_package(hip REQUIRED) + +# Set compiler and linker +set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_AR ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_BUILD_TYPE Release) + +# Turn static library generation ON +option(BUILD_SHARED_LIBS "Build as a shared library" OFF) + +set(CPP_SOURCES hipOptLibrary.cpp) + +# Generate static lib libHipOptLibrary.a. +add_library(HipOptLibrary STATIC ${CPP_SOURCES}) + +# Set-up the correct flags to generate the static library. +target_link_libraries(HipOptLibrary PRIVATE --emit-static-lib) +target_include_directories(HipOptLibrary PRIVATE /opt/rocm/hsa/include) +get_property(link_libraries TARGET HipOptLibrary PROPERTY LINK_LIBRARIES) +string (REPLACE ";" " " LINK_PROPS "${link_libraries}") +set(CMAKE_CXX_ARCHIVE_CREATE " -o ${LINK_PROPS} ") + +# Create test executable that uses libHipOptLibrary.a +set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain1.cpp) + +add_executable(test_opt_static ${TEST_SOURCES}) +add_dependencies(test_opt_static HipOptLibrary) +target_link_libraries(test_opt_static PRIVATE -lHipOptLibrary -L${CMAKE_BINARY_DIR}) +target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr hsa-runtime64::hsa-runtime64) + diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/Makefile b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/Makefile new file mode 100644 index 0000000000..2bbc26727a --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/Makefile @@ -0,0 +1,38 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif + +HIPCC=$(HIP_PATH)/bin/hipcc +GXX=g++ + +EMIT_STATIC_LIB_SRC=hipOptLibrary.cpp +EMIT_STATIC_LIB=./libHipOptLibrary.a +EMIT_STATIC_MAIN_SRC=hipMain1.cpp +HIPCC_EXE=./test_emit_static_hipcc_linker.out +HOST_EXE=./test_emit_static_host_linker.out + +.PHONY: test + +all: $(HIPCC_EXE) $(HOST_EXE) test + +$(EMIT_STATIC_LIB): + $(HIPCC) $(EMIT_STATIC_LIB_SRC) --emit-static-lib -fPIC -o $@ + +# Compiles hipMain1 with hipcc and links with libHipOptLibrary.a which contains host function. +$(HIPCC_EXE): $(EMIT_STATIC_LIB) + $(HIPCC) $(EMIT_STATIC_MAIN_SRC) -L. -lHipOptLibrary -o $@ + +# Compiles hipMain1 with g++ and links with libHipOptLibrary.a which contains host function. +$(HOST_EXE): $(EMIT_STATIC_LIB) + $(GXX) $(EMIT_STATIC_MAIN_SRC) -L. -lHipOptLibrary -L$(HIP_PATH)/lib -lamdhip64 -o $@ + +test: $(HIPCC_EXE) $(HOST_EXE) + $(HIPCC_EXE) + $(HOST_EXE) + +clean: + rm -f $(HIPCC_EXE) + rm -f $(HOST_EXE) + rm -f $(EMIT_STATIC_LIB) + rm -f *.o diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipMain1.cpp b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipMain1.cpp new file mode 100644 index 0000000000..5bb6a1e3fb --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipMain1.cpp @@ -0,0 +1,28 @@ +/* + * Copyright (c) 2020-present 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. + * */ + +extern void run_test1(); + +int main(){ + // Run test that generates static lib with -emit-static-lib + run_test1(); +} diff --git a/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipOptLibrary.cpp b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipOptLibrary.cpp new file mode 100644 index 0000000000..b969109f13 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/15_static_library/host_functions/hipOptLibrary.cpp @@ -0,0 +1,62 @@ +/* + * Copyright (c) 2020-present 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 HIP_ASSERT(status) assert(status == hipSuccess) +#define LEN 512 + +__global__ void copy(uint32_t* A, uint32_t* B) { + size_t tid = threadIdx.x + blockIdx.x * blockDim.x; + B[tid] = A[tid]; +} + +void run_test1() { + uint32_t *A_h, *B_h, *A_d, *B_d; + size_t valbytes = LEN * sizeof(uint32_t); + + A_h = (uint32_t*)malloc(valbytes); + B_h = (uint32_t*)malloc(valbytes); + for (uint32_t i = 0; i < LEN; i++) { + A_h[i] = i; + B_h[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&A_d, valbytes)); + HIP_ASSERT(hipMalloc((void**)&B_d, valbytes)); + + HIP_ASSERT(hipMemcpy(A_d, A_h, valbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(copy, dim3(LEN/64), dim3(64), 0, 0, A_d, B_d); + HIP_ASSERT(hipMemcpy(B_h, B_d, valbytes, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + assert(A_h[i] == B_h[i]); + } + + HIP_ASSERT(hipFree(A_d)); + HIP_ASSERT(hipFree(B_d)); + free(A_h); + free(B_h); + std::cout << "Test Passed!\n"; +}