Section author: Mia Doričić, Vedran Miletić

rocPRIM: ROCm Parallel Primitives

Kod: https://github.com/ROCmSoftwarePlatform/rocPRIM

Dokumentacija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/

rocPRIM implementira osnovne funkcije za paralelno računanje na grafičkim procesorima Radeon.

Terminologija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/group__rocprim__glossary.html (TODO prevesti i objasniti)

Paralelne primitive su alati pomoću kojih možemo implementirati željene paralelne algoritme. Za svaku od funkcija postoji više varijanti ovisno o parametrima.

Njih dijelimo na:

Block-wide

Dokumentacija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/group__blockmodule.html

Neki koji se često koriste su:

  • load() vrši učitavanje podataka iz memorije

  • reduce() vrši operaciju smanjivanja na razne načine (npr. minimum, maksimum, zbroj, produkt)

  • scan() vrši uključno i isključno skeniranje (tzv. prefiksni zbroj)

  • sort() vrši sortiranje (ključeva, parova ključeva…)

  • store() vrši pohranu skupova podataka u memoriju

Device-wide (mogu koristiti čitav GPU)

Dokumentacija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/group__devicemodule.html

Neki koji se često koriste su:

  • histogram_even() računa histogram

  • merge() izvodi sjedinjavanje

  • partition() stvara particije

  • reduce(...) vrši operaciju smanjivanja na razne načine (npr. minimum, maksimum, zbroj, produkt)

  • select(...) odabire kontretnog primitivca

Warp-wide (grupa niti koje se izvode na jednom CU)

Dokumentacija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/group__warpmodule.html

  • reduce() vrši operaciju smanjivanja na razne načine (npr. minimum, maksimum, zbroj, produkt)

  • scan() vrši uključno i isključno skeniranje (tzv. prefiksni zbroj)

  • sort() vrši sortiranje (ključeva, parova ključeva…)

Pomoćni alati

Dokumentacija: https://codedocs.xyz/ROCmSoftwarePlatform/rocPRIM/group__utilsmodule__functional.html

  • less gleda manju vrijednost

  • less_equal gleda vrijednost manje ili jednako

  • greater gleda veću vrijednost

  • greater_equal gleda vrijednost veće ili jednako

  • equal_to gleda vrijednost jednaku danom elementu

  • not_equal_to gleda vrijednost različitu od danog elementa

  • plus vrši zbrajanje

  • minus vrši oduzimanje

  • multiplies vrši množenje

  • maximum postavlja najveću moguću vrijednost

  • minimum postavlja najmanju moguću vrijednost

  • identity ostavlja sve elemente istima, ne mijenja ništa

Primjeri

Službeni primjer example/example_temporary_storage.cpp: https://github.com/ROCmSoftwarePlatform/rocPRIM/blob/develop/example/example_temporary_storage.cpp

Program pokazuje korištenje 4 vrste memorije za izvođenje redukcije (scan) zbrajanjem.

Promotrimo funkciju main() od koje program kreće:

int main()
{
  // Initializing HIP device
  hipDeviceProp_t device_properties;
  HIP_CHECK(hipGetDeviceProperties(&device_properties, 0));

  // Show device info
  printf("Selected device:         %s  \n", device_properties.name              );
  printf("Available global memory: %lu \n", device_properties.totalGlobalMem    );
  printf("Shared memory per block: %lu \n", device_properties.sharedMemPerBlock );
  printf("Warp size:               %d  \n", device_properties.warpSize          );
  printf("Max threads per block:   %d  \n", device_properties.maxThreadsPerBlock);

  // Running kernels
  run_example_global_memory_storage<int>(1024);
  run_example_shared_memory<int>(1024);
  run_example_union_storage_types<int>(1024);
  run_example_dynamic_shared_memory<int>(1024);
}

U funkcije se prvo inicijalizira HIP uređaj i provjerava da je inicijalizacija uspješna.

