Засунуть внутрь написанных пользователем ядер
Я новичок в Thrust. Я вижу, что все презентации и примеры Thrust показывают только код хоста.
Я хотел бы знать, могу ли я передать device_vector в свое собственное ядро? Как? Если да, какие операции разрешены для него внутри кода ядра / устройства?
4 ответа
Как было изначально написано, Thrust - это просто абстракция на стороне хоста. Его нельзя использовать внутри ядер. Вы можете передать память устройства, инкапсулированную внутри thrust::device_vector к вашему собственному ядру, как это:
thrust::device_vector< Foo > fooVector;
// Do something thrust-y with fooVector
Foo* fooArray = thrust::raw_pointer_cast( &fooVector[0] );
// Pass raw array and its size to kernel
someKernelCall<<< x, y >>>( fooArray, fooVector.size() );
и вы также можете использовать память устройства, не распределенную по тяге в алгоритмах тяги, путем создания экземпляра thrust::device_ptr с указателем памяти устройства cuda.
Отредактировано четыре с половиной года спустя, чтобы добавить, что согласно ответу @JackOLantern, Thrust 1.8 добавляет политику последовательного выполнения, которая означает, что вы можете запускать однопоточные версии аллогритов Thrust на устройстве. Обратите внимание, что до сих пор не представляется возможным непосредственно передать вектор толкающее устройство к векторам ядра и устройства не могут быть использованы непосредственно в коде устройства.
Обратите внимание, что также можно использовать thrust::device В некоторых случаях политика выполнения должна иметь параллельное выполнение тяги, запускаемое ядром как дочерняя сетка. Это требует отдельной компиляции / связывания устройства и аппаратного обеспечения, которое поддерживает динамический параллелизм. Я не уверен, поддерживается ли это на самом деле во всех алгоритмах тяги или нет, но, безусловно, работает с некоторыми.
Это обновление к моему предыдущему ответу.
Начиная с Thrust 1.8.1, примитивы CUDA Thrust можно комбинировать с thrust::device политика выполнения для параллельного запуска в одном потоке CUDA, использующем динамический параллелизм CUDA. Ниже приведен пример.
#include <stdio.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
#include "TimingGPU.cuh"
#include "Utilities.cuh"
#define BLOCKSIZE_1D 256
#define BLOCKSIZE_2D_X 32
#define BLOCKSIZE_2D_Y 32
/*************************/
/* TEST KERNEL FUNCTIONS */
/*************************/
__global__ void test1(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) {
const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;
if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::seq, d_data + tid * Ncols, d_data + (tid + 1) * Ncols);
}
__global__ void test2(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) {
const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;
if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::device, d_data + tid * Ncols, d_data + (tid + 1) * Ncols);
}
/********/
/* MAIN */
/********/
int main() {
const int Nrows = 64;
const int Ncols = 2048;
gpuErrchk(cudaFree(0));
// size_t DevQueue;
// gpuErrchk(cudaDeviceGetLimit(&DevQueue, cudaLimitDevRuntimePendingLaunchCount));
// DevQueue *= 128;
// gpuErrchk(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, DevQueue));
float *h_data = (float *)malloc(Nrows * Ncols * sizeof(float));
float *h_results = (float *)malloc(Nrows * sizeof(float));
float *h_results1 = (float *)malloc(Nrows * sizeof(float));
float *h_results2 = (float *)malloc(Nrows * sizeof(float));
float sum = 0.f;
for (int i=0; i<Nrows; i++) {
h_results[i] = 0.f;
for (int j=0; j<Ncols; j++) {
h_data[i*Ncols+j] = i;
h_results[i] = h_results[i] + h_data[i*Ncols+j];
}
}
TimingGPU timerGPU;
float *d_data; gpuErrchk(cudaMalloc((void**)&d_data, Nrows * Ncols * sizeof(float)));
float *d_results1; gpuErrchk(cudaMalloc((void**)&d_results1, Nrows * sizeof(float)));
float *d_results2; gpuErrchk(cudaMalloc((void**)&d_results2, Nrows * sizeof(float)));
gpuErrchk(cudaMemcpy(d_data, h_data, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice));
timerGPU.StartCounter();
test1<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("Timing approach nr. 1 = %f\n", timerGPU.GetCounter());
gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<Nrows; i++) {
if (h_results1[i] != h_results[i]) {
printf("Approach nr. 1; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]);
return 0;
}
}
timerGPU.StartCounter();
test2<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("Timing approach nr. 2 = %f\n", timerGPU.GetCounter());
gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<Nrows; i++) {
if (h_results1[i] != h_results[i]) {
printf("Approach nr. 2; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]);
return 0;
}
}
printf("Test passed!\n");
}
Вышеприведенный пример выполняет сокращение строк матрицы в том же смысле, что и Сокращение строк матрицы с помощью CUDA, но это делается не так, как в предыдущем посте, а именно путем вызова примитивов CUDA Thrust непосредственно из написанных пользователем ядер. Кроме того, приведенный выше пример служит для сравнения производительности одних и тех же операций при выполнении с двумя политиками выполнения, а именно: thrust::seq а также thrust::device, Ниже приведены некоторые графики, показывающие разницу в производительности.
Производительность была оценена на Kepler K20c и Maxwell GeForce GTX 850M.
Я хотел бы предоставить обновленный ответ на этот вопрос.
Начиная с Thrust 1.8, примитивы CUDA Thrust можно комбинировать с thrust::seq политика выполнения для запуска последовательно в одном потоке CUDA (или последовательно в одном потоке ЦП). Ниже приведен пример.
Если вы хотите параллельное выполнение внутри потока, то вы можете рассмотреть возможность использования CUB, который предоставляет процедуры сокращения, которые могут быть вызваны из блока потока, при условии, что ваша карта обеспечивает динамический параллелизм.
Вот пример с Thrust
#include <stdio.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
__global__ void test(float *d_A, int N) {
float sum = thrust::reduce(thrust::seq, d_A, d_A + N);
printf("Device side result = %f\n", sum);
}
int main() {
const int N = 16;
float *h_A = (float*)malloc(N * sizeof(float));
float sum = 0.f;
for (int i=0; i<N; i++) {
h_A[i] = i;
sum = sum + h_A[i];
}
printf("Host side result = %f\n", sum);
float *d_A; gpuErrchk(cudaMalloc((void**)&d_A, N * sizeof(float)));
gpuErrchk(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));
test<<<1,1>>>(d_A, N);
}
Если вы хотите использовать данные, выделенные / обработанные с помощью thrust yes, вы можете просто получить необработанный указатель на выделенные данные.
int * raw_ptr = thrust::raw_pointer_cast(dev_ptr);
Если вы хотите распределить векторы тяги в ядре, я никогда не пытался, но я не думаю, что это будет работать, а также, если это работает, я не думаю, что это даст какую-то выгоду.