Cuda migration (by amp and amohan)

Memory allocations must be handles

Author: vancemiller

cava/nightwatch/generator/c/callee.py

@@ -390,7 +390,7 @@ def buffer_case():
                 return ""
 
         def default_case():
-            return (Expr(type.transfer).equals("NW_HANDLE")).if_then_else(
+            return (Expr(type.transfer).one_of(["NW_HANDLE", "NW_OPAQUE"])).if_then_else(
                 Expr(not type.deallocates).if_then_else(
                     assign_record_replay_functions(param_value, type).then(record_call_metadata(param_value, type)),
                     expunge_calls(param_value, type),
@@ -399,9 +399,10 @@ def default_case():
 
         if type.fields:
             return for_all_elements(values, type, depth=depth, original_type=original_type, **other)
-        return type.is_simple_buffer().if_then_else(
-            simple_buffer_case, Expr(type.transfer).equals("NW_BUFFER").if_then_else(buffer_case, default_case)
-        )
+        return Expr(type.transfer).equals("NW_BUFFER").if_then_else(
+                buffer_case,
+                default_case
+            )
 
     with location(f"at {term.yellow(str(arg.name))}", arg.location):
         conv = convert_result_value(

cava/samples/cudart/cudart.cpp

@@ -58,6 +58,8 @@ typedef struct {
     int num_funcs;
     struct fatbin_function *func; /* for functions */
 
+    size_t buffer_size;
+
     /* global states */
     CUmodule cur_module;
 
@@ -442,21 +444,51 @@ cudaFreeHost(void *ptr)
 }
 ava_end_replacement;
 
+
+/// Migration: @mem_extract_tag
+ava_utility void* object_extract(void* obj, size_t* length) {
+    //called from host
+    void* buffer;
+    printf("object_replay: object=%lx\n", (uintptr_t)obj);
+    cudaDeviceSynchronize();
+
+    *length = ava_metadata(obj)->buffer_size;
+    buffer = malloc(*length);
+    cudaMemcpy(buffer, obj, *length, cudaMemcpyDeviceToHost);
+    return buffer;
+}
+
+ava_utility void object_replace(void* obj, void* data, size_t length) {
+    printf("object_replace: object=%lx, len=%lu\n", (uintptr_t)obj, length);
+    assert(length != 0);
+    cudaMemcpy(obj, data, length, cudaMemcpyHostToDevice);
+}
+
 __host__ __cudart_builtin__ cudaError_t CUDARTAPI
 cudaMalloc(void **devPtr, size_t size)
 {
     ava_argument(devPtr) {
+
         ava_out; ava_buffer(1);
-        ava_element ava_opaque;
+        ava_element{
+            ava_allocates;
+            ava_handle;
+            ava_object_explicit_state_functions(object_extract, object_replace);
+            ava_object_record;
+        }
     }
+
+    ava_execute();
+    ava_metadata(*devPtr)->buffer_size = size;
 }
 
 __host__ cudaError_t CUDARTAPI
 cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
 {
     ava_argument(dst) {
         if (kind == cudaMemcpyHostToDevice) {
-            ava_opaque;
+            ava_handle;
+            ava_object_record;
         }
         else if (kind == cudaMemcpyDeviceToHost) {
             ava_out; ava_buffer(count);
@@ -468,15 +500,18 @@ cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
             ava_in; ava_buffer(count);
         }
         else if (kind == cudaMemcpyDeviceToHost) {
-            ava_opaque;
+            ava_handle;
         }
     }
 }
 
 __host__ __cudart_builtin__ cudaError_t CUDARTAPI
 cudaFree(void *devPtr)
 {
-    ava_argument(devPtr) ava_opaque;
+    ava_argument(devPtr){
+        ava_handle;
+        ava_object_record;
+    }
 }
 
 /* Rich set of APIs */