DeviceCUDAkernel.cu

Go to the documentation of this file.
00001 #include <stdio.h>
00002 #include "CudaUtils.h"
00003 
00004 #ifdef NAMD_CUDA
00005 #include <cuda.h>
00006 
00007 __global__ void read_CUDA_ARCH_kernel(int *cuda_arch) {
00008   if (threadIdx.x == 0) {
00009 
00010 #if __CUDA_ARCH__ == 100
00011     *cuda_arch = 100;
00012 #elif __CUDA_ARCH__ == 110
00013     *cuda_arch = 110;
00014 #elif __CUDA_ARCH__ == 120
00015     *cuda_arch = 120;
00016 #elif __CUDA_ARCH__ == 130
00017     *cuda_arch = 130;
00018 #elif __CUDA_ARCH__ == 200
00019     *cuda_arch = 200;
00020 #elif __CUDA_ARCH__ == 210
00021     *cuda_arch = 210;
00022 #elif __CUDA_ARCH__ == 300
00023     *cuda_arch = 300;
00024 #elif __CUDA_ARCH__ == 350
00025     *cuda_arch = 350;
00026 #elif __CUDA_ARCH__ == 500
00027     *cuda_arch = 500;
00028 #else
00029     *cuda_arch = 500;
00030 #endif
00031 
00032   }
00033 }
00034 
00035 //
00036 // Reads the value of __CUDA_ARCH__ from device code
00037 //
00038 int read_CUDA_ARCH() {
00039   int *d_cuda_arch;
00040   int h_cuda_arch;
00041   allocate_device<int>(&d_cuda_arch, 1);
00042   
00043   read_CUDA_ARCH_kernel <<< 1, 1 >>> (d_cuda_arch);
00044   cudaError_t err = cudaGetLastError();
00045   if (err != cudaSuccess) {
00046     char str[1024];
00047     sprintf(str, "Error executing CUDA kernel read_CUDA_ARCH_kernel in file %s\nError string: %s\nPossible cause: Device compute capability is less than the compute capability the code was compiled for.\n",
00048       __FILE__,cudaGetErrorString(err));
00049     cudaDie(str);
00050   }
00051   cudaCheck(cudaDeviceSynchronize());
00052 
00053   copy_DtoH_sync<int>(d_cuda_arch, &h_cuda_arch, 1);
00054   deallocate_device<int>(&d_cuda_arch);
00055 
00056   return h_cuda_arch;
00057 }
00058 
00059 #endif // NAMD_CUDA

Generated on Sat Nov 18 01:17:13 2017 for NAMD by  doxygen 1.4.7