AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-436821 Update hip samples Readme files

Browse Source 
Change-Id: I6bf3a72eac4a4242cb2dbf4e6eee73e0e1bef2ef
This commit is contained in:
Rahul Manocha
2023-12-11 03:43:30 +00:00
parent f5deb13ce3
commit 76dd8ea569
29 changed files with 471 additions and 237 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/2_Cookbook/15_static_library/README.md
+60 -158
View File
@@ -1,179 +1,81 @@
# 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++:
```
ROCM_PATH is the path where ROCM is installed. default path is /opt/rocm.
g++ hipMain1.cpp -L. -lHipOptLibrary -L<ROCM_PATH>/hip/lib -lamdhip64 -o test_emit_static_host_linker.out
# Compile to assembly and create an executable from modified asm
This sample shows how to generate the assembly code for a simple HIP source application, then re-compiling it and generating a valid HIP executable.
This sample uses a previous HIP application sample, please see [0_Intro/square](https://github.com/ROCm-Developer-Tools/HIP/blob/master/samples/0_Intro/square).
## Compiling the HIP source into assembly
Using HIP flags `-c -S` will help generate the host x86_64 and the device AMDGCN assembly code when paired with `--cuda-host-only` and `--cuda-device-only` respectively. In this sample we use these commands:
```
<ROCM_PATH>/hip/bin/hipcc -c -S --cuda-host-only -target x86_64-linux-gnu -o square_host.s square.cpp
<ROCM_PATH>/hip/bin/hipcc -c -S --cuda-device-only --offload-arch=gfx900 --offload-arch=gfx906 --offload-arch=gfx908 --offload-arch=gfx1010 --offload-arch=gfx1030 --offload-arch=gfx1100 --offload-arch=gfx1101 --offload-arch=gfx1102 --offload-arch=gfx1103 square.cpp
```
## Static libraries with device functions
The device assembly will be output into two separate files:
- square-hip-amdgcn-amd-amdhsa-gfx900.s
- square-hip-amdgcn-amd-amdhsa-gfx906.s
- square-hip-amdgcn-amd-amdhsa-gfx908.s
- square-hip-amdgcn-amd-amdhsa-gfx1010.s
- square-hip-amdgcn-amd-amdhsa-gfx1030.s
- square-hip-amdgcn-amd-amdhsa-gfx1100.s
- square-hip-amdgcn-amd-amdhsa-gfx1101.s
- square-hip-amdgcn-amd-amdhsa-gfx1102.s
- square-hip-amdgcn-amd-amdhsa-gfx1103.s
### 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).
You may modify `--offload-arch` flag to build other archs and choose to enable or disable xnack and sram-ecc.
hipDevice.cpp:
**Note:** At this point, you may evaluate the assembly code, and make modifications if you are familiar with the AMDGCN assembly language and architecture.
## Compiling the assembly into a valid HIP executable
If valid, the modified host and device assembly may be compiled into a HIP executable. The host assembly can be compiled into an object using this command:
```
#include <hip/hip_runtime.h>
__device__ int square_me(int A) {
return A*A;
}
<ROCM_PATH>/hip/bin/hipcc -c square_host.s -o square_host.o
```
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:
However, the device assembly code will require a few extra steps. The device assemblies needs to be compiled into device objects, then offload-bundled into a HIP fat binary using the clang-offload-bundler, then llvm-mc embeds the binary inside of a host object using the MC directives provided in `hip_obj_gen.mcin`. The output is a host object with an embedded device object. Here are the steps for device side compilation into an object:
```
hipcc hipDevice.cpp -c -fgpu-rdc -fPIC -o hipDevice.o
ar rcsD libHipDevice.a hipDevice.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx900 square-hip-amdgcn-amd-amdhsa-gfx900.s -o square-hip-amdgcn-amd-amdhsa-gfx900.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 square-hip-amdgcn-amd-amdhsa-gfx906.s -o square-hip-amdgcn-amd-amdhsa-gfx906.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx908 square-hip-amdgcn-amd-amdhsa-gfx908.s -o square-hip-amdgcn-amd-amdhsa-gfx908.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1010 square-hip-amdgcn-amd-amdhsa-gfx1010.s -o square-hip-amdgcn-amd-amdhsa-gfx1010.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1030 square-hip-amdgcn-amd-amdhsa-gfx1030.s -o square-hip-amdgcn-amd-amdhsa-gfx1030.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1100 square-hip-amdgcn-amd-amdhsa-gfx1100.s -o square-hip-amdgcn-amd-amdhsa-gfx1100.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1101 square-hip-amdgcn-amd-amdhsa-gfx1101.s -o square-hip-amdgcn-amd-amdhsa-gfx1101.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1102 square-hip-amdgcn-amd-amdhsa-gfx1102.s -o square-hip-amdgcn-amd-amdhsa-gfx1102.o
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx1103 square-hip-amdgcn-amd-amdhsa-gfx1103.s -o square-hip-amdgcn-amd-amdhsa-gfx1103.o
<ROCM_PATH>/llvm/bin/clang-offload-bundler -type=o -bundle-align=4096 -targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa-gfx900,hip-amdgcn-amd-amdhsa-gfx906,hip-amdgcn-amd-amdhsa-gfx908,hip-amdgcn-amd-amdhsa-gfx1010,hip-amdgcn-amd-amdhsa-gfx1030,hip-amdgcn-amd-amdhsa-gfx1100,hip-amdgcn-amd-amdhsa-gfx1101,hip-amdgcn-amd-amdhsa-gfx1102,hip-amdgcn-amd-amdhsa-gfx1103 -inputs=/dev/null,square-hip-amdgcn-amd-amdhsa-gfx900.o,square-hip-amdgcn-amd-amdhsa-gfx906.o,square-hip-amdgcn-amd-amdhsa-gfx908.o,square-hip-amdgcn-amd-amdhsa-gfx1010.o,square-hip-amdgcn-amd-amdhsa-gfx1030.o,square-hip-amdgcn-amd-amdhsa-gfx1100.o,square-hip-amdgcn-amd-amdhsa-gfx1101.o,square-hip-amdgcn-amd-amdhsa-gfx1102.o,square-hip-amdgcn-amd-amdhsa-gfx1103.o -outputs=offload_bundle.hipfb
<ROCM_PATH>/llvm/bin/llvm-mc -triple x86_64-unknown-linux-gnu hip_obj_gen.mcin -o square_device.o --filetype=obj
```
### 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:
**Note:** Using option `-bundle-align=4096` only works on ROCm 4.0 and newer compilers. Also, the architecture must match the same arch as when compiling to assembly.
hipMain2.cpp:
Finally, using the system linker, hipcc, or clang, link the host and device objects into an executable:
```
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#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();
}
<ROCM_PATH>/hip/bin/hipcc square_host.o square_device.o -o square_asm.out
```
To link to the static library:
## How to build and run this sample:
- Build the sample using cmake
```
hipcc libHipDevice.a hipMain2.cpp -fgpu-rdc -o test_device_static_hipcc.out
$ mkdir build; cd build
$ cmake .. -DCMAKE_PREFIX_PATH=/opt/rocm
$ make
```
## 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.
- Execute sample
```
cd device_functions
mkdir -p build
cd build
cmake ..
make
./test_*.out
$ ./square_asm.out
info: running on device AMD Radeon Graphics
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!
```
It is recommended to use Visual Studio's command prompt for this sample due to requirement of MS Librarian tool - LIB.exe on windows platform.
Override CMAKE_C_COMPILER and CMAKE_CXX_COMPILER to hipcc as Visual Studio's compiler would use cl.exe as default compiler.
i.e. cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=%HIP_PATH% -DCMAKE_C_COMPILER=%HIP_PATH%/bin/hipcc.bat -DCMAKE_CXX_COMPILER=%HIP_PATH%/bin/hipcc.bat ..
## For More Infomation, please refer to the HIP FAQ.
**Note:** Currently, defined arch is `gfx900`, `gfx906`, `gfx908`, `gfx1010`,`gfx1030`,`gfx1100`,`gfx1101`,`gfx1102` and `gfx1103`. Any undefined arch can be modified with make argument `GPU_ARCHxx`.
## For More Information, please refer to the HIP FAQ.