axpy_partition_gpu.cu 1.4 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354
  1. #include <starpu.h>
  2. #include "axpy_partition_gpu.h"
  3. #include <stdio.h>
  4. //This code demonstrates how to transform a kernel to execute on a given set of GPU SMs.
  5. // Original kernel
  6. __global__ void saxpy(int n, float a, float *x, float *y)
  7. {
  8. int i = blockIdx.x*blockDim.x + threadIdx.x;
  9. if (i<n) y[i] = a*x[i] + y[i];
  10. }
  11. // Transformed kernel
  12. __global__ void saxpy_partitioned(__P_KARGS, int n, float a, float *x, float *y)
  13. {
  14. __P_BEGIN;
  15. __P_LOOPX;
  16. int i = blockid.x*blockDim.x + threadIdx.x; // note that blockIdx is replaced.
  17. if (i<n) y[i] = a*x[i] + y[i];
  18. __P_LOOPEND;
  19. }
  20. extern "C" void cuda_axpy(void *descr[], void *_args)
  21. {
  22. float a = *((float *)_args);
  23. unsigned n = STARPU_VECTOR_GET_NX(descr[0]);
  24. float *x = (float *)STARPU_VECTOR_GET_PTR(descr[0]);
  25. float *y = (float *)STARPU_VECTOR_GET_PTR(descr[1]);
  26. int SM_mapping_start = -1;
  27. int SM_mapping_end = -1;
  28. int SM_allocation = -1;
  29. cudaStream_t stream = starpu_cuda_get_local_stream();
  30. int workerid = starpu_worker_get_id();
  31. starpu_sched_ctx_get_sms_interval(workerid, &SM_mapping_start, &SM_mapping_end);
  32. SM_allocation = SM_mapping_end - SM_mapping_start;
  33. int dimensions = 512;
  34. //partitioning setup
  35. // int SM_mapping_start = 0;
  36. // int SM_allocation = 13;
  37. __P_HOSTSETUP(saxpy_partitioned,dim3(dimensions,1,1),dimensions,0,SM_mapping_start,SM_allocation,stream);
  38. saxpy_partitioned<<<width,dimensions,0,stream>>>(__P_HKARGS,n,a,x,y);
  39. }