Use read only cache instead of texture cache for CC >= 3.5

SimFlowCFD · Jan 5, 2020 · 6b40d01 · 6b40d01
1 parent b1772a3
commit 6b40d01
Show file tree

Hide file tree

Showing 5 changed files with 95 additions and 40 deletions.
diff --git a/src/OpenFOAM/Make/files b/src/OpenFOAM/Make/files
@@ -105,6 +105,7 @@ $(ranges)/labelRange/labelRanges.C
 $(ranges)/scalarRange/scalarRange.C
 $(ranges)/scalarRange/scalarRanges.C
 
+containers/Lists/gpuList/gpuConfig.C
 containers/HashTables/HashTable/HashTableCore.C
 containers/HashTables/StaticHashTable/StaticHashTableCore.C
 containers/Lists/PackedList/PackedListCore.C

diff --git a/src/OpenFOAM/containers/Lists/gpuList/gpuConfig.C b/src/OpenFOAM/containers/Lists/gpuList/gpuConfig.C
@@ -0,0 +1,42 @@
+#include "gpuConfig.H"
+
+namespace Foam {
+
+    int deviceCount()
+    {
+        int num_devices;
+        CUDA_CALL(cudaGetDeviceCount(&num_devices));
+        return num_devices;
+    }
+
+    int currentDevice()
+    {
+        int device;
+        CUDA_CALL(cudaGetDevice(&device));
+        return device;
+    }
+
+    void setCurrentDevice(int device)
+    {
+        CUDA_CALL(cudaSetDevice(device));
+    }
+
+    int deviceComputeCapability(int device)
+    {
+        cudaDeviceProp deviceProp;
+        CUDA_CALL(cudaGetDeviceProperties(&deviceProp, device));
+        return 10*deviceProp.major + deviceProp.minor;
+    }
+
+    int currentComputeCapability()
+    {
+        return deviceComputeCapability(currentDevice());
+    }
+
+    bool needTextureBind()
+    {
+        static bool needBind = currentComputeCapability() < 35;
+        return needBind;
+    }
+
+}
diff --git a/src/OpenFOAM/containers/Lists/gpuList/gpuConfig.H b/src/OpenFOAM/containers/Lists/gpuList/gpuConfig.H
@@ -12,6 +12,7 @@
 #include <thrust/iterator/permutation_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/discard_iterator.h>
 
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
@@ -37,26 +38,19 @@ namespace gpu_api = thrust;
 
 #define GPU_ERROR_CHECK()                                  \
  cudaDeviceSynchronize();                                  \
- CUDA_CALL( cudaPeekAtLastError());    
+ CUDA_CALL( cudaPeekAtLastError());
 
 #define GPU_ERROR_CHECK_ASYNC()                            \
- CUDA_CALL(cudaPeekAtLastError());    
+ CUDA_CALL(cudaPeekAtLastError());
 
 namespace Foam
 {
-
-inline int getGpuDeviceCount()
-{
-    int num_devices;
-    CUDA_CALL(cudaGetDeviceCount(&num_devices));
-    return num_devices;
-}
-
-inline void setGpuDevice(int device)
-{
-   CUDA_CALL(cudaSetDevice(device));
-}
-
+    int deviceCount();
+    int currentDevice();
+    void setCurrentDevice(int device);
+    int deviceComputeCapability(int device);
+    int currentComputeCapability();
+    bool needTextureBind();
 }
 
 #else

diff --git a/src/OpenFOAM/containers/Lists/gpuList/textures.H b/src/OpenFOAM/containers/Lists/gpuList/textures.H
@@ -13,46 +13,45 @@ private:
     const T* data;
 
     inline void initResourceDesc(cudaResourceDesc& resDesc);
-    void init(int n, T* data_);
+    void bind(int n);
 
 public:
     textures(int n, T* _data):
         tex(0),
         data(_data)
     {
-        init(n,_data);
+        if(needTextureBind())
+            bind(n);
     }
 
     textures(const gpuList<T>& list):
-        tex(0)
+        tex(0),
+        data(list.data())
     {
-        init(list.size(),const_cast<T*>(list.data()));
+        if(needTextureBind())
+            bind(list.size());
     }
 
     inline __device__ T operator[](const int& i) const;
 
     void destroy()
     {
-        cudaDestroyTextureObject(tex);
+        if(needTextureBind())
+            cudaDestroyTextureObject(tex);
     }
 };
 
 template<class T>
-inline void textures<T>::initResourceDesc(cudaResourceDesc& resDesc)
-{
-}
-
-template<class T>
-inline void textures<T>::init(int n, T* _data)
+inline void textures<T>::bind(int n)
 {
     cudaResourceDesc resDesc;
     memset(&resDesc, 0, sizeof(cudaResourceDesc));
-   
+
     resDesc.resType = cudaResourceTypeLinear;
-    resDesc.res.linear.devPtr = _data;
+    resDesc.res.linear.devPtr = const_cast<T*>(data);
     resDesc.res.linear.sizeInBytes = n*sizeof(T);
 
-    initResourceDesc(resDesc); 
+    initResourceDesc(resDesc);
 
     cudaTextureDesc texDesc;
     memset(&texDesc, 0, sizeof(cudaTextureDesc));
@@ -61,29 +60,46 @@ inline void textures<T>::init(int n, T* _data)
     cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);
 }
 
+
+template<class T>
+inline void textures<T>::initResourceDesc(cudaResourceDesc& resDesc)
+{
+}
+
+
 template<>
 inline void textures<float>::initResourceDesc(cudaResourceDesc& resDesc)
 {
     resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
-    resDesc.res.linear.desc.x = 32; 
+    resDesc.res.linear.desc.x = 32;
 }
 
 
 template<>
 inline void textures<int>::initResourceDesc(cudaResourceDesc& resDesc)
 {
     resDesc.res.linear.desc.f = cudaChannelFormatKindSigned;
-    resDesc.res.linear.desc.x = 32; 
+    resDesc.res.linear.desc.x = 32;
 }
 
 template<>
 inline void textures<double>::initResourceDesc(cudaResourceDesc& resDesc)
 {
     resDesc.res.linear.desc.f = cudaChannelFormatKindSigned;
-    resDesc.res.linear.desc.x = 32; 
-    resDesc.res.linear.desc.y = 32; 
+    resDesc.res.linear.desc.x = 32;
+    resDesc.res.linear.desc.y = 32;
+}
+
+#if __CUDA_ARCH__ >= 350
+
+template<class T>
+inline __device__ T textures<T>::operator[](const int& i) const
+{
+    return __ldg(data + i);
 }
 
+#else
+
 template<>
 inline __device__ float textures<float>::operator[](const int& i) const
 {
@@ -103,4 +119,6 @@ inline __device__ double textures<double>::operator[](const int& i) const
     return __hiloint2double(v.y, v.x);
 }
 
+#endif
+
 }
diff --git a/src/OpenFOAM/global/argList/argList.C b/src/OpenFOAM/global/argList/argList.C
@@ -586,7 +586,7 @@ void Foam::argList::parse
     // Case is a single processor run unless it is running parallel
     int nProcs = 1;
 
-    int deviceCount = getGpuDeviceCount();
+    int nDeviceCount = deviceCount();
 
     // Roots if running distributed
     fileNameList roots;
@@ -787,27 +787,27 @@ void Foam::argList::parse
             {
                 int device = devices[Pstream::myProcNo()];
 
-                if(device < 0 || device >= deviceCount)
+                if(device < 0 || device >= nDeviceCount)
                 {
                     FatalError
                         <<"Invalid device ID: "<<device
                         <<" for processor "<<Pstream::myProcNo()<<endl;
                     FatalError.exit();
                 }
 
-                setGpuDevice(device);
+                setCurrentDevice(device);
             }
         }
         else
         {
-            if(Pstream::myProcNo() >= deviceCount)
+            if(Pstream::myProcNo() >= nDeviceCount)
             {
                 FatalError
                     <<"Specify device IDs with 'devices' argument"<<endl;
                 FatalError.exit();
             }
 
-            setGpuDevice(Pstream::myProcNo());
+            setCurrentDevice(Pstream::myProcNo());
         }
     }
     else
@@ -820,14 +820,14 @@ void Foam::argList::parse
         {
             int device = optionRead<int>("device");
 
-            if(device < 0 || device >= deviceCount)
+            if(device < 0 || device >= nDeviceCount)
             {
                 FatalError
                     <<"Invalid device ID: "<<device<<endl;
                 FatalError.exit();
             }
 
-            setGpuDevice(device);
+            setCurrentDevice(device);
 
         }
     }