获得CUDA Thrust以使用您选择的CUDA流

刺柏

查看CUDA Thrust代码中的内核启动,看来它们始终使用默认流。我可以让Thrust使用我选择的流吗?我在API中缺少什么吗?

活力

我想在Thrust 1.8版本发布后更新talonmies提供的答案,它引入了将CUDA执行流指示为

thrust::cuda::par.on(stream)

也可以看看

推力释放1.8.0

在下面,我将重铸示例

Fermi架构的虚假依赖问题

就CUDA Thrust API而言。

#include <iostream>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include <thrust\device_vector.h>
#include <thrust\execution_policy.h>

#include "Utilities.cuh"

using namespace std;

#define NUM_THREADS 32
#define NUM_BLOCKS 16
#define NUM_STREAMS 3

struct BinaryOp{ __host__ __device__ int operator()(const int& o1,const int& o2) { return o1 * o2; } };

int main()
{
    const int N = 6000000;

    // --- Host side input data allocation and initialization. Registering host memory as page-locked (required for asynch cudaMemcpyAsync).
    int *h_in = new int[N]; for(int i = 0; i < N; i++) h_in[i] = 5;
    gpuErrchk(cudaHostRegister(h_in, N * sizeof(int), cudaHostRegisterPortable));

    // --- Host side input data allocation and initialization. Registering host memory as page-locked (required for asynch cudaMemcpyAsync).
    int *h_out = new int[N]; for(int i = 0; i < N; i++) h_out[i] = 0;
    gpuErrchk(cudaHostRegister(h_out, N * sizeof(int), cudaHostRegisterPortable));

    // --- Host side check results vector allocation and initialization
    int *h_checkResults = new int[N]; for(int i = 0; i < N; i++) h_checkResults[i] = h_in[i] * h_in[i];

    // --- Device side input data allocation.
    int *d_in = 0;              gpuErrchk(cudaMalloc((void **)&d_in, N * sizeof(int)));

    // --- Device side output data allocation. 
    int *d_out = 0;             gpuErrchk( cudaMalloc((void **)&d_out, N * sizeof(int)));

    int streamSize = N / NUM_STREAMS;
    size_t streamMemSize = N * sizeof(int) / NUM_STREAMS;

    // --- Set kernel launch configuration
    dim3 nThreads       = dim3(NUM_THREADS,1,1);
    dim3 nBlocks        = dim3(NUM_BLOCKS, 1,1);
    dim3 subKernelBlock = dim3((int)ceil((float)nBlocks.x / 2));

    // --- Create CUDA streams
    cudaStream_t streams[NUM_STREAMS];
    for(int i = 0; i < NUM_STREAMS; i++)
        gpuErrchk(cudaStreamCreate(&streams[i]));

    /**************************/
    /* BREADTH-FIRST APPROACH */
    /**************************/

    for(int i = 0; i < NUM_STREAMS; i++) {
        int offset = i * streamSize;
        cudaMemcpyAsync(&d_in[offset], &h_in[offset], streamMemSize, cudaMemcpyHostToDevice,     streams[i]);
    }

    for(int i = 0; i < NUM_STREAMS; i++)
    {
        int offset = i * streamSize;

        thrust::transform(thrust::cuda::par.on(streams[i]), thrust::device_pointer_cast(&d_in[offset]), thrust::device_pointer_cast(&d_in[offset]) + streamSize/2, 
                                                            thrust::device_pointer_cast(&d_in[offset]), thrust::device_pointer_cast(&d_out[offset]), BinaryOp());
        thrust::transform(thrust::cuda::par.on(streams[i]), thrust::device_pointer_cast(&d_in[offset + streamSize/2]), thrust::device_pointer_cast(&d_in[offset + streamSize/2]) + streamSize/2, 
                                                            thrust::device_pointer_cast(&d_in[offset + streamSize/2]), thrust::device_pointer_cast(&d_out[offset + streamSize/2]), BinaryOp());

    }

    for(int i = 0; i < NUM_STREAMS; i++) {
        int offset = i * streamSize;
        cudaMemcpyAsync(&h_out[offset], &d_out[offset], streamMemSize, cudaMemcpyDeviceToHost,   streams[i]);
    }

    for(int i = 0; i < NUM_STREAMS; i++)
        gpuErrchk(cudaStreamSynchronize(streams[i]));

    gpuErrchk(cudaDeviceSynchronize());

    // --- Release resources
    gpuErrchk(cudaHostUnregister(h_in));
    gpuErrchk(cudaHostUnregister(h_out));
    gpuErrchk(cudaFree(d_in));
    gpuErrchk(cudaFree(d_out));

    for(int i = 0; i < NUM_STREAMS; i++)
        gpuErrchk(cudaStreamDestroy(streams[i]));

    cudaDeviceReset();  

    // --- GPU output check
    int sum = 0;
    for(int i = 0; i < N; i++) {     
        //printf("%i %i\n", h_out[i], h_checkResults[i]);
        sum += h_checkResults[i] - h_out[i];
    }

    cout << "Error between CPU and GPU: " << sum << endl;

    delete[] h_in;
    delete[] h_out;
    delete[] h_checkResults;

    return 0;
}

运行此示例所需Utilities.cuUtilities.cuh文件在此github页面上维护

Visual Profiler时间线显示CUDA Thrust操作和内存传输的并发

在此处输入图片说明

本文收集自互联网,转载请注明来源。

如有侵权,请联系[email protected] 删除。

编辑于
0

我来说两句

0条评论
登录后参与评论

相关文章

来自分类Dev

TypeError:预期的CPU(获得CUDA)

来自分类Dev

如何获得CUDA的printf以打印到任意流?

来自分类Dev

如何获得CUDA内核的汇编代码?

来自分类Dev

使用Thrust,CUDA进行慢速排序

来自分类Dev

在主机上顺序使用CUDA Thrust算法

来自分类Dev

使用CUDA / Thrust使关键出现的次数相等

来自分类Dev

在cuda Thrust中排序

来自分类Dev

在cuda Thrust中排序

来自分类Dev

使用CUDA / Thrust对多个数组进行排序

来自分类Dev

如何使用CUDA选择GPU?

来自分类Dev

如何使用CUDA选择GPU?

来自分类Dev

如何使用CUDA Thrust执行策略来覆盖Thrust的低级设备内存分配器

来自分类Dev

CUDA流压缩算法

来自分类Dev

同步多个Cuda流

来自分类Dev

CUDA / Thrust:如何对交错数组的列求和?

来自分类Dev

仅对向量CUDA / THRUST的正元素求和

来自分类Dev

CUDA / Thrust:如何对交错数组的列求和?

来自分类Dev

在CUDA中,我无法获得有效值

来自分类Dev

Cuda中的选择排序

来自分类Dev

Tensorflow如何支持Cuda流?

来自分类Dev

使用CUDA Thrust确定2个最大元素及其在每个矩阵行中的位置

来自分类Dev

在现有的C项目中使用CUDA Thrust:编译错误

来自分类Dev

使用CUDA / thrust在数组中设置每个float4的一个元素

来自分类Dev

CUDA Thrust如何使用矢量将对象传输到设备

来自分类Dev

使用cuda时可以使用桌面或gui界面吗?

来自分类Dev

可以使用__syncthreads()合并单独的CUDA内核吗?

来自分类Dev

使用CUDA调试时的CUDA堆栈大小

来自分类Dev

运行时错误:后端 CUDA 的预期对象,但为参数 #3 'index' 获得了后端 CPU

来自分类Dev

使用多个CUDA GPU