Advertisement
Guest User

Untitled

a guest
Apr 21st, 2014
134
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. $ cat t395.cu
  2. #include <stdio.h>
  3. #define INCS 10
  4. #define nTPB 256
  5. #define nBLKS 10
  6. #define USE_ATOMICS
  7. #define cudaCheckErrors(msg) \
  8. do { \
  9. cudaError_t __err = cudaGetLastError(); \
  10. if (__err != cudaSuccess) { \
  11. fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
  12. msg, cudaGetErrorString(__err), \
  13. __FILE__, __LINE__); \
  14. fprintf(stderr, "*** FAILED - ABORTING\n"); \
  15. exit(1); \
  16. } \
  17. } while (0)
  18.  
  19. __global__ void mykernel(volatile int *data){
  20.  
  21. for (int i = 0; i < INCS; i++){
  22. #ifdef USE_ATOMICS
  23. atomicAdd((int *)data,1);
  24. #else
  25. (*data)++;
  26. #endif
  27. }
  28. }
  29.  
  30. int main(){
  31. // simple test to demonstrate reading progress data from kernel
  32. volatile int *d_data, *h_data;
  33. cudaSetDeviceFlags(cudaDeviceMapHost);
  34. cudaCheckErrors("cudaSetDeviceFlags error");
  35. cudaHostAlloc((void **)&h_data, sizeof(int), cudaHostAllocMapped);
  36. cudaCheckErrors("cudaHostAlloc error");
  37. cudaHostGetDevicePointer((int **)&d_data, (int *)h_data, 0);
  38. cudaCheckErrors("cudaHostGetDevicePointer error");
  39. *h_data = 0;
  40. printf("kernel starting\n");
  41. mykernel<<<nBLKS,nTPB>>>(d_data);
  42. cudaDeviceSynchronize();
  43. cudaCheckErrors("kernel fail");
  44. if ((INCS*nBLKS*nTPB) != *h_data) {printf("fail! should be: %d, was: %d\n", (INCS*nBLKS*nTPB), *h_data); return 1;}
  45. printf("Success!\n");
  46. return 0;
  47. }
  48.  
  49. $ nvcc -arch=sm_20 -o t395 t395.cu
  50. $ ./t395
  51. kernel starting
  52. Success!
  53. $
Advertisement
RAW Paste Data Copied
Advertisement