René's URL Explorer Experiment

Title: [GPU/Full System Test] GPU standalone benchmarking by mconcas · Pull Request #6484 · AliceO2Group/AliceO2 · GitHub

Open Graph Title: [GPU/Full System Test] GPU standalone benchmarking by mconcas · Pull Request #6484 · AliceO2Group/AliceO2

X Title: [GPU/Full System Test] GPU standalone benchmarking by mconcas · Pull Request #6484 · AliceO2Group/AliceO2

Description: Hi @davidrohr, In this draft I propose a possible implementation of dedicated executables to run generic GPU benchmark. Notably this PR has: Two dedicated binaries o2-gpu-memory-benchmark-{cuda,hip} automatically produced based on "GPU autodetection" Automatic (re-)generation of HIP code on-the-fly by mean of hipify-perl script, upon changes in .cu CUDA corresponding files Automatic scratching based on fraction (default 95%) of the free GPU resident memory Templated benchmarking class to use desired class type backend and templated benchmarking function A dummy example of reading kernel (to be heavily revised, but just to create a full benchmarking "workflow") Configurable CLI params Store results on root file Overall, for the moment it's very simple but I wanted to be sure I am going in the right direction. Please, let me know what do you think. Cheers, Matteo Summary Kernels are always launched in one-dimensional fashion, assigning gridDim.x to the number of multiprocessors and blockDim.x to the number of available threads per block. seq, single-block seq, multi-block conc, single-block conc, multi-block read ✔️ ✔️ ✔️ ✔️ write ✔️ ✔️ ✔️ ✔️ copy ✔️ ✔️ ✔️ ✔️ random read TODO TODO TODO TODO random write TODO TODO TODO TODO random copy TODO TODO TODO TODO Names can be imprecise and misleading, I try to explain what I mean: seq: kernels are launched one after the other finishes, one per different 1GB(default) partition of scratch; so that each one run with no others in parallel on other partitions single-block: each block is pinned (by its Id.x) to a different 1GB(default) partition of scratch. To ergodically span over it it uses iteration increase by a stride equal to blockDim.x. multi-block: all blocks work on each segment but in a ordered and strided way (stride = blockDim * gridDim ) conc: kernels run at the same time on different slices of scratch. Benchmarks measure per-slice performance single-block: each block is pinned (by its Id.x) to a different sub-buffer (regardless partitions) of scratch. To ergodically span over it it uses iteration increase by a stride equal to blockDim.x. multi-block: The idea is to partition the pool of blocks so that each subset can access a sub-buffer of scratch. All blocks work on each segment but in a ordered and strided way (stride = blockDim * gridDim )

Open Graph Description: Hi @davidrohr, In this draft I propose a possible implementation of dedicated executables to run generic GPU benchmark. Notably this PR has: Two dedicated binaries o2-gpu-memory-benchmark-{cuda,hi...

X Description: Hi @davidrohr, In this draft I propose a possible implementation of dedicated executables to run generic GPU benchmark. Notably this PR has: Two dedicated binaries o2-gpu-memory-benchmark-{cuda,hi...

Opengraph URL: https://github.com/AliceO2Group/AliceO2/pull/6484

X: @github

direct link

Domain: github.com

Links:

Viewport: width=device-width