<html><head><meta http-equiv="Content-Type" content="text/html; charset=utf-8"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class="">Hey Ees, Johannes,<div class=""><br class=""></div><div class="">AFAIK the latest version of the code should be publicly accessible here (it just isn’t maintained):</div><div class=""><a href="https://github.com/UniHD-CEG/mekong-cuda" class="">https://github.com/UniHD-CEG/mekong-cuda</a></div><div class=""><br class=""></div><div class="">Johannes low level polyhedral analysis is in “llvm-patches/”, which you should be able to apply to the LLVM trunk around Jan 8, 2018 (I don’t remember the commit, but it’s unlikely there are changes breaking the relevant APIs just randomly on that day).</div><div class="">It already comes with support to recognize the standard “threadIdx.{} + blockIdx.{} * blockDim.{}” (with {} being either x, y, or z), making the whole expression a “constant”, removing the non-linearity and enabling polyhedral analysis.</div><div class=""><br class=""></div><div class="">The post-processing we’re applying to the raw analysis result is in “lib/MeKernelAnalysis.cpp”.</div><div class="">I’m not sure how helpful that will be for you, because we were going in a different direction and combining all accesses of a thread, but it might be helpful as a fairly simple example of how to use polyhedral analysis from within LLVM (and of course it was developed shooting from the hip, without tests).</div><div class=""><br class=""></div><div class="">P.S.: I like your idea of inlining everything, then making the CUDA intrinsics constants and applying constant folding. It sounds simple enough to get something working quickly that can then be expanded upon</div><div class=""><br class=""></div><div class="">Cheers,</div><div class="">Alex</div><div class=""><div><br class=""><blockquote type="cite" class=""><div class="">On 23. Aug 2020, at 21:24, Ees <<a href="mailto:kayesg42@gmail.com" class="">kayesg42@gmail.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div class="">Hello Johannes,<br class=""><br class="">Thank you very much for the material. I will have a look and get back to you (possibly with questions if you don't mind :) ).<br class="">I would also appreciate the code if that's available.<br class=""><br class="">- Ees<br class=""><br class="">On 23-08-2020 18:47, Johannes Doerfert wrote:<br class=""><blockquote type="cite" class="">Hi Ees,<br class=""><br class="">a while back we started a project with similar scope.<br class="">Unfortunately the development slowed down and the plans to revive it this summer got tanked by the US travel restrictions.<br class=""><br class="">Anyway, there is some some existing code that might be useful, though in a prototype stage. While I'm obviously biased, I would suggest we continue from there.<br class=""><br class="">@Alex @Holger can we put the latest version on github or some other place to share it, I'm unsure if the code I (might have) access to is the latest.<br class=""><br class="">@Ees I attached a recent paper and you might find the following links useful:<br class=""><br class=""> * 2017 LLVM Developers’ Meeting: J. Doerfert “Polyhedral Value & Memory Analysis ” <a href="https://youtu.be/xSA0XLYJ-G0" class="">https://youtu.be/xSA0XLYJ-G0</a><br class=""><br class=""> * "Automated Partitioning of Data-Parallel Kernels using Polyhedral Compilation.", P2S2 2020 (slides and video <a href="https://www.mcs.anl.gov/events/workshops/p2s2/2020/program.php" class="">https://www.mcs.anl.gov/events/workshops/p2s2/2020/program.php</a>)<br class=""><br class=""><br class="">Let us know what you think :)<br class=""><br class="">~ Johannes<br class=""><br class=""><br class=""><br class=""><br class="">On 8/22/20 9:38 AM, Ees Kee via llvm-dev wrote:<br class="">> Hi all,<br class="">><br class="">> As part of my research I want to investigate the relation between the<br class="">> grid's geometry and the memory accesses of a kernel in common gpu<br class="">> benchmarks (e.g Rodinia, Polybench etc). As a first step i want to<br class="">> answer the following question:<br class="">><br class="">> - Given a kernel function with M possible memory accesses. For how many of<br class="">> those M accesses we can statically infer its location given concrete values<br class="">> for the grid/block and executing thread?<br class="">><br class="">> (Assume CUDA only for now)<br class="">><br class="">> My initial idea is to replace all uses of dim-related values, e.g:<br class="">> __cuda_builtin_blockDim_t::__fetch_builtin_x()<br class="">> __cuda_builtin_gridDim_t::__fetch_builtin_x()<br class="">><br class="">> and index related values, e.g:<br class="">> __cuda_builtin_blockIdx_t::__fetch_builtin_x()<br class="">> __cuda_builtin_threadIdx_t::__fetch_builtin_x()<br class="">><br class="">> with ConstantInts. Then run constant folding on the result and check how<br class="">> many GEPs have constant values.<br class="">><br class="">> Would something like this work or are there complications I am not thinking<br class="">> of? I'd appreciate any suggestions.<br class="">><br class="">> P.S i am new to LLVM<br class="">><br class="">> Thanks in advance,<br class="">> Ees<br class="">><br class="">><br class="">> _______________________________________________<br class="">> LLVM Developers mailing list<br class="">> <a href="mailto:llvm-dev@lists.llvm.org" class="">llvm-dev@lists.llvm.org</a><br class="">> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" class="">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br class=""><br class=""></blockquote></div></div></blockquote></div><br class=""></div></body></html>