Added support for AMD backend

.gitignore

@@ -2,6 +2,7 @@
 # Lines that start with '#' are comments.
 
 # Editor and misc backup files - anywhere
+.vscode
 *~
 .*~
 .*.swp

.gitlab-ci.yml

@@ -6,10 +6,11 @@ default:
   tags:
     - cuda
 
-build-job:   
+build-job:
   stage: build
   script:
     - echo "Compiling the code..."
+    - export GPUFOAM_BACKEND_NVIDIA=1
     - ./Allwmake -j 8
     - echo "Compile complete."
   after_script:
@@ -24,12 +25,12 @@ build-job:
     when: always
     expire_in: 8 hours
 
-unit-test-job:  
-  stage: test  
+unit-test-job:
+  stage: test
   before_script:
     - mkdir -p build/bin build/lib $FOAM_USER_APPBIN $FOAM_USER_LIBBIN
     - cp -r build/bin/. $FOAM_USER_APPBIN
     - cp -r build/lib/. $FOAM_USER_LIBBIN
   script:
-    - testGpuChemistry
+    - Test-gpuChemistry [CPU]
     - echo "Tests done."

Allwmake

@@ -2,6 +2,17 @@
 cd ${0%/*} || exit 1    # Run from this directory
 
 
+if [ -n "${GPUFOAM_BACKEND_NVIDIA}" ]
+then
+    echo "Compiling gpuFoam for Nvidia backend"
+elif [ -n "${GPUFOAM_BACKEND_AMD}" ]
+then
+    echo "Compiling gpuFoam for AMD backend"
+else
+    echo "You need to set either GPUFOAM_BACKEND_NVIDIA=1 or GPUFOAM_BACKEND_AMD=1 to compile gpu models."
+    exit
+fi
+
 # Compile gpu utilities
 gpu_utils/Allwmake
 

README.md

@@ -6,14 +6,55 @@
 For more details check the README section in each of the subfolders listed above.
 
 ## Requirements
-* C++17 supporting host compiler 
-* Nvidia [cuda compiler](https://developer.nvidia.com/hpc-sdk) (nvcc) version 10+
+* C++17 supporting host compiler
+* Either Nvidia [cuda compiler](https://developer.nvidia.com/hpc-sdk) (nvcc) version 10+
+* Or AMD [hip compiler](https://rocm.docs.amd.com/projects/HIP/en/docs-6.0.0/how_to_guides/install.html) (hipcc) version 5.6+
 * Latest [OpenFOAM foundation development release](https://openfoam.org/version/dev/)
 
+## Compilation
+Everything in this folder should compile with the Allwmake script. Before running it, ensure that you have a valid OpenFOAM installation and compilers available.
+
+* #### For Nvidia
+
+    To check that you have valid compilers for the Nvidia backend
+    ```
+    foamRun -help
+    g++ --version
+    nvcc --version
+    ```
+
+    Then check the name of your graphics card (for example with nvidia-smi) on and go to [this page](https://developer.nvidia.com/cuda-gpus) to see which compute capability it has. Modify the file ```./nvcpp``` accordingly if necessary.
+
+
+
+    Finally, run the commands
+    ```
+    export GPUFOAM_NVIDIA_BACKEND=1
+    ./Allwmake
+    ```
+
+* #### For AMD
+
+    To check that you have valid compilers for the AMD backend
+    ```
+    foamRun -help
+    g++ --version
+    hipcc --version
+    ```
+    Then check the name of your graphics card and go to [this page](https://llvm.org/docs/AMDGPUUsage.html#processors) to see the matching architecture keyword. Modify the file ```./hipcc``` if necessary.
+
+    Finally, run the commands
+    ```
+    export GPUFOAM_AMD_BACKEND=1
+    ./Allwmake
+    ```
+
+
+
 ## Credits
 This project contains source code taken from the following projects
 
-* [mdspan](https://github.com/kokkos/mdspan) 
-* [variant](https://github.com/bryancatanzaro/variant) 
+* [mdspan](https://github.com/kokkos/mdspan)
+* [variant](https://github.com/bryancatanzaro/variant)
 
 The licenses of the external projects can be found from their respective folders.

gpu_chemistry/Allwmake

@@ -9,7 +9,11 @@ wmakeLnInclude src
 wmake src/gpuKernelEvaluator
 wmake src/gpuChemistryModels
 wmake catchMain
-wmake all unittest
-#wmake all benchmark
+wmake unittest/testHelpers
+wmake unittest/OpenFOAMReferenceKernels
+wmake unittest/cpuTestKernels
+wmake unittest/gpuTestKernels
+wmake unittest/tests
+wmake all benchmark
 
 # ----------------------------------------------------------------- end-of-file

gpu_chemistry/README.md

@@ -1,13 +1,5 @@
 # gpu_chemistry
 
-## Compilation
-Everything in this folder such compile with the Allwmake script. Before running it, ensure that you have a valid OpenFOAM installation and compilers available by typing the following commands
-
-```
-foamRun -help
-g++ --version
-nvcc --version
-```
 
 ## Performance
 

gpu_chemistry/benchmark/others/Benchmark-others.cu → gpu_chemistry/benchmark/Benchmark-gpuChemistry.cu

@@ -8,10 +8,67 @@
 #include "ludecompose.H"
 #include "gpuODESystem.H"
 #include "makeGpuOdeSolver.H"
-
+#include "gpuMemoryResource.H"
 
 
 static constexpr TestData::Mechanism mech = TestData::GRI;
+using memoryResource_t = FoamGpu::gpuMemoryResource;
+
+
+template<class T, class R>
+__global__ void on_device(T t, R* r)
+{
+    *r = t();
+}
+
+
+
+#ifdef __NVIDIA_BACKEND__
+
+    template<class T>
+    static inline gScalar eval(T t)
+    {
+
+        gScalar *d_result;
+        gpuErrorCheck(cudaMalloc(&d_result, sizeof(gScalar)));
+        on_device<<<1,1>>>(t, d_result);
+        gpuErrorCheck(cudaGetLastError());
+        gpuErrorCheck(cudaDeviceSynchronize());
+        gScalar h_result;
+        gpuErrorCheck(cudaMemcpy(&h_result, d_result, sizeof(gScalar), cudaMemcpyDeviceToHost));
+        gpuErrorCheck(cudaDeviceSynchronize());
+        gpuErrorCheck(cudaFree(d_result));
+        gpuErrorCheck(cudaDeviceSynchronize());
+        return h_result;
+
+    }
+
+    //AMD-backend
+    #else
+
+    template<class T>
+    static inline gScalar eval(T t)
+    {
+
+        gScalar *d_result;
+        gpuErrorCheck(hipMalloc(&d_result, sizeof(gScalar)));
+        hipLaunchKernelGGL
+        (
+            on_device, dim3(1), dim3(1), 0, 0, t, d_result
+        );
+        gpuErrorCheck(hipGetLastError());
+        gpuErrorCheck(hipDeviceSynchronize());
+        gScalar h_result;
+        gpuErrorCheck(hipMemcpy(&h_result, d_result, sizeof(gScalar), hipMemcpyDeviceToHost));
+        gpuErrorCheck(hipDeviceSynchronize());
+        gpuErrorCheck(hipFree(d_result));
+        gpuErrorCheck(hipDeviceSynchronize());
+        return h_result;
+
+    }
+
+#endif
+
 
 TEST_CASE("LU")
 {
@@ -29,9 +86,9 @@ TEST_CASE("LU")
 
     device_vector<gScalar> matrix(vals.begin(), vals.end());
 
-    device_vector<gLabel> pivot(nEqns, 0);
-    device_vector<gScalar> v(nEqns, 0);
-    device_vector<gScalar> source(nEqns, 0);
+    device_vector<gLabel> pivot(nEqns);
+    device_vector<gScalar> v(nEqns);
+    device_vector<gScalar> source(nEqns);
 
     auto m_span = make_mdspan(matrix, extents<2>{nEqns, nEqns});
     auto p_span = make_mdspan(pivot, extents<1>{nEqns});
@@ -129,6 +186,7 @@ TEST_CASE("gpuODESystem"){
 
 }
 
+/*
 TEST_CASE("gpuReactionRate"){
 
     using namespace FoamGpu;
@@ -138,8 +196,10 @@ TEST_CASE("gpuReactionRate"){
     const gScalar T = 900.0;
 
 
-    device_vector<gScalar> c(nSpecie, 0.43);
-    device_vector<gScalar> ddc(nSpecie, 0.43);
+    device_vector<gScalar> c = toDeviceVector(TestData::get_concentration_vector(mech));
+    device_vector<gScalar> ddc = c;
+
+
 
     //SECTION("gpuArrheniusReactionRate")
     {
@@ -386,8 +446,8 @@ TEST_CASE("gpuReactionRate"){
             for (int i = 0; i < nReactions; ++i){
 
                 auto pair = reactions[i].k_.everything(p, T, c, ddc);
-                ret += std::get<0>(pair);
-                ret += std::get<1>(pair);
+                ret += pair[0];
+                ret += pair[1];
                 ret += ddc[0] + ddc[5] + ddc[7];
 
             }
@@ -412,10 +472,12 @@ TEST_CASE("gpuReactionRate"){
 
             return eval(op4);
         };
+
     }
 
 
 }
+*/
 
 TEST_CASE("gpuReaction"){
 
@@ -429,17 +491,12 @@ TEST_CASE("gpuReaction"){
     const gScalar p = 1E5;
     const gScalar T = 900.0;
 
-    device_vector<gScalar> c = [=](){
-        std::vector<gScalar> vals(nSpecie);
-        assign_test_concentration(vals, mech);
-        //fill_random(vals);
-        return device_vector<gScalar>(vals.begin(), vals.end());
-    }();
+    device_vector<gScalar> c = toDeviceVector(TestData::get_concentration_vector(mech));
+
+
+
+    device_vector<gScalar> dndt(nEqns);
 
-    device_vector<gScalar> dndt = [=](){
-        std::vector<gScalar> vals(nEqns);
-        return device_vector<gScalar>(vals.begin(), vals.end());
-    }();
 
     BENCHMARK("dNdtByV"){
 

gpu_chemistry/benchmark/Make/files

@@ -0,0 +1,5 @@
+Benchmark-gpuChemistry.cu
+
+EXE = $(FOAM_USER_APPBIN)/Benchmark-gpuChemistry
+
+

gpu_chemistry/benchmark/others/Make/options → gpu_chemistry/benchmark/Make/options

@@ -1,12 +1,10 @@
 EXE_INC = \
-    -I../../catchMain/lnInclude \
-    -I../../src/lnInclude \
-    -I../../src/gpuChemistryModels/lnInclude \
-    -I../../src/gpuKernelEvaluator/lnInclude \
-    -I../../../gpu_utils/common/lnInclude \
-    -I../../unittest/testHelpers/ \
-    -I../
-    -std=c++17
+    -I../catchMain/lnInclude \
+    -I../src/lnInclude \
+    -I../src/gpuChemistryModels/lnInclude \
+    -I../src/gpuKernelEvaluator/lnInclude \
+    -I../../gpu_utils/common/lnInclude \
+    -I../unittest/testHelpers/
 
 EXE_LIBS = \
     -L$(FOAM_USER_LIBBIN) \
@@ -23,5 +21,8 @@ EXE_LIBS = \
     -lfiniteVolume \
     -lchemistryModel
 
-
-include ../../nvcc
+ifeq ($(GPUFOAM_BACKEND_NVIDIA),1)
+    include ../../nvcc
+else
+    include ../../hipcc
+endif