Skip to content

Commit

Permalink
Fix review comments
Browse files Browse the repository at this point in the history
  • Loading branch information
MKKnorr committed Oct 1, 2024
1 parent 835a788 commit 958ca19
Show file tree
Hide file tree
Showing 2 changed files with 21 additions and 10 deletions.
19 changes: 10 additions & 9 deletions docs/faq.rst
Original file line number Diff line number Diff line change
Expand Up @@ -76,6 +76,8 @@ As a result, the OpenCL syntax is a lot different from HIP, and porting tools
have to perform complex transformations, especially when it comes to templates
or other C++ features in kernels.

An overview over the different syntaxes can be seen `:doc:here<reference/terms>`.

Can I install both CUDA and ROCm on the same machine?
-----------------------------------------------------

Expand All @@ -84,8 +86,7 @@ Yes. Beware, that you still need a compatible GPU to actually run the compiled c
HIP detected my platform (``amd`` vs ``nvidia``) incorrectly. What should I do?
-------------------------------------------------------------------------------

If HIP is unable to detect the correct platform, or you want to manually specify
it, set the ``HIP_PLATFORM`` environment variable to either ``amd`` or ``nvidia``.
See the `:ref:HIP porting guide<porting_identify_hip_runtime>` under the section "Identifying HIP Runtime".

On NVIDIA platforms, can I mix HIP code with CUDA code?
-------------------------------------------------------
Expand All @@ -111,7 +112,7 @@ What is HIP-Clang?
------------------

HIP-Clang is a Clang/LLVM based compiler to compile HIP programs for AMD
platforms. It's executable is ``amdclang++``.
platforms. Its executable is called ``amdclang++`` on Linux and ``clang++`` on Windows.

Can I link HIP device code with host code compiled with another compiler such as gcc, icc, or clang?
-----------------------------------------------------------------------------------------------------------
Expand Down Expand Up @@ -139,7 +140,7 @@ The following is the HIP device code, assumed to be saved in ``device.hip``:
const int x = threadIdx.x + blockIdx.x * blockDim.x;
if(x < size){array[x] = x;}
};

extern "C"{
hipError_t callKernel(int blocks, int threadsPerBlock, double* array, size_t size){
kernel<<<blocks, threadsPerBlock, 0, hipStreamDefault>>>(array, size);
Expand All @@ -154,32 +155,32 @@ The following is the host code, written in C, saved in ``host.c``:
#include <hip/hip_runtime_api.h>
#include <stdio.h>
#include <stdlib.h>
#define HIP_CHECK(c) { \
if (c != hipSuccess){ \
printf("HIP Error : %s", hipGetErrorString(c)); \
printf(" %s %d\n", __FILE__, __LINE__); \
exit(c); \
} \
}
// Forward declaration - the implementation needs to be compiled with
// a device compiler like hipcc or amdclang++
hipError_t callKernel(int blocks, int threadsPerBlock, double* array, size_t size);
int main(int argc, char** argv) {
int blocks = 1024;
int threadsPerBlock = 256;
size_t arraySize = blocks * threadsPerBlock;
double* d_array;
double* h_array;
h_array = (double*)malloc(arraySize * sizeof(double));
HIP_CHECK(hipMalloc((void**)&d_array, arraySize * sizeof(double)));
HIP_CHECK(callKernel(blocks, threadsPerBlock, d_array, arraySize));
HIP_CHECK(hipMemcpy(h_array, d_array, arraySize * sizeof(double), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(d_array));
free(h_array);
return 0;
}
Expand Down
12 changes: 11 additions & 1 deletion docs/how-to/hip_porting_guide.md
Original file line number Diff line number Diff line change
Expand Up @@ -247,6 +247,8 @@ Makefiles can use the following syntax to conditionally provide a default HIP_PA
HIP_PATH ?= $(shell hipconfig --path)
```
(porting_identify_hip_runtime)=
## Identifying HIP Runtime
HIP can depend on rocclr, or CUDA as runtime
Expand All @@ -257,7 +259,15 @@ ROCclr is a virtual device interface that HIP runtimes interact with different b
* NVIDIA platform
On NVIDIA platform, HIP is just a thin layer on top of CUDA.
On non-AMD platform, HIP runtime determines if CUDA is available and can be used. If available, HIP_PLATFORM is set to `nvidia` and underneath CUDA path is used.
The environment variable `HIP_PLATFORM` can be used to specify, what runtime
should be used. The platform is detected automatically by HIP. When an AMD
graphics driver and an AMD GPU is detected, `HIP_PLATFORM` is set to `amd`.
If both runtimes are installed, and a specific one should be used, or HIP is
unable to detect the runtime, setting the environment variable manually tells
`hipcc` what compilation path to choose.
To use the CUDA compilation path, set the environment variable
to `HIP_PLATFORM=nvidia`.
## `hipLaunchKernelGGL`
Expand Down

0 comments on commit 958ca19

Please sign in to comment.