找回密码
 立即注册
nanhuier 该用户已被删除
发表于 2019-11-20 09:07:44
15891
  1. 显卡是Quadro P1000 最新驱动,CUDA 8, 代码如下
复制代码


使用道具 举报 回复
发表于 2019-11-20 09:11:36

  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"

  3. #include <stdio.h>
  4. //#include <cmath>
  5. #include<stdint.h>
  6. //typedef BYTE  uint16_t;
  7. //typedef int  uint16_t;
  8. #include "CudaKernelInfo.h"
  9. #include <iostream>
  10. //#include <iostream>
  11. using namespace std;

  12. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

  13. __global__ void addKernel(int *c, const int *a, const int *b)
  14. {
  15.     int i = threadIdx.x;
  16.     c[i] = a[i] + b[i];
  17. }


  18. /*__device__ inline float lerp(float v0, float v1, float t)
  19. {
  20.                
  21.         return fmaf(t, v1, fmaf(-t, v0, v0));
  22. }*/

  23. __device__ float lerp(float v0, float v1, float t)
  24. {

  25.         return fmaf(t, v1, fmaf(-t, v0, v0));
  26. }


  27. __global__ void VolumeProcessing_resizeAndMaskKernel(
  28.         uint16_t * out_ptr, const int out_stride,
  29.         const float * in_ptr, const int in_stride,
  30.         const int out_size_x, const int out_size_y,
  31.         const int in_size_x, const int in_size_y, const int in_size_z,
  32.         int slice_index, float resample_step, float radius_sqr,
  33.         const int max_voxel_value)
  34. {
  35.         int ox = blockIdx.x*blockDim.x + threadIdx.x;
  36.         int oy = blockIdx.y*blockDim.y + threadIdx.y;

  37.         //std::cout << ox << " " << oy;
  38.         if (ox >= out_size_x || oy >= out_size_y)
  39.         {
  40.                 //cout << "exceed limit";
  41.                 return;
  42.         }
  43.        
  44.        
  45.         float dx = fmaf(0.5f, out_size_x, float(-ox) - 0.5f);
  46.         float dy = fmaf(0.5f, out_size_y, float(-oy) - 0.5f);
  47.         float d = fmaf(dx, dx, dy*dy);
  48.        

  49.         float value = 0.0f;

  50.         if (d <= radius_sqr)
  51.         {
  52.                 // clamp to edge
  53.                 float ix = resample_step * ox;
  54.                 float iy = resample_step * oy;
  55.                 float iz = resample_step * slice_index;

  56.                 // 0: first, 1: next voxel
  57.                 int x0 = min(int(ix) + 0, in_size_x - 1);
  58.                 int x1 = min(int(ix) + 1, in_size_x - 1);
  59.                 int y0 = min(int(iy) + 0, in_size_y - 1);
  60.                 int y1 = min(int(iy) + 1, in_size_y - 1);
  61.                 int z0 = min(int(iz) + 0, in_size_z - 1);
  62.                 int z1 = min(int(iz) + 1, in_size_z - 1);

  63.                 // weight of next voxel
  64.                 float t = min(ix - x0, 1.0f);
  65.                 float u = min(iy - y0, 1.0f);
  66.                 float v = min(iz - z0, 1.0f);



  67.                 float xy0 = lerp(
  68.                         lerp(in_ptr[x0 + (y0 + z0 * in_size_y) * in_stride]
  69.                                 , in_ptr[x1 + (y0 + z0 * in_size_y) * in_stride], t),
  70.                         lerp(in_ptr[x0 + (y1 + z0 * in_size_y) * in_stride]
  71.                                 , in_ptr[x1 + (y1 + z0 * in_size_y) * in_stride], t), u);

  72.                 float xy1 = lerp(
  73.                         lerp(in_ptr[x0 + (y0 + z1 * in_size_y) * in_stride]
  74.                                 , in_ptr[x1 + (y0 + z1 * in_size_y) * in_stride], t),
  75.                         lerp(in_ptr[x0 + (y1 + z1 * in_size_y) * in_stride]
  76.                                 , in_ptr[x1 + (y1 + z1 * in_size_y) * in_stride], t), u);

  77.                 value = lerp(xy0, xy1, v);
  78.         }

  79.         out_ptr[ox + oy * out_stride] = uint16_t(min(max(0.5f, value + 0.5f), 0.5f + max_voxel_value));
  80. }



  81. int main()
  82. {
  83.    /* const int arraySize = 5;
  84.     const int a[arraySize] = { 1, 2, 3, 4, 5 };
  85.     const int b[arraySize] = { 10, 20, 30, 40, 50 };
  86.     int c[arraySize] = { 0 };

  87.     // Add vectors in parallel.
  88.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
  89.     if (cudaStatus != cudaSuccess) {
  90.         fprintf(stderr, "addWithCuda failed!");
  91.         return 1;
  92.     }

  93.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
  94.         c[0], c[1], c[2], c[3], c[4]); */

  95.     // cudaDeviceReset must be called before exiting in order for profiling and
  96.     // tracing tools such as Nsight and Visual Profiler to show complete traces.

  97.         // test resample
  98.         // add by yyy
  99.        
  100.         cudaError_t cudaStatus;

  101.         cudaStatus = cudaSetDevice(0);
  102.         if (cudaStatus != cudaSuccess) {
  103.                 fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
  104.                
  105.         }

  106.         uint16_t *out_ptr;
  107.         const int out_stride = 448;
  108.         //const float *in_ptr;
  109.         float *in_ptr;
  110.         //void *v_in_ptr;
  111.         const int in_stride = 480;
  112.         const int out_size_x = 420;
  113.         const int out_size_y = 420;
  114.         const int in_size_x = 420;
  115.         const int in_size_y = 420;
  116.         const int in_size_z = 250;
  117.         int slice_index = 0;
  118.         float resample_step = 1;
  119.         float radius_sqr = 209.5;
  120.         const int max_voxel_value = 8191;

  121.         //CUDACHECK
  122.         //cudaMemset
  123.         out_ptr = new uint16_t[420 * 420];
  124.         //float *t1 = new float[420 * 420 * 250];

  125.       [color=Red] // 执行到这里报错 an illegal memory access was encountered[/color]
  126.         [color=Red]cudaStatus=cudaMalloc((void**)&in_ptr, 420 * 420 * 250 * sizeof(float));[/color]
  127.        
  128.         if (cudaStatus != cudaSuccess) {
  129.                 fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
  130.         }

  131.        
  132.         //cudaMemcpy(v_in_ptr, t1, 420 * 420 * 250 * sizeof(float), cudaMemcpyHostToDevice);
  133.         //cudaMemcpy()
  134.         //CudaKernelInfo launch(420, 420);
  135.         //CudaKernelInfo launch(slice.sizeX(), slice.sizeY());
  136.        
  137.         // block size
  138.         CudaKernelInfo launch(420, 420);
  139.        
  140.         //cout << launch.gridSize() << " " << launch.threadBlockSize();
  141.        
  142.     VolumeProcessing_resizeAndMaskKernel<<<launch.gridSize(),launch.threadBlockSize()>>>(out_ptr, out_stride, in_ptr, in_stride,
  143.                 out_size_x, out_size_y, in_size_x, in_size_y, in_size_z, slice_index, resample_step, radius_sqr, max_voxel_value);


  144.         cudaStatus = cudaDeviceSynchronize();
  145.         if (cudaStatus != cudaSuccess) {
  146.                 fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
  147.         }

  148.     /*cudaStatus = cudaDeviceReset();
  149.     if (cudaStatus != cudaSuccess) {
  150.         fprintf(stderr, "cudaDeviceReset failed!");
  151.         return 1;
  152.     }*/
  153.         //cudaFreeArray(in_ptr);
  154.         delete[]out_ptr;
  155.         cudaFree(in_ptr);
  156.        
  157.     return 0;
  158. }

  159. // Helper function for using CUDA to add vectors in parallel.
  160. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
  161. {
  162.     int *dev_a = 0;
  163.     int *dev_b = 0;
  164.     int *dev_c = 0;
  165.     cudaError_t cudaStatus;

  166.     // Choose which GPU to run on, change this on a multi-GPU system.
  167.     cudaStatus = cudaSetDevice(0);
  168.     if (cudaStatus != cudaSuccess) {
  169.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
  170.         goto Error;
  171.     }

  172.     // Allocate GPU buffers for three vectors (two input, one output)    .
  173.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
  174.     if (cudaStatus != cudaSuccess) {
  175.         fprintf(stderr, "cudaMalloc failed!");
  176.         goto Error;
  177.     }

  178.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
  179.     if (cudaStatus != cudaSuccess) {
  180.         fprintf(stderr, "cudaMalloc failed!");
  181.         goto Error;
  182.     }

  183.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
  184.     if (cudaStatus != cudaSuccess) {
  185.         fprintf(stderr, "cudaMalloc failed!");
  186.         goto Error;
  187.     }

  188.     // Copy input vectors from host memory to GPU buffers.
  189.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
  190.     if (cudaStatus != cudaSuccess) {
  191.         fprintf(stderr, "cudaMemcpy failed!");
  192.         goto Error;
  193.     }

  194.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
  195.     if (cudaStatus != cudaSuccess) {
  196.         fprintf(stderr, "cudaMemcpy failed!");
  197.         goto Error;
  198.     }

  199.     // Launch a kernel on the GPU with one thread for each element.
  200.     addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

  201.     // Check for any errors launching the kernel
  202.     cudaStatus = cudaGetLastError();
  203.     if (cudaStatus != cudaSuccess) {
  204.         fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
  205.         goto Error;
  206.     }
  207.    
  208.     // cudaDeviceSynchronize waits for the kernel to finish, and returns
  209.     // any errors encountered during the launch.
  210.     cudaStatus = cudaDeviceSynchronize();
  211.     if (cudaStatus != cudaSuccess) {
  212.         fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
  213.         goto Error;
  214.     }

  215.     // Copy output vector from GPU buffer to host memory.
  216.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
  217.     if (cudaStatus != cudaSuccess) {
  218.         fprintf(stderr, "cudaMemcpy failed!");
  219.         goto Error;
  220.     }

  221. Error:
  222.     cudaFree(dev_c);
  223.     cudaFree(dev_a);
  224.     cudaFree(dev_b);
  225.    
  226.     return cudaStatus;
  227. }
复制代码
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册

zzczczxczxczx