zoukankan      html  css  js  c++  java
  • 数组逆序=全局内存版 VS 共享内存版

    全局内存版

     1 #include <stdio.h>
     2 #include <assert.h>
     3 #include "cuda.h"
     4 #include "cuda_runtime.h"
     5 #include "device_launch_parameters.h"
     6 //检查CUDA运行时是否有错误
     7 void checkCUDAError(const char* msg);
     8 // Part3: 在全局内存执行内核
     9 /*
    10 blockDim块内的线程数
    11 blockIdx网格内的块索引
    12 gridDim网格内块个数
    13 threadIdx块内线程索引
    14 */
    15 __global__ void reverseArrayBlock(int *d_out, int *d_in)
    16 {
    17     int inOffset = blockDim.x * blockIdx.x;
    18     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
    19     int in = inOffset + threadIdx.x;
    20     int out = outOffset + (blockDim.x - 1 - threadIdx.x);
    21     d_out[out] = d_in[in];
    22 }
    23 /////////////////////////////////////////////////////////////////////
    24 //主函数
    25 /////////////////////////////////////////////////////////////////////
    26 int main(int argc, char** argv)
    27 {
    28     //指向主机的内存空间和大小
    29     int *h_a;
    30     int dimA = 256 * 1024; // 256K elements (1MB total)
    31     //指向设备的指针和大小
    32     int *d_b, *d_a;
    33     //定义网格和块大小,每个块的线程数量
    34     int numThreadsPerBlock = 256;
    35 
    36     /*
    37     根据数组大小和预设的块大小来计算需要的块数
    38     */
    39     int numBlocks = dimA / numThreadsPerBlock;
    40     //申请主机及设备上的存储空间
    41     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
    42     //主机上的大小
    43     h_a = (int *)malloc(memSize);
    44     //设备上的大小
    45     cudaMalloc((void **)&d_a, memSize);
    46     cudaMalloc((void **)&d_b, memSize);
    47     //在主机上初始化输入数组
    48     for (int i = 0; i < dimA; ++i)
    49     {
    50         h_a[i] = i;
    51     }
    52     //将主机数组拷贝到设备上,h_a-->d_a
    53     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);
    54     //启动内核
    55     dim3 dimGrid(numBlocks);
    56     dim3 dimBlock(numThreadsPerBlock);
    57     reverseArrayBlock <<< dimGrid,    dimBlock >>>(d_b, d_a);
    58     //阻塞,一直到设备完成计算
    59     cudaThreadSynchronize();
    60     //检查是否设备产生了错误
    61     //检查任何CUDA错误
    62     checkCUDAError("kernel invocation");
    63     //将结果从设备拷贝到主机,d_b-->h_a
    64     cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);
    65     //检查任何CUDA错误
    66     checkCUDAError("memcpy");
    67     //核对返回到主机上的结果是否正确
    68     for (int i = 0; i < dimA; i++)
    69     {
    70         assert(h_a[i] == dimA - 1 - i);
    71     }
    72     //释放设备内存
    73     cudaFree(d_a);
    74     cudaFree(d_b);
    75     //释放主机内存
    76     free(h_a);
    77     printf("Correct!
    ");
    78     return 0;
    79 }
    80 void checkCUDAError(const char *msg)
    81 {
    82     cudaError_t err = cudaGetLastError();
    83     if (cudaSuccess != err)
    84     {
    85         fprintf(stderr, "Cuda error: %s: %s.
    ", msg,cudaGetErrorString(err));
    86         exit(EXIT_FAILURE);
    87     }
    88 }

    共享内存版

      1 #include <stdio.h>  
      2 #include <assert.h>
      3 #include "cuda.h"
      4 #include "cuda_runtime.h"
      5 #include "device_launch_parameters.h"
      6 #include <device_functions.h>
      7 //检查CUDA运行时是否有错误
      8 void checkCUDAError(const char* msg);
      9 // Part 2 of 2: 使用共享内存执行内核
     10 __global__ void reverseArrayBlock(int *d_out, int *d_in)
     11 {
     12     extern __shared__ int s_data[];
     13     int inOffset = blockDim.x * blockIdx.x;
     14     int in = inOffset + threadIdx.x;
     15     // Load one element per thread from device memory and store it   
     16     // *in reversed order* into temporary shared memory  
     17     /*
     18     每个线程从设备内存加载一个数据元素并按逆序存储在共享存储器上
     19     */
     20     s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
     21     /*
     22     阻塞,一直到所有线程将他们的数据都写入到共享内存中
     23     */
     24     __syncthreads();
     25     // write the data from shared memory in forward order,   
     26     // but to the reversed block offset as before  
     27     /*
     28     将共享内存中的数据s_data写入到d_out中,按照前序
     29     */
     30     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
     31     int out = outOffset + threadIdx.x;
     32     d_out[out] = s_data[threadIdx.x];
     33 }
     34 ////////////////////////////////////////////////////////////////////  
     35 //主函数
     36 ////////////////////////////////////////////////////////////////////  
     37 int main(int argc, char** argv)
     38 {
     39     //指向主机的内存空间和大小
     40     int *h_a;
     41     int dimA = 256 * 1024; // 256K elements (1MB total)  
     42     // pointer for device memory  
     43     int *d_b, *d_a;
     44     //指向设备的指针和大小
     45     int numThreadsPerBlock = 256;
     46     
     47     /*
     48     根据数组大小和预设的块大小来计算需要的块数
     49     */
     50     int numBlocks = dimA / numThreadsPerBlock;
     51     /*
     52     Part 1 of 2:
     53     计算共享内存所需的内存空间大小,这在下面的内核调用时被使用
     54     */
     55     int sharedMemSize = numThreadsPerBlock * sizeof(int);
     56     //申请主机及设备上的存储空间
     57     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
     58     //主机上的大小
     59     h_a = (int *)malloc(memSize);
     60     //设备上的大小
     61     cudaMalloc((void **)&d_a, memSize);
     62     cudaMalloc((void **)&d_b, memSize);
     63     //在主机上初始化输入数组
     64     for (int i = 0; i < dimA; ++i)
     65     {
     66         h_a[i] = i;
     67     }
     68     //将主机数组拷贝到设备上,h_a-->d_a
     69     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);
     70     //启动内核
     71     dim3 dimGrid(numBlocks);
     72     dim3 dimBlock(numThreadsPerBlock);
     73     reverseArrayBlock << < dimGrid, dimBlock, sharedMemSize >> >(d_b, d_a);
     74     //阻塞,一直到设备完成计算
     75     cudaThreadSynchronize();
     76     //检查是否设备产生了错误
     77     //检查任何CUDA错误
     78     checkCUDAError("kernel invocation");
     79     //将结果从设备拷贝到主机,d_b-->h_a
     80     cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);
     81     //检查任何CUDA错误
     82     checkCUDAError("memcpy");
     83     //核对返回到主机上的结果是否正确
     84     for (int i = 0; i < dimA; i++)
     85     {
     86         assert(h_a[i] == dimA - 1 - i);
     87     }
     88     //释放设备内存
     89     cudaFree(d_a);
     90     cudaFree(d_b);
     91     //释放主机内存
     92     free(h_a);
     93     printf("Correct!
    ");
     94     return 0;
     95 }
     96 
     97 void checkCUDAError(const char *msg)
     98 {
     99     cudaError_t err = cudaGetLastError();
    100     if (cudaSuccess != err)
    101     {
    102         fprintf(stderr, "Cuda error: %s: %s.
    ", msg, cudaGetErrorString(err));
    103         exit(EXIT_FAILURE);
    104     }
    105 }

     两个全部是数组逆序的实验,可以仔细观察其中更多而不同。

    项目下载链接

  • 相关阅读:
    hdu 4521 小明系列问题——小明序列(线段树 or DP)
    hdu 1115 Lifting the Stone
    hdu 5476 Explore Track of Point(2015上海网络赛)
    Codeforces 527C Glass Carving
    hdu 4414 Finding crosses
    LA 5135 Mining Your Own Business
    uva 11324 The Largest Clique
    hdu 4288 Coder
    PowerShell随笔3 ---别名
    PowerShell随笔2---初始命令
  • 原文地址:https://www.cnblogs.com/liangliangdetianxia/p/3981072.html
Copyright © 2011-2022 走看看