Zatim se objekt device_properties koristi za ispis svojstava uređaja kao što su ime, dostupna globalna memorija (na suvremnim grafičkim karticama to je GDDR6 ili HBM memorija koja se navodi u specifikacijama), dijeljena memorija po bloku…

Kao zadnji korak, pozivaju se sve funkcije za pokretanje svih zrna u ovom programu s parametrima koji definiraju veličinu vektora. U nastavku analiziramo svaku od pojedinih funkcija.

Funkcija za pokretanje zrna koje koristi globalnu memoriju je oblika:

template<class T>
void run_example_global_memory_storage(size_t size)
{
  constexpr unsigned int block_size = 256;
  // Make sure size is a multiple of block_size
  auto grid_size = (size + block_size - 1) / block_size;
  size = block_size * grid_size;

  // Generate input on host and copy it to device
  std::vector<T> host_input = get_random_data<T>(size, 0, 1000);
  // Generating expected output for kernel
  std::vector<T> host_expected_output = get_expected_output<T>(host_input, block_size);
  // For reading device output
  std::vector<T> host_output(size);

  // Device memory allocation
  T * device_input;
  T * device_output;
  HIP_CHECK(hipMalloc(&device_input, host_input.size() * sizeof(typename decltype(host_input)::value_type)));
  HIP_CHECK(hipMalloc(&device_output, host_output.size() * sizeof(typename decltype(host_output)::value_type)));

  // Writing input data to device memory
  hip_write_device_memory<T>(device_input, host_input);

  // Allocating temporary storage in global memory
  using storage_type = typename rocprim::block_scan<T, block_size>::storage_type;
  storage_type *global_storage;
  HIP_CHECK(hipMalloc(&global_storage, (grid_size * sizeof(storage_type))));

  // Launching kernel example_shared_memory
  hipLaunchKernelGGL(
    HIP_KERNEL_NAME(example_global_memory_storage<block_size, T>),
    dim3(grid_size), dim3(block_size),
    0, 0,
    device_input, device_output, global_storage
  );

  // Reading output from device
  hip_read_device_memory<T>(host_output, device_output);

  // Validating output
  OUTPUT_VALIDATION_CHECK(
    validate_device_output(host_output, host_expected_output)
  );

  HIP_CHECK(hipFree(device_input));
  HIP_CHECK(hipFree(device_output));
  HIP_CHECK(hipFree(global_storage));

  std::cout << "Kernel run_example_global_memory_storage run was successful!" << std::endl;
}

U deklaraciji ove funkcije možemo vidjeti da je inicijalizirana varijabla size tipa size_t, koju ćemo koristiti u sljedećim koracima.

Na početku funkcije možemo vidjeti da je veličina bloka postavljena na 265:

constexpr unsigned int block_size = 256;

Kako nam je kasnije u kodu potrebna veličina mreže, moramo ju saznati, a to ćemo postići na sljedeći način, pritom koristeći prije spomenutu varijablu size koja mora biti umnožak od veličine bloka i mreže:

auto grid_size = (size + block_size - 1) / block_size;
size = block_size * grid_size;

