Hipify
- Convert CUDA to portable C++,
hipify
- C++ kernel language ( C++11/14/17 features )
- C runtime API
- same performance as native CUDA
__global__ void add_kernel(const T * a,
const T * b,
T * c){
const int i = blockDim.x * blockIdx.x + threadIdx.x;
c[i] = a[i] + b[i];}
void CUDAStream<T>::add(){
add_kernel<<<array_size/TBSIZE, TBSIZE>>>(d_a, d_b, d_c);
check_error(); //..
}
__global__ void add_kernel(hipLaunchParm lp,
const T * a, const T * b,
T * c){
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
c[i] = a[i] + b[i];
}
void HIPStream<T>::add(){
hipLaunchKernel(HIP_KERNEL_NAME(add_kernel),
dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0,
d_a, d_b, d_c); check_error(); //...
}
HIP summary
very interesting tool to get started with production or legacy code
still low-level CUDA programming
HIP library eco-system available: hipBlas, hipFFT, hipRNG, machine learning acceleration, ...
HC
- C++ parallel runtime and API
- based on C++AMP in
hc namespace plus C++14
- (asynchronous) copy commands for host-device i/o
- explicit pointer-based memory allocation (am_alloc / am_free)
- hc::accelerator_view, hc::array_view, hc::completion_future
- device specific 'instrinsics' (wavefront shuffle, bit extraction, atomics)
Very similar to thrust, boost.compute, sycl.
HC API Overview
#include "Stream.h"
#include "hc.hpp"
template <class T>
class HCStream : public Stream<T>
{
protected:
unsigned int array_size;
hc::array<T,1> d_a;
hc::array<T,1> d_b;
hc::array<T,1> d_c;
//...
template <class T>
void HCStream<T>::init_arrays(T _a, T _b, T _c)
{
hc::array_view<T,1> view_a(this->d_a);
hc::parallel_for_each(hc::extent<1>(array_size)
, [=](hc::index<1> i) [[hc]] {
view_a[i] = _a;
});
//...
template <class T>
void HCStream<T>::add()
{
hc::array_view<T,1> view_a(this->d_a);
hc::array_view<T,1> view_b(this->d_b);
hc::array_view<T,1> view_c(this->d_c);
hc::parallel_for_each(hc::extent<1>(array_size)
, [=](hc::index<1> i) [[hc]] {
view_c[i] = view_a[i]+view_b[i];
});
Let's compare the results
Comparing to HBM2 and GDDR5X
HC can be low-level too (WIP)
hc::parallel_for_each(tiled_ex,
[=,
&view_a,
&view_b,
&partial](const hc::tiled_index<1>& tidx) [[hc]] {
auto gidx = tidx.global[0];
T r = T{0}; // Assumes reduction op is addition.
while (gidx < view_a.get_extent().size()) {
r += view_a[gidx] * view_b[gidx]; //dot-product
gidx += domain_sz;
}
tile_static T tileData[TBSIZE];
tileData[tidx.local[0]] = r;
tidx.barrier.wait_with_tile_static_memory_fence();
for (auto h = TBSIZE / 2; h; h /= 2) {
if (tidx.local[0] < h) {
tileData[tidx.local[0]] += tileData[tidx.local[0] + h];
}
tidx.barrier.wait_with_tile_static_memory_fence();
}
if (tidx.global == tidx.tile_origin) partial[tidx.tile] = tileData[0];
Concurrency constructs
- asynchronous operations (memory copies, kernel launches) return completion future
std::vector<float> payload (/*pick a number*/);
hc::array<float,1> d_payload(payload.size());
hc::completion_future when_done = hc::async_copy(payload.begin(),
payload.end(),
d_payload);
when_done.then(call_kernel_functor); //continuation function!
- for me hc::completion_future::then API not production ready yet:
template<typename functor >
void then (const functor &func);//just a callback for now
Concurrency TS?
std::vector<hc::completion_future> streams(n);
for(hc::completion_future when_done : streams){
when_done = hc::async_copy(payload_begin_itr,
payload_end_itr,
d_payload_view);
when_done.then(parallel_for_each(/*do magic*/))
.then(hc::async_copy(d_payload_view,result_begin_itr));
}
hc::when_all(streams);
concurrency constructs are the glue code of host-device interactions!
(see when_all, co_await and friends)
