HIP

https://rocm.docs.amd.com/projects/HIP/en/latest
Admin
The Heterogeneous-computing Interface for Portability (HIP) is a C++ runtime API and kernel language...
Tokens:718,065
Snippets:2,585
Trust Score:10
Update:1 year ago
Show doc for...
Tokens:
Raw
### Setup Visual Studio Dev Environment on Windows

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This PowerShell script sets up the Developer Command Prompt environment for Visual Studio on Windows. It locates the latest Visual Studio installation, imports necessary modules, and enters the developer shell with specified architecture and arguments. It also appends the HIP binary path to the system's PATH environment variable.

```powershell
$InstallationPath = Get-CimInstance MSFT_VSInstance | Sort-Object -Property Version -Descending | Select-Object -First 1 -ExpandProperty InstallLocation
Import-Module $InstallationPath\Common7\Tools\Microsoft.VisualStudio.DevShell.dll
Enter-VsDevShell -InstallPath $InstallationPath -SkipAutomaticLocation -Arch amd64 -HostArch amd64 -DevCmdArguments '-no_logo'
$env:PATH = "${env:HIP_PATH}bin;${env:PATH}"
```

--------------------------------

### Setup Visual Studio Dev Environment on Windows (NVIDIA)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This PowerShell script configures the command line environment for Visual Studio on Windows, specifically for NVIDIA toolchains. It finds the latest Visual Studio installation, loads the developer shell module, and enters the developer environment. This setup is crucial for using CUDA tools and compilers.

```powershell
$InstallationPath = Get-CimInstance MSFT_VSInstance | Sort-Object -Property Version -Descending | Select-Object -First 1 -ExpandProperty InstallLocation
Import-Module $InstallationPath\Common7\Tools\Microsoft.VisualStudio.DevShell.dll
Enter-VsDevShell -InstallPath $InstallationPath -SkipAutomaticLocation -Arch amd64 -HostArch amd64 -DevCmdArguments '-no_logo'
```

--------------------------------

### Verify AMD Compiler Installation on Linux

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command verifies that the AMD compiler (amdclang++) is installed and accessible via the command line after updating the PATH. It checks the compiler's version, confirming successful setup.

```bash
amdclang++ --version
```

--------------------------------

### Verify Compiler Installation on Windows (clang++)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command verifies that the clang++ compiler is installed and accessible after setting up the Visual Studio developer environment. It checks the compiler's version, confirming successful configuration for HIP development.

```powershell
clang++ --version
```

--------------------------------

### Cloning HIP Examples Repository

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Instructions to clone the ROCm HIP examples repository from GitHub using Git. This repository contains sample code, including a SAXPY implementation.

```shell
git clone https://github.com/amd/rocm-examples.git

```

--------------------------------

### Verify NVIDIA Compiler Installation on Linux

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command verifies that the NVIDIA CUDA compiler (nvcc) is installed and accessible via the command line. It checks the compiler's version, confirming that the CUDA toolkit is properly set up and discoverable.

```bash
nvcc --version
```

--------------------------------

### Configure and Build HIP Runtime for AMD (Shell)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/build

Configures the HIP build for AMD using CMake, specifying common directories, platform, installation prefix, and build options. It then compiles the HIP runtime using 'make' and installs it with 'sudo make install'.

```shell
cd "$CLR_DIR"
mkdir -p build; cd build
cmake -DHIP_COMMON_DIR=$HIP_DIR -DHIP_PLATFORM=amd -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=$PWD/install -DHIP_CATCH_TEST=0 -DCLR_BUILD_HIP=ON -DCLR_BUILD_OCL=OFF ..

make -j$(nproc)
sudo make install
```

--------------------------------

### Add ROCm Bin to PATH on Linux

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command exports the ROCm bin directory to the system's PATH variable on Linux. This allows the system to find and execute ROCm-related tools like the compiler from any directory. Ensure the path '/opt/rocm/bin' is correct for your installation.

```bash
export PATH=/opt/rocm/bin:${PATH}
```

--------------------------------

### Install ROCm LLVM - apt-get

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/build

Installs the ROCm LLVM development package using apt-get. This provides the necessary LLVM tools for building HIP on an AMD ROCm platform.

```shell
apt-get install rocm-llvm-dev
```

--------------------------------

### HIP Kernel Launch Example

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_driver_api

A comprehensive example demonstrating HIP kernel launch. It includes device initialization, memory allocation and transfer, module loading, function retrieval, argument preparation, and kernel execution. Conditional compilation is used for NVIDIA platform-specific initialization and cleanup.

```HIP
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>

#include <vector>

int main() {

    size_t elements = 64*1024;
    size_t size_bytes = elements * sizeof(float);

    std::vector<float> A(elements), B(elements);

    // On NVIDIA platforms the driver runtime needs to be initiated
    #ifdef __HIP_PLATFORM_NVIDIA__
    hipInit(0);
    hipDevice_t device;
    hipCtx_t context;
    HIPCHECK(hipDeviceGet(&device, 0));
    HIPCHECK(hipCtxCreate(&context, 0, device));
    #endif

    // Allocate device memory
    hipDeviceptr_t d_A, d_B;
    HIPCHECK(hipMalloc(&d_A, size_bytes));
    HIPCHECK(hipMalloc(&d_B, size_bytes));

    // Copy data to device
    HIPCHECK(hipMemcpyHtoD(d_A, A.data(), size_bytes));
    HIPCHECK(hipMemcpyHtoD(d_B, B.data(), size_bytes));

    // Load module
    hipModule_t Module;
    // For AMD the module file has to contain architecture specific object codee
    // For NVIDIA the module file has to contain PTX, found in e.g. "vcpy_isa.ptx"
    HIPCHECK(hipModuleLoad(&Module, "vcpy_isa.co"));
    // Get kernel function from the module via its name
    hipFunction_t Function;
    HIPCHECK(hipModuleGetFunction(&Function, Module, "hello_world"));

    // Create buffer for kernel arguments
    std::vector<void*> argBuffer{&d_A, &d_B};
    size_t arg_size_bytes = argBuffer.size() * sizeof(void*);

    // Create configuration passed to the kernel as arguments
    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, argBuffer.data(),
                      HIP_LAUNCH_PARAM_BUFFER_SIZE, &arg_size_bytes, HIP_LAUNCH_PARAM_END};

    int threads_per_block = 128;
    int blocks = (elements + threads_per_block - 1) / threads_per_block;

    // Actually launch kernel
    HIPCHECK(hipModuleLaunchKernel(Function, blocks, 1, 1, threads_per_block, 1, 1, 0, 0, NULL, config));

    HIPCHECK(hipMemcpyDtoH(A.data(), d_A, elements));
    HIPCHECK(hipMemcpyDtoH(B.data(), d_B, elements));

    #ifdef __HIP_PLATFORM_NVIDIA__
    HIPCHECK(hipCtxDetach(context));
    #endif

    HIPCHECK(hipFree(d_A));
    HIPCHECK(hipFree(d_B));

    return 0;
}

```

--------------------------------

### Example AMDGPU Assembly Snippet

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

A snippet of AMDGPU assembly code, likely from a HIP kernel. It demonstrates basic operations like vector moves, additions, loads, stores, and fused multiply-add instructions, typical for GPU computations.

```assembly
v_mov_b32_e32 v3, s3
v_add_co_u32_e32 v0, vcc, s2, v0
v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
global_load_dword v3, v[0:1], off
s_load_dword s0, s[6:7], 0x0
s_waitcnt vmcnt(0) lgkmcnt(0)
v_fmac_f32_e32 v3, s0, v2
global_store_dword v[0:1], v3, off
.LBB0_2:
 s_endpgm
```

--------------------------------

### HIP Kernel Launch Configuration

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Illustrates how to launch a HIP kernel on the device using the triple chevron syntax. This specifies the grid size, block size, shared memory allocation, and the stream to enqueue the operation on.

```c++
saxpy_kernel<<<dim3(grid_size), dim3(block_size), 0, hipStreamDefault>>>(a, d_x, d_y, size);

```

--------------------------------

### Get HIP Path for Makefiles

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide

This snippet demonstrates how to use the `hipconfig` tool to dynamically retrieve the HIP installation path, which can then be used in Makefiles to set the HIP_PATH variable. This ensures that Makefiles can locate HIP headers and libraries.

```makefile
HIP_PATH ?= $(shell hipconfig --path)

```

--------------------------------

### Get HIP Compiler Configuration for Standard C++ Compilers

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide

This example shows how to use `hipconfig` to obtain the necessary compiler flags (like defines and include paths) for compiling HIP code with a standard C++ compiler (e.g., gcc, icc). This is useful for host code that only calls HIP APIs without defining or launching kernels.

```bash
hipconfig --cpp_config

```

```bash
-D__HIP_PLATFORM_AMD__= -I/opt/rocm/include

```

--------------------------------

### HIP Kernel Launch Example using <<<>>> Syntax

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_cpp_language_extensions

Demonstrates launching a HIP kernel (example_kernel) using the standard triple chevron syntax. It includes memory allocation on the GPU, kernel configuration (grid and block dimensions, shared memory, stream), kernel execution, synchronization, and memory deallocation. Error checking is performed using a HIP_CHECK macro.

