diff --git a/src/cuda_queue.cu b/src/cuda_queue.cu index 07e28fe510c367f72709a5ec0b74e1f6f30e4427..4477f61be656a653bb98c98b8d1dd67e68a5d15d 100644 --- a/src/cuda_queue.cu +++ b/src/cuda_queue.cu @@ -28,6 +28,12 @@ __device__ struct queue_cuda cuda_queues[ cuda_numqueues ]; __constant__ int cuda_nrqueues; __constant__ int cuda_queue_size; +__device__ struct task *tasks_cuda; +__device__ int *locks_cuda; +__device__ int *uses_cuda; +__device__ int *deps_cuda; +__device__ struct res *res_cuda; + /** * @brief Get a task ID from the given queue. * @@ -140,3 +146,38 @@ __device__ int runner_cuda_gettask ( struct queue_cuda *q ) { #endif + + + +int qsched_prepare_cuda ( struct qsched *s , ) { + +/*Copy the qsched data to the device*/ +if( cudaMalloc( &tasks_cuda , sizeof(struct task) * s->count != cudaSuccess ) + error("Failed to allocate task array on the device."); + +if( cudaMemcpy( ) + +if( cudaMalloc( &locks_cuda , sizeof(int) * s->count_lock != cudaSuccess ) + error("Failed to allocate locks array on the device."); + +if( cudaMalloc( &deps_cuda , sizeof(int) * s->count_deps != cudaSuccess ) + error("Failed to allocate deps array on the device."); + +if( cudaMalloc( &uses_cuda , sizeof(int) * s->count_uses != cudaSuccess ) + error("Failed to allocate use array on the device."); + +if( cudaMalloc( &res_cuda , sizeof( res) * s->count_res != cudaSuccess ) + error("Failed to allocated on the device."); + +/* Initialize the queues. */ +int nr_queues= 2, k, qsize; +int *data; +struct queue_cuda queues[ cuda_maxqueues ]; + + + + + +} + + diff --git a/src/cuda_queue.h b/src/cuda_queue.h index d86df771b754c59571d8c559633cd0f1b9553d2e..a546786ba406fc17ef584e50c4fe44ae19f8737a 100644 --- a/src/cuda_queue.h +++ b/src/cuda_queue.h @@ -39,27 +39,3 @@ struct queue_cuda { volatile int *rec_data; }; - - -/** Struct for each task. */ -struct task_cuda { - - /** Task type and subtype. */ - short int type, subtype; - - /** Wait counters. */ - volatile int wait; - - /** Task flags. */ - int flags; - - /** Indices of the cells involved. */ - int i, j; - - /** Nr of task that this task unlocks. */ - int nr_unlock; - - /** List of task that this task unlocks (dependencies). */ - int unlock[ cuda_max_unlock ]; - - }; diff --git a/src/lock.h b/src/lock.h index 60329d9709efbd68a2936c52021a847a51edbb81..abdfc5ddf1df0ee03438ca3aa99385c9ccf6f3ad 100644 --- a/src/lock.h +++ b/src/lock.h @@ -47,7 +47,7 @@ int res = atomicCAS ( l , 1 , 0 ) != 1; return res; } - #define lock_unlock_blink( l ) atomicCAS( l , 1 , 0 ) + #define lock_unlock_blind( l ) atomicCAS( &l , 1 , 0 ) #elif defined (PTHREAD_LOCK) #define lock_type pthread_spinlock_t #define lock_init( l ) ( pthread_spin_init( l , PTHREAD_PROCESS_PRIVATE ) != 0 ) diff --git a/src/qsched.c b/src/qsched.c index e95cfdae39dee8adf37afd8ae702a82220e1ae1f..a183dd24611a3b9e8f6b7c8ac95e5b9a168bb148 100644 --- a/src/qsched.c +++ b/src/qsched.c @@ -998,11 +998,11 @@ void qsched_prepare ( struct qsched *s ) { * * @param s Pointer to the #qsched. * @param parent ID of the parent resource or #qsched_res_none if none. - * + * @param size Size of the data in bytes. * @return The ID of the new shared resource. */ -int qsched_addres ( struct qsched *s , int owner , int parent ) { +int qsched_addres ( struct qsched *s , int owner , int parent, void* data, int size ) { struct res *res_new; int id; @@ -1040,6 +1040,13 @@ int qsched_addres ( struct qsched *s , int owner , int parent ) { s->res[ id ].hold = 0; s->res[ id ].owner = owner; s->res[ id ].parent = parent; + s->res[ id ].data = data; + s->res[ id ].size = size; + + #ifdef WITH_CUDA + cudaMalloc( &s->res[ id ].cuda_data, size ); + #endif + /* Unlock the sched. */ lock_unlock_blind( &s->lock ); diff --git a/src/qsched.h b/src/qsched.h index a1dd360243adfd85eff64e01eedcfe7a5dacc6d5..ee08fe41f7a58d530e8af0277a499aafa2167b4e 100644 --- a/src/qsched.h +++ b/src/qsched.h @@ -177,7 +177,7 @@ void qsched_enqueue ( struct qsched *s , struct task *t ); /* External functions. */ void qsched_init ( struct qsched *s , int nr_queues , int flags ); -qsched_res_t qsched_addres ( struct qsched *s , int owner , qsched_res_t parent ); +qsched_res_t qsched_addres ( struct qsched *s , int owner , qsched_res_t parent , void* data, int size); void qsched_addlock ( struct qsched *s , qsched_task_t t , qsched_res_t res ); void qsched_addunlock ( struct qsched *s , qsched_task_t ta , qsched_task_t tb ); qsched_task_t qsched_addtask ( struct qsched *s , int type , unsigned int flags , void *data , int data_size , int cost ); diff --git a/src/res.h b/src/res.h index 9e5a81efa5aa4e7d3ec331699eb2f6c72e517321..bb2b3043eb974d09193a6dd06b153b111bfe198c 100644 --- a/src/res.h +++ b/src/res.h @@ -33,9 +33,15 @@ struct res { /* The resource's parent. */ int parent; - #ifdef WITH_CUDA - void* data, gpu_data; + /* Pointer to data on the CPU */ + void* data; + + /*Size of the data associated with this resource */ int size; + + #ifdef WITH_CUDA + /* Pointer to the data on the GPU */ + void* gpu_data; #endif };