Extend Support for CUDA and HIP kernel invocations

This MR extends the support for CUDA and HIP kernel invocations and API reflection in pystencils-sfg, reflecting the concepts introduced in pystencils!449 (merged).

• Move GPU-related stuff into new composer mix-in SfgGpuComposer
• Replace cuda_invoke() by gpu_invoke()
• Implement three separate signatures of gpu_invoke(), depending on the kernel indexing scheme selected by the user
• Replace extensions.gpu by lang.gpu
• Add explanations on GPU kernels to user manual
• Fix some errors and clarify some interface descriptions in the user manual

This MR will already be using Target.HIP and therefore will have to wait for pystencils!458 (merged).

tests/generator_scripts/source/CudaKernels.harness.cpp

@@ -42,7 +42,8 @@ int main(void)
     std::mt19937 gen{rd()};
     std::uniform_real_distribution<double> distrib{-1.0, 1.0};
 
-    auto check = [&](std::function< void () > invoke) {
+    auto check = [&](std::function<void()> invoke)
+    {
         for (size_t i = 0; i < items; ++i)
         {
             data_src[i] = distrib(gen);
@@ -62,24 +63,33 @@ int main(void)
         }
     };
 
-    check([&]() {
+    check([&]()
+          {
         /* Linear3D Dynamic */
         dim3 blockSize{64, 8, 1};
         cudaStream_t stream;
         checkCudaError(cudaStreamCreate(&stream));
         gen::linear3d::scaleKernel(blockSize, dst, src, stream);
-        checkCudaError(cudaStreamSynchronize(stream));
-    });
+        checkCudaError(cudaStreamSynchronize(stream)); });
+
+    check([&]()
+          {
+        /* Linear3D Automatic */
+        cudaStream_t stream;
+        checkCudaError(cudaStreamCreate(&stream));
+        gen::linear3d::scaleKernel(dst, src, stream);
+        checkCudaError(cudaStreamSynchronize(stream)); });
 
-    check([&]() {
+    check([&]()
+          {
         /* Blockwise4D Automatic */
         cudaStream_t stream;
         checkCudaError(cudaStreamCreate(&stream));
         gen::blockwise4d::scaleKernel(dst, src, stream);
-        checkCudaError(cudaStreamSynchronize(stream));
-    });
+        checkCudaError(cudaStreamSynchronize(stream)); });
 
-    check([&]() {
+    check([&]()
+          {
         /* Linear3D Manual */
         dim3 blockSize{32, 8, 1};
         dim3 gridSize{5, 4, 23};
@@ -87,18 +97,17 @@ int main(void)
         cudaStream_t stream;
         checkCudaError(cudaStreamCreate(&stream));
         gen::linear3d_manual::scaleKernel(blockSize, dst, gridSize, src, stream);
-        checkCudaError(cudaStreamSynchronize(stream));
-    });
+        checkCudaError(cudaStreamSynchronize(stream)); });
 
-    check([&]() {
+    check([&]()
+          {
         /* Blockwise4D Manual */
         dim3 blockSize{132, 1, 1};
         dim3 gridSize{25, 23, 1};
         cudaStream_t stream;
         checkCudaError(cudaStreamCreate(&stream));
         gen::blockwise4d_manual::scaleKernel(blockSize, dst, gridSize, src, stream);
-        checkCudaError(cudaStreamSynchronize(stream));
-    });
+        checkCudaError(cudaStreamSynchronize(stream)); });
 
     checkCudaError(cudaFree(data_src));
     checkCudaError(cudaFree(data_dst));