```C++
#include <hip/hip_runtime.h>
#include <iostream>

#define HIP_CHECK(expression)                                \ 
{                                                          \ 
    const hipError_t err = expression;                       \ 
    if(err != hipSuccess){\ 
        std::cerr << "HIP error: " << hipGetErrorString(err) \ 
            << " at " << __LINE__ << "\n";               \ 
    }                                                      \ 
}

// Performs a simple initialization of an array with the thread's index variables.
// This function is only available in device code.
__device__ void init_array(float * const a, const unsigned int arraySize){\ 
  // globalIdx uniquely identifies a thread in a 1D launch configuration.
  const int globalIdx = threadIdx.x + blockIdx.x * blockDim.x;\ 
  // Each thread initializes a single element of the array.\ 
  if(globalIdx < arraySize){\ 
    a[globalIdx] = globalIdx;\ 
  }\n}

// Rounds a value up to the next multiple.
// This function is available in host and device code.
__host__ __device__ constexpr int round_up_to_nearest_multiple(int number, int multiple){\ 
  return (number + multiple - 1)/multiple;\ 
}

__global__ void example_kernel(float * const a, const unsigned int N)\ 
{\ 
  // Initialize array.\ 
  init_array(a, N);\ 
  // Perform additional work:\ 
  // - work with the array\ 
  // - use the array in a different kernel\ 
  // - ...\ 
}

int main()\ 
{\ 
  constexpr int N = 100000000; // problem size\ 
  constexpr int blockSize = 256; //configurable block size\ 

  //needed number of blocks for the given problem size\ 
  constexpr int gridSize = round_up_to_nearest_multiple(N, blockSize);\ 

  float *a;\ 
  // allocate memory on the GPU\ 
  HIP_CHECK(hipMalloc(&a, sizeof(*a) * N));\ 

  std::cout << "Launching kernel." << std::endl;\ 
  example_kernel<<<dim3(gridSize), dim3(blockSize), 0/*example doesn't use shared memory*/, 0/*default stream*/>>>(a, N);\ 
  // make sure kernel execution is finished by synchronizing. The CPU can also\ 
  // execute other instructions during that time\ 
  HIP_CHECK(hipDeviceSynchronize());\ 
  std::cout << "Kernel execution finished." << std::endl;\ 

  HIP_CHECK(hipFree(a));\ 
}
```

--------------------------------

### Compile HIP Application (Windows AMD)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Compiles and links a single-file HIP application on Windows for AMD GPUs using clang++. It specifies include paths, libraries, and uses an environment variable for the HIP path.

```bash
clang++ .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I .\Common -lamdhip64 -L ${env:HIP_PATH}lib -O2

```

--------------------------------

### Full HIPRTC Kernel Compilation and Loading Example

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_rtc

A complete example demonstrating the HIPRTC workflow: defining a kernel, creating and compiling a program, checking for errors, retrieving compilation logs, getting the compiled code, and loading it into a module and function handle. Includes necessary headers and error checking macros.

```c++
#include <hip/hip_runtime.h>
#include <hip/hiprtc.h>

#include <iostream>
#include <string>
#include <vector>

#define CHECK_RET_CODE(call, ret_code) \
  { \
    if ((call) != ret_code) { \
      std::cout << "Failed in call: " << #call << std::endl; \
      std::abort(); \
    } \
  }
#define HIP_CHECK(call) CHECK_RET_CODE(call, hipSuccess)
#define HIPRTC_CHECK(call) CHECK_RET_CODE(call, HIPRTC_SUCCESS)

// source code for hiprtc
static constexpr auto kernel_source{
    R"(
    extern "C"
    __global__ void vector_add(float* output, float* input1, float* input2, size_t size) {
      int i = threadIdx.x;
      if (i < size) {
        output[i] = input1[i] + input2[i];
      }
    }
)"};

int main() {
  hiprtcProgram prog;
  auto rtc_ret_code = hiprtcCreateProgram(&prog,            // HIPRTC program handle
                                          kernel_source,    // kernel source string
                                          "vector_add.cpp", // Name of the file
                                          0,                // Number of headers
                                          NULL,             // Header sources
                                          NULL);            // Name of header file

  if (rtc_ret_code != HIPRTC_SUCCESS) {
    std::cout << "Failed to create program" << std::endl;
    std::abort();
  }

  hipDeviceProp_t props;
  int device = 0;
  HIP_CHECK(hipGetDeviceProperties(&props, device));
  std::string sarg = std::string("--gpu-architecture=") +
      props.gcnArchName;  // device for which binary is to be generated

  // Compilation step would follow here
  // hiprtcCompileProgram(...);
  // Error checking and log retrieval would follow
  // hiprtcGetProgramLogSize(...), hiprtcGetProgramLog(...)
  // Getting code size and data would follow
  // hiprtcGetCodeSize(...), hiprtcGetCode(...)
  // Loading module and function would follow
  // hipModuleLoadData(...), hipModuleGetFunction(...)

  // Cleanup
  hiprtcDestroyProgram(&prog);

  return 0;
}

```

