Использование типов CUDA в pyCUDA - PullRequest
Новая
       16

Использование типов CUDA в pyCUDA

2 голосов
/ 02 апреля 2019

Давайте рассмотрим код CUDA на Mersenne Twister CUDA для произвольного числа потоков и предположим, что я хочу преобразовать его в приложение pyCUDA.

Я знаю, что яможно использовать ctypes и CDLL, а именно 

cudart = CDLL("/usr/local/cuda/lib64/libcudart.so")

для использования подпрограмм cudart.

Однако мне также необходимо выделить, например, curandStateMtgp32 массив, определение которого находится в curand_mtgp32.h, или вызовите 

curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, devKernelParams);

и используйте mtgp32dc_params_fast_11213, определение которого находится в curand_mtgp32_host.h.

Как работать с CUDA определениями типаи значения в pyCUDA?

1 Ответ

1 голос
/ 09 апреля 2019

Я решил проблему со ссылкой на API на стороне устройства следующим образом:

  1. Я создал .dll, содержащий две функции: MTGP32Setup() для настройки генератора Мерсенна Твистера и MTGP32Generation() для генерациислучайные числа;
  2. Я вызывал вышеупомянутые функции, используя ctypes.

Исходный код для .dll

// --- Generate random numbers with cuRAND's Mersenne Twister

#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <time.h>

#include <cuda.h>
#include <curand_kernel.h>
/* include MTGP host helper functions */
#include <curand_mtgp32_host.h>

#define BLOCKSIZE   256
#define GRIDSIZE    64

curandStateMtgp32 *devMTGPStates;

/********************/
/* CUDA ERROR CHECK */
/********************/
// --- Credit to http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api
void gpuAssert(cudaError_t code, const 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); }
    }
}

void gpuErrchk(cudaError_t ans) { gpuAssert((ans), __FILE__, __LINE__); }

/*************************/
/* CURAND ERROR CHECKING */
/*************************/
static const char *_curandGetErrorEnum(curandStatus_t error)
{
    switch (error)
    {
    case CURAND_STATUS_SUCCESS:
        return "CURAND_SUCCESS";

    case CURAND_STATUS_VERSION_MISMATCH:
        return "CURAND_STATUS_VERSION_MISMATCH";

    case CURAND_STATUS_NOT_INITIALIZED:
        return "CURAND_STATUS_NOT_INITIALIZED";

    case CURAND_STATUS_ALLOCATION_FAILED:
        return "CURAND_STATUS_ALLOCATION_FAILED";

    case CURAND_STATUS_TYPE_ERROR:
        return "CURAND_STATUS_TYPE_ERROR";

    case CURAND_STATUS_OUT_OF_RANGE:
        return "CURAND_STATUS_OUT_OF_RANGE";

    case CURAND_STATUS_LENGTH_NOT_MULTIPLE:
        return "CURAND_STATUS_LENGTH_NOT_MULTIPLE";

    case CURAND_STATUS_DOUBLE_PRECISION_REQUIRED:
        return "CURAND_STATUS_DOUBLE_PRECISION_REQUIRED";

    case CURAND_STATUS_LAUNCH_FAILURE:
        return "CURAND_STATUS_LAUNCH_FAILURE";

    case CURAND_STATUS_PREEXISTING_FAILURE:
        return "CURAND_STATUS_PREEXISTING_FAILURE";

    case CURAND_STATUS_INITIALIZATION_FAILED:
        return "CURAND_STATUS_INITIALIZATION_FAILED";

    case CURAND_STATUS_ARCH_MISMATCH:
        return "CURAND_STATUS_ARCH_MISMATCH";

    case CURAND_STATUS_INTERNAL_ERROR:
        return "CURAND_STATUS_INTERNAL_ERROR";

    }

    return "<unknown>";
}

inline void __curandSafeCall(curandStatus_t err, const char *file, const int line)
{
    if (CURAND_STATUS_SUCCESS != err) {
        fprintf(stderr, "CURAND error in file '%s', line %d, error: %s \nterminating!\n", __FILE__, __LINE__, \
            _curandGetErrorEnum(err)); \
            assert(0); \
    }
}

void curandSafeCall(curandStatus_t err) { __curandSafeCall(err, __FILE__, __LINE__); }

/*******************/
/* iDivUp FUNCTION */
/*******************/
__host__ __device__ int iDivUp(int a, int b) { return ((a % b) != 0) ? (a / b + 1) : (a / b); }

/*********************/
/* GENERATION KERNEL */
/*********************/
__global__ void generate_kernel(curandStateMtgp32 * __restrict__ state, float * __restrict__ result, const int N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    for (int k = tid; k < N; k += blockDim.x * gridDim.x)
        result[k] = curand_uniform(&state[blockIdx.x]);
}

extern "C" {

    /**************************/
    /* MERSENNE TWISTER SETUP */
    /**************************/
    __declspec(dllexport)
    void MTGP32Setup() {
        // --- Setup the pseudorandom number generator
        gpuErrchk(cudaMalloc(&devMTGPStates, GRIDSIZE * sizeof(curandStateMtgp32)));
        mtgp32_kernel_params *devKernelParams; gpuErrchk(cudaMalloc(&devKernelParams, sizeof(mtgp32_kernel_params)));
        curandSafeCall(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, devKernelParams));
        curandSafeCall(curandMakeMTGP32KernelState(devMTGPStates, mtgp32dc_params_fast_11213, devKernelParams, GRIDSIZE, time(NULL)));
    }

    /*******************************/
    /* MERSENNE TWISTER GENERATION */
    /*******************************/
    __declspec(dllexport)
    void MTGP32Generation(float * __restrict__ devResults, const int N) {
        // --- Generate pseudo-random sequence and copy to the host
        generate_kernel << <GRIDSIZE, BLOCKSIZE >> > (devMTGPStates, devResults, N);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());
    }
} //

Исходный код для PyCUDA абонента 

import os
import sys
import numpy              as     np
import ctypes
from   ctypes             import *

import pycuda.driver      as     drv
import pycuda.gpuarray    as     gpuarray
import pycuda.autoinit

lib = cdll.LoadLibrary('D:\\Project\\cuRAND\\mersenneTwisterDLL\\x64\\Release\\mersenneTwisterDLL.dll')

N = 10

d_x = gpuarray.zeros((N, 1), dtype = np.float32)

lib.MTGP32Setup()
lib.MTGP32Generation(ctypes.cast(d_x.ptr, POINTER(c_float)), N)

print(d_x)

API на стороне хоста можно обрабатывать аналогично Вызову host функций в PyCUDA.

...