Сборка пакета R с кодом перемещаемого устройства

Question

Я пишу пакет R, который использует Thrust как для обработки выделения памяти, так и для того, чтобы не писать собственные ядра CUDA.

В некоторых случаях я вызываю подпрограммы cuBLAS из кода устройства, а не из кода хоста. Это изменяет требования к компиляции. Хотя код компилируется с использованием команд nvcc ниже, может быть желательно явно вызвать главный компоновщик (g++). Как я могу изменить текущий процесс сборки, чтобы выполнить это?

шагов, которые я использую являются:

1. Компиляция выходного файла (max.o), содержащим устройство перемещаемого кода, используя -dc переключатель

2. Создать библиотеку (libmax.a), чтобы связать с

3. Скомпилировать выходной файл, содержащий функции обертки (somePackage.o) с использованием -c переключатель

4. Создать разделяемую библиотеку (somePackage.so), которая ссылается на libmax.a используя -shared переключатель

Рабочий пример приведен ниже:

iterator.h: Это определяет некоторые типы, в том числе strideAccessor.

max.h: Декларация функции в max.cu

max.cu: Определяет функцию, которая находит индекс максимального элемента в каждом из массивов n каскадных размерности d.

somePackage.cu: Оболочка обработки/C++ интерфейс R

$ cat iterator.h 
#ifndef ITER_H 
#define ITER_H 

#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/permutation_iterator.h> 
#include <thrust/tuple.h> 
#include <thrust/iterator/zip_iterator.h> 

typedef thrust::device_vector<int> ivec_d; 
typedef thrust::device_vector<double> fvec_d; 
typedef thrust::device_vector<int>::iterator intIter; 
typedef thrust::device_vector<double>::iterator realIter; 
typedef thrust::host_vector<int> ivec_h; 
typedef thrust::host_vector<double> fvec_h; 

typedef thrust::counting_iterator<int> countIter; 

//Used for generating rep((1:len)*incr, times=infinity) 
struct stride: public thrust::unary_function<int, int>{ 

    int incr; 

    __host__ __device__ stride(int incr=1): incr(incr){} 

    __host__ __device__ int operator()(int x){ 

    return x*incr; 
    } 
}; 

typedef thrust::transform_iterator<stride, countIter> strideIter; 
typedef thrust::permutation_iterator<realIter, strideIter> strideAccessor; 


#endif 

$ cat max.h 
#include "iterator.h" 

void cublas_max(fvec_d &x, ivec_d &result, int n, int d); 

$ cat max.cu 
#include "iterator.h" 
#include <thrust/functional.h> 
#include <thrust/transform.h> 
#include <cublas_v2.h> 
#include <iostream> 

struct whichMax : thrust::unary_function<double, int>{ 
    int dim; 

    __host__ __device__ whichMax(int dim): dim(dim){} 

    __host__ __device__ int operator()(double &vec){ 

    cublasHandle_t handle; 
    cublasCreate_v2(&handle); 
    int incx=1, n = dim, result =0; 
    double *vec_ptr = thrust::raw_pointer_cast(&vec); 

    //find the first index of a maximal element 
    cublasIdamax(handle, n, vec_ptr, incx, &result); 
    cublasDestroy_v2(handle); 
    return result; 
    } 
}; 

void cublas_max(fvec_d &x, ivec_d &result, int n, int d){ 

    stride f(d); 
    strideIter siter = thrust::transform_iterator<stride, countIter>(thrust::make_counting_iterator<int>(0), f); 
    strideAccessor stridex = thrust::permutation_iterator<realIter, strideIter>(x.begin(), siter); 

    whichMax g(d); 

    //find the index of maximum for each of n subvectors 
    thrust::copy(result.begin(), result.end(), std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
    thrust::transform(stridex, stridex + n, result.begin(), g); 
    thrust::copy(result.begin(), result.end(), std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
} 

$ cat somePackage.cu 
#include "iterator.h" 
#include "max.h" 
#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <R.h> 
#include <Rinternals.h> 
#include <Rmath.h> 
#include <iostream> 

extern "C" SEXP Rcublas_max(SEXP x, SEXP n, SEXP dim){ 

    double *xptr = REAL(x); 
    int N = INTEGER(n)[0], D = INTEGER(n)[0]; 

    fvec_d dx(xptr, xptr+N*D); 
    ivec_d dresult(N); 

    cublas_max(dx, dresult, N, D); 

    ivec_h hresult(N); 
    thrust::copy(dresult.begin(), dresult.end(), hresult.begin()); 

    SEXP indices = PROTECT(allocVector(INTSXP, N)); 

    for(int i=0; i<N; ++i) 
    INTEGER(indices)[i] = hresult[i]; 

    UNPROTECT(1); 
    return indices; 
} 

$ make 
nvcc -dc -arch=sm_35 -Xcompiler -fPIC -lcublas_device -lcublas_device max.cu -o max.o 
nvcc -lib -arch=sm_35 -Xcompiler -fPIC -lcublas_device -lcublas_device max.o -o libmax.a 
nvcc -c -arch=sm_35 -Xcompiler -fPIC -lcublas_device somePackage.cu -lmax -I/home/emittman/src/R-3.3.1/builddir/include -I. -o somePackage.o 
nvcc -shared -arch=sm_35 -Xcompiler -fPIC -lcublas_device somePackage.o -I/home/emittman/src/R-3.3.1/builddir/include -I. -L. -lcublas_device -lmax -o somePackage.so 
ptxas info : 'device-function-maxrregcount' is a BETA feature 

Answer 1

Я создал R пакет с Rcpp, вызывая некоторые Экстерн функции из C++ общей библиотеке, которая затем вызвать ядра CUDA для выполнения вычислений обязательный.

Что вы пытаетесь сделать здесь, это скомпилировать ваш код CUDA в статическую библиотеку, а затем связать его с вашим R-пакетом (который сам будет скомпилирован в общую библиотеку). Мой подход отличается от вашего, и я предоставляю описания для моего метода, чтобы дать вам другую мысль.

Вот упрощенный пример.

kernels.cu разделяемой библиотеки, содержащей код CUDA:

__global__ 
void my_cuda_kernel(...) { 
    // ...... 
} 

main.cu разделяемой библиотеки, содержащей код CUDA:

extern "C" { 
    void do_cuda_work(...) { 
     thrust::copy(...); 
     my_cuda_kernel <<<...>>> (...); 
    } 
} 

пакет.каст в R упаковке:

extern void do_cuda_work(...); 

// [[Rcpp::export]] 
void call_cuda_code(...) { 
    do_cuda_work(...); 
} 

компилировать код CUDA в разделяемую библиотеку, вам нужно использовать:

nvcc -arch=sm_35 -dc ... kernels.cu -o kernels.o 
nvcc -arch=sm_35 -dc ... main.cu -o main.o 
nvcc --shared -arch=sm_35 ... kernels.o main.o ... libMyCUDALibrary.so 

Пожалуйста, обратите внимание, что для раздельной компиляции для работы необходимо указать -arch=sm_35 для обоих компилятор и компоновщик и -dc для компилятора. Как только вы успешно создали общую библиотеку, привязка вашего R-пакета к ней довольно прямолинейна. (Тем не менее, вам, возможно, потребуется создать Makevars файл под src папке R пакет, чтобы указать, включают и библиотеку путей, и, возможно, RPATH а):

CXX_STD= CXX11 
PKG_CPPFLAGS= -I../../../CPP/include 
PKG_LIBS= -L../../../CPP/bin/Release -lMyCUDALibrary -Wl,-rpath=$$HOME/MyCUDALibrary/CPP/bin/Release `$(R_HOME)/bin/Rscript -e "Rcpp::LdFlags()"`