By extension, this means that we'll also have to move cuda-runtime related calls (or wrappers) into a .hpp file and separate out the device-side utilities into a .cuh file. This will further aid in renaming certain .cu's to .cpp's.

@chaithyagr, how about this one?

Makes sense, will give a good dive into the codebase.
This was also on my mind.

So I could finally get code building!!
I see that code base is divied up into src and src_prims for easier usage.
For now I thought of starting the process from src_prims (A kind of bottom up approach)

Going by the rules described in #1675 (comment) , we would have a lot of file renaming, do we want this?

For now, I havent looked at any helpers for doing this, as I want to have a look into the library as I do this. But do let me know if there is any obvious fast way of doing this. Something which I thought could be used (not yet sure) is https://github.com/include-what-you-use/include-what-you-use to analyze the usage patterns

Glad to know that the build worked for you! In case you faced any issues during the build process, please do write about those in issue #1503 . We want to know such difficulties that the devs are having.

I'm pretty sure about all the rules I described above, except number 2, since that is mostly my personal preference. I'd like to have @JohnZed, @cjnolet and @dantegd to chime in here and get their opinion before proceeding any further.

BTW, @chaithyagr, do note that the rule number 1 above works the other way too. Meaning, let's assume that there is a .cuh which doesn't have any kernels or device methods in it. Even if it contains cuda runtime API calls, it should fall under rule number 2. In other words, we need nvcc only for compiling the kernels or device methods. Rest everything can be compiled via host compiler itself.

IWYU can be an interesting addition to our CI process. Can you file an issue about it against cuML?

yes, and it is better to start from src_prims. In fact I'd say start with one of the smaller subfolders in there, just to get a hang of this process and file a PR for just that change. It is better to get this change in small installments (or else everyone will face merge conflicts!)

Sure, I will work on src_prims first.
However, I ended up making a major modification as I changed utils and cuda_utils first which has changed a lot of files, although the build works so I didnt break anything yet :P

However, I think I agree with you on having this PR in multiple parts, I will keep this change aside and work on smaller groups of folders and submit them rather than have a bulky PR.

I will go through the files and try to come up with a logical order of things and how to proceed in free time

@chaithyagr any updates on this?

Oops, I seem to have forgotten to update here.
So I made one major change of changing the base cuda files as mentioned and I am about to make a more targetted change for small batched PRs.
However, for now, I seem to be limited by some other priorities and might get at this a little later, especially with the lockdown due to Corona.
Do we have any expected deadline, so that I can give an idea of what I can do before?
I can get the logical order in my opinion by end of tomorrow.

I just realized that it is better for me to be sure of the renaming rules:

Just to ensure that I am on the right track, here are some example renames that I think needs to be done :

1. src_prims/cuda_utils.h and src_prims/utils.h
   I am pretty sure that cuda_utils.h should be renamed to cuda_utils.cuh. The same applies to utiils.h as they do have device kernel calls
   However, I just feel this will cause massive change as I described in #1675 (comment) (esp from Rule 4)

That brings me to a question, so right now all files are compiled using nvcc irrespective?

2. This could be a little stupid, but I was going through src_prims/sparse and I find that almost all files should become cuh. Am I right in saying this? Somehow I am getting uncomfortable with the sheer amount of changes that even a minute update could bring out.

Am I missing something here or have something grossly wrong?

@chaithyagr

re: ETA

the sooner you can get atleast some set of changes filed for PR the better it is for us. This change could also help us in reducing compilation time. Hence the need to address this issue faster.

... Somehow I am getting uncomfortable with the sheer amount of changes that even a minute update could bring out.

That's why it is probably better not to start with core set of files like utils.h or cuda_utils.h. For example: pick up an algorithm inside cpp/src, just mentally apply the above said rules on such core header files. Now start renaming the headers in this algo's folder appropriately. In other words, doing a top-down approach (instead of going bottom-up) is easier for you to contain the changes in small PR's. At the end, most of the files need to be touched anyways, because of updating these core header files. But by then, you should have gotten used to the process and have become effective in doing this.

I am pretty sure that cuda_utils.h should be renamed to cuda_utils.cuh...

That's correct, as it follows rule no 1 above.

The same applies to utiils.h as they do have device kernel calls

utils.h does not have any of the items mentioned in rule no 1. However, the name is too generic to cause conflicts in future! So, I'd probably rename this as cpp/src_prims/common/cudart_utils.hpp.

.. but I was going through src_prims/sparse and I find that almost all files should become cuh...

That's correct. except for cusparse_wrappers.h which will stay as it is.

@chaithyagr I just updated rule no. 2 to help minimize amount of changes. In short, if a header needs to be renamed as a .hpp, but it is already a .h, feel free to leave it as it is. Example: cpp/src_prims/sparse/cusparse_wrappers.h, as per rule no. 2, should be changed to a .hpp. But, in order to minimize the changes, this renaming will cause, let us just leave this file as it is.

Yes @teju85 , I think have bottoms up approach was a bad idea initially.

Now, I started by understand the directory structure and also the organization of codes in some subdirectories of src

I am sorry, but I seem to be getting a bit confused here, and I couldnt find good pointers to what is right definitions at this point, so do let me know if there is anything basic that I am missing here..

I will commit my changes in this bunch and hopefully start a PR tomorrow, I hope with the answers I will be better equipped to work on some aspect over the weekend.

@chaithyagr

1. kmeans - your analysis is correct
2. kalmanfilter - I'd ignore this one for now, as this might be deprecated soon
3. randomforest - please ignore this one too for now. RF requires a major refactor that's been happening internally
4. solver
   • cd.h -> should be renamed to cd.cuh (reasoning: your point about a future cuda_utils.cuh renaming and also because many of the LinAlg::* namespace usages are actually kernel calls)
   • learning_rate.h -> no change
   • sgd.h -> should be renamed to sgd.cuh (reasoning: same as cd.cuh renaming)
   • shuffle.h -> no change
   • solver.cu -> no change

Hope this helps.

Thanks @teju85 , I have made the first PR, with the changes in solver, dbscan and svm.
I have not done a complete test yet, as I need to clean up some space, but given that it is just renaming of headers, I expect things to be fine, as I got a clean build.
Nevertheless, I will build local gpuCI for future testing as needed

@teju85 , Shall I proceed with more changes, or shall I wait for the first PR to get a clearer picture of how PR is done?

Hey @chaithyagr, give me few more days to look at that PR. Currently busy with some other stuffs.

Sure sure!

@chaithyagr I noticed a silly mistake from me in prims folder! It is: cpp/src_prims/common/seive.cuh

1. It should be renamed to seive.hpp
2. It doesn't require cuda_utils.h
3. It's corresponding test under cpp/test/prims/seive.cu should be renamed to cpp/test/prims/seive.cpp.

Just pointing it out explicitly so that we don't miss it.

I have opened what seems to be the last PR for src/ directory files. I have ignored randomForest and kalman filter for now as you suggested above.

Moving towards src_prims, it is weird but I am a bit stuck on how to make changes without a major change, I think the movement of cuda_utils.h to cuda_utils.cuh must be the last change as it will have the maximum impact. Also, I am still trying to ease in minimum batches of changes to keep PR tight.

This is good progress. Thanks @chaithyagr for doing this!

Sadly, for the src_prims, I think the only way is all-or-nothing, because it gets used at multiple places inside src. :( So, I'm ok if you create a monolithic PR for this one. I'll somehow manage the review for it.

Quick question, I found a wrapper: cpp/src_prims/common/cub_wrapper.h, which is primarily a .h file, as needed.
However, it includes cub.cuh , so by the rules, it must be a .cuh file.
It feels weird to make this change, and we use sort functions from cub.cuh

@chaithyagr I noticed a silly mistake from me in prims folder! It is: cpp/src_prims/common/seive.cuh

1. It should be renamed to seive.hpp
2. It doesn't require cuda_utils.h
3. It's corresponding test under cpp/test/prims/seive.cu should be renamed to cpp/test/prims/seive.cpp.

Just pointing it out explicitly so that we don't miss it.

I'm afraid, we need to use cuda_utils.h for ceildiv

ah... good catch regarding ceildiv. Thanks @chaithyagr

@chaithyagr regarding cub_wrapper, it must be named as cub_wrapper.cuh itself, as it calls cub device routines. In other words, files that include and use device-side methods from thrust and also from cub should just be named with .cuh extension (if they are headers) and .cu (if they are source files)

I think the majority of the changes are done (once the last PR merges)
As a wrap-up and in my opinion, a final step:

Although we could the changes, I feel there is a need to check the above rules in an actual CI script test.
I can compile some simple scripts I used in a bash file. This script, while doesn't completely model the actual set of rules stated, as I am not sure how to exactly be sure if a file has device functions (other than simple heuristics).
This would ensure that we don't need to do such similar cleanup again.

@teju85 shall we try to get this done as a part of this issue (effectively at high priority) or try to add this to #1895 as an extra task as they both are related?

In my opinion, its better to do a complete wrap for good :P

PS: I am not very well aware of the CI methodology here, I can work on a script, but placing it in right place, I might need some help.

@chaithyagr that's a good idea (I was planning to ask you about adding such a "checker" after the prims header rename PR was done!). Yes, adding such a check (inside cpp/scripts/) will certainly help guide devs to write cleaner code. I can help you with getting it enabled in our CI.

However, I'm not sure if IWYU and this check need to be in the same PR. I'd prefer if we did these 2 separately.

I tried to add in most of checks in this one simple code:
find -name '*.hpp' -or -name '*.h' -or -name '*.cpp' -or -name '*.c' | grep -v 'cpp\/build' | xargs grep "#include.*.cuh\|__device__\|threadIdx\|blockIdx"

Currently here is the output, which lists all possibly bad cases:

./cpp/src/decisiontree/memory.h:#include "memory.cuh"
./cpp/src_prims/linalg/transpose.h:                   [=] __device__(int idx) {
./cpp/src_prims/distance/distance_tile_traits.h:        threadIdx.x / Threads::kW * kStrideH * Iterations::kH;
./cpp/src_prims/distance/distance_tile_traits.h:      int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
./cpp/include/cuml/common/cubAllocatorAdapter.hpp:#include <cub/util_allocator.cuh>
./cpp/include/cuml/tsa/arima_common.h:                     counting + batch_size, [=] __device__(int bid) {
./cpp/include/cuml/tsa/arima_common.h:                     counting + batch_size, [=] __device__(int bid) {

The second part of checks, that is check if threadIdx and blockIdx exists is a very crude way to look for device kernels, does it help? This hopes that no one used this as variable name or even coded this specifically in comments.

Why not use the __global__ keyword reserved for kernel definitions, instead of threadIdx/blockIdx?

Ahh, I forgot to add __global__.

Perhaps threadIdx and blockIdx could be used more so for warning.