Skip to content
Permalink
Browse files

Adding in more changes for New and BCL integration, including meta lo…

…oking, assembly reading, etc.
  • Loading branch information...
kaby76 committed Apr 14, 2018
1 parent 17fea96 commit 987209c6e9b852c18c6bc08dbbd6e4778b309b32
@@ -163,18 +163,17 @@ public IntPtr AddDataStructure(object to_gpu)
{
IntPtr result = IntPtr.Zero;
var type = to_gpu.GetType();
var btype = type;

result = _allocated_objects.Where(p => p.Key == to_gpu).FirstOrDefault().Value;
var find_object = _allocated_objects.Where(p => p.Key == to_gpu);

// Allocate new buffer for object on GPU.
if (result == IntPtr.Zero)
if (!find_object.Any())
{
if (Campy.Utils.Options.IsOn("copy_trace"))
System.Console.WriteLine("Allocating GPU buf " + to_gpu);
result = New(BUFFERS.SizeOf(btype));
result = New(BUFFERS.SizeOf(type));
_allocated_objects[to_gpu] = result;
}
else result = find_object.First().Value;

// Copy to GPU if it hasn't been done before.
if (!_copied_to_gpu.Contains(result))
@@ -1293,44 +1292,65 @@ public IntPtr New(int bytes)
//}
}

[global::System.Runtime.InteropServices.DllImport(
@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint =
"?GfsAddFile@@YAXPEAX0_K0@Z")]
public static extern System.IntPtr BclHeapAlloc(
[MarshalAs(UnmanagedType.LPStr)]string assemblyName,
[MarshalAs(UnmanagedType.LPStr)]string nameSpace,
[MarshalAs(UnmanagedType.LPStr)]string name);
[global::System.Runtime.InteropServices.DllImport(@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint = "?BclHeapAlloc@@YAPEAXPEAX@Z")]
public static extern System.IntPtr BclHeapAlloc(System.IntPtr bcl_type);

[global::System.Runtime.InteropServices.DllImport(@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint = "?BclGetMetaOfType@@YAPEAXPEAD00PEAX@Z")]
public static extern System.IntPtr BclGetMetaOfType(
[MarshalAs(UnmanagedType.LPStr)] string assemblyName,
[MarshalAs(UnmanagedType.LPStr)] string nameSpace,
[MarshalAs(UnmanagedType.LPStr)] string name,
System.IntPtr nested);

private IntPtr GetBclType(Type type)
{
Stack<Type> chain = new Stack<Type>();
while (type != null)
{
chain.Push(type);
type = type.DeclaringType;
}

System.IntPtr result = System.IntPtr.Zero;

while (chain.Any())
{
type = chain.Pop();
var tr = type.ToMonoTypeReference();
tr = RUNTIME.RewriteType(tr);
var assembly_name = tr.Module.Name;
var name_space = tr.Namespace;
var name = tr.Name;
// Make sure assembly is placed in GPU BCL file system.
JITER.Singleton.AddAssemblyToFileSystem(tr.Module);
result = BclGetMetaOfType(assembly_name, name_space, name, result);
}

return result;
}

public IntPtr New(Type type)
{
var assembly_name = type.Assembly.FullName;
var name_space = type.Namespace;
var name = type.FullName;
return BclHeapAlloc(assembly_name, name_space, name);
var bcl_type = GetBclType(type);
return BclHeapAlloc(bcl_type);
}


[global::System.Runtime.InteropServices.DllImport(
@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint =
"?BclArrayAlloc@@YAPEAXPEAD00H@Z")]
[global::System.Runtime.InteropServices.DllImport(@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint = "?BclArrayAlloc@@YAPEAXPEAXHPEAI@Z")]
public static extern System.IntPtr BclArrayAlloc(
[MarshalAs(UnmanagedType.LPStr)]string assemblyName,
[MarshalAs(UnmanagedType.LPStr)]string nameSpace,
[MarshalAs(UnmanagedType.LPStr)]string name,
int length);
System.IntPtr bcl_type,
int rank,
uint[] lengths);

public IntPtr New(Array array)
{
Type type = array.GetType().GetElementType();
var bcl_type = GetBclType(type);

var tr = type.ToMonoTypeReference();
tr = RUNTIME.RewriteType(tr);

var assembly_name = tr.Module.Name;
var name_space = tr.Namespace;
var name = tr.Name;
uint[] lengths = new uint[array.Rank];
for (int i = 0; i < array.Rank; ++i) lengths[i] = (uint)array.GetLength(i);

return BclArrayAlloc(assembly_name, name_space, name, array.Length);
return BclArrayAlloc(bcl_type, array.Rank, lengths);
}

public void Free(IntPtr pointer)
@@ -447,8 +447,18 @@ public class JITER
private static bool init;
private Dictionary<MethodInfo, IntPtr> method_to_image;
private bool done_init;
private static JITER _singleton;

public JITER()
public static JITER Singleton
{
get
{
if (_singleton == null) _singleton = new JITER();
return _singleton;
}
}

private JITER()
{
global_llvm_module = default(ModuleRef);
all_llvm_modules = new List<ModuleRef>();
@@ -693,8 +703,7 @@ private CFG.Vertex Eval(CFG.Vertex current, Dictionary<Tuple<TypeReference, Gene
return current;
}

private bool TypeUsingGeneric()
{ return false; }
private bool TypeUsingGeneric() { return false; }

public List<CFG.Vertex> InstantiateGenerics(IEnumerable<CFG.Vertex> change_set, List<System.Type> list_of_data_types_used, List<Mono.Cecil.TypeReference> list_of_mono_data_types_used)
{
@@ -2008,8 +2017,6 @@ public CUfunction GetCudaFunction(MethodInfo kernel_method, IntPtr image)
return helloWorld;
}

private bool all_new = true;

[global::System.Runtime.InteropServices.DllImport(@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint = "?InitTheBcl@@YAXPEAX_KH0@Z")]
public static extern void InitTheBcl(System.IntPtr a1, long a2, int a3, System.IntPtr a4);

@@ -2025,6 +2032,28 @@ public CUfunction GetCudaFunction(MethodInfo kernel_method, IntPtr image)
[global::System.Runtime.InteropServices.DllImport(@"C:\Users\kenne\Documents\Campy2\ConsoleApp4\bin\Debug\Campy.Runtime.Wrapper.dll", EntryPoint = "?InitializeBCL2@@YAXXZ")]
public static extern void InitializeBCL2();


public void AddAssemblyToFileSystem(Mono.Cecil.ModuleDefinition module)
{
// Set up corlib.dll in file system.
string full_path_assem = module.FileName;
string assem = Path.GetFileName(full_path_assem);
Stream stream = new FileStream(full_path_assem, FileMode.Open, FileAccess.Read, FileShare.Read);
var corlib_bytes_handle_len = stream.Length;
var corlib_bytes = new byte[corlib_bytes_handle_len];
stream.Read(corlib_bytes, 0, (int)corlib_bytes_handle_len);
var corlib_bytes_handle = GCHandle.Alloc(corlib_bytes, GCHandleType.Pinned);
var corlib_bytes_intptr = corlib_bytes_handle.AddrOfPinnedObject();
stream.Close();
stream.Dispose();
var ptrx = Marshal.StringToHGlobalAnsi(assem);
BUFFERS buffers = new BUFFERS();
IntPtr pointer1 = buffers.New(assem.Length + 1);
BUFFERS.Cp(pointer1, ptrx, assem.Length + 1);
var pointer4 = buffers.New(sizeof(int));
GfsAddFile(pointer1, corlib_bytes_intptr, corlib_bytes_handle_len, pointer4);
}

public void InitBCL(CUmodule mod)
{
if (!done_init)
@@ -2045,27 +2074,30 @@ public void InitBCL(CUmodule mod)
{
InitFileSystem();

// Set up corlib.dll in file system.
string full_path_assem = RUNTIME.FindCoreLib();
string assem = Path.GetFileName(full_path_assem);
Stream stream = new FileStream(full_path_assem, FileMode.Open, FileAccess.Read, FileShare.Read);
var corlib_bytes_handle_len = stream.Length;
var corlib_bytes = new byte[corlib_bytes_handle_len];
stream.Read(corlib_bytes, 0, (int) corlib_bytes_handle_len);
var corlib_bytes_handle = GCHandle.Alloc(corlib_bytes, GCHandleType.Pinned);
var corlib_bytes_intptr = corlib_bytes_handle.AddrOfPinnedObject();
stream.Close();
stream.Dispose();
var ptrx = Marshal.StringToHGlobalAnsi(assem);
BUFFERS buffers = new BUFFERS();
IntPtr pointer1 = buffers.New(assem.Length + 1);
BUFFERS.Cp(pointer1, ptrx, assem.Length + 1);
var pointer4 = buffers.New(sizeof(int));

GfsAddFile(pointer1, corlib_bytes_intptr, corlib_bytes_handle_len, pointer4);
{
// Set up corlib.dll in file system.
string full_path_assem = RUNTIME.FindCoreLib();
string assem = Path.GetFileName(full_path_assem);
Stream stream = new FileStream(full_path_assem, FileMode.Open, FileAccess.Read, FileShare.Read);
var corlib_bytes_handle_len = stream.Length;
var corlib_bytes = new byte[corlib_bytes_handle_len];
stream.Read(corlib_bytes, 0, (int)corlib_bytes_handle_len);
var corlib_bytes_handle = GCHandle.Alloc(corlib_bytes, GCHandleType.Pinned);
var corlib_bytes_intptr = corlib_bytes_handle.AddrOfPinnedObject();
stream.Close();
stream.Dispose();
var ptrx = Marshal.StringToHGlobalAnsi(assem);
BUFFERS buffers = new BUFFERS();
IntPtr pointer1 = buffers.New(assem.Length + 1);
BUFFERS.Cp(pointer1, ptrx, assem.Length + 1);
var pointer4 = buffers.New(sizeof(int));
GfsAddFile(pointer1, corlib_bytes_intptr, corlib_bytes_handle_len, pointer4);
}

InitializeBCL1();

InitializeBCL2();

done_init = true;
}
}
@@ -171,6 +171,7 @@ function_space_specifier PTR MetaData_GetTypeMethodField(tMetaData *pMetaData, I

function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromName(tMetaData *pMetaData, STRING nameSpace, STRING name, tMD_TypeDef *pInNestedClass);
function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromFullName(STRING assemblyName, STRING nameSpace, STRING name);
function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromFullNameAndNestedType(STRING assemblyName, STRING nameSpace, STRING name, tMD_TypeDef* nested);
function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromDefRefOrSpec(tMetaData *pMetaData, IDX_TABLE token, tMD_TypeDef **ppClassTypeArgs, tMD_TypeDef **ppMethodTypeArgs);
function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromMethodDef(tMD_MethodDef *pMethodDef);
function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromFieldDef(tMD_FieldDef *pFieldDef);
@@ -201,6 +201,16 @@ function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromFullName(STRING ass
return MetaData_GetTypeDefFromName(pTypeMetaData, nameSpace, name, NULL);
}

function_space_specifier tMD_TypeDef* MetaData_GetTypeDefFromFullNameAndNestedType(STRING assemblyName, STRING nameSpace, STRING name, tMD_TypeDef* nested) {
tMetaData *pTypeMetaData;

pTypeMetaData = CLIFile_GetMetaDataForAssembly(assemblyName);

// Note that this cannot get a nested class, as this final parameter is always NULL
return MetaData_GetTypeDefFromName(pTypeMetaData, nameSpace, name, nested);
}


__global__
void Bcl_MetaData_GetTypeDefFromDefRefOrSpec(tMetaData *pMetaData, IDX_TABLE token, tMD_TypeDef **ppClassTypeArgs, tMD_TypeDef **ppMethodTypeArgs, tMD_TypeDef** result)
{
@@ -234,7 +234,8 @@ function_space_specifier tAsyncCall* System_Array_Resize(PTR pThis_, PTR pParams
newSize = *(U32*)p++;

pOldArray = (tSystemArray*)*ppArray_;
U32 len = *((&(pOldArray->rank)) + 1);
U32 rank = *((&(pOldArray->rank)) + 1);
int len = *((&(pOldArray->rank)) + 2);;
oldSize = len;

if (oldSize == newSize) {
@@ -243,7 +244,7 @@ function_space_specifier tAsyncCall* System_Array_Resize(PTR pThis_, PTR pParams
}

pArrayTypeDef = Heap_GetType(*ppArray_);
pHeap = SystemArray_NewVector(pArrayTypeDef, newSize);
pHeap = SystemArray_NewVector(pArrayTypeDef, rank, &newSize);
pNewArray = (tSystemArray*)pHeap;
*ppArray_ = pHeap;
PTR beginning_of_elements = pNewArray->ptr_elements;
@@ -284,21 +285,26 @@ function_space_specifier tAsyncCall* System_Array_Reverse(PTR pThis_, PTR pParam
return NULL;
}

function_space_specifier HEAP_PTR SystemArray_NewVector(tMD_TypeDef *pArrayTypeDef, U32 length) {
function_space_specifier HEAP_PTR SystemArray_NewVector(tMD_TypeDef *pArrayTypeDef, U32 rank, U32* lengths) {
U32 heapSize;
tSystemArray *pArray;
// The size of an array depends on the rank.
heapSize = sizeof(void*); // ptr to first element.
int next = sizeof(INT64); // size of rank = 1
int next = sizeof(INT64); // for rank
heapSize += next;
next = sizeof(INT64) * 1; // assume rank = 1.
next = sizeof(INT64) * rank;
heapSize += next;
next = length * pArrayTypeDef->pArrayElementType->arrayElementSize;
next = 1;
for (int i = 0; i < rank; ++i) next *= lengths[i];
next = next * pArrayTypeDef->pArrayElementType->arrayElementSize;
heapSize += next;
pArray = (tSystemArray*)Heap_Alloc(pArrayTypeDef, heapSize);
pArray->ptr_elements = (PTR)((&(pArray->rank)) + 2);
pArray->rank = 1;
*((&(pArray->rank)) + 1) = length;
pArray->ptr_elements = (PTR)((&(pArray->rank)) + 1 + rank);
pArray->rank = rank;
for (int i = 0; i < rank; ++i)
{
*((&(pArray->rank)) + 1 + i) = *lengths;
}
return (HEAP_PTR)pArray;
}

@@ -31,7 +31,7 @@ function_space_specifier tAsyncCall* System_Array_Internal_Copy(PTR pThis_, PTR
function_space_specifier tAsyncCall* System_Array_Resize(PTR pThis_, PTR pParams, PTR pReturnValue);
function_space_specifier tAsyncCall* System_Array_Reverse(PTR pThis_, PTR pParams, PTR pReturnValue);

function_space_specifier HEAP_PTR SystemArray_NewVector(tMD_TypeDef *pArrayTypeDef, U32 length);
function_space_specifier HEAP_PTR SystemArray_NewVector(tMD_TypeDef *pArrayTypeDef, U32 rank, U32* lengths);
#define SystemArray_GetLength(pArray) (*(U32*)(pArray))
function_space_specifier void SystemArray_StoreElement(HEAP_PTR pThis_, U32 index, PTR value);
function_space_specifier void SystemArray_LoadElement(HEAP_PTR pThis_, U32 index, PTR value);
@@ -42,8 +42,9 @@ function_space_specifier tAsyncCall* System_Enum_Internal_GetInfo(PTR pThis_, PT
HEAP_PTR names, values;

// An enum type always has just one non-literal field, with all other fields being the values.
names = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_STRING], pEnumType->numFields - 1);
values = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_INT32], pEnumType->numFields - 1);
U32 v = pEnumType->numFields - 1;
names = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_STRING], 1, &v);
values = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_INT32], 1, &v);

for (i=0, retIndex=0; i<pEnumType->numFields; i++) {
tMD_FieldDef *pField = pEnumType->ppFields[i];
@@ -134,7 +134,8 @@ function_space_specifier tAsyncCall* System_RuntimeType_GetGenericArguments(PTR
}
}

ret = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_TYPE], argCount);
U32 v = argCount;
ret = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_TYPE], 1, &v);
// Allocate to return value straight away, so it cannot be GCed
*(HEAP_PTR*)pReturnValue = ret;

@@ -50,7 +50,8 @@ function_space_specifier tAsyncCall* System_ValueType_GetFields(PTR pThis_, PTR
}
}

ret = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_OBJECT], numInstanceFields << ((o2 == NULL)?0:1));
U32 v = numInstanceFields << ((o2 == NULL) ? 0 : 1);
ret = SystemArray_NewVector(_bcl_->types[TYPE_SYSTEM_ARRAY_OBJECT], 1, &v);

retOfs = 0;
for (i=0; i<pType->numFields; i++) {
@@ -295,18 +295,11 @@ global_space_specifier void Initialize_BCL2()
Finalizer_Init();
}

function_space_specifier void* Bcl_Heap_Alloc(STRING assemblyName, STRING nameSpace, STRING name)
{
tMD_TypeDef* type_def = MetaData_GetTypeDefFromFullName(assemblyName, nameSpace, name);
void * result = Heap_AllocType(type_def);
return result;
}

function_space_specifier void* Bcl_Array_Alloc(STRING assemblyName, STRING nameSpace, STRING name, int length)
function_space_specifier void* Bcl_Array_Alloc(tMD_TypeDef* element_type_def, int rank, unsigned int* lengths)
{
tMD_TypeDef* type_def = MetaData_GetTypeDefFromFullName(assemblyName, nameSpace, name);
type_def = Type_GetArrayTypeDef(type_def, NULL, NULL);
return (void*)SystemArray_NewVector(type_def, length);
tMD_TypeDef* array_type_def = Type_GetArrayTypeDef(element_type_def, NULL, NULL);
return (void*)SystemArray_NewVector(array_type_def, rank, lengths);
}

function_space_specifier int get_kernel_base_index()
@@ -304,7 +304,9 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)"</Command>
</CudaCompile>
<CudaCompile Include="..\Native\Filesystem.c" />
<CudaCompile Include="kernel.cu" />
<CudaCompile Include="_wrapper.cu" />
<CudaCompile Include="_wrapper.cu">
<Include Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">..\Native</Include>
</CudaCompile>
</ItemGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
<ImportGroup Label="ExtensionTargets">
Oops, something went wrong.

0 comments on commit 987209c

Please sign in to comment.
You can’t perform that action at this time.