--------------------------------

### CUDA __launch_bounds__ Qualifier Example

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide

Shows the syntax for CUDA's `__launch_bounds__` qualifier, which specifies the maximum number of threads per block and the minimum number of blocks per multiprocessor required for optimal occupancy.

```cuda
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MULTIPROCESSOR)

```

--------------------------------

### HIP Initialization and Version Functions

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/group___driver

This section provides documentation for functions related to HIP runtime initialization, device retrieval, and version information.

```APIDOC
## hipInit

### Description
Explicitly initializes the HIP runtime.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
unsigned int flags;
hipError_t err = hipInit(flags);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` on successful initialization.

#### Response Example
```c
hipSuccess
```

---

## hipDriverGetVersion

### Description
Returns the approximate HIP driver version.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
int driverVersion;
hipError_t err = hipDriverGetVersion(&driverVersion);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `driverVersion` with the driver version.

#### Response Example
```c
hipSuccess
```

---

## hipRuntimeGetVersion

### Description
Returns the approximate HIP Runtime version.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
int runtimeVersion;
hipError_t err = hipRuntimeGetVersion(&runtimeVersion);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `runtimeVersion` with the runtime version.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGet

### Description
Returns a handle to a compute device.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
hipDevice_t device;
int ordinal = 0;
hipError_t err = hipDeviceGet(&device, ordinal);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `device` with the handle to the specified device.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceComputeCapability

### Description
Returns the compute capability of the device.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
int major, minor;
hipDevice_t device = 0;
hipError_t err = hipDeviceComputeCapability(&major, &minor, device);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `major` and `minor` with the compute capability.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGetName

### Description
Returns an identifier string for the device.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
char name[256];
int len = 256;
hipDevice_t device = 0;
hipError_t err = hipDeviceGetName(name, len, device);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `name` with the device identifier string.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGetUuid

### Description
Returns a UUID for the device. [BETA]

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
hipUUID uuid;
hipDevice_t device = 0;
hipError_t err = hipDeviceGetUuid(&uuid, device);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `uuid` with the device UUID.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGetP2PAttribute

### Description
Returns a value for an attribute of the link between two devices.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
int value;
hipDeviceP2PAttr attr = hipDevP2PAttrAccessSharedRecon;
int srcDevice = 0;
int dstDevice = 1;
hipError_t err = hipDeviceGetP2PAttribute(&value, attr, srcDevice, dstDevice);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `value` with the requested attribute value.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGetPCIBusId

### Description
Returns a PCI Bus Id string for the device, overloaded to take an int device ID.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
char pciBusId[256];
int len = 256;
int device = 0;
hipError_t err = hipDeviceGetPCIBusId(pciBusId, len, device);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `pciBusId` with the PCI Bus Id string.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceGetByPCIBusId

### Description
Returns a handle to a compute device given its PCI Bus Id.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
int device;
const char *pciBusId = "0000:01:00.0";
hipError_t err = hipDeviceGetByPCIBusId(&device, pciBusId);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `device` with the handle to the device.

#### Response Example
```c
hipSuccess
```

---

## hipDeviceTotalMem

### Description
Returns the total amount of memory on the device.

### Method
N/A (Function call)

### Endpoint
N/A

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c
size_t bytes;
hipDevice_t device = 0;
hipError_t err = hipDeviceTotalMem(&bytes, device);
```

### Response
#### Success Response (hipSuccess)
Returns `hipSuccess` and populates `bytes` with the total memory in bytes.

#### Response Example
```c
hipSuccess
```
```

--------------------------------

### HIP Get Proc Address for Runtime Function

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_driver_api

Shows how to retrieve the address of a HIP runtime function, specifically `hipInit`, using `hipGetProcAddress()`. This example initializes the HIP runtime, obtains the function pointer, and then calls the function through the obtained pointer. It requires `hip_runtime.h` and `hip_runtime_api.h`.

