diff --git a/README.md b/README.md index 4b69485ba3..f9f653f2ba 100644 --- a/README.md +++ b/README.md @@ -21,10 +21,69 @@ New projects can be developed directly in the portable HIP C++ language and can - [clang-hipify](clang-hipify/README.md) - [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) + ## How do I get set up? See the [Installation](INSTALL.md) notes. +## Simple Example +The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree. +Programmers familiar with CUDA will also be able to quickly learn and start coding with the HIP API. +Compute kernels are launched with the "hipLaunchKernel" macro call. Here is simple example showing a +snippet of HIP API code: + +```cpp +hipMalloc(&A_d, Nbytes)); +hipMalloc(&C_d, Nbytes)); + +hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice); + +const unsigned blocks = 512; +const unsigned threadsPerBlock = 256; +hipLaunchKernel(vector_square, /* compute kernel*/ + dim3(blocks), dim3(threadsPerBlock), 0/*dynamic shared*/, 0/*stream*/, /* launch config*/ + C_d, A_d, N); /* arguments to the compute kernel */ + +hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost); +``` + + +The HIP kernel language defines builtins for determining grid and block coordinates, math functions, short vectors, +atomics, and timer functions. It also specifies additional defines and keywords for function types, address spaces, and +optimization controls. (See the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) for a full description). +Here's an example of defining a simple 'vector_square' kernel. + + + +```cpp +template +__global__ void +vector_square(T *C_d, const T *A_d, size_t N) +{ + size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + size_t stride = hipBlockDim_x * hipGridDim_x ; + + for (size_t i=offset; i