feat: Add cudaStreamSynchronize

To support cudaStreamSynchronize API

Author: coldfunction

qcu-device/hw/misc/virtio-qcuda.c

@@ -923,6 +923,18 @@ static void qcu_cudaStreamDestroy(VirtioQCArg *arg)
 	memset(&cudaStream[idx], 0, sizeof(cudaStream_t));
 }
 
+static void qcu_cudaStreamSynchronize(VirtioQCArg *arg)
+{
+	cudaError_t err;
+	uint32_t idx;
+	idx = arg->pA;
+
+	cudaError((err = cudaStreamSynchronize(cudaStream[idx])));
+	
+	arg->cmd = err;
+}
+
+
 #endif // CONFIG_CUDA
 
 static int qcu_cmd_write(VirtioQCArg *arg)
@@ -1252,6 +1264,10 @@ static void virtio_qcuda_cmd_handle(VirtIODevice *vdev, VirtQueue *vq)
 				qcu_cudaStreamDestroy(arg);	
 				break;
 
+			case VIRTQC_cudaStreamSynchronize:
+				qcu_cudaStreamSynchronize(arg);
+				break;
+
 			// Event Management (runtime API)
 			case VIRTQC_cudaEventCreate:
 				qcu_cudaEventCreate(arg);

qcu-driver/qcuda_common.h

@@ -78,6 +78,7 @@ enum
 	//stream
 	VIRTQC_cudaStreamCreate,
 	VIRTQC_cudaStreamDestroy,
+	VIRTQC_cudaStreamSynchronize,
 };
 
 typedef struct VirtioQCArg   VirtioQCArg;

qcu-driver/qcuda_driver.c

@@ -1044,6 +1044,11 @@ void qcu_cudaStreamDestroy(VirtioQCArg *arg)
 	qcu_misc_send_cmd(arg);
 }
 
+void qcu_cudaStreamSynchronize(VirtioQCArg *arg)
+{
+	qcu_misc_send_cmd(arg);
+}
+
 
 // @_cmd: device command
 // @_arg: argument of cuda function
@@ -1216,6 +1221,10 @@ static long qcu_misc_ioctl(struct file *filp, unsigned int _cmd, unsigned long _
 			qcu_cudaStreamDestroy(arg);
 			break;
 
+		case VIRTQC_cudaStreamSynchronize:
+			qcu_cudaStreamSynchronize(arg);
+			break;
+
 		default:
 			error("unknow cmd= %d\n", arg->cmd);
 			break;

qcu-library/libcudart.c

@@ -821,7 +821,6 @@ cudaError_t cudaStreamCreate(cudaStream_t *pStream)
 cudaError_t cudaStreamDestroy(cudaStream_t stream) 
 {
 	VirtioQCArg arg;
-
 	memset(&arg, 0, sizeof(VirtioQCArg));
 
 	ptr( arg.pA, stream, 0);
@@ -831,6 +830,20 @@ cudaError_t cudaStreamDestroy(cudaStream_t stream)
 
 }	
 
+cudaError_t cudaStreamSynchronize(cudaStream_t 	stream) 	
+{
+	VirtioQCArg arg;
+	memset(&arg, 0, sizeof(VirtioQCArg));
+
+	uint64_t mystream = (stream==NULL)?(uint64_t)-1:(uint64_t)stream;
+
+	ptr( arg.pA, mystream, 0);
+
+	send_cmd_to_device( VIRTQC_cudaStreamSynchronize, &arg);
+
+	return (cudaError_t)arg.cmd;
+}
+
 // Macro to aligned up to the memory size in question                                                         
 #define MEMORY_ALIGNMENT  4096
 #define ALIGN_UP(x,size) ( ((size_t)x+(size-1))&(~(size-1)) )
@@ -869,3 +882,5 @@ cudaError_t cudaFreeHost(void *ptr)
 }
 
 
+
+