```C++
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>

#include <iostream>

typedef hipError_t (*hipInit_t)(unsigned int);

int main() {
    // Initialize the HIP runtime
    hipError_t res = hipInit(0);
    if (res != hipSuccess) {
        std::cerr << "Failed to initialize HIP runtime." << std::endl;
        return 1;
    }

    // Get the address of the hipInit function
    hipInit_t hipInitFunc;
    int hipVersion = HIP_VERSION; // Use the HIP version defined in hip_runtime_api.h
    uint64_t flags = 0; // No special flags
    hipDriverProcAddressQueryResult symbolStatus;

    res = hipGetProcAddress("hipInit", (void**)&hipInitFunc, hipVersion, flags, &symbolStatus);
    if (res != hipSuccess) {
        std::cerr << "Failed to get address of hipInit()." << std::endl;
        return 1;
    }

    // Call the hipInit function using the obtained address
    res = hipInitFunc(0);
    if (res == hipSuccess) {
        std::cout << "HIP runtime initialized successfully using hipGetProcAddress()." << std::endl;
    } else {
        std::cerr << "Failed to initialize HIP runtime using hipGetProcAddress()." << std::endl;
    }

    return 0;
}

```

--------------------------------

### Device Initialization and Configuration

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__runtime__api_8h

Functions for initializing the HIP runtime, setting device flags, choosing devices, and retrieving device information.

```APIDOC
## HIP Device Initialization and Configuration

### hipInit

**Description**: Explicitly initializes the HIP runtime.

**Method**: `hipError_t hipInit(unsigned int flags)`

**Parameters**:

*   **flags** (unsigned int) - Flags to control initialization behavior.

### hipSetValidDevices

**Description**: Sets a list of devices that can be used.

**Method**: `hipError_t hipSetValidDevices(int *device_arr, int len)`

**Parameters**:

*   **device_arr** (int *) - Array of device IDs.
*   **len** (int) - The number of devices in the array.

### hipChooseDevice

**Description**: Returns a device that matches the specified properties.

**Method**: `hipError_t hipChooseDevice(int *device, const hipDeviceProp_t *prop)`

**Parameters**:

*   **device** (int *) - Pointer to store the chosen device ID.
*   **prop** (const hipDeviceProp_t *) - Pointer to a structure with desired device properties.

### hipGetDeviceFlags

**Description**: Gets the flags set for the current device.

**Method**: `hipError_t hipGetDeviceFlags(unsigned int *flags)`

**Parameters**:

*   **flags** (unsigned int *) - Pointer to store the device flags.

### hipSetDeviceFlags

**Description**: Changes the current device behavior according to the flags passed.

**Method**: `hipError_t hipSetDeviceFlags(unsigned flags)`

**Parameters**:

*   **flags** (unsigned) - Flags to set for the current device.

### hipDeviceGetTexture1DLinearMaxWidth

**Description**: Gets the maximum width for 1D linear textures on the specified device.

**Method**: `hipError_t hipDeviceGetTexture1DLinearMaxWidth(size_t *max_width, const hipChannelFormatDesc *desc, int device)`

**Parameters**:

*   **max_width** (size_t *) - Pointer to store the maximum width.
*   **desc** (const hipChannelFormatDesc *) - Description of the channel format.
*   **device** (int) - The device ID.

### hipDeviceGetLimit

**Description**: Gets resource limits of the current device.

**Method**: `hipError_t hipDeviceGetLimit(size_t *pValue, enum hipLimit_t limit)`

**Parameters**:

*   **pValue** (size_t *) - Pointer to store the limit value.
*   **limit** (enum hipLimit_t) - The type of limit to query.

### hipDeviceSetLimit

**Description**: Sets resource limits of the current device.

**Method**: `hipError_t hipDeviceSetLimit(enum hipLimit_t limit, size_t value)`

**Parameters**:

*   **limit** (enum hipLimit_t) - The type of limit to set.
*   **value** (size_t) - The value to set for the limit.

### hipDeviceGetSharedMemConfig

**Description**: Returns the bank width of shared memory for the current device.

**Method**: `hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig *pConfig)`

**Parameters**:

*   **pConfig** (hipSharedMemConfig *) - Pointer to store the shared memory configuration.

### hipDeviceSetSharedMemConfig

**Description**: Sets the bank width of shared memory on the current device.

**Method**: `hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config)`

**Parameters**:

*   **config** (hipSharedMemConfig) - The shared memory configuration to set.

### hipExtGetLinkTypeAndHopCount

**Description**: Returns the link type and hop count between two devices.

**Method**: `hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t *linktype, uint32_t *hopcount)`

**Parameters**:

*   **device1** (int) - The ID of the first device.
*   **device2** (int) - The ID of the second device.
*   **linktype** (uint32_t *) - Pointer to store the link type.
*   **hopcount** (uint32_t *) - Pointer to store the hop count.

```

