| 41 | #define DEVICE __forceinline__ __device__ |
| 42 | |
| 43 | int get_sm_count() { |
| 44 | int device_id; |
| 45 | cudaError_t result = cudaGetDevice(&device_id); |
| 46 | if (result != cudaSuccess) { |
| 47 | std::cerr << "cudaGetDevice() returned error " << cudaGetErrorString(result) |
| 48 | << std::endl; |
| 49 | return 1; |
| 50 | } |
| 51 | int multiprocessor_count; |
| 52 | result = cudaDeviceGetAttribute(&multiprocessor_count, |
| 53 | cudaDevAttrMultiProcessorCount, device_id); |
| 54 | if (result != cudaSuccess) { |
| 55 | std::cerr << "cudaDeviceGetAttribute() returned error " |
| 56 | << cudaGetErrorString(result) << std::endl; |
| 57 | return 1; |
| 58 | } |
| 59 | return multiprocessor_count; |
| 60 | } |
| 61 | |
| 62 | static constexpr int MAX_CLUSTER_SIZE = 16; |
| 63 | static constexpr int WARP_GROUP_SIZE = 128; |
nothing calls this directly
no outgoing calls
no test coverage detected