NAMD
DeviceCUDAkernel.cu
Go to the documentation of this file.
1 #ifdef NAMD_CUDA
2 #include <cuda.h>
3 #endif // NAMD_CUDA
4 
5 #include <stdio.h>
6 #include "CudaUtils.h"
7 
8 #ifdef NAMD_CUDA
9 __global__ void read_CUDA_ARCH_kernel(int *cuda_arch) {
10  if (threadIdx.x == 0) {
11 
12 #if __CUDA_ARCH__ == 100
13  *cuda_arch = 100;
14 #elif __CUDA_ARCH__ == 110
15  *cuda_arch = 110;
16 #elif __CUDA_ARCH__ == 120
17  *cuda_arch = 120;
18 #elif __CUDA_ARCH__ == 130
19  *cuda_arch = 130;
20 #elif __CUDA_ARCH__ == 200
21  *cuda_arch = 200;
22 #elif __CUDA_ARCH__ == 210
23  *cuda_arch = 210;
24 #elif __CUDA_ARCH__ == 300
25  *cuda_arch = 300;
26 #elif __CUDA_ARCH__ == 350
27  *cuda_arch = 350;
28 #elif __CUDA_ARCH__ == 500
29  *cuda_arch = 500;
30 #else
31  *cuda_arch = 500;
32 #endif
33 
34  }
35 }
36 
37 //
38 // Reads the value of __CUDA_ARCH__ from device code
39 //
40 int read_CUDA_ARCH() {
41  int *d_cuda_arch;
42  int h_cuda_arch;
43  allocate_device<int>(&d_cuda_arch, 1);
44 
45  read_CUDA_ARCH_kernel <<< 1, 1 >>> (d_cuda_arch);
46  cudaError_t err = cudaGetLastError();
47  if (err != cudaSuccess) {
48  char str[1024];
49  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",
50  __FILE__,cudaGetErrorString(err));
51  cudaDie(str);
52  }
53  cudaCheck(cudaDeviceSynchronize());
54 
55  copy_DtoH_sync<int>(d_cuda_arch, &h_cuda_arch, 1);
56  deallocate_device<int>(&d_cuda_arch);
57 
58  return h_cuda_arch;
59 }
60 
61 #endif // NAMD_CUDA
Generated by   doxygen 1.8.14