Add CUDA support

This MR adds cuda support

remaining to-dos

• Fix compilation problems and add required NVCC flags
• Add launch bound option

requested review from @hoenig

assigned to @holzer

marked the checklist item Fix compilation problems and add required NVCC flags as completed

marked the checklist item Fix compilation problems and add required NVCC flags as incomplete

added 1 commit

• a0a17570 - Working CUDA benchmark Version. CUDA needs '.cu' files, otherwise it doesn't work?

Compare with previous version

marked the checklist item Fix compilation problems and add required NVCC flags as completed

added 16 commits

• a0a17570...ae05d616 - 11 commits from branch master
• c51eee43 - Added CUDA benchmarks
• 96c63098 - Working CUDA benchmark Version. CUDA needs '.cu' files, otherwise it doesn't work?
• ac4b31c1 - Updated import to new pystencils api
• 39b4029c - Exposing the cuda block size option to the generate_benchmark function
• 0028ee62 - Merge branch 'CUDA' of i10git.cs.fau.de:pycodegen/pystencils-benchmark into CUDA

Compare with previous version

approved this merge request

added 2 commits

• 3300460d - made the gpu test more streamlined with the cpu tests
• 3e930f35 - removed some code duplication between benchmark and benchmark_gpu

Compare with previous version

added 2 commits

• 24f81cf6 - added submodules from cpu and gpu benchmark generation
• 857f1848 - removed the mutable default argument from the _kernel_header and

Compare with previous version

resolved all threads

The code looks really nice now and easily extendible to other platforms/compilers. Is the second TODO in the PRs description also done?

marked this merge request as ready

Thank you and thank you for your review I am not really sure what it is meant there. In principle it is possible to configure the cuda_block_size within the generate_benchmark call. But as I see now this parameter is not used at all

It means this:

__global__ void
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP)
fooKernel(int *inArr, int *outArr)
{
    // ... Computation of kernel
}

So that you can add __launch_bounds__ as an optional argument. It was pretty important for example on AMD GPUs to limit register usage with LBM kernels.

In most cases, you don't need the second argument. However, when using pystencils standalone it can add the launch bounds option via cupy, but since this is not what we are doing here we need to add it manually as a tuning parameter.

approved this merge request

Would it make sense to also add ROCm support? Mostly this would be just a renaming for example: #include <cuda_runtime.h> --> #include <hip_runtime.h> .

I'm not sure if it is better to add this in a second MR or directly here?