Advertisement
Guest User

Untitled

a guest
May 11th, 2016
77
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 3.19 KB | None | 0 0
  1. #include <iostream>
  2. #include <utility>
  3. #include <cuda_runtime.h>
  4. using namespace std;
  5.  
  6. /*
  7.  * Fake destructor for C-like resources, to cope with exceptions. E.g.:
  8.  * void* array = malloc(123);
  9.  * destructor([=]{ free(array); });
  10.  */
  11.  
  12. #define destructor(f) destructor_helper_macro_1(f, __LINE__)
  13.  
  14. template<typename F> struct destructor_helper{
  15.     F f;
  16.     ~destructor_helper(){ f(); }
  17. };
  18. template<typename F> destructor_helper<F> make_destructor_helper(F&& f){
  19.     return destructor_helper<F>{std::move(f)};
  20. }
  21. #define destructor_helper_macro_2(f, l) auto destructor_ ## l = make_destructor_helper(f)
  22. #define destructor_helper_macro_1(f, l) destructor_helper_macro_2(f, l)
  23.  
  24. #define cudacheckpoint_(file, line) { cudaDeviceSynchronize() assertcu; std::cerr<<"cudacheckpoint "<<file<<":"<<line<<std::endl; }
  25. #define cudacheckpoint cudacheckpoint_(__FILE__, __LINE__)
  26.  
  27. /*
  28.  * Cuda errors to exceptions translation.
  29.  * Append to cuda/cufft calls assertcu/assertcufft, e.g. cudaMemcpy(...) assertcu;
  30.  */
  31.  
  32. #define assertcufft >>assertcufft_helper()
  33. #define assertcu >>assertcu_helper()
  34.  
  35. #include <stdexcept>
  36. #include <string>
  37.  
  38. struct cuda_error : std::runtime_error{
  39.     cuda_error(cudaError err): std::runtime_error(cudaGetErrorString(err)) {}
  40. protected:
  41.     cuda_error(std::string&& err): std::runtime_error(std::move(err)) {}
  42. };
  43.  
  44. struct assertcu_helper{};
  45. inline int operator>>(cudaError ret, assertcu_helper&&){
  46.     return ret == cudaSuccess ? cudaSuccess : throw cuda_error(ret);
  47. }
  48.  
  49. template<typename T> T get_device_object(const T& on_device, cudaStream_t stream){
  50.     T on_host;
  51.     cudaMemcpyFromSymbolAsync((void*)&on_host, (const void*)&on_device, sizeof(T), 0, cudaMemcpyDeviceToHost, stream) assertcu;
  52.     cudaStreamSynchronize(stream) assertcu;
  53.     return on_host;
  54. }
  55.  
  56.  
  57. #define block_size 256
  58.  
  59. __global__ void fillmem(int* ints){
  60.     size_t offset = size_t(blockIdx.x) * blockDim.x + threadIdx.x;
  61.     ints[offset] = 42;
  62. }
  63.  
  64. __device__ bool ok = true;
  65.  
  66. __global__ void validate(const int* ints){
  67.     size_t offset = size_t(blockIdx.x) * blockDim.x + threadIdx.x;
  68.     if(ints[offset] != 42) ok = false;
  69. }
  70.  
  71. int main(int argc, char *argv[]){
  72.     int offset = 2;
  73.         if (argc == 2) {
  74.             offset = atoi(argv[1]);
  75.         }
  76.     unsigned int total_ints = (((1024ULL * 1024ULL * 1024ULL * 2) / 4)+offset);
  77.     cudaDeviceProp deviceProp;
  78.     cudaGetDeviceProperties(&deviceProp, 0) assertcu;
  79.     if(deviceProp.major < 3) throw runtime_error("required sm >= 3.0");     //for large grid size in x dimension
  80.     cudaStream_t stream = 0;
  81.     destructor([&]{ cudaStreamDestroy(stream); });
  82.     cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) assertcu;
  83.     int* ints = 0;
  84.     destructor([&]{ cudaFree(ints); });
  85.     cudaMalloc(&ints, total_ints * 4) assertcu;
  86.     destructor(cudaDeviceSynchronize);
  87.     cudaMemsetAsync(ints, 0, total_ints * 4, stream) assertcu;
  88.     cudaDeviceSynchronize();
  89.     fillmem<<<total_ints / block_size, block_size, 0, stream>>>(ints);
  90.     cudaGetLastError() assertcu;
  91.     validate<<<total_ints / block_size, block_size, 0, stream>>>(ints);
  92.     cudaGetLastError() assertcu;
  93.     if(!get_device_object(ok, stream)) {
  94.         std::cout << offset << " failed\n";
  95.         //throw logic_error("validation failed");
  96.     } else {
  97.         std::cout << offset << " ok\n";
  98.     }
  99. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement