Dnia 2011-04-24, nie o godzinie 17:50 +0200, Tomasz Rybak pisze: > Dnia 2011-04-24, nie o godzinie 01:22 -0400, Andreas Kloeckner pisze: > > On Sun, 17 Apr 2011 17:10:30 +0200, "=?UTF-8?B?VG9tYXN6IFJ5YmFr?=" > > <[email protected]> wrote: > > > Dnia 2011-04-16 00:48 Andreas Kloeckner napisał(a): > > > > > > >Hi Tomasz, > > > > > > > >any progress on the things below? Has maybe another email gone missing? > > > >I'd like to release PyCUDA 2011.1 soon. > > > > > > > > > I have send the patch to the mailing list on 2011-03-18 and 2011-03-24. > > > Please check the archives. > > > > Weird. I'm really not sure what happened there--I could've sworn I never > > saw these emails. > > > > In any case, I've merged your patch into the > > curand-wrapper-v2-from-tomasz branch, which I've also brought up to > > current master. > > Thanks! > > > > > Looks like we're almost done here--only initialization is missing. > > I hope to have it done by the end of the next week.
I have troubles with transferring data between C and Python. I want to: 1. call curandGetDirectionVectors32 which returns pointer to the 32 int32 2. transfer this data to the device memory 3. create GPUArray from it 4. call prepare() kernel which will pass appropriate direction vectors to the curand_init() kernel I have tried two approaches (curand-hostptr.diff and curand-memcpy.diff). The former, when compiled, causes _curand module to misbehave - e.g. it get_curand_version() returns None instead of (3, 2, 0) or (0, 0, 0) The latter compiles but cuMemcpy call fails because of parameter type mismatch. Can someone give me some direction how to transfer raw pointer from C to Python? Best regards. -- Tomasz Rybak <[email protected]> GPG/PGP key ID: 2AD5 9860 Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860 http://member.acm.org/~tomaszrybak
diff --git a/pycuda/curandom.py b/pycuda/curandom.py
index 73ca18d..f9a2da2 100644
--- a/pycuda/curandom.py
+++ b/pycuda/curandom.py
@@ -526,6 +526,15 @@ if get_curand_version() >= (3, 2, 0):
# {{{ Sobol32 RNG
+# TODO
+def generate_direction_vectors(count):
+ result = numpy.empty((count, 32), dtype=numpy.int32)
+ for i in range(count):
+ a = get_direction_vectors32()
+ b = numpy.fromstring(a.data(), dtype=numpy.int32, count=32)
+ result[i, :] = b
+ return pycuda.GPUArray.to_gpu(result)
+
sobol32_random_source = """
extern "C" {
__global__ void prepare(curandStateSobol32 *s, const int n, unsigned int **v,
@@ -589,6 +598,8 @@ if get_curand_version() >= (3, 2, 0):
if has_stack:
drv.Context.set_limit(drv.limit.STACK_SIZE, prev_stack_size)
+
+
# }}}
# }}}
diff --git a/src/cpp/curand.hpp b/src/cpp/curand.hpp
index e025b0c..a9f09bd 100644
--- a/src/cpp/curand.hpp
+++ b/src/cpp/curand.hpp
@@ -6,6 +6,23 @@
#include <curand.h>
#endif
+#ifdef CUDAPP_TRACE_CUDA
+ #define CURAND_PRINT_ERROR_TRACE(NAME, CODE) \
+ if (CODE != CURAND_STATUS_SUCCESS) \
+ std::cerr << NAME << " failed with code " << CODE << std::endl;
+#else
+ #define CURAND_PRINT_ERROR_TRACE(NAME, CODE) /*nothing*/
+#endif
+// TODO: add printing as in PRINT_ERROR
+#define CURAND_CALL_GUARDED(NAME, ARGLIST) \
+ { \
+ CUDAPP_PRINT_CALL_TRACE(#NAME); \
+ curandStatus_t cu_status_code; \
+ cu_status_code = NAME ARGLIST; \
+ CURAND_PRINT_ERROR_TRACE(#NAME, cu_status_code); \
+ if (cu_status_code != CURAND_STATUS_SUCCESS) \
+ throw pycuda::error(#NAME, CUDA_SUCCESS);\
+ }
namespace pycuda { namespace curandom {
@@ -22,10 +39,14 @@ namespace pycuda { namespace curandom {
}
#if CUDAPP_CUDA_VERSION >= 3020
- void py_curand_get_direction_vectors32(curandDirectionVectors32_t *vectors[],
+ boost::shared_ptr<host_pointer> py_curand_get_direction_vectors32(
curandDirectionVectorSet_t set)
-// TODO: checking; cannot use CUDAPP_CALL_GUARDED because function returns CURAND enum
- { curandGetDirectionVectors32(vectors, set); }
+ {
+ curandDirectionVectors32_t *vectors;
+ CURAND_CALL_GUARDED(curandGetDirectionVectors32, (&vectors, set));
+ boost::shared_ptr<host_pointer> result(new host_pointer((void *)vectors));
+ return result;
+ }
#endif
} }
diff --git a/src/wrapper/wrap_cudadrv.cpp b/src/wrapper/wrap_cudadrv.cpp
index de19bba..30d2ca2 100644
--- a/src/wrapper/wrap_cudadrv.cpp
+++ b/src/wrapper/wrap_cudadrv.cpp
@@ -902,8 +902,6 @@ BOOST_PYTHON_MODULE(_driver)
.def("_param_setv", function_param_setv)
.DEF_SIMPLE_METHOD(param_set_texref)
- .def("launch_kernel", &cl::launch_kernel)
-
.def("_launch", &cl::launch)
.def("_launch_grid", &cl::launch_grid,
py::args("grid_width", "grid_height"))
@@ -916,6 +914,9 @@ BOOST_PYTHON_MODULE(_driver)
#if CUDAPP_CUDA_VERSION >= 3000 && defined(CUDAPP_POST_30_BETA)
.DEF_SIMPLE_METHOD(set_cache_config)
#endif
+#if CUDAPP_CUDA_VERSION >= 4000
+ .def("launch_kernel", &cl::launch_kernel)
+#endif
;
}
diff --git a/src/wrapper/wrap_curand.cpp b/src/wrapper/wrap_curand.cpp
index 71adf57..d970b13 100644
--- a/src/wrapper/wrap_curand.cpp
+++ b/src/wrapper/wrap_curand.cpp
@@ -25,7 +25,8 @@ void pycuda_expose_curand()
py::def("get_curand_version", py_curand_version);
#if CUDAPP_CUDA_VERSION >= 3020
- py::def("get_direction_vectors32", py_curand_get_direction_vectors32, (arg("vectors"), arg("set")));
+ py::def("get_direction_vectors32", py_curand_get_direction_vectors32,
+ arg("set"));
#endif
}
diff --git a/pycuda/curandom.py b/pycuda/curandom.py
index 73ca18d..57362d8 100644
--- a/pycuda/curandom.py
+++ b/pycuda/curandom.py
@@ -526,9 +526,16 @@ if get_curand_version() >= (3, 2, 0):
# {{{ Sobol32 RNG
+# TODO
+def generate_direction_vectors(count):
+ result = pycuda.GPUArray.empty((count, 32), dtype=numpy.int32)
+ for i in range(count):
+ get_direction_vectors32(direction_vector_set.VECTOR_32, result.gpudata, i)
+ return result
+
sobol32_random_source = """
extern "C" {
-__global__ void prepare(curandStateSobol32 *s, const int n, unsigned int **v,
+__global__ void prepare(curandStateSobol32 *s, const int n, curandDirectionVectors32_t *v,
const unsigned int o)
{
const int id = blockIdx.x*blockDim.x+threadIdx.x;
@@ -589,6 +596,8 @@ if get_curand_version() >= (3, 2, 0):
if has_stack:
drv.Context.set_limit(drv.limit.STACK_SIZE, prev_stack_size)
+
+
# }}}
# }}}
diff --git a/src/cpp/curand.hpp b/src/cpp/curand.hpp
index e025b0c..fc19670 100644
--- a/src/cpp/curand.hpp
+++ b/src/cpp/curand.hpp
@@ -6,6 +6,23 @@
#include <curand.h>
#endif
+#ifdef CUDAPP_TRACE_CUDA
+ #define CURAND_PRINT_ERROR_TRACE(NAME, CODE) \
+ if (CODE != CURAND_STATUS_SUCCESS) \
+ std::cerr << NAME << " failed with code " << CODE << std::endl;
+#else
+ #define CURAND_PRINT_ERROR_TRACE(NAME, CODE) /*nothing*/
+#endif
+// TODO: add printing as in PRINT_ERROR
+#define CURAND_CALL_GUARDED(NAME, ARGLIST) \
+ { \
+ CUDAPP_PRINT_CALL_TRACE(#NAME); \
+ curandStatus_t cu_status_code; \
+ cu_status_code = NAME ARGLIST; \
+ CURAND_PRINT_ERROR_TRACE(#NAME, cu_status_code); \
+ if (cu_status_code != CURAND_STATUS_SUCCESS) \
+ throw pycuda::error(#NAME, CUDA_SUCCESS);\
+ }
namespace pycuda { namespace curandom {
@@ -22,10 +39,13 @@ namespace pycuda { namespace curandom {
}
#if CUDAPP_CUDA_VERSION >= 3020
- void py_curand_get_direction_vectors32(curandDirectionVectors32_t *vectors[],
- curandDirectionVectorSet_t set)
-// TODO: checking; cannot use CUDAPP_CALL_GUARDED because function returns CURAND enum
- { curandGetDirectionVectors32(vectors, set); }
+ void py_curand_get_direction_vectors32(
+ curandDirectionVectorSet_t set, int array, int position)
+ {
+ curandDirectionVectors32_t *vectors;
+ CURAND_CALL_GUARDED(curandGetDirectionVectors32, (&vectors, set));
+ CUDAPP_CALL_GUARDED(cuMemcpyHtoD, (array+position*32*4, vectors, 32*4));
+ }
#endif
} }
diff --git a/src/wrapper/wrap_cudadrv.cpp b/src/wrapper/wrap_cudadrv.cpp
index de19bba..30d2ca2 100644
--- a/src/wrapper/wrap_cudadrv.cpp
+++ b/src/wrapper/wrap_cudadrv.cpp
@@ -902,8 +902,6 @@ BOOST_PYTHON_MODULE(_driver)
.def("_param_setv", function_param_setv)
.DEF_SIMPLE_METHOD(param_set_texref)
- .def("launch_kernel", &cl::launch_kernel)
-
.def("_launch", &cl::launch)
.def("_launch_grid", &cl::launch_grid,
py::args("grid_width", "grid_height"))
@@ -916,6 +914,9 @@ BOOST_PYTHON_MODULE(_driver)
#if CUDAPP_CUDA_VERSION >= 3000 && defined(CUDAPP_POST_30_BETA)
.DEF_SIMPLE_METHOD(set_cache_config)
#endif
+#if CUDAPP_CUDA_VERSION >= 4000
+ .def("launch_kernel", &cl::launch_kernel)
+#endif
;
}
diff --git a/src/wrapper/wrap_curand.cpp b/src/wrapper/wrap_curand.cpp
index 71adf57..085c307 100644
--- a/src/wrapper/wrap_curand.cpp
+++ b/src/wrapper/wrap_curand.cpp
@@ -25,7 +25,8 @@ void pycuda_expose_curand()
py::def("get_curand_version", py_curand_version);
#if CUDAPP_CUDA_VERSION >= 3020
- py::def("get_direction_vectors32", py_curand_get_direction_vectors32, (arg("vectors"), arg("set")));
+ py::def("get_direction_vectors32", py_curand_get_direction_vectors32,
+ (arg("set"), arg("array"), arg("position")));
#endif
}
signature.asc
Description: This is a digitally signed message part
_______________________________________________ PyCUDA mailing list [email protected] http://lists.tiker.net/listinfo/pycuda