--------------------------------

### Verify HIP Installation

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/install

Executes the hipconfig command with the --full flag to display detailed information about the HIP installation. This is used to verify that HIP is installed correctly and to check its configuration, especially the installation path.

```bash
/opt/rocm/bin/hipconfig --full
```

--------------------------------

### Setup Kernel Argument (hipSetupArgument)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/group___clang

Sets up a kernel argument by providing a pointer to the host memory, the size of the argument, and its offset on the argument stack. This is a prerequisite for launching kernels with specific arguments.

```c
hipError_t hipSetupArgument(const void *arg, size_t size, size_t offset)
```

--------------------------------

### Execute Compiled HIP Saxpy Sample

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command executes the compiled HIP Saxpy sample program. The output shows the calculation being performed and the first 10 elements of the resulting array, demonstrating the program's functionality.

```Shell
.\saxpy.exe
```

--------------------------------

### Query Device Properties at Runtime with HIP API

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide

Query GPU architecture feature flags at runtime using hipGetDeviceProperties() or hipDeviceGetAttribute() in host code. This example demonstrates checking for shared int32 atomic operations support using hipDeviceProp_t. A HIP_CHECK macro is included for error handling.

```c++
#include <hip/hip_runtime.h>
#include <cstdlib>
#include <iostream>

#define HIP_CHECK(expression) {                           \ 
  const hipError_t err = expression;                      \ 
  if (err != hipSuccess){                                 \ 
    std::cout << "HIP Error: " << hipGetErrorString(err)) \ 
              << " at line " << __LINE__ << std::endl;      \ 
    std::exit(EXIT_FAILURE);                              \ 
  }                                                       \ 
}

int main(){
  int deviceCount;
  HIP_CHECK(hipGetDeviceCount(&deviceCount));

  int device = 0; // Query first available GPU. Can be replaced with any
                  // integer up to, not including, deviceCount
  hipDeviceProp_t deviceProp;
  HIP_CHECK(hipGetDeviceProperties(&deviceProp, device));

  std::cout << "The queried device ";
  if (deviceProp.arch.hasSharedInt32Atomics) // portable HIP feature query
    std::cout << "supports";
  else
    std::cout << "does not support";
  std::cout << " shared int32 atomic operations" << std::endl;
}

```

--------------------------------

### Install HIP Packages on NVIDIA Systems

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/install

Installs the necessary HIP packages for NVIDIA systems using apt-get. This command installs the CUDA SDK and the HIP porting layer, enabling HIP functionality on NVIDIA GPUs.

```bash
apt-get install hip-runtime-nvidia hip-dev
```

--------------------------------

### hipExtModuleLaunchKernel

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__ext_8h

Launches a kernel with specified parameters, shared memory, and stream. It supports optional start and stop events for timing and flags for additional options.

```APIDOC
## hipExtModuleLaunchKernel

### Description
Launches kernel with parameters and shared memory on stream with arguments passed to kernel params or extra arguments.

### Method
`hipError_t`

### Endpoint
N/A (Runtime API Function)

### Parameters
#### Path Parameters
N/A

#### Query Parameters
N/A

#### Request Body
N/A

### Request Example
N/A

### Response
#### Success Response (200)
`hipError_t`: Success or error code.

#### Response Example
N/A
```

--------------------------------

### Compile HIP Application (Windows NVIDIA)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Compiles and links a single-file HIP application on Windows for NVIDIA GPUs using nvcc. It uses environment variables for the HIP path and includes common include paths.

```bash
nvcc .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I ${env:HIP_PATH}include -I .\Common -O2 -x cu

```

--------------------------------

### Programmatic Dependent Kernel Launch using Streams and Events

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_runtime_api/asynchronous

This example shows how to achieve programmatic dependent kernel launches in HIP using streams and events, similar to CUDA's functionality. `hipStreamWaitEvent` is used to ensure that a kernel launched in `stream2` does not start until the specified `event` (recorded in `stream1`) has completed, effectively creating a dependency.

```hip
hipEvent_t event;
// ... create and record event in stream1 ...

// Wait for the event before launching kernel in stream2
hipStreamWaitEvent(stream2, event);

// Launch kernel in stream2 after event completion
myKernel<<<blocks, threads, 0, stream2>>>(...);
```

--------------------------------

### HIP Peer-to-Peer Memory Access Example

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_runtime_api/multi_device

This C++ code demonstrates enabling and utilizing peer-to-peer memory access between two GPUs using the HIP runtime API. It includes memory allocation, kernel execution on each device, and inter-GPU memory copying. Error checking is performed using a HIP_CHECK macro. Ensure HIP runtime is installed and accessible.

```cpp
#include <hip/hip_runtime.h>
#include <iostream>

#define HIP_CHECK(expression)                        \
{                                                    \
    const hipError_t status = expression;            \
    if (status != hipSuccess) {                      \
        std::cerr << "HIP error " << status          \
                << ": " << hipGetErrorString(status) \
                << " at " << __FILE__ << ":"         \
                << __LINE__ << std::endl;            \
        exit(status);                                \
    }                                                \
}

__global__ void simpleKernel(double *data)
{
    int idx   = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] = idx * 2.0;
}

int main()
{
    double* deviceData0;
    double* deviceData1;
    size_t  size = 1024 * sizeof(*deviceData0);

    int deviceId0 = 0;
    int deviceId1 = 1;

    // Enable peer access to the memory (allocated and future) on the peer device.
    // Ensure the device is active before enabling peer access.
    hipSetDevice(deviceId0);
    hipDeviceEnablePeerAccess(deviceId1, 0);

    hipSetDevice(deviceId1);
    hipDeviceEnablePeerAccess(deviceId0, 0);

    // Set device 0 and perform operations
    HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
    HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
    simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
    HIP_CHECK(hipDeviceSynchronize());

    // Set device 1 and perform operations
    HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
    HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
    simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
    HIP_CHECK(hipDeviceSynchronize());

    // Use peer-to-peer access
    hipSetDevice(deviceId0);

    // Now device 0 can access memory allocated on device 1
    hipMemcpy(deviceData0, deviceData1, size, hipMemcpyDeviceToDevice);

    // Copy result from device 0
    double hostData0[1024];
    HIP_CHECK(hipSetDevice(deviceId0));
    HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));

    // Copy result from device 1
    double hostData1[1024];
    HIP_CHECK(hipSetDevice(deviceId1));
    HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));

    // Display results from both devices
    std::cout << "Device 0 data: " << hostData0[0] << std::endl;
    std::cout << "Device 1 data: " << hostData1[0] << std::endl;

    // Free device memory
    HIP_CHECK(hipFree(deviceData0));
    HIP_CHECK(hipFree(deviceData1));

    return 0;
}

```

--------------------------------

### hipHccModuleLaunchKernel (Deprecated)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__ext_8h

This HIP API is deprecated and users should migrate to `hipExtModuleLaunchKernel()` for launching kernels.

```APIDOC
## hipHccModuleLaunchKernel (Deprecated)

### Description
This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.

### Method
`hipError_t`

### Endpoint
N/A (Runtime API Function)

### Parameters
#### Path Parameters
N/A

#### Query Parameters
N/A

#### Request Body
N/A

### Request Example
N/A

### Response
#### Success Response (200)
`hipError_t`: Success or error code.

#### Response Example
N/A
```

--------------------------------

### Compile HIP Application (Linux AMD)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Compiles and links a single-file HIP application on Linux for AMD GPUs using amdclang++. It specifies include paths, libraries, and optimization levels.

```bash
amdclang++ ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -lamdhip64 -L /opt/rocm/lib -O2

```

--------------------------------

### Get HIP Device

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__runtime__api_8h_source

Gets the device handle for a specified ordinal.

```APIDOC
## hipDeviceGet

### Description
Gets the device handle for a specified ordinal.

### Method
N/A (Function Call)

### Endpoint
N/A

### Parameters
#### Path Parameters
N/A

#### Query Parameters
N/A

#### Request Body
N/A

### Request Example
```c
hipDevice_t device;
int ordinal = 0;
hipError_t err = hipDeviceGet(&device, ordinal);
```

### Response
#### Success Response (hipSuccess)
- hipDevice_t* device - Pointer to a `hipDevice_t` that will be populated with the device handle.
- int ordinal - The ordinal of the device to retrieve.

#### Response Example
```json
{
  "device": "0x...",
  "ordinal": 0
}
```
```

--------------------------------

### HIP Kernel Argument Setup (C)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__runtime__api_8h_source

This C function is used to set up kernel arguments one by one. Arguments are specified by their value, size, and offset within the argument buffer. This is typically used in conjunction with a kernel launch that reads arguments sequentially.

```c
hipError_t hipSetupArgument(const void *arg, size_t size, size_t offset);
```

--------------------------------

### Inspect Binary Artifacts with dumpbin (Windows/AMD)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

On Windows with AMD, the HIP SDK lacks native binary inspection tools. The `dumpbin` utility from the Windows SDK can extract raw data from the `.hip_fat` section of executables to identify embedded binary formats and target triples.

```powershell
dumpbin.exe /nologo /section:.hip_fat /rawdata:8 .\saxpy.exe | select -Skip 20 -First 12
```

--------------------------------

### Compile HIP Application (Linux NVIDIA)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Compiles and links a single-file HIP application on Linux for NVIDIA GPUs using nvcc. It includes necessary include paths and optimization flags.

```bash
nvcc ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -I /opt/rocm/include -O2 -x cu

```

--------------------------------

### Get Device by PCI Bus ID

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__runtime__api_8h_source

Gets the device ordinal for a given PCI bus ID string.

```APIDOC
## hipDeviceGetByPCIBusId

### Description
Gets the device ordinal for a given PCI bus ID string.

### Method
N/A (Function Call)

### Endpoint
N/A

### Parameters
#### Path Parameters
N/A

#### Query Parameters
N/A

#### Request Body
N/A

### Request Example
```c
int device;
const char* pciBusId = "0000:01:00.0";
hipError_t err = hipDeviceGetByPCIBusId(&device, pciBusId);
```

### Response
#### Success Response (hipSuccess)
- int* device - Pointer to an integer to store the device ordinal.
- const char* pciBusId - The PCI bus ID string.

#### Response Example
```json
{
  "device": 0,
  "pciBusId": "0000:01:00.0"
}
```
```

--------------------------------

### Install Python 3 - apt-get

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/build

Installs Python 3 using the apt-get package manager. This is a prerequisite for building HIP from source.

```shell
apt-get install python3
```

--------------------------------

### Start Profiler Recording (HIP) [Deprecated]

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/reference/hip_runtime_api/modules/profiler_control

Starts the recording of profiling information. This function is deprecated and users should use roctracer/rocTX instead. It returns hipErrorNotSupported.

```c
hipError_t hipProfilerStart()
```

--------------------------------

### Install CppHeaderParser - pip3

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/install/build

Installs the CppHeaderParser Python package using pip3. This package is required for handling C++ headers during the HIP build process.

```shell
pip3 install CppHeaderParser
```

--------------------------------

### List Temporary Files

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Lists all temporary files generated during the compilation process of `main.hip` when the `--save-temps` flag is used. This helps in locating the dumped device binaries for inspection.

```bash
ls main-hip-amdgcn-amd-amdhsa-*

```

--------------------------------

### Compile HIP Program with Specific SM Architectures using nvcc

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

This command compiles a HIP program (`main.hip`) using the `nvcc` compiler. It specifies the include paths for HIP and common libraries, enables optimization (`-O2`), sets the CUDA language (`-x cu`), and targets specific SM architectures (7.0 and 8.6).

```Shell
nvcc .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I ${env:HIP_PATH}include -I .\Common -O2 -x cu -arch=sm_70,sm_86
```

--------------------------------

### HIP Logging: Example ClPrint Usage

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/logging

An example demonstrating how to use the ClPrint macro to log an informational message during the initialization of the HSA stack. This specific call logs initialization events.

```cpp
ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Initializing HSA stack.");
```

--------------------------------

### hipLaunchCooperativeKernelMultiDevice

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/hip__runtime__api_8h

Launches cooperative kernels on multiple devices using a list of launch parameters.

```APIDOC
## hipError_t hipLaunchCooperativeKernelMultiDevice (hipLaunchParams *launchParamsList, int numDevices, unsigned int flags)

### Description
Launches kernels on multiple devices where thread blocks can cooperate and synchronize as they execute.

### Method
N/A

### Endpoint
N/A

### Parameters
N/A

### Request Example
N/A

### Response
N/A
```

--------------------------------

### Disassemble Offload Bundle (Linux AMD)

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/tutorial/saxpy

Extracts and disassembles a specific offload bundle (e.g., for a particular GPU architecture) from an executable using `llvm-objdump`. The output is redirected to a file for human readability.

```bash
llvm-objdump --disassemble saxpy.0.hipv4-amdgcn-amd-amdhsa--gfx942 > saxpy.s

```

--------------------------------

### Get HIPRTC Compilation Log

Source: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_rtc

Retrieves the compilation log for a HIPRTC program. If compilation fails or produces warnings, this function can be used to get the log messages. The log size is first queried, and then the log content is retrieved.

```c++
size_t logSize;
hiprtcGetProgramLogSize(prog, &logSize);

if (logSize) {
  string log(logSize, '\0');
  hiprtcGetProgramLog(prog, &log[0]);
  // Corrective action with logs
  std::cout << "Compiler log: " << log << std::endl;
}

```
Tokens:
Raw