Skip to content
Browse files

Created Module class and cuModuleLoad

  • Loading branch information...
1 parent dd3633a commit 33b4489afe6ebb2eb90a708d2f3b7b2f50d80a13 Jitu Das committed
Showing with 144 additions and 10 deletions.
  1. +4 −2 src/bindings.cpp
  2. +2 −4 src/mem.cpp
  3. +2 −1 src/mem.hpp
  4. +39 −0 src/module.cpp
  5. +26 −0 src/module.hpp
  6. +7 −0 test/test.cu
  7. +6 −2 test/test.js
  8. +57 −0 test/test.ptx
  9. +1 −1 wscript
View
6 src/bindings.cpp
@@ -1,7 +1,8 @@
#include "bindings.hpp"
-#include "device.hpp"
#include "ctx.hpp"
+#include "device.hpp"
#include "mem.hpp"
+#include "module.hpp"
void init (Handle<Object> target) {
HandleScope scope;
@@ -14,9 +15,10 @@ void init (Handle<Object> target) {
target->SetAccessor(String::New("deviceCount"), GetDeviceCount);
// Initialize driver api bindings
- Device::Initialize(target);
Ctx::Initialize(target);
+ Device::Initialize(target);
Mem::Initialize(target);
+ Module::Initialize(target);
}
Handle<Value> GetDriverVersion(Local<String> property, const AccessorInfo &info) {
View
6 src/mem.cpp
@@ -34,9 +34,8 @@ Handle<Value> Mem::Alloc(const Arguments& args) {
size_t bytesize = args[0]->Uint32Value();
CUresult error = cuMemAlloc(&(pmem->m_devicePtr), bytesize);
- // TODO: Throw error if any
-
result->Set(String::New("size"), Integer::NewFromUnsigned(bytesize));
+ result->Set(String::New("error"), Integer::New(error));
return scope.Close(result);
}
@@ -52,10 +51,9 @@ Handle<Value> Mem::AllocPitch(const Arguments& args) {
size_t Height = args[1]->Uint32Value();
CUresult error = cuMemAllocPitch(&(pmem->m_devicePtr), &pPitch, WidthInBytes, Height, ElementSizeBytes);
- // TODO: Throw error if any
-
result->Set(String::New("size"), Integer::NewFromUnsigned(pPitch * Height));
result->Set(String::New("pitch"), Integer::NewFromUnsigned(pPitch));
+ result->Set(String::New("error"), Integer::New(error));
return scope.Close(result);
}
View
3 src/mem.hpp
@@ -11,7 +11,6 @@ class Mem : public ObjectWrap {
protected:
static Persistent<FunctionTemplate> constructor_template;
- static Handle<Value> New(const Arguments& args);
static Handle<Value> Alloc(const Arguments& args);
static Handle<Value> AllocPitch(const Arguments& args);
static Handle<Value> Free(const Arguments& args);
@@ -21,6 +20,8 @@ class Mem : public ObjectWrap {
~Mem() {}
private:
+ static Handle<Value> New(const Arguments& args);
+
CUdeviceptr m_devicePtr;
};
View
39 src/module.cpp
@@ -0,0 +1,39 @@
+#include "module.hpp"
+
+Persistent<FunctionTemplate> Module::constructor_template;
+
+void Module::Initialize(Handle<Object> target) {
+ HandleScope scope;
+
+ Local<FunctionTemplate> t = FunctionTemplate::New(Module::New);
+ constructor_template = Persistent<FunctionTemplate>::New(t);
+ constructor_template->InstanceTemplate()->SetInternalFieldCount(1);
+ constructor_template->SetClassName(String::NewSymbol("CudaModule"));
+
+ // Module objects can only be created by load functions
+ NODE_SET_METHOD(target, "moduleLoad", Module::Load);
+}
+
+Handle<Value> Module::New(const Arguments& args) {
+ HandleScope scope;
+
+ Module *pmem = new Module();
+ pmem->Wrap(args.This());
+
+ return args.This();
+}
+
+Handle<Value> Module::Load(const Arguments& args) {
+ HandleScope scope;
+ Local<Object> result = constructor_template->InstanceTemplate()->NewInstance();
+ Module *pmodule = ObjectWrap::Unwrap<Module>(result);
+
+ String::AsciiValue fname(args[0]);
+ CUresult error = cuModuleLoad(&(pmodule->m_module), *fname);
+
+ result->Set(String::New("fname"), args[0]);
+ result->Set(String::New("error"), Integer::New(error));
+
+ return scope.Close(result);
+}
+
View
26 src/module.hpp
@@ -0,0 +1,26 @@
+#ifndef MODULE_HPP
+#define MODULE_HPP
+
+#include <cuda.h>
+#include "bindings.hpp"
+
+class Module : public ObjectWrap {
+public:
+ static void Initialize(Handle<Object> target);
+
+protected:
+ static Persistent<FunctionTemplate> constructor_template;
+
+ static Handle<Value> Load(const Arguments& args);
+
+ Module() : ObjectWrap(), m_module(0) {}
+
+ ~Module() {}
+
+private:
+ static Handle<Value> New(const Arguments& args);
+
+ CUmodule m_module;
+};
+
+#endif
View
7 test/test.cu
@@ -0,0 +1,7 @@
+#include <cstdio>
+
+__global__ void helloWorld() {
+#if __CUDA_ARCH__ >= 200
+ printf("Hello, world!");
+#endif
+}
View
8 test/test.js
@@ -31,7 +31,7 @@ var cuMem = cu.memAlloc(100);
console.log("Allocated 100 bytes:", cuMem);
//cuMemFree
-error = cuMem.free();
+var error = cuMem.free();
console.log("Mem Free with error code: " + error);
//cuMemAllocPitch
@@ -39,10 +39,14 @@ var cuMem = cu.memAllocPitch(100, 100, 8);
console.log("Allocated 100x100 array of doubles:", cuMem);
//cuMemFree
-error = cuMem.free();
+var error = cuMem.free();
console.log("Mem Free with error code: " + error);
+//cuModuleLoad
+var error = cu.moduleLoad("test/test.ptx");
+console.log("Loaded module:", error);
+
//cuCtxDestroy
error = cuCtx.destroy();
console.log("Context destroyed with error code: " + error);
View
57 test/test.ptx
@@ -0,0 +1,57 @@
+ .version 1.4
+ .target sm_10, map_f64_to_f32
+ // compiled with /usr/local/cuda/open64/lib//be
+ // nvopencc 4.0 built on 2011-05-12
+
+ //-----------------------------------------------------------
+ // Compiling /tmp/tmpxft_000050ad_00000000-9_test.cpp3.i (/tmp/ccBI#.dEPtVw)
+ //-----------------------------------------------------------
+
+ //-----------------------------------------------------------
+ // Options:
+ //-----------------------------------------------------------
+ // Target:ptx, ISA:sm_10, Endian:little, Pointer Size:64
+ // -O3 (Optimization level)
+ // -g0 (Debug level)
+ // -m2 (Report advisories)
+ //-----------------------------------------------------------
+
+ .file 1 "<command-line>"
+ .file 2 "/tmp/tmpxft_000050ad_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"
+ .file 6 "/usr/local/cuda/bin/../include/builtin_types.h"
+ .file 7 "/usr/local/cuda/bin/../include/device_types.h"
+ .file 8 "/usr/local/cuda/bin/../include/driver_types.h"
+ .file 9 "/usr/local/cuda/bin/../include/surface_types.h"
+ .file 10 "/usr/local/cuda/bin/../include/texture_types.h"
+ .file 11 "/usr/local/cuda/bin/../include/vector_types.h"
+ .file 12 "/usr/local/cuda/bin/../include/device_launch_parameters.h"
+ .file 13 "/usr/local/cuda/bin/../include/crt/storage_class.h"
+ .file 14 "/usr/include/bits/types.h"
+ .file 15 "/usr/include/time.h"
+ .file 16 "test.cu"
+ .file 17 "/usr/local/cuda/bin/../include/common_functions.h"
+ .file 18 "/usr/local/cuda/bin/../include/math_functions.h"
+ .file 19 "/usr/local/cuda/bin/../include/math_constants.h"
+ .file 20 "/usr/local/cuda/bin/../include/device_functions.h"
+ .file 21 "/usr/local/cuda/bin/../include/sm_11_atomic_functions.h"
+ .file 22 "/usr/local/cuda/bin/../include/sm_12_atomic_functions.h"
+ .file 23 "/usr/local/cuda/bin/../include/sm_13_double_functions.h"
+ .file 24 "/usr/local/cuda/bin/../include/sm_20_atomic_functions.h"
+ .file 25 "/usr/local/cuda/bin/../include/sm_20_intrinsics.h"
+ .file 26 "/usr/local/cuda/bin/../include/surface_functions.h"
+ .file 27 "/usr/local/cuda/bin/../include/texture_fetch_functions.h"
+ .file 28 "/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h"
+
+
+ .entry _Z10helloWorldv
+ {
+ .loc 16 3 0
+$LDWbegin__Z10helloWorldv:
+ .loc 16 7 0
+ exit;
+$LDWend__Z10helloWorldv:
+ } // _Z10helloWorldv
+
View
2 wscript
@@ -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/device.cpp src/ctx.cpp src/mem.cpp"
+ obj.source = "src/bindings.cpp src/ctx.cpp src/device.cpp src/mem.cpp src/module.cpp"
obj.uselib = "CUDA"

0 comments on commit 33b4489

Please sign in to comment.
Something went wrong with that request. Please try again.