add CuContext/MContextwrapped-handles

1 files changed, 3 insertions, 4 deletions

diff --git a/nerv/lib/matrix/generic/cukernel.cu b/nerv/lib/matrix/generic/cukernel.cu
index 51e3b6a..0e09cfa 100644
--- a/nerv/lib/matrix/generic/cukernel.cu
+++ b/nerv/lib/matrix/generic/cukernel.cu
@@ -445,13 +445,12 @@ extern "C" {
         cudaStreamSynchronize(0);
     }
 
-    extern curandGenerator_t curand_gen;
-    void cudak_(cuda_rand_uniform)(const Matrix *a) {
+    void cudak_(cuda_rand_uniform)(const Matrix *a, CuContext *context) {
         #ifdef MATRIX_USE_FLOAT
-        curandGenerateUniform(curand_gen, MATRIX_ELEM_PTR(a), a->nrow * a->stride / sizeof(MATRIX_ELEM));
+        curandGenerateUniform(context->curand_gen, MATRIX_ELEM_PTR(a), a->nrow * a->stride / sizeof(MATRIX_ELEM));
         #endif
         #ifdef MATRIX_USE_DOUBLE
-        curandGenerateUniformDouble(curand_gen, MATRIX_ELEM_PTR(a), a->nrow * a->stride / sizeof(MATRIX_ELEM));
+        curandGenerateUniformDouble(context->curand_gen, MATRIX_ELEM_PTR(a), a->nrow * a->stride / sizeof(MATRIX_ELEM));
         #endif
     }