Получение CUDA Thrust для использования потока CUDA по вашему выбору

Если посмотреть на запуск ядра в коде CUDA Thrust, кажется, что они всегда используют поток по умолчанию. Могу ли я заставить Thrust использовать поток по своему выбору? Я что-то упустил в API?

2 ответа

Я хочу обновить ответ, предоставленный talonmies после выпуска Thrust 1.8, который вводит возможность указания потока выполнения CUDA как

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

смотрите также

Thrust Release 1.8.0.

В следующем примере я повторяю пример в

Проблема ложной зависимости для архитектуры Fermi

в терминах API CUDA Thrust.

#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;
}
</thrust\execution_policy.h></thrust\device_vector.h></stdio.h></iostream>

Файлы Utilities.cu и Utilities.cuh, необходимые для запуска такого примера, сохраняются на этой странице github.

Временная шкала Visual Profiler показывает параллелизм операций CUDA Thrust и передачи памяти


Нет, вам ничего не хватает (по крайней мере, до моментального снимка, который поставляется с CUDA 6.0).

Первоначальная система отправки на основе трюков намеренно абстрагирует все базовые API API CUDA, жертвуя некоторой производительностью для удобства использования и согласованности (имейте в виду, что у тяги есть другие, кроме CUDA). Если вам нужен такой уровень гибкости, вам нужно попробовать другую библиотеку (например, CUB).

В версиях, начиная с моментального снимка CUDA 7.0, стало возможным установить поток выбора для операций тяги с помощью политики выполнения и функции отправки.

licensed under cc by-sa 3.0 with attribution.