U sljedećem koraku funkcije generiraju se unos, očekivani ispis i dobiveni ispis. Oni su svi tipa std::vector<T>. (Za više informacija pogledajte: https://en.cppreference.com/w/cpp/container/vector):

std::vector<T> host_input = get_random_data<T>(size, 0, 1000);
std::vector<T> host_expected_output = get_expected_output<T>(host_input, block_size);
std::vector<T> host_output(size);

Slijedi alokacija memorije uređaja. Potrebno je alocirati memoriju i za unos i za ispis. Putem HIP_CHECK provjeravamo hoće li alokacija biti uspješna. Alokaciju memorije za unos i ispis vršimo putem funkcije hipMalloc na sljedeći način (za više informacija o hipMalloc možete pogledati ovdje: https://rocmdocs.amd.com/en/latest/ROCm_API_References/HIP_API/Memory-Management.html) :

T * device_input;
T * device_output;
HIP_CHECK(hipMalloc(&device_input, host_input.size() * sizeof(typename decltype(host_input)::value_type)));
HIP_CHECK(hipMalloc(&device_output, host_output.size() * sizeof(typename decltype(host_output)::value_type)));

Koristimo prethodno definirane varijable device_input i host_input za upisivanje unosa na za to predodređenu memoriju na uređaju.

Slijedi alokacija privremene pohrane podataka u globalnoj memoriji:

using storage_type = typename rocprim::block_scan<T, block_size>::storage_type;
storage_type *global_storage;
HIP_CHECK(hipMalloc(&global_storage, (grid_size * sizeof(storage_type))));

Ovdje se javlja prethodno definirana funkcija rocPRIMa, block_scan. Za više informacija pogledajte podnaslov Block-wide. Ponovno se, slično prethodnim koracima, putem HIP_CHECK traži provjera uspješne alokacije memorije hipMalloc.

Pokreće se zrno koje koristi globalnu memoriju za pohranu (za više informacija o funkciji hipLaunchKernellGGL koja se koristi u ovom primjeru možete saznati ovdje: https://rocmdocs.amd.com/en/latest/Programming_Guides/HIP-GUIDE.html) :

hipLaunchKernelGGL(
  HIP_KERNEL_NAME(example_global_memory_storage<block_size, T>),
  dim3(grid_size), dim3(block_size),
  0, 0,
  device_input, device_output, global_storage
);

Prije nego što pogledamo kojeg je zrno oblika, postoji još par koraka pri kraju ove funkcije. Sljedeće što je potrebno je čitanje ispisa sa uređaja:

hip_read_device_memory<T>(host_output, device_output);

Taj je ispis potrebno validirati odnosno pokrenuti provjeru da se utvrdi je li sve u redu sa ispisom:

OUTPUT_VALIDATION_CHECK(
  validate_device_output(host_output, host_expected_output)
);

Kada su svi koraci zadovoljeni, sada je vrijeme za oslobađanje memorije koju smo prethodno alocirali, i pritom provjeravamo hoće li ta akcija proći uspješno, a to činimo na sljedeći način:

HIP_CHECK(hipFree(device_input));
HIP_CHECK(hipFree(device_output));
HIP_CHECK(hipFree(global_storage));

Ako je sve prošlo po planu, funkcija će ispisati da je pokretanje zrna bilo uspješno.

Uspješno smo prošli kroz prvi primjer pokretanja zrna.

Naposlijetku, zrno koje smo pokretali je oblika:

// Kernel 4 - Using global memory for storage
template<
  const unsigned int BlockSize,
  class T
>
__global__
__launch_bounds__(BlockSize, ROCPRIM_DEFAULT_MIN_WARPS_PER_EU)
void example_global_memory_storage(
  const T *input,
  T *output,
  typename rocprim::block_scan<T, BlockSize>::storage_type *global_storage)
{
  // Indexing for  this block
  unsigned int index = (hipBlockIdx_x * BlockSize) + hipThreadIdx_x;
  // specialize block_scan for type T and block of 256 threads
  using block_scan_type = rocprim::block_scan<T, BlockSize>;
  // Variables required for performing a scan
  T input_value, output_value;

  // execute inclusive scan
  input_value = input[index];

  block_scan_type()
    .inclusive_scan(
       input_value, output_value,
       global_storage[hipBlockIdx_x],
       rocprim::plus<T>()
    );

  output[index] = output_value;
}

Ponovno u početnom dijelu koda stoji funkcija block_scan, nakon čega slijedi indeksiranje niti za blokove.

Indeks niti računamo na način:

unsigned int index = (hipBlockIdx_x * BlockSize) + hipThreadIdx_x;

Usmjerimo block_scan funkciju na T i block koji smo prije postavili na 256 niti.

using block_scan_type = rocprim::block_scan<T, BlockSize>;

Zatim postavljamo varijable koje su nam potrebne za provođenje skeniranja putem funkcije block_scan_type (specificirana funkcija sa parametrom type, kako bi se odredilo koji algoritam funkcija treba pratiti).

Skeniranje započinje, te se podatci pohranjuju u memoriju određenu za ispis.

Prošli smo jedno zrno i njegovu funkciju za pokretanje. Nastavimo dalje promatrati sljedeće funkcije koje se pozivaju u prije pokazanoj main() funkciji.

Sljedeća funkcija koju ćemo promatrati je oblika:

template<class T>
void run_example_shared_memory(size_t size)
{
  constexpr unsigned int block_size = 256;
  // Make sure size is a multiple of block_size
  unsigned int grid_size = (size + block_size - 1) / block_size;
  size = block_size * grid_size;

  // Generate input on host and copy it to device
  std::vector<T> host_input = get_random_data<T>(size, 0, 1000);
  // Generating expected output for kernel
  std::vector<T> host_expected_output = get_expected_output<T>(host_input, block_size);
  // For reading device output
  std::vector<T> host_output(size);

  // Device memory allocation
  T * device_input;
  T * device_output;
  HIP_CHECK(hipMalloc(&device_input, host_input.size() * sizeof(typename decltype(host_input)::value_type)));
  HIP_CHECK(hipMalloc(&device_output, host_output.size() * sizeof(typename decltype(host_output)::value_type)));

  // Writing input data to device memory
  hip_write_device_memory<T>(device_input, host_input);

  // Launching kernel example_shared_memory
  hipLaunchKernelGGL(
    HIP_KERNEL_NAME(example_shared_memory<block_size, T>),
    dim3(grid_size), dim3(block_size),
    0, 0,
    device_input, device_output
  );

  // Reading output from device
  hip_read_device_memory<T>(host_output, device_output);

  // Validating output
  OUTPUT_VALIDATION_CHECK(
    validate_device_output(host_output, host_expected_output)
  );

  HIP_CHECK(hipFree(device_input));
  HIP_CHECK(hipFree(device_output));

  std::cout << "Kernel run_example_shared_memory run was successful!" << std::endl;
}

Usporedite funkciju run_example_global_memory_storage s funkcijom ``run_example_shared_memory`:

run_example_global_memory_storage                                                        run_example_shared_memory

 hip_write_device_memory<T>(device_input, host_input);                                   hip_write_device_memory<T>(device_input, host_input);

 using storage_type = typename rocprim::block_scan<T, block_size>::storage_type;
 storage_type *global_storage;
 HIP_CHECK(hipMalloc(&global_storage, (grid_size * sizeof(storage_type))));

hipLaunchKernelGGL(                                                                      hipLaunchKernelGGL(
  HIP_KERNEL_NAME(example_global_memory_storage<block_size, T>),                           HIP_KERNEL_NAME(example_shared_memory<block_size, T>),
  dim3(grid_size), dim3(block_size),                                                       dim3(grid_size), dim3(block_size),
  0, 0,                                                                                    0, 0,
  device_input, device_output, global_storage                                              device_input, device_output
  );                                                                                       );

Primjetiti ćete da se razlikuju jedino u dijelu gdje se “run_example_global_memory_storage” bavi globalnom memorijom. Sve ostalo je postavljeno na jednak način.

Naime, zrno koje se pokreće u funkciji run_example_shared_memory() je oblika:

template<
  const unsigned int BlockSize,
  class T
>
__global__
__launch_bounds__(BlockSize, ROCPRIM_DEFAULT_MIN_WARPS_PER_EU)
void example_shared_memory(const T *input, T *output)
{
  // Indexing for  this block
  unsigned int index = (hipBlockIdx_x * BlockSize) + hipThreadIdx_x;

  // Allocating storage in shared memory for the block
  using block_scan_type = rocprim::block_scan<T, BlockSize>;
  __shared__ typename block_scan_type::storage_type storage;

  // Variables required for performing a scan
  T input_value, output_value;

  // Execute inclusive plus scan
  input_value = input[index];

  block_scan_type()
    .inclusive_scan(
       input_value,
       output_value,
       storage,
       rocprim::plus<T>()
  );

  output[index] = output_value;
}

Ovo zrno možete usporediti sa zrnom “example_global_memory_storage” koje smo prvo promatrali. Primjetiti ćete da su razlike minimalne, upravo iz razloga jer su oba zrna namijenjena za istu svrhu, no koriste dva različita tipa memorije.

Konkretno ovo zrno namjenjeno je za dijeljenu memoriju, dok je “example_global_memory_storage” namijenjeno za globalnu memoriju.

Funkcija sljedeća na redu pozivanja u main() funkciji je oblika:

template<class T>
void run_example_union_storage_types(size_t size)
{
  constexpr unsigned int block_size = 256;
  constexpr unsigned int items_per_thread = 4;
  // Make sure size is a multiple of block_size
  auto grid_size = (size + block_size - 1) / block_size;
  size = block_size * grid_size;

  // Generate input on host and copy it to device
  std::vector<T> host_input = get_random_data<T>(size, 0, 1000);
  // Generating expected output for kernel
  std::vector<T> host_expected_output = get_expected_output<T>(host_input, block_size, items_per_thread);
  // For reading device output
  std::vector<T> host_output(size);

  // Device memory allocation
  T * device_input;
  T * device_output;
  HIP_CHECK(hipMalloc(&device_input, host_input.size() * sizeof(typename decltype(host_input)::value_type)));
  HIP_CHECK(hipMalloc(&device_output, host_output.size() * sizeof(typename decltype(host_output)::value_type)));

  // Writing input data to device memory
  hip_write_device_memory<T>(device_input, host_input);

  // Launching kernel example_union_storage_types
  hipLaunchKernelGGL(
     HIP_KERNEL_NAME(example_union_storage_types<block_size, items_per_thread, int>),
     dim3(grid_size), dim3(block_size),
     0, 0,
     device_input, device_output
  );

  // Reading output from device
  hip_read_device_memory<T>(host_output, device_output);

  // Validating output
  OUTPUT_VALIDATION_CHECK(
     validate_device_output(host_output, host_expected_output)
  );

  HIP_CHECK(hipFree(device_input));
  HIP_CHECK(hipFree(device_output));

  std::cout << "Kernel run_example_union_storage_types run was successful!" << std::endl;
}

Vidjeli smo do sada dvije funkcije run_example... koje su podosta slične, razlikuju se u detaljima oko globalne memorije. Ova funkcija je isto tako slična, jedine razlike koje postoje su:

...

constexpr unsigned int block_size = 256;
constexpr unsigned int items_per_thread = 4;
...

std::vector<T> host_expected_output = get_expected_output<T>(host_input, block_size, items_per_thread);
...

hipLaunchKernelGGL(
    HIP_KERNEL_NAME(example_union_storage_types<block_size, items_per_thread, int>),
    dim3(grid_size), dim3(block_size),
    0, 0,
    device_input, device_output
);
...

Ako pogledate malo pažljivije, ovdje je uvedena nova int varijabla; items_per_thread. S obzirom da se radi o unijama memorija, potrebna nam je ta varijabla za zrno koje pozivamo putem te funkcije, a ono je oblika:

template<
  const unsigned int BlockSize,
  const unsigned int ItemsPerThread,
  class T
>
__global__
__launch_bounds__(BlockSize, ROCPRIM_DEFAULT_MIN_WARPS_PER_EU)
void example_union_storage_types(const T *input, T *output)
{
  // Specialize primitives
  using block_scan_type = rocprim::block_scan<
     T, BlockSize, rocprim::block_scan_algorithm::using_warp_scan
  >;
  using block_load_type = rocprim::block_load<
     T, BlockSize, ItemsPerThread, rocprim::block_load_method::block_load_transpose
  >;
  using block_store_type = rocprim::block_store<
     T, BlockSize, ItemsPerThread, rocprim::block_store_method::block_store_transpose
  >;
  // Allocate storage in shared memory for both scan and sort operations

  __shared__ union
  {
     typename block_scan_type::storage_type scan;
     typename block_load_type::storage_type load;
     typename block_store_type::storage_type store;
  } storage;

  constexpr int items_per_block = BlockSize * ItemsPerThread;
  int block_offset = (hipBlockIdx_x * items_per_block);

  // Input/output array for block scan primitive
  T values[ItemsPerThread];

  // Loading data for this thread
  block_load_type().load(
     input + block_offset,
     values,
     storage.load
  );
  rocprim::syncthreads();

  // Perform scan
  block_scan_type()
     .inclusive_scan(
         values, // as input
         values, // as output
         storage.scan,
         rocprim::plus<T>()
     );
  rocprim::syncthreads();

  // Save elements to output
  block_store_type().store(
     output + block_offset,
     values,
     storage.store
  );
}

U slučaju ovog zrna, koristiti će se operacije već definirane u samom rocPRIMu:

using block_scan_type = rocprim::block_scan<
    T, BlockSize, rocprim::block_scan_algorithm::using_warp_scan
>;
using block_load_type = rocprim::block_load<
    T, BlockSize, ItemsPerThread, rocprim::block_load_method::block_load_transpose
>;
using block_store_type = rocprim::block_store<
    T, BlockSize, ItemsPerThread, rocprim::block_store_method::block_store_transpose
>;

Ovo su tzv. primitive (za više informacija o svakoj od ovih operacija pogledajte početak ovog teksta podnaslov Block-wide).

Slijedi alokacija prostora za pohranu u dijeljenoj memoriji za ove operacije:

__shared__ union
 {
     typename block_scan_type::storage_type scan;
     typename block_load_type::storage_type load;
     typename block_store_type::storage_type store;
 } storage;

 constexpr int items_per_block = BlockSize * ItemsPerThread;
 int block_offset = (hipBlockIdx_x * items_per_block);

Stvara se input/output polje za skeniranje prije navedenih blokova, nakon čega se učitavaju dobiveni podatci u ovoj sekvenci naredbi. Pritom se pokreće i skeniranje:

// Input/output array for block scan primitive
 T values[ItemsPerThread];

// Loading data for this thread
block_load_type().load(
    input + block_offset,
    values,
    storage.load
);
rocprim::syncthreads();

// Perform scan
block_scan_type()
    .inclusive_scan(
        values, // as input
        values, // as output
        storage.scan,
        rocprim::plus<T>()
    );
rocprim::syncthreads();

// Save elements to output
block_store_type().store(
    output + block_offset,
    values,
    storage.store
);

}

Na kraju skeniranja podatci se pohranjuju u memoriju za ispis.

Terminologija

Warp

Odnosi se na grupu niti koje se izvrše na SIMT (Single Instruction, Multiple Thread) način. Također, za njih se kaže da su “valna fronta” na AMD GPU-ima.

Hardware Warp Size

Odnosi se na broj niti u warpu, i taj broj je definiran hardverom. Na Nvidijinim grafičkim procesorima veličina warpa iznosi 32, a na AMD-ovim grafičkim karticama iznosi 64.

Logical Warp Size

Odnosi se na broj niti u warpu definiran sa strane korisnika, koji može biti jednak ili manji od broja niti definiranim hardverom.

Lane ID

Odnosi se na identifikator niti unutar warpa. Logički lane ID se odnosi na identifikator niti u logičkom warpu, koji može biti manji od broja niti definiranim hardverom.

Warp ID

Odnosi se na identifikator hardverskog/logičkog warpa u bloku. On je jedinstven za svaki warp.

Block

Odnosi se na grupu niti koje se izvršavaju na jednakoj računskoj jedinici. Ove niti mogu biti indeksirane korištenjem jedne dimenzije (X), dvije dimenzije (X, Y) ili 3 dimenzije (X, Y, Z). Blok se sastoji od više warpova.

Tile

Odnosi se na blok, ali u C++AMP/HCC nominklaturi.