Permalink
Browse files

Kernel launching works! Added Function class and cuLaunchKernel.

  • Loading branch information...
1 parent 33b4489 commit 0c93a6e4f892aebc1d187b5500a1fa05f9d493c3 Jitu Das committed Mar 24, 2012
Showing with 161 additions and 18 deletions.
  1. +7 −3 src/bindings.cpp
  2. +4 −0 src/bindings.hpp
  3. +2 −0 src/ctx.cpp
  4. +5 −0 src/ctx.hpp
  5. +2 −0 src/device.cpp
  6. +4 −0 src/device.hpp
  7. +52 −0 src/function.cpp
  8. +33 −0 src/function.hpp
  9. +20 −0 src/module.cpp
  10. +5 −0 src/module.hpp
  11. +8 −4 test/test.cu
  12. BIN test/test.cubin
  13. +10 −2 test/test.js
  14. +8 −8 test/test.ptx
  15. +1 −1 wscript
View
@@ -1,12 +1,15 @@
#include "bindings.hpp"
#include "ctx.hpp"
#include "device.hpp"
+#include "function.hpp"
#include "mem.hpp"
#include "module.hpp"
+using namespace NodeCuda;
+
void init (Handle<Object> target) {
HandleScope scope;
-
+
// Initiailze the cuda driver api
cuInit(0);
@@ -17,18 +20,19 @@ void init (Handle<Object> target) {
// Initialize driver api bindings
Ctx::Initialize(target);
Device::Initialize(target);
+ NodeCuda::Function::Initialize(target);
Mem::Initialize(target);
Module::Initialize(target);
}
-Handle<Value> GetDriverVersion(Local<String> property, const AccessorInfo &info) {
+Handle<Value> NodeCuda::GetDriverVersion(Local<String> property, const AccessorInfo &info) {
HandleScope scope;
int driverVersion = 0;
cuDriverGetVersion(&driverVersion);
return scope.Close(Integer::New(driverVersion));
}
-Handle<Value> GetDeviceCount(Local<String> property, const AccessorInfo &info) {
+Handle<Value> NodeCuda::GetDeviceCount(Local<String> property, const AccessorInfo &info) {
HandleScope scope;
int count = 0;
cuDeviceGetCount(&count);
View
@@ -7,7 +7,11 @@
using namespace v8;
using namespace node;
+namespace NodeCuda {
+
static Handle<Value> GetDriverVersion(Local<String> property, const AccessorInfo &info);
static Handle<Value> GetDeviceCount(Local<String> property, const AccessorInfo &info);
+}
+
#endif
View
@@ -1,6 +1,8 @@
#include "ctx.hpp"
#include "device.hpp"
+using namespace NodeCuda;
+
Persistent<FunctionTemplate> Ctx::constructor_template;
void Ctx::Initialize(Handle<Object> target) {
View
@@ -4,6 +4,8 @@
#include <cuda.h>
#include "bindings.hpp"
+namespace NodeCuda {
+
class Ctx : public ObjectWrap {
public:
static void Initialize(Handle<Object> target);
@@ -27,4 +29,7 @@ class Ctx : public ObjectWrap {
private:
CUcontext m_context;
};
+
+}
+
#endif
View
@@ -1,5 +1,7 @@
#include "device.hpp"
+using namespace NodeCuda;
+
Persistent<FunctionTemplate> Device::constructor_template;
void Device::Initialize(Handle<Object> target) {
View
@@ -4,6 +4,8 @@
#include <cuda.h>
#include "bindings.hpp"
+namespace NodeCuda {
+
class Device : public ObjectWrap {
public:
static void Initialize(Handle<Object> target);
@@ -29,4 +31,6 @@ class Device : public ObjectWrap {
friend class Ctx;
};
+}
+
#endif
View
@@ -0,0 +1,52 @@
+#include "function.hpp"
+#include <iostream>
+using namespace std;
+using namespace NodeCuda;
+
+Persistent<FunctionTemplate> NodeCuda::Function::constructor_template;
+
+void NodeCuda::Function::Initialize(Handle<Object> target) {
+ HandleScope scope;
+
+ Local<FunctionTemplate> t = FunctionTemplate::New(NodeCuda::Function::New);
+ constructor_template = Persistent<FunctionTemplate>::New(t);
+ constructor_template->InstanceTemplate()->SetInternalFieldCount(1);
+ constructor_template->SetClassName(String::NewSymbol("CudaFunction"));
+
+ NODE_SET_PROTOTYPE_METHOD(constructor_template, "launch", NodeCuda::Function::LaunchKernel);
+
+ // Function objects can only be created by cuModuleGetFunction
+}
+
+Handle<Value> NodeCuda::Function::New(const Arguments& args) {
+ HandleScope scope;
+
+ NodeCuda::Function *pfunction = new NodeCuda::Function();
+ pfunction->Wrap(args.This());
+
+ return args.This();
+}
+
+Handle<Value> NodeCuda::Function::LaunchKernel(const Arguments& args) {
+ HandleScope scope;
+ Function *pfunction = ObjectWrap::Unwrap<Function>(args.This());
+
+ Local<Array> gridDim = Local<Array>::Cast(args[0]);
+ unsigned int gridDimX = gridDim->Get(0)->Uint32Value();
+ unsigned int gridDimY = gridDim->Get(1)->Uint32Value();
+ unsigned int gridDimZ = gridDim->Get(2)->Uint32Value();
+
+ Local<Array> blockDim = Local<Array>::Cast(args[1]);
+ unsigned int blockDimX = blockDim->Get(0)->Uint32Value();
+ unsigned int blockDimY = blockDim->Get(1)->Uint32Value();
+ unsigned int blockDimZ = blockDim->Get(2)->Uint32Value();
+
+ cout << "Grid " << gridDimX << "," << gridDimY << "," << gridDimZ << endl;
+ cout << "Block " << blockDimX << "," << blockDimY << "," << blockDimZ << endl;
+
+ CUresult error = cuLaunchKernel(pfunction->m_function,
+ gridDimX, gridDimY, gridDimZ,
+ blockDimX, blockDimY, blockDimZ,
+ 0, 0, NULL, NULL);
+ return scope.Close(Number::New(error));
+}
View
@@ -0,0 +1,33 @@
+#ifndef FUNCTION_HPP
+#define FUNCTION_HPP
+
+#include <cuda.h>
+#include "bindings.hpp"
+#include "module.hpp"
+
+namespace NodeCuda {
+
+class Function : public ObjectWrap {
+public:
+ static void Initialize(Handle<Object> target);
+
+protected:
+ static Persistent<FunctionTemplate> constructor_template;
+
+ static Handle<Value> LaunchKernel(const Arguments& args);
+
+ Function() : ObjectWrap(), m_function(0) {}
+
+ ~Function() {}
+
+private:
+ static Handle<Value> New(const Arguments& args);
+
+ CUfunction m_function;
+
+ friend Handle<Value> Module::GetFunction(const Arguments&);
+};
+
+}
+
+#endif
View
@@ -1,4 +1,7 @@
#include "module.hpp"
+#include "function.hpp"
+
+using namespace NodeCuda;
Persistent<FunctionTemplate> Module::constructor_template;
@@ -12,6 +15,8 @@ void Module::Initialize(Handle<Object> target) {
// Module objects can only be created by load functions
NODE_SET_METHOD(target, "moduleLoad", Module::Load);
+
+ NODE_SET_PROTOTYPE_METHOD(constructor_template, "getFunction", Module::GetFunction);
}
Handle<Value> Module::New(const Arguments& args) {
@@ -37,3 +42,18 @@ Handle<Value> Module::Load(const Arguments& args) {
return scope.Close(result);
}
+Handle<Value> Module::GetFunction(const Arguments& args) {
+ HandleScope scope;
+ Local<Object> result = NodeCuda::Function::constructor_template->InstanceTemplate()->NewInstance();
+ Module *pmodule = ObjectWrap::Unwrap<Module>(args.This());
+ NodeCuda::Function *pfunction = ObjectWrap::Unwrap<NodeCuda::Function>(result);
+
+ String::AsciiValue name(args[0]);
+ CUresult error = cuModuleGetFunction(&(pfunction->m_function), pmodule->m_module, *name);
+
+ result->Set(String::New("name"), args[0]);
+ result->Set(String::New("error"), Integer::New(error));
+
+ return scope.Close(result);
+}
+
View
@@ -4,9 +4,12 @@
#include <cuda.h>
#include "bindings.hpp"
+namespace NodeCuda {
+
class Module : public ObjectWrap {
public:
static void Initialize(Handle<Object> target);
+ static Handle<Value> GetFunction(const Arguments& args);
protected:
static Persistent<FunctionTemplate> constructor_template;
@@ -23,4 +26,6 @@ class Module : public ObjectWrap {
CUmodule m_module;
};
+}
+
#endif
View
@@ -1,7 +1,11 @@
#include <cstdio>
-__global__ void helloWorld() {
-#if __CUDA_ARCH__ >= 200
- printf("Hello, world!");
-#endif
+extern "C" {
+ __global__ void helloWorld() {
+ #if __CUDA_ARCH__ >= 200
+ printf("Hello, world! I'm thread (%d,%d,%d) in block (%d,%d,%d).\n",
+ threadIdx.x, threadIdx.y, threadIdx.z,
+ blockIdx.x, blockIdx.y, blockIdx.z);
+ #endif
+ }
}
View
Binary file not shown.
View
@@ -44,8 +44,16 @@ console.log("Mem Free with error code: " + error);
//cuModuleLoad
-var error = cu.moduleLoad("test/test.ptx");
-console.log("Loaded module:", error);
+var cuModule = cu.moduleLoad("test/test.cubin");
+console.log("Loaded module:", cuModule);
+
+//cuModuleGetFunction
+var cuFunction = cuModule.getFunction("helloWorld");
+console.log("Got function:", cuFunction);
+
+//cuLaunchKernel
+var error = cuFunction.launch([3,1,1],[2,2,2]);
+console.log("Launched kernel:", error);
//cuCtxDestroy
error = cuCtx.destroy();
View
@@ -4,7 +4,7 @@
// nvopencc 4.0 built on 2011-05-12
//-----------------------------------------------------------
- // Compiling /tmp/tmpxft_000050ad_00000000-9_test.cpp3.i (/tmp/ccBI#.dEPtVw)
+ // Compiling /tmp/tmpxft_00005ddb_00000000-9_test.cpp3.i (/tmp/ccBI#.hpZn7e)
//-----------------------------------------------------------
//-----------------------------------------------------------
@@ -17,7 +17,7 @@
//-----------------------------------------------------------
.file 1 "<command-line>"
- .file 2 "/tmp/tmpxft_000050ad_00000000-8_test.cudafe2.gpu"
+ .file 2 "/tmp/tmpxft_00005ddb_00000000-8_test.cudafe2.gpu"
.file 3 "/usr/lib/gcc/x86_64-redhat-linux/4.4.5/include/stddef.h"
.file 4 "/usr/local/cuda/bin/../include/crt/device_runtime.h"
.file 5 "/usr/local/cuda/bin/../include/host_defines.h"
@@ -46,12 +46,12 @@
.file 28 "/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h"
- .entry _Z10helloWorldv
+ .entry helloWorld
{
- .loc 16 3 0
-$LDWbegin__Z10helloWorldv:
- .loc 16 7 0
+ .loc 16 4 0
+$LDWbegin_helloWorld:
+ .loc 16 8 0
exit;
-$LDWend__Z10helloWorldv:
- } // _Z10helloWorldv
+$LDWend_helloWorld:
+ } // helloWorld
View
@@ -20,5 +20,5 @@ def configure(conf):
def build(bld):
obj = bld.new_task_gen('cxx', 'shlib', 'node_addon')
obj.target = "cuda"
- obj.source = "src/bindings.cpp src/ctx.cpp src/device.cpp src/mem.cpp src/module.cpp"
+ obj.source = "src/bindings.cpp src/ctx.cpp src/device.cpp src/function.cpp src/mem.cpp src/module.cpp"
obj.uselib = "CUDA"

0 comments on commit 0c93a6e

Please sign in to comment.