private static void TestResult(CUresult r)
{
if (r != CUresult.CUDA_SUCCESS)
{
throw new ApplicationException("CUDA error: " + r.ToString());
}
}
/// <summary> see CUDA doc; </summary>
public static void ModuleGetFunction(out CUfunction func, CUmodule module, string name)
{
IntPtr _name = Marshal.StringToHGlobalAnsi(name);
CUresult res = my.cuModuleGetFunction(out func, module, _name);
Marshal.FreeHGlobal(_name);
TestResult(res);
}
/// <summary> see CUDA doc; </summary>
public static void ModuleLoad(out CUmodule module, string fname)
{
IntPtr _fname = Marshal.StringToHGlobalAnsi(fname);
CUresult res = my.cuModuleLoad(out module, _fname);
Marshal.FreeHGlobal(_fname);
TestResult(res);
}
public static void CheckCudaError(CUresult res)
{
if (res.Value != cudaError_enum.CUDA_SUCCESS)
{
IntPtr pStr = IntPtr.Zero;
Functions.cuGetErrorString(res, ref pStr);
var cuda_error = Marshal.PtrToStringAnsi(pStr);
throw new Exception("CUDA error: " + cuda_error);
}
}
public static void Check(CUresult error, string message)
{
if (error != CUresult.CUDA_SUCCESS)
{
message = string.Format("{0}: {1}", error.ToString(), message);
Exception exception = new CudaException(message);
exception.Data.Add("CUresult", error);
throw exception;
}
}
private static void testResult(CUresult r)
{
//Console.WriteLine("thr = '" + Thread.CurrentThread.Priority + "'");
//{
// StackFrame fr = new StackFrame(1, true);
// _MethodBase m = fr.GetMethod();
// string _name = m.DeclaringType.FullName + "." + m.Name;
// Console.WriteLine(_name);
//}
if (r != CUresult.CUDA_SUCCESS)
{
throw new ApplicationException("CUDA error: " + r.ToString());
}
}
public static CUmodule InitializeModule(IntPtr cubin)
{
if (cached_modules.TryGetValue(cubin, out CUmodule value))
{
return(value);
}
uint num_ops = 0;
var op = new CUjit_option[num_ops];
ulong[] op_values = new ulong[num_ops];
var op_values_link_handle = GCHandle.Alloc(op_values, GCHandleType.Pinned);
var op_values_link_intptr = op_values_link_handle.AddrOfPinnedObject();
CUresult res = Cuda.cuModuleLoadDataEx(out CUmodule module, cubin, 0, op, op_values_link_intptr);
CudaHelpers.CheckCudaError(res);
cached_modules[cubin] = module;
return(module);
}
public CudaException(CUresult error, Exception innerException)
: this(error, null, innerException)
{
}
public CudaException(CUresult error)
: this(error, null, null)
{
}
/// <summary> see CUDA doc; </summary>
public static void ModuleLoadData(out CUmodule module, IntPtr img)
{
CUresult res = my.cuModuleLoadData(out module, img);
TestResult(res);
}
public CudaException(CUresult error, Exception innerException)
: this(error, null, innerException)
{
}
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 = new CUresult(cudaError_enum.CUDA_SUCCESS);
Campy.Utils.CudaHelpers.MakeLinearTiling(1,
out Campy.Utils.CudaHelpers.dim3 tile_size, out Campy.Utils.CudaHelpers.dim3 tiles);
res = Functions.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 = Functions.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 = new CUresult(cudaError_enum.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 = Functions.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 CudaException(CUresult error, String message, Exception innerException)
: this((CudaError)error, message, innerException)
{
}
public CudaException(CUresult error, String message)
: this(error, message, null)
{
}
public CudaException(CUresult error, String message)
: this(error, message, null)
{
}
public CudaException(CUresult error, String message, Exception innerException)
: this((CudaError)error, message, innerException)
{
}
public CudaException(CUresult error)
: this(error, null, null)
{
}
