--- /dev/null
+#include <iostream>\r
+#include <stdio.h>\r
+#include <stdlib.h>\r
+#include <assert.h>\r
+\r
+#include "Device.h"\r
+#include "cudaTools.h"\r
+#include "Chronos.h"\r
+#include "omp.h"\r
+#include "StringTools.h"\r
+\r
+using std::string;\r
+using std::cout;\r
+using std::cerr;\r
+using std::endl;\r
+using std::flush;\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |* Declaration *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |* Public *|\r
+ \*-------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |* Private *|\r
+ \*-------------------------------------*/\r
+\r
+static void cout01(int isTrue);\r
+static void coutBool(bool isFlag);\r
+static int dim(const dim3& dim);\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |* Implementation *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |* Public *|\r
+ \*-------------------------------------*/\r
+\r
+/*--------------*\\r
+|* wrapper *|\r
+ \*-------------*/\r
+\r
+void Device::synchronize(void)\r
+ {\r
+ cudaDeviceSynchronize();\r
+ }\r
+\r
+/*--------------*\\r
+|* Tools *|\r
+ \*-------------*/\r
+\r
+void Device::checkKernelError(const char *message)\r
+ {\r
+ cudaError_t error = cudaGetLastError();\r
+ if (error != cudaSuccess)\r
+ {\r
+ if (message != NULL)\r
+ {\r
+ fprintf(stderr, "\n[CUDA ERROR] : Kernel Execution Failed : %s: %s\n\n", message, cudaGetErrorString(error));\r
+ }\r
+ else\r
+ {\r
+ fprintf(stderr, "\n[CUDA ERROR] : Kernel Execution Failed : %s\n\n", cudaGetErrorString(error));\r
+ }\r
+ exit (EXIT_FAILURE);\r
+ }\r
+ }\r
+\r
+void Device::checkDimError(const dim3& dg, const dim3& db)\r
+ {\r
+ assertDim(dg, db);\r
+ checkDimOptimiser(dg, db);\r
+ }\r
+\r
+void Device::checkDimOptimiser(const dim3& dg, const dim3& db)\r
+ {\r
+ int mpCount = Device::getMPCount();\r
+ int warpSize = Device::getWarpSize();\r
+\r
+ int nbBlock = dim(dg);\r
+ int nbThreadBlock = dim(db);\r
+\r
+ // grid\r
+ if (nbBlock < mpCount || nbBlock % mpCount != 0)\r
+ {\r
+ string messageGrid = "nbBlock = " + StringTools::toString(nbBlock) + " : Heuristic : For best performance, for nbBlock, try a multiple of #MP="\r
+ + StringTools::toString(mpCount);\r
+\r
+ cout << "dg" << "(" << dg.x << "," << dg.y << "," << dg.z << ") " << messageGrid << endl;\r
+ }\r
+\r
+ // block\r
+ if (nbThreadBlock < warpSize || nbThreadBlock % warpSize != 0) // TODO kepler\r
+ {\r
+ string messageBlock = "nbThreadBlock = " + StringTools::toString(nbThreadBlock)\r
+ + " : Heuristic : For best performance, for nbThreadBlock, try a multiple of |warp|=" + StringTools::toString(warpSize);\r
+\r
+ cout << "db" << "(" << db.x << "," << db.y << "," << db.z << ") " << messageBlock << endl;\r
+ }\r
+ }\r
+\r
+void Device::assertDim(const dim3& dg, const dim3& db)\r
+ {\r
+ // grid\r
+ {\r
+ dim3 dimGridMax = Device::getMaxGridDim();\r
+\r
+ assert(dg.x <= dimGridMax.x);\r
+ assert(dg.y <= dimGridMax.y);\r
+ assert(dg.z <= dimGridMax.z);\r
+ }\r
+\r
+ // block\r
+ {\r
+ dim3 dimBlockMax = Device::getMaxBlockDim();\r
+\r
+ assert(db.x <= dimBlockMax.x);\r
+ assert(db.y <= dimBlockMax.y);\r
+ assert(db.z <= dimBlockMax.z);\r
+ }\r
+\r
+ // Thread per block\r
+ assert(dim(db) <= getMaxThreadPerBlock());\r
+ }\r
+\r
+void Device::print(const dim3& dg, const dim3& db)\r
+ {\r
+ cout << "dg" << "(" << dg.x << "," << dg.y << "," << dg.z << ") " << endl;\r
+ cout << "db" << "(" << db.x << "," << db.y << "," << db.z << ") " << endl;\r
+ }\r
+\r
+int Device::dim(const dim3& dim)\r
+ {\r
+ return dim.x * dim.y * dim.z;\r
+ }\r
+\r
+int Device::nbThread(const dim3& dg, const dim3& db)\r
+ {\r
+ return dim(dg) * dim(db);\r
+ }\r
+\r
+/*--------------*\\r
+|* get *|\r
+ \*-------------*/\r
+\r
+int Device::getDeviceId(void)\r
+ {\r
+ int deviceId;\r
+ HANDLE_ERROR(cudaGetDevice(&deviceId));\r
+\r
+ return deviceId;\r
+ }\r
+\r
+int Device::getDeviceCount(void)\r
+ {\r
+ int nbDevice;\r
+ HANDLE_ERROR(cudaGetDeviceCount(&nbDevice));\r
+\r
+ return nbDevice;\r
+ }\r
+\r
+cudaDeviceProp Device::getDeviceProp(int idDevice)\r
+ {\r
+ cudaDeviceProp prop;\r
+ HANDLE_ERROR(cudaGetDeviceProperties(&prop, idDevice));\r
+ return prop;\r
+ }\r
+\r
+cudaDeviceProp Device::getDeviceProp(void)\r
+ {\r
+ return getDeviceProp(getDeviceId());\r
+ }\r
+\r
+/*--------------------------------------*\\r
+ |* Secondaire *|\r
+ \*-------------------------------------*/\r
+\r
+/*--------------*\\r
+|* get *|\r
+ \*-------------*/\r
+\r
+string Device::getNameSimple(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).name;\r
+ }\r
+\r
+string Device::getNameSimple()\r
+ {\r
+ return getNameSimple(getDeviceId());\r
+ }\r
+\r
+string Device::getName(int idDevice)\r
+ {\r
+ string id = StringTools::toString(idDevice);\r
+ string a = StringTools::toString(getCapacityMajor(idDevice));\r
+ string b = StringTools::toString(getCapacityMinor(idDevice));\r
+\r
+ return "[" + getNameSimple(idDevice) + "] id = " + id + " sm=" + a + "" + b;\r
+ }\r
+\r
+string Device::getName()\r
+ {\r
+ return getName(getDeviceId());\r
+ }\r
+\r
+dim3 Device::getMaxGridDim(int idDevice)\r
+ {\r
+ cudaDeviceProp prop = getDeviceProp(idDevice);\r
+\r
+ return dim3(prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);\r
+ }\r
+\r
+dim3 Device::getMaxGridDim()\r
+ {\r
+ return getMaxGridDim(getDeviceId());\r
+ }\r
+\r
+dim3 Device::getMaxBlockDim(int idDevice)\r
+ {\r
+ cudaDeviceProp prop = getDeviceProp(idDevice);\r
+\r
+ return dim3(prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);\r
+ }\r
+\r
+dim3 Device::getMaxBlockDim()\r
+ {\r
+ return getMaxBlockDim(getDeviceId());\r
+ }\r
+\r
+int Device::getMaxThreadPerBlock(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).maxThreadsPerBlock;\r
+ }\r
+\r
+int Device::getMaxThreadPerBlock()\r
+ {\r
+ return getMaxThreadPerBlock(getDeviceId());\r
+ }\r
+\r
+int Device::getWarpSize(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).warpSize;\r
+ }\r
+\r
+int Device::getWarpSize(void)\r
+ {\r
+ return getWarpSize(getDeviceId());\r
+ }\r
+\r
+int Device::getMPCount(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).multiProcessorCount;\r
+ }\r
+\r
+int Device::getMPCount(void)\r
+ {\r
+ return getMPCount(getDeviceId());\r
+ }\r
+\r
+int Device::getCapacityMajor(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).major;\r
+ }\r
+\r
+int Device::getCapacityMajor()\r
+ {\r
+ return getCapacityMajor(getDeviceId());\r
+ }\r
+\r
+int Device::getCapacityMinor(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).minor;\r
+ }\r
+\r
+int Device::getCapacityMinor()\r
+ {\r
+ return getCapacityMinor(getDeviceId());\r
+ }\r
+\r
+int Device::getRuntimeVersion()\r
+ {\r
+ int version = -1;\r
+ cudaRuntimeGetVersion(&version);\r
+ return version;\r
+\r
+ // return CUDART_VERSION;\r
+ }\r
+\r
+int Device::getDriverVersion(void)\r
+ {\r
+ int version = -1;\r
+ cudaDriverGetVersion(&version);\r
+ return version;\r
+ }\r
+\r
+int Device::getAsyncEngineCount(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).asyncEngineCount;\r
+ }\r
+\r
+int Device::getAsyncEngineCount()\r
+ {\r
+ return getAsyncEngineCount(getDeviceId());\r
+ }\r
+\r
+/*--------------*\\r
+|* is *|\r
+ \*-------------*/\r
+\r
+bool Device::isCuda(void)\r
+ {\r
+ return getDeviceCount() >= 1;\r
+ }\r
+\r
+bool Device::isFermi(int idDevice)\r
+ {\r
+ int c = getCapacityMajor(idDevice);\r
+\r
+ return c >= 2 && c < 3;\r
+ }\r
+\r
+bool Device::isFermi()\r
+ {\r
+ return isFermi(getDeviceId());\r
+ }\r
+\r
+bool Device::isKepler(int idDevice)\r
+ {\r
+ int c = getCapacityMajor(idDevice);\r
+\r
+ return c >= 3 && c < 4;\r
+ }\r
+\r
+bool Device::isKepler()\r
+ {\r
+ return isKepler(getDeviceId());\r
+ }\r
+\r
+/**\r
+ * 64 bits only\r
+ */\r
+bool Device::isUVAEnable(int idDevice)\r
+ {\r
+ return getCapacityMajor() >= 2.0 && getRuntimeVersion() >= 4000;\r
+ }\r
+\r
+/**\r
+ * 64 bits only\r
+ */\r
+bool Device::isUVAEnable()\r
+ {\r
+ return isUVAEnable(getDeviceId());\r
+ }\r
+\r
+bool Device::isAtomicShareMemoryEnable(int idDevice)\r
+ {\r
+ return (getCapacityMajor(idDevice) == 1 && Device::getCapacityMinor(idDevice) >= 2) || getCapacityMajor(idDevice) >= 2;\r
+ }\r
+\r
+bool Device::isAtomicShareMemoryEnable()\r
+ {\r
+ return isAtomicShareMemoryEnable(getDeviceId());\r
+ }\r
+\r
+bool Device::isHostMapMemoryEnable(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).canMapHostMemory;\r
+ }\r
+\r
+bool Device::isHostMapMemoryEnable()\r
+ {\r
+ return isHostMapMemoryEnable(getDeviceId());\r
+ }\r
+\r
+bool Device::isECCEnable(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).ECCEnabled;\r
+ }\r
+\r
+bool Device::isECCEnable(void)\r
+ {\r
+ return isECCEnable(getDeviceId());\r
+ }\r
+\r
+bool Device::isAsyncEngine(int idDevice)\r
+ {\r
+ return getDeviceProp(idDevice).deviceOverlap;\r
+ }\r
+\r
+bool Device::isAsyncEngine(void)\r
+ {\r
+ return isAsyncEngine(getDeviceId());\r
+ }\r
+\r
+/*--------------------------------------*\\r
+ |* Print *|\r
+ \*-------------------------------------*/\r
+\r
+void Device::print(int idDevice)\r
+ {\r
+ cudaDeviceProp prop = getDeviceProp(idDevice);\r
+\r
+ cout << endl;\r
+ cout << "===========================================" << endl;\r
+ cout << " " << prop.name << " : id=" << idDevice << " : sm=" << prop.major << "" << prop.minor << " :" << endl;\r
+ cout << "===========================================" << endl;\r
+\r
+ cout << endl;\r
+ cout << "Device id : " << idDevice << endl;\r
+ cout << "Name : " << prop.name << endl;\r
+ cout << "GPU capability : " << prop.major << "." << prop.minor << "" << endl;\r
+ cout << "is Fermi : ";\r
+ coutBool(isFermi(idDevice));\r
+ cout << "is Kepler : ";\r
+ coutBool(isKepler(idDevice));\r
+ cout << "Clock rate : " << prop.clockRate / 1000 << " MHZ" << endl;\r
+ cout << "GPU integrated on MB : ";\r
+ cout01(prop.integrated);\r
+ // cout << "ComputeMode : " << prop.computeMode << endl;\r
+\r
+ cout << endl;\r
+ cout << "Kernel : " << endl;\r
+ cout << "Limit execution (timeout): ";\r
+ cout01(prop.kernelExecTimeoutEnabled);\r
+\r
+ cout << endl;\r
+ cout << "Memory : " << endl;\r
+ cout << "Global Memory : " << prop.totalGlobalMem / 1024 / 1024 << " MB" << endl;\r
+ cout << "Constant Memory : " << prop.totalConstMem / 1024 << " KB" << endl;\r
+ cout << "Texture1D max size : (" << prop.maxTexture1D << ")" << endl;\r
+ cout << "Texture2D max size : (" << prop.maxTexture2D[0] << "," << prop.maxTexture2D[1] << ")" << endl;\r
+ cout << "Texture3D max size : (" << prop.maxTexture3D[0] << "," << prop.maxTexture3D[1] << "," << prop.maxTexture3D[2] << ")" << endl;\r
+ //cout << "Texture2D Array max Size : (" << ptrProp.maxTexture2DArray[0] << "," << ptrProp.maxTexture2DArray[1] << "," << ptrProp.maxTexture2DArray[2] << ")"<< endl;\r
+ cout << "Texture Alignment : " << prop.textureAlignment << " B" << endl;\r
+ cout << "Max mem pitch : " << prop.memPitch << endl;\r
+\r
+ cout << endl;\r
+ cout << "Multiprocesseur(MP) : " << endl;\r
+ cout << "MP count : " << prop.multiProcessorCount << endl;\r
+ cout << "Shared memory per block : " << prop.sharedMemPerBlock / 1024 << " KB " << endl;\r
+ cout << "Register memory per block: " << prop.regsPerBlock / 1024 << " KB " << endl;\r
+ cout << "Max threads per block : " << prop.maxThreadsPerBlock << endl;\r
+ cout << "Max block dim : (" << prop.maxThreadsDim[0] << "," << prop.maxThreadsDim[1] << "," << prop.maxThreadsDim[2] << ")" << endl;\r
+ cout << "Max grid dim : (" << prop.maxGridSize[0] << "," << prop.maxGridSize[1] << "," << prop.maxGridSize[2] << ")" << endl;\r
+ cout << "Threads in warp : " << prop.warpSize << endl;\r
+\r
+ cout << endl;\r
+ cout << "GPU Capacity : " << endl;\r
+ cout << "MapHostMemory : ";\r
+ cout01(isHostMapMemoryEnable());\r
+ cout << "AtomicOperation sharedMemory : ";\r
+ cout01(isAtomicShareMemoryEnable());\r
+\r
+ cout << "UVA (Unified Virtual Addressing) : ";\r
+ cout01(isUVAEnable());\r
+\r
+ cout << "ECCEnabled : ";\r
+ cout01(prop.ECCEnabled);\r
+\r
+ cout << "Concurrent bidirectional copy (DMA only!) : ";\r
+ cout01(prop.deviceOverlap);\r
+ if (prop.deviceOverlap)\r
+ {\r
+ cout << "Concurrent bidirectional copy : host->device // device->host : " << getAsyncEngineCount(idDevice) << endl;\r
+ }\r
+\r
+ cout << "Concurrent Kernels (fermi 16x) : ";\r
+ cout01(prop.concurrentKernels);\r
+\r
+ cout << endl;\r
+ if (Device::getDeviceCount() >= 2)\r
+ {\r
+ printP2PmatrixCompatibility();\r
+ }\r
+\r
+ cout << "Cuda Runtime version : " << getRuntimeVersion() << endl;\r
+ cout << "Cuda Driver version : " << getDriverVersion() << endl;\r
+ cout << endl;\r
+ // cout << "=================== end ========================" << endl;\r
+\r
+// cout << endl;\r
+// cout << "===========================================" << endl;\r
+// cout << "Cuda Runtime version : " << getRuntimeVersion() << endl;\r
+// cout << "===========================================" << endl;\r
+// cout << endl;\r
+ }\r
+\r
+void Device::print()\r
+ {\r
+ print(getDeviceId());\r
+ }\r
+\r
+void Device::printAll()\r
+ {\r
+ cout << "\nList of all GPU available :" << endl;\r
+ printAllSimple();\r
+\r
+ cout << endl << "Details :" << endl;\r
+ int deviceCount = getDeviceCount();\r
+\r
+ for (int id = 0; id < deviceCount; id++)\r
+ {\r
+ print(id);\r
+ }\r
+\r
+ printCurrent();\r
+ }\r
+\r
+void Device::printAllSimple()\r
+ {\r
+ cout << endl;\r
+ cout << "==============================================================" << endl;\r
+ cout << "[CUDA] : List GPU Available : cuda version = " << getRuntimeVersion() <<endl;\r
+ cout << "==============================================================" << endl;\r
+ cout << endl;\r
+\r
+ int current=getDeviceId();\r
+ for (int id = 0; id < getDeviceCount(); id++)\r
+ {\r
+ cout << getName(id) ;\r
+ if(id == current)\r
+ {\r
+ cout<< " : [CURRENT]";\r
+ }\r
+ cout << endl;\r
+ }\r
+\r
+ cout << endl;\r
+ }\r
+\r
+void Device::printCurrent()\r
+ {\r
+ cout << "==============================================================" << endl;\r
+ cout << "[Cuda] : current device : " << getName() << endl;\r
+ cout << "==============================================================" << endl;\r
+ }\r
+\r
+/*--------------*\\r
+|* load *|\r
+ \*-------------*/\r
+\r
+/**\r
+ * Linux : nvidia-smi -pm 1 utile? TODO\r
+ * marche pas pour opengl\r
+ */\r
+void Device::loadCudaDriver(int deviceID, bool isMapMemoryAsk)\r
+ {\r
+ Chronos chrono;\r
+ cout << "\nDevice(" << deviceID << ") : Load Driver ";\r
+\r
+ int* ptrBidon;\r
+\r
+ HANDLE_ERROR(cudaSetDevice(deviceID));\r
+\r
+ if (isHostMapMemoryEnable() && isMapMemoryAsk)\r
+ {\r
+ HANDLE_ERROR(cudaSetDeviceFlags(cudaDeviceMapHost));\r
+ cout << "(HostMapMemory activate) : "<<flush;\r
+ }\r
+ else if (!isHostMapMemoryEnable() && isMapMemoryAsk)\r
+ {\r
+ cerr << "(HostMapMemory not enable) : "<<flush;\r
+ }\r
+\r
+ HANDLE_ERROR(cudaMalloc((void** ) &ptrBidon, sizeof(int)));\r
+ HANDLE_ERROR(cudaFree(ptrBidon));\r
+\r
+ chrono.stop();\r
+ cout << chrono.getDeltaTime() << " (s)" << endl<<endl;\r
+ }\r
+\r
+void Device::loadCudaDriver(bool isMapMemoryEnable)\r
+ {\r
+ loadCudaDriver(getDeviceId(), isMapMemoryEnable);\r
+ }\r
+\r
+void Device::loadCudaDriverAll(bool isMapMemoryEnable)\r
+ {\r
+ cout << "\nLoad Cuda Driver : start ..." << endl;\r
+ Chronos chrono;\r
+\r
+ int k = Device::getDeviceCount();\r
+ //omp_set_num_threads(k);\r
+ //#pragma omp parallel for\r
+ for (int i = 0; i < k; i++)\r
+ {\r
+ loadCudaDriver(i, isMapMemoryEnable);\r
+ }\r
+\r
+ chrono.stop();\r
+ cout << "Load Cuda Driver : end : " << chrono.getDeltaTime() << " (s)\n" << endl;\r
+ }\r
+\r
+/*--------------*\\r
+|* p2p *|\r
+ \*-------------*/\r
+\r
+void Device::printP2PmatrixCompatibility()\r
+ {\r
+ int* matrixP2PCompatibility = p2pMatrixCompatibility();\r
+ int* ptrMatrixP2PCompatibility = matrixP2PCompatibility;\r
+\r
+ int n = Device::getDeviceCount();\r
+ cout << "P2P compatibility : symetric matrix (" << n << "x" << n << "):" << endl;\r
+ for (int i = 0; i < n; i++)\r
+ {\r
+ for (int j = 0; j < n; j++)\r
+ {\r
+ if (i != j)\r
+ {\r
+ cout << *ptrMatrixP2PCompatibility << " ";\r
+ }\r
+ else\r
+ {\r
+ cout << " ";\r
+ }\r
+\r
+ ptrMatrixP2PCompatibility++;\r
+ }\r
+ cout << endl;\r
+ }\r
+ cout << endl;\r
+\r
+ delete[] matrixP2PCompatibility;\r
+ }\r
+\r
+void Device::p2pEnableALL()\r
+ {\r
+ int n = Device::getDeviceCount();\r
+\r
+ int* matrixP2PCompatibility = p2pMatrixCompatibility();\r
+ int* ptrMatrixP2PCompatibility = matrixP2PCompatibility;\r
+\r
+ if (n >= 2)\r
+ {\r
+ cout << "P2P enable : symetric matrix (" << n << "x" << n << "):" << endl;\r
+ }\r
+\r
+ for (int i = 0; i < n; i++)\r
+ {\r
+ for (int j = 0; j < n; j++)\r
+ {\r
+ if (i != j)\r
+ {\r
+ cout << *ptrMatrixP2PCompatibility << " ";\r
+ }\r
+ else\r
+ {\r
+ cout << " ";\r
+ }\r
+\r
+ if (*ptrMatrixP2PCompatibility)\r
+ {\r
+ int flaginutile = 0;\r
+ HANDLE_ERROR(cudaSetDevice(i));\r
+ HANDLE_ERROR(cudaDeviceEnablePeerAccess(j, flaginutile));\r
+ }\r
+ ptrMatrixP2PCompatibility++;\r
+ }\r
+ }\r
+\r
+ delete[] matrixP2PCompatibility;\r
+ }\r
+\r
+int* Device::p2pMatrixCompatibility()\r
+ {\r
+ int n = Device::getDeviceCount();\r
+\r
+ int* matrixP2PCompatibility = new int[n * n];\r
+ int* ptrMatrixP2PCompatibility = matrixP2PCompatibility;\r
+\r
+ for (int i = 0; i < n; i++)\r
+ {\r
+ for (int j = 0; j < n; j++)\r
+ {\r
+ int isP2PAutorized01;\r
+ HANDLE_ERROR(cudaDeviceCanAccessPeer(&isP2PAutorized01, i, j));\r
+\r
+ *ptrMatrixP2PCompatibility++ = isP2PAutorized01;\r
+ }\r
+ }\r
+\r
+ return matrixP2PCompatibility;\r
+ }\r
+\r
+/*--------------------------------------*\\r
+ |* Private *|\r
+ \*-------------------------------------*/\r
+\r
+void cout01(int isTrue)\r
+ {\r
+ if (isTrue)\r
+ cout << "True" << endl;\r
+ else\r
+ cout << "False" << endl;\r
+ }\r
+\r
+void coutBool(bool isFlag)\r
+ {\r
+ if (isFlag)\r
+ cout << "True" << endl;\r
+ else\r
+ cout << "False" << endl;\r
+ }\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |* End *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
projects / GPU.git / blobdiff