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?

I think it would be a bit cleaner to do that in another MR

Alright, fine for me

added 1 commit

• 9140da63 - Added a parameter to insert a launch bounds to the kernel

Compare with previous version

added 1 commit

• 4ee400e9 - added the new packages to the setup.cfg and the new templates to the

Compare with previous version

added 1 commit

• 4b1f3f53 - fixed the _add_launch_bounds and also added some small tests

Compare with previous version

added 1 commit

• 6e88d389 - Using cuda as a base for the docker container to also test the gpu

Compare with previous version

added 1 commit

• 82ce1d7d - Fix the missing constants for the gpu main file and added a kernel with

Compare with previous version

added 1 commit

• 1e542f17 - Skipping compiling and running cuda kernels if cuda or gpu is not

Compare with previous version

This is dangerous, as it makes the test pass, although it was not tested completely. I'd use pytest.fail, pytest.skip, or pytest.exit, depending on your intent how to handle a GPU test, without having a GPU compiler/runtime.

Good point, I think skipping is the the most reasonable in that case

changed this line in version 12 of the diff

added 1 commit

• d38a9324 - using pytest skip if there is no nvcc or gpu available

Compare with previous version

marked the checklist item Add launch bound option as completed

added 1 commit

• 879ee872 - removed the unused `cuda_block_size` for the `gpu.generate_benchmark`

Compare with previous version

merged

mentioned in commit fcfbef80