private static void SetBaseIndex(int base_index)
{
unsafe
{
IntPtr[] x1 = new IntPtr[] { new IntPtr(base_index) };
GCHandle handle2 = GCHandle.Alloc(x1, GCHandleType.Pinned);
var parm1 = handle2.AddrOfPinnedObject();
IntPtr[] kp = new IntPtr[] { parm1 };
CUmodule module = RUNTIME.RuntimeModule;
CudaHelpers.CheckCudaError(Cuda.cuModuleGetFunction(out CUfunction function, module, "_Z21set_kernel_base_indexi"));
Campy.Utils.CudaHelpers.MakeLinearTiling(1,
out Campy.Utils.CudaHelpers.dim3 tile_size,
out Campy.Utils.CudaHelpers.dim3 tiles);
CUresult res;
fixed(IntPtr *kernelParams = kp)
{
res = Cuda.cuLaunchKernel(
function,
tiles.x, tiles.y, tiles.z, // grid has one block.
tile_size.x, tile_size.y, tile_size.z, // n threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
Utils.CudaHelpers.CheckCudaError(res);
res = Cuda.cuCtxSynchronize(); // Make sure it's copied back to host.
Utils.CudaHelpers.CheckCudaError(res);
}
}
public static void For(int number_of_threads, SimpleKernel simpleKernel)
{
if (Campy.Utils.Options.IsOn("import-only"))
{
JustImport(simpleKernel);
return;
}
GCHandle handle1 = default(GCHandle);
GCHandle handle2 = default(GCHandle);
try
{
unsafe
{
System.Reflection.MethodInfo method_info = simpleKernel.Method;
String kernel_assembly_file_name = method_info.DeclaringType.Assembly.Location;
Mono.Cecil.ModuleDefinition md = Campy.Meta.StickyReadMod.StickyReadModule(
kernel_assembly_file_name, new ReaderParameters {
ReadSymbols = true
});
MethodReference method_reference = md.ImportReference(method_info);
CUfunction ptr_to_kernel = default(CUfunction);
CUmodule module = default(CUmodule);
Campy.Utils.TimePhase.Time("compile ", () =>
{
IntPtr image = Singleton._compiler.Compile(method_reference, simpleKernel.Target);
module = Singleton._compiler.SetModule(method_reference, image);
Singleton._compiler.StoreJits(module);
ptr_to_kernel = Singleton._compiler.GetCudaFunction(method_reference, module);
});
RUNTIME.BclCheckHeap();
BUFFERS buffer = Singleton.Buffer;
IntPtr kernel_target_object = IntPtr.Zero;
Campy.Utils.TimePhase.Time("deep copy ", () =>
{
int count = simpleKernel.Method.GetParameters().Length;
var bb = Singleton._compiler.GetBasicBlock(method_reference);
if (bb.HasThis)
{
count++;
}
if (!(count == 1 || count == 2))
{
throw new Exception("Expecting at least one parameter for kernel.");
}
if (bb.HasThis)
{
kernel_target_object = buffer.AddDataStructure(simpleKernel.Target);
}
});
Campy.Utils.TimePhase.Time("kernel cctor set up", () =>
{
// For each cctor, run on GPU.
// Construct dependency graph of methods.
List <MethodReference> order_list = COMPILER.Singleton.ConstructCctorOrder();
// Finally, call cctors.
foreach (var bb in order_list)
{
if (Campy.Utils.Options.IsOn("trace-cctors"))
{
System.Console.WriteLine("Executing cctor "
+ bb.FullName);
}
var cctor = Singleton._compiler.GetCudaFunction(bb, module);
var res = CUresult.CUDA_SUCCESS;
Campy.Utils.CudaHelpers.MakeLinearTiling(1,
out Campy.Utils.CudaHelpers.dim3 tile_size, out Campy.Utils.CudaHelpers.dim3 tiles);
res = Cuda.cuLaunchKernel(
cctor,
tiles.x, tiles.y, tiles.z, // grid has one block.
tile_size.x, tile_size.y, tile_size.z, // n threads.
0, // no shared memory
default(CUstream),
(IntPtr)IntPtr.Zero,
(IntPtr)IntPtr.Zero
);
CudaHelpers.CheckCudaError(res);
res = Cuda.cuCtxSynchronize(); // Make sure it's copied back to host.
CudaHelpers.CheckCudaError(res);
}
});
if (Campy.Utils.Options.IsOn("trace-cctors"))
{
System.Console.WriteLine("Done with cctors");
}
Campy.Utils.TimePhase.Time("kernel call ", () =>
{
IntPtr[] parm1 = new IntPtr[1];
IntPtr[] parm2 = new IntPtr[1];
parm1[0] = kernel_target_object;
parm2[0] = buffer.New(BUFFERS.SizeOf(typeof(int)));
IntPtr[] x1 = parm1;
handle1 = GCHandle.Alloc(x1, GCHandleType.Pinned);
IntPtr pointer1 = handle1.AddrOfPinnedObject();
IntPtr[] x2 = parm2;
handle2 = GCHandle.Alloc(x2, GCHandleType.Pinned);
IntPtr pointer2 = handle2.AddrOfPinnedObject();
IntPtr[] kp = new IntPtr[] { pointer1, pointer2 };
var res = CUresult.CUDA_SUCCESS;
fixed(IntPtr * kernelParams = kp)
{
Campy.Utils.CudaHelpers.MakeLinearTiling(number_of_threads,
out Campy.Utils.CudaHelpers.dim3 tile_size, out Campy.Utils.CudaHelpers.dim3 tiles);
//MakeLinearTiling(1, out dim3 tile_size, out dim3 tiles);
res = Cuda.cuLaunchKernel(
ptr_to_kernel,
tiles.x, tiles.y, tiles.z, // grid has one block.
tile_size.x, tile_size.y, tile_size.z, // n threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
public static void For(int number_of_threads, SimpleKernel simpleKernel)
{
GCHandle handle1 = default(GCHandle);
GCHandle handle2 = default(GCHandle);
try
{
unsafe
{
//////// COMPILE KERNEL INTO GPU CODE ///////
/////////////////////////////////////////////
var stopwatch_cuda_compile = new Stopwatch();
stopwatch_cuda_compile.Start();
IntPtr image = Singleton()._converter.Compile(simpleKernel.Method, simpleKernel.Target);
CUfunction ptr_to_kernel = Singleton()._converter.GetCudaFunction(simpleKernel.Method, image);
var elapse_cuda_compile = stopwatch_cuda_compile.Elapsed;
RUNTIME.CheckHeap();
//////// COPY DATA INTO GPU /////////////////
/////////////////////////////////////////////
var stopwatch_deep_copy_to = new Stopwatch();
stopwatch_deep_copy_to.Reset();
stopwatch_deep_copy_to.Start();
BUFFERS buffer = Singleton().Buffer;
// Set up parameters.
int count = simpleKernel.Method.GetParameters().Length;
var bb = Singleton()._converter.GetBasicBlock(simpleKernel.Method);
if (bb.HasThis)
{
count++;
}
if (!(count == 1 || count == 2))
{
throw new Exception("Expecting at least one parameter for kernel.");
}
IntPtr[] parm1 = new IntPtr[1];
IntPtr[] parm2 = new IntPtr[1];
IntPtr ptr = IntPtr.Zero;
// The method really should have a "this" because it's a closure
// object.
if (bb.HasThis)
{
RUNTIME.CheckHeap();
ptr = buffer.AddDataStructure(simpleKernel.Target);
parm1[0] = ptr;
}
{
Type btype = typeof(int);
var s = BUFFERS.SizeOf(btype);
var ptr2 = buffer.New(s);
// buffer.DeepCopyToImplementation(index, ptr2);
parm2[0] = ptr2;
}
stopwatch_deep_copy_to.Start();
var elapse_deep_copy_to = stopwatch_cuda_compile.Elapsed;
var stopwatch_call_kernel = new Stopwatch();
stopwatch_call_kernel.Reset();
stopwatch_call_kernel.Start();
IntPtr[] x1 = parm1;
handle1 = GCHandle.Alloc(x1, GCHandleType.Pinned);
IntPtr pointer1 = handle1.AddrOfPinnedObject();
IntPtr[] x2 = parm2;
handle2 = GCHandle.Alloc(x2, GCHandleType.Pinned);
IntPtr pointer2 = handle2.AddrOfPinnedObject();
RUNTIME.CheckHeap();
IntPtr[] kp = new IntPtr[] { pointer1, pointer2 };
var res = CUresult.CUDA_SUCCESS;
fixed(IntPtr *kernelParams = kp)
{
Campy.Utils.CudaHelpers.MakeLinearTiling(number_of_threads, out Campy.Utils.CudaHelpers.dim3 tile_size, out Campy.Utils.CudaHelpers.dim3 tiles);
//MakeLinearTiling(1, out dim3 tile_size, out dim3 tiles);
res = Cuda.cuLaunchKernel(
ptr_to_kernel,
tiles.x, tiles.y, tiles.z, // grid has one block.
tile_size.x, tile_size.y, tile_size.z, // n threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
CudaHelpers.CheckCudaError(res);
res = Cuda.cuCtxSynchronize(); // Make sure it's copied back to host.
CudaHelpers.CheckCudaError(res);
stopwatch_call_kernel.Stop();
var elapse_call_kernel = stopwatch_call_kernel.Elapsed;
if (Campy.Utils.Options.IsOn("jit_trace"))
{
System.Console.WriteLine("cuda compile " + elapse_cuda_compile);
System.Console.WriteLine("deep copy in " + elapse_deep_copy_to);
System.Console.WriteLine("cuda kernel " + elapse_call_kernel);
}
{
var stopwatch_deep_copy_back = new Stopwatch();
stopwatch_deep_copy_back.Reset();
RUNTIME.CheckHeap();
stopwatch_deep_copy_back.Start();
buffer.SynchDataStructures();
stopwatch_deep_copy_back.Stop();
RUNTIME.CheckHeap();
var elapse_deep_copy_back = stopwatch_deep_copy_back.Elapsed;
if (Campy.Utils.Options.IsOn("jit_trace"))
{
System.Console.WriteLine("deep copy out " + elapse_deep_copy_back);
}
}
}
}
catch (Exception e)
{
Console.WriteLine(e);
throw e;
}
finally
{
if (default(GCHandle) != handle1)
{
handle1.Free();
}
if (default(GCHandle) != handle2)
{
handle2.Free();
}
}
}
static unsafe void Part1()
{
Cuda.cuInit(0);
var res = Cuda.cuDeviceGet(out int device, 0);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetPCIBusId(out string pciBusId, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetName(out string name, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuCtxCreate_v2(out CUcontext cuContext, 0, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
string assembly_directory = AssemblyDirectory;
StreamReader sr = new StreamReader(assembly_directory + @"/../../Project2/x64/Debug/stuff.ptx");
string kernel = sr.ReadToEnd();
IntPtr ptr = Marshal.StringToHGlobalAnsi(kernel);
int[] option_values = new int[] { };
GCHandle option_values_handle = GCHandle.Alloc(option_values, GCHandleType.Pinned);
IntPtr options_values_ptr = option_values_handle.AddrOfPinnedObject();
CUjit_option[] options = new CUjit_option[] { };
res = Cuda.cuModuleLoadDataEx(out CUmodule cuModule, ptr, (uint)0, options, options_values_ptr);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuModuleGetFunction(out CUfunction helloWorld, cuModule, "_Z5helloPc");
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
byte[] v = { (byte)'G', (byte)'d', (byte)'k', (byte)'k', (byte)'n', 31, (byte)'v', (byte)'n', (byte)'q', (byte)'k', (byte)'c', 0 };
GCHandle handle = GCHandle.Alloc(v, GCHandleType.Pinned);
IntPtr pointer = IntPtr.Zero;
pointer = handle.AddrOfPinnedObject();
res = Cuda.cuMemAlloc_v2(out IntPtr dptr, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyHtoD_v2(dptr, pointer, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
IntPtr[] x = new IntPtr[] { dptr };
GCHandle handle2 = GCHandle.Alloc(x, GCHandleType.Pinned);
IntPtr pointer2 = IntPtr.Zero;
pointer2 = handle2.AddrOfPinnedObject();
IntPtr[] kp = new IntPtr[] { pointer2 };
fixed(IntPtr *kernelParams = kp)
{
res = Cuda.cuLaunchKernel(helloWorld,
1, 1, 1, // grid has one block.
11, 1, 1, // block has 11 threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyDtoH_v2(pointer, dptr, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
for (int i = 0; i < 11; ++i)
{
System.Console.Write((char)v[i]);
}
System.Console.WriteLine();
Cuda.cuCtxDestroy_v2(cuContext);
}
static unsafe void Part2()
{
Cuda.cuInit(0);
var res = Cuda.cuDeviceGet(out int device, 0);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetPCIBusId(out string pciBusId, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetName(out string name, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuCtxCreate_v2(out CUcontext cuContext, 0, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
string assembly_directory = AssemblyDirectory;
FileStream fs = new FileStream(assembly_directory + @"/../../Project2/x64/Debug/stuff.cu.obj", FileMode.Open);
var len = fs.Length;
var gpu_obj = new byte[len];
fs.Read(gpu_obj, 0, (int)len);
uint num_ops_link = 5;
var op_link = new CUjit_option[num_ops_link];
ulong[] op_values_link = new ulong[num_ops_link];
int size = 1024 * 100;
op_link[0] = CUjit_option.CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
op_values_link[0] = (ulong)size;
op_link[1] = CUjit_option.CU_JIT_INFO_LOG_BUFFER;
byte[] info_log_buffer = new byte[size];
var info_log_buffer_handle = GCHandle.Alloc(info_log_buffer, GCHandleType.Pinned);
var info_log_buffer_intptr = info_log_buffer_handle.AddrOfPinnedObject();
op_values_link[1] = (ulong)info_log_buffer_intptr;
op_link[2] = CUjit_option.CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
op_values_link[2] = (ulong)size;
op_link[3] = CUjit_option.CU_JIT_ERROR_LOG_BUFFER;
byte[] error_log_buffer = new byte[size];
var error_log_buffer_handle = GCHandle.Alloc(error_log_buffer, GCHandleType.Pinned);
var error_log_buffer_intptr = error_log_buffer_handle.AddrOfPinnedObject();
op_values_link[3] = (ulong)error_log_buffer_intptr;
op_link[4] = CUjit_option.CU_JIT_LOG_VERBOSE;
op_values_link[4] = (ulong)1;
var op_values_link_handle = GCHandle.Alloc(op_values_link, GCHandleType.Pinned);
var op_values_link_intptr = op_values_link_handle.AddrOfPinnedObject();
res = Cuda.cuLinkCreate_v2(num_ops_link, op_link, op_values_link_intptr, out CUlinkState linkState);
{
string info = Marshal.PtrToStringAnsi(info_log_buffer_intptr);
System.Console.WriteLine(info);
string error = Marshal.PtrToStringAnsi(error_log_buffer_intptr);
System.Console.WriteLine(error);
}
uint num_ops = 0;
CUjit_option[] op = new CUjit_option[0];
ulong[] op_values = new ulong[0];
var op_values_handle = GCHandle.Alloc(op_values, GCHandleType.Pinned);
var op_values_intptr = op_values_handle.AddrOfPinnedObject();
var kernel_handle = GCHandle.Alloc(gpu_obj, GCHandleType.Pinned);
IntPtr gpu_bcl_obj_intptr = kernel_handle.AddrOfPinnedObject();
res = Cuda.cuLinkAddData_v2(linkState, CUjitInputType.CU_JIT_INPUT_OBJECT,
gpu_bcl_obj_intptr, (uint)len,
"", num_ops, op, op_values_intptr);
{
string info = Marshal.PtrToStringAnsi(info_log_buffer_intptr);
System.Console.WriteLine(info);
string error = Marshal.PtrToStringAnsi(error_log_buffer_intptr);
System.Console.WriteLine(error);
}
IntPtr image;
res = Cuda.cuLinkComplete(linkState, out image, out ulong sz);
res = Cuda.cuModuleLoadDataEx(out CUmodule module, image, 0, op, op_values_link_intptr);
res = Cuda.cuModuleGetFunction(out CUfunction helloWorld, module, "_Z5helloPc");
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
byte[] v = { (byte)'G', (byte)'d', (byte)'k', (byte)'k', (byte)'n', 31, (byte)'v', (byte)'n', (byte)'q', (byte)'k', (byte)'c', 0 };
GCHandle handle = GCHandle.Alloc(v, GCHandleType.Pinned);
IntPtr pointer = IntPtr.Zero;
pointer = handle.AddrOfPinnedObject();
res = Cuda.cuMemAlloc_v2(out IntPtr dptr, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyHtoD_v2(dptr, pointer, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
IntPtr[] x = new IntPtr[] { dptr };
GCHandle handle2 = GCHandle.Alloc(x, GCHandleType.Pinned);
IntPtr pointer2 = IntPtr.Zero;
pointer2 = handle2.AddrOfPinnedObject();
IntPtr[] kp = new IntPtr[] { pointer2 };
fixed(IntPtr *kernelParams = kp)
{
res = Cuda.cuLaunchKernel(helloWorld,
1, 1, 1, // grid has one block.
11, 1, 1, // block has 11 threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyDtoH_v2(pointer, dptr, 11 * sizeof(byte));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
for (int i = 0; i < 11; ++i)
{
System.Console.Write((char)v[i]);
}
System.Console.WriteLine();
Cuda.cuCtxDestroy_v2(cuContext);
}
static unsafe void Main(string[] args)
{
System.Console.WriteLine("1");
Cuda.cuInit(0);
System.Console.WriteLine("2");
// Device api.
var res = Cuda.cuDeviceGet(out int device, 0);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetPCIBusId(out string pciBusId, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuDeviceGetName(out string name, 100, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
System.Console.WriteLine("3");
res = Cuda.cuCtxCreate_v2(out CUcontext cuContext, 0, device);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
string cu_kernel = @"
#include <stdio.h>
__global__
void kern(int * ar)
{
int i = threadIdx.x;
if (i < 11)
ar[i] = ar[i] + 1;
}
";
string compile_string = @"
nvcc --ptx --gpu-architecture=sm_20 -ccbin ""C:\Program Files(x86)\Microsoft Visual Studio 14.0\VC\bin"" y.cu";
string kernel = @"
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21373419
// Cuda compilation tools, release 8.0, V8.0.55
// Based on LLVM 3.4svn
//
.version 5.0
.target sm_20
.address_size 64
// .globl _Z4kernPi
.visible .entry _Z4kernPi(
.param .u64 _Z4kernPi_param_0
)
{
.reg .pred %p<2>;
.reg .b32 %r<4>;
.reg .b64 %rd<5>;
ld.param.u64 %rd1, [_Z4kernPi_param_0];
mov.u32 %r1, %tid.x;
setp.gt.s32 %p1, %r1, 10;
@%p1 bra BB0_2;
cvta.to.global.u64 %rd2, %rd1;
mul.wide.s32 %rd3, %r1, 4;
add.s64 %rd4, %rd2, %rd3;
ld.global.u32 %r2, [%rd4];
add.s32 %r3, %r2, 1;
st.global.u32 [%rd4], %r3;
BB0_2:
ret;
}
";
IntPtr ptr = Marshal.StringToHGlobalAnsi(kernel);
res = Cuda.cuModuleLoadData(out CUmodule cuModule, ptr);
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuModuleGetFunction(out CUfunction helloWorld, cuModule, "_Z4kernPi");
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
int[] v = { 'G', 'd', 'k', 'k', 'n', (char)31, 'v', 'n', 'q', 'k', 'c' };
GCHandle handle = GCHandle.Alloc(v, GCHandleType.Pinned);
IntPtr pointer = IntPtr.Zero;
pointer = handle.AddrOfPinnedObject();
res = Cuda.cuMemAlloc_v2(out IntPtr dptr, 11 * sizeof(int));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyHtoD_v2(dptr, pointer, 11 * sizeof(int));
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
IntPtr[] x = new IntPtr[] { dptr };
GCHandle handle2 = GCHandle.Alloc(x, GCHandleType.Pinned);
IntPtr pointer2 = IntPtr.Zero;
pointer2 = handle2.AddrOfPinnedObject();
IntPtr[] kp = new IntPtr[] { pointer2 };
fixed(IntPtr *kernelParams = kp)
{
res = Cuda.cuLaunchKernel(helloWorld,
1, 1, 1, // grid has one block.
11, 1, 1, // block has 11 threads.
0, // no shared memory
default(CUstream),
(IntPtr)kernelParams,
(IntPtr)IntPtr.Zero
);
}
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
res = Cuda.cuMemcpyDtoH_v2(pointer, dptr, 11 * sizeof(int));
for (int i = 0; i < 11; ++i)
{
System.Console.Write(Convert.ToChar(v[i]));
}
System.Console.WriteLine();
if (res != CUresult.CUDA_SUCCESS)
{
throw new Exception();
}
Cuda.cuCtxDestroy_v2(cuContext);
}
