updated CUDA hybridAPI and added a CUDA hello example
authorLukasz Wesolowski <wesolwsk@illinois.edu>
Wed, 4 Jun 2008 23:14:20 +0000 (23:14 +0000)
committerLukasz Wesolowski <wesolwsk@illinois.edu>
Wed, 4 Jun 2008 23:14:20 +0000 (23:14 +0000)
19 files changed:
examples/charm++/cuda/hello/Makefile [new file with mode: 0644]
examples/charm++/cuda/hello/README [new file with mode: 0644]
examples/charm++/cuda/hello/hello.C [new file with mode: 0644]
examples/charm++/cuda/hello/hello.ci [new file with mode: 0644]
examples/charm++/cuda/hello/helloCUDA.cu [new file with mode: 0644]
examples/converse/pingpong/Makefile
examples/converse/pingpong/pingpong.C
src/arch/cuda/hybridAPI/Makefile
src/arch/cuda/hybridAPI/cuda-hybrid-api.cu
src/arch/cuda/hybridAPI/cuda-hybrid-api.h
src/arch/cuda/hybridAPI/wr.h [new file with mode: 0644]
src/arch/cuda/hybridAPI/wrqueue.cu
src/arch/cuda/hybridAPI/wrqueue.h
src/arch/net-linux-amd64-cuda/conv-mach.sh
src/arch/net-linux-amd64-cuda/special.sh
src/conv-core/convcore.c
src/conv-core/converse.h
src/scripts/Makefile
tests/charm++/pingpong/Makefile

diff --git a/examples/charm++/cuda/hello/Makefile b/examples/charm++/cuda/hello/Makefile
new file mode 100644 (file)
index 0000000..5f44014
--- /dev/null
@@ -0,0 +1,35 @@
+CHARMC=../../../../bin/charmc $(OPTS)
+
+OBJS = hello.o helloCUDA.o
+
+NVCC = /usr/local/cuda/bin/nvcc
+NVCC_FLAGS = -c -use_fast_math #-device-debug -deviceemu 
+NVCC_INC = -I/usr/local/cuda/include -I../../../../../NVIDIA_CUDA_SDK/common/inc
+
+LDLIBS += -lcuda -lcudart -lGL -lGLU -lcutil
+
+export LD_RUN_PATH = /usr/local/cuda/lib
+
+all: hello
+
+hello: $(OBJS)
+       $(CHARMC) -language charm++ -o hello $(OBJS) $(LD_LIBS)
+
+hello.decl.h: hello.ci
+       $(CHARMC)  hello.ci
+
+clean:
+       rm -f *.decl.h *.def.h conv-host *.o wr.h hello charmrun
+
+hello.o: hello.C hello.decl.h
+       $(CHARMC) -g -c hello.C
+
+helloCUDA.o: helloCUDA.cu
+       cp ../../../../src/arch/cuda/hybridAPI/wr.h .
+       $(NVCC) -g $(NVCC_FLAGS) $(NVCC_INC) helloCUDA.cu
+
+test: all
+       ./charmrun hello +p4 10
+
+bgtest: all
+       ./charmrun hello +p4 10 +x2 +y2 +z1
diff --git a/examples/charm++/cuda/hello/README b/examples/charm++/cuda/hello/README
new file mode 100644 (file)
index 0000000..59a60ed
--- /dev/null
@@ -0,0 +1,8 @@
+Hello World in Charm with CUDA using the hybrid API.
+
+This example passes a Hi message along the array from index 0 to N. Each chare 
+executes an empty kernel on the GPU when it receives the message. 
+When the kernel returns, the runtime system executes the specified callback 
+function which passes the message to the subsequent chare in the array. 
+
+
diff --git a/examples/charm++/cuda/hello/hello.C b/examples/charm++/cuda/hello/hello.C
new file mode 100644 (file)
index 0000000..60552ce
--- /dev/null
@@ -0,0 +1,78 @@
+#include <stdio.h>
+#include "hello.decl.h"
+
+/*readonly*/ CProxy_Main mainProxy;
+/*readonly*/ int nElements;
+/*readonly*/ CProxy_Hello arr; 
+
+int currentIndex; 
+
+extern void kernelSetup(); 
+
+void kernelReturn() {
+  printf("Kernel returned\n"); 
+  arr[currentIndex].SendHi(); 
+  //  sendHi(); 
+}
+
+/*mainchare*/
+class Main : public CBase_Main
+{
+public:
+  Main(CkArgMsg* m)
+  {
+    //Process command-line arguments
+    nElements=5;
+    if(m->argc >1 ) nElements=atoi(m->argv[1]);
+    delete m;
+
+    //Start the computation
+    CkPrintf("Running Hello on %d processors for %d elements\n",
+            CkNumPes(),nElements);
+    mainProxy = thisProxy;
+
+    arr = CProxy_Hello::ckNew(nElements);
+
+    currentIndex = 0; 
+
+    arr[0].SayHi();
+  };
+
+  void done(void)
+  {
+    CkPrintf("All done\n");
+    CkExit();
+  };
+};
+
+/*array [1D]*/
+class Hello : public CBase_Hello
+{
+public:
+  Hello()
+  {
+    CkPrintf("Hello %d created\n",thisIndex);
+  }
+
+  Hello(CkMigrateMessage *m) {}
+
+  void SayHi()
+  {
+    currentIndex = thisIndex; 
+    CkPrintf("Hi from element %d\n", thisIndex);
+    if (thisIndex < nElements-1)
+      kernelSetup(); 
+    else 
+      //We've been around once-- we're done.
+      mainProxy.done();
+  }
+
+  void SendHi() {
+    //Pass the hello on:
+    CkPrintf("Executing sendHi\n"); 
+    thisProxy[thisIndex+1].SayHi();
+  }
+
+};
+
+#include "hello.def.h"
diff --git a/examples/charm++/cuda/hello/hello.ci b/examples/charm++/cuda/hello/hello.ci
new file mode 100644 (file)
index 0000000..5ce047b
--- /dev/null
@@ -0,0 +1,16 @@
+mainmodule hello {
+  readonly CProxy_Main mainProxy;
+  readonly int nElements;
+  readonly CProxy_Hello arr;
+
+  mainchare Main {
+    entry Main(CkArgMsg *m);
+    entry void done(void);
+  };
+
+  array [1D] Hello {
+    entry Hello(void);
+    entry void SayHi(void);
+    entry void SendHi(void); 
+  };           
+};
diff --git a/examples/charm++/cuda/hello/helloCUDA.cu b/examples/charm++/cuda/hello/helloCUDA.cu
new file mode 100644 (file)
index 0000000..895d90f
--- /dev/null
@@ -0,0 +1,53 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include "wr.h"
+
+extern workRequestQueue* wrQueue; 
+extern void kernelReturn(); 
+
+__global__ void helloKernel() { 
+
+}
+
+void kernelSetup() {
+  workRequest *wr; 
+  wr = (workRequest*) malloc(sizeof(workRequest)); 
+
+  wr->dimGrid.x = 1; 
+  wr->dimBlock.x = 1; 
+  wr->smemSize = 0;
+  
+  wr->readWriteDevicePtr = NULL;
+  wr->readWriteHostPtr = NULL; 
+  wr->readWriteLen = 0; 
+
+  wr->readOnlyDevicePtr = NULL;
+  wr->readOnlyHostPtr = NULL;
+  wr->readOnlyLen = 0; 
+
+  wr->writeOnlyDevicePtr = NULL;
+  wr->writeOnlyHostPtr = NULL;
+  wr->writeOnlyLen = 0; 
+
+  wr->callbackFn = kernelReturn; 
+
+  wr->id = 0; 
+
+  wr->executing = 0; 
+
+  enqueue(wrQueue, wr); 
+
+}
+
+void kernelSelect(workRequest *wr) {
+  printf("inside kernelSelect\n"); 
+  switch (wr->id) {
+  case 0: 
+    printf("calling kernel\n"); 
+    helloKernel<<<wr->dimGrid,wr->dimBlock,wr->smemSize>>>();
+    break;
+  default:
+    printf("error: id %d not valid\n", wr->id); 
+    break; 
+  }
+}
index a08c1fe4ec2764ebdf71ac0147f1d81bedd79044..87422190f2b70e7b4cd0c64a0e4a0c753e9369b3 100644 (file)
@@ -3,10 +3,10 @@ CHARMC=../../../bin/charmc $(OPTS)
 all: pingpong
 
 pingpong: pingpong.o
-       $(CHARMC) -language converse++ -o pingpong pingpong.o
+       $(CHARMC) -language converse++ -pg -o pingpong pingpong.o
 
 pingpong.o: pingpong.C
-       $(CHARMC) -language converse++ -c pingpong.C
+       $(CHARMC) -language converse++ -pg -c pingpong.C
 
 test: pingpong
        ./charmrun ./pingpong +p2 $(TESTOPTS)
index e7977296f70006d9c0c00ba7d9ca0925306630dc..a56b5011d24f6aa6a6ec2140ce5da3c9d2e0fe5f 100644 (file)
@@ -82,7 +82,7 @@ CmiHandler node0HandlerFunc(char *msg)
         CmiSetHandler(msg,CpvAccess(node1Handler));
         *((int *)(msg+CmiMsgHeaderSizeBytes)) = CpvAccess(msgSize);
         
-        CmiSyncSendAndFree(1,CpvAccess(msgSize),msg);
+        CmiSyncSendAndFree(0,CpvAccess(msgSize),msg);
     }
     return 0;
 }
index a05e6b2a7e499246cbcba4579317802a0f35a930..e3ebd852142dba1668f7a90607e8332a7ec02bd3 100644 (file)
@@ -1,25 +1,30 @@
 CHARMDIR=~/charm
 CHARMC=$(CHARMDIR)/bin/charmc
 
-#LDFLAGS += -L/usr/local/cuda/lib -L~/NVIDIA_CUDA_SDK/lib
-#LDLIBS += -lcuda -lcudart -lGL -lGLU -lcutil
-
-
 NVCC = /usr/local/cuda/bin/nvcc
-NVCC_FLAGS = -c -use_fast_math #-device-debug -deviceemu
-NVCC_INC = -I/usr/local/cuda/include -I/~/NVIDIA_CUDA_SDK/common/inc
+NVCC_FLAGS = -g -c -use_fast_math #-device-debug -deviceemu
+NVCC_INC = -I/usr/local/cuda/include -I~/NVIDIA_CUDA_SDK/common/inc
 
 RM = rm -f
 
-install: cudahybridapi.a
-       cp cudahybridapi.a $(CHARMDIR)/lib
+all: libs
+       cp libcudahybridapi.a $(CHARMDIR)/lib
+
+libs: libcudahybridapi.a
+
+install: libcudahybridapi.a
+       cp libcudahybridapi.a $(CHARMDIR)/lib
+#      cp hybridapi.o wrqueue.o $(CHARMDIR)/tmp
 
-cudahybridapi.a: hybridapi.o wrqueue.o
-       ar q cudahybridapi.a hybridapi.o wrqueue.o 
+libcudahybridapi.a: hybridapi.o wrqueue.o
+       -rm $@
+       ar q $@ hybridapi.o wrqueue.o 
 
-hybridapi.o: cuda-hybrid-api.h cuda-hybrid-api.cu
+hybridapi.o: cuda-hybrid-api.cu cuda-hybrid-api.h
        $(NVCC) $(NVCC_FLAGS) $(NVCC_INC) -o hybridapi.o cuda-hybrid-api.cu
 
-wrqueue.o: wrqueue.cu wrqueue.h
+wrqueue.o: wrqueue.cu wrqueue.h wr.h 
        $(NVCC) $(NVCC_FLAGS) $(NVCC_INC) -o wrqueue.o wrqueue.cu 
 
+clean:
+       rm -f *.o *.a 
index ee4fb9916bb54fb99331af4b5d9ed951cc6c85d2..d58a28148e15597dc1194faf83258d165fef0bc5 100644 (file)
@@ -17,7 +17,7 @@
 #include "wrqueue.h"
 #include "cuda-hybrid-api.h"
 
-workRequestQueue *wrQueue; 
+workRequestQueue *wrQueue = NULL
 
 /*
   TO DO
@@ -26,6 +26,8 @@ workRequestQueue *wrQueue;
   stream 3 - memory copies
 */
 
+/* setupMemory
+   set up memory on the gpu for this kernel's execution */
 void setupMemory(workRequest *wr) {
 
   cudaMalloc((void **)&(wr->readWriteDevicePtr), wr->readWriteLen);
@@ -38,6 +40,8 @@ void setupMemory(workRequest *wr) {
                  cudaMemcpyHostToDevice); 
 } 
 
+/* cleanupMemory
+   free memory no longer needed on the gpu */ 
 void cleanupMemory(workRequest *wr) {
 
   cudaMemcpy(wr->readWriteHostPtr, wr->readWriteDevicePtr, wr->readWriteLen, cudaMemcpyDeviceToHost); 
@@ -49,24 +53,38 @@ void cleanupMemory(workRequest *wr) {
 
 }
 
+/* kernelSelect
+   a switch statement defined by the user to allow the library to execute
+   the correct kernel */ 
 void kernelSelect(workRequest *wr);
 
+/* initHybridAPI
+   initializes the work request queue
+*/
 void initHybridAPI() {
-  init_wrqueue(wrQueue); 
+  initWRqueue(&wrQueue); 
 }
 
+/* gpuProgressFn
+   called periodically to check if the current kernel has completed,
+   and invoke subsequent kernel */
 void gpuProgressFn() {
-  while (!isEmpty(wrQueue)) {
+  if (wrQueue == NULL) {
+    return; 
+  }
 
+  while (!isEmpty(wrQueue)) {
     workRequest *wr = head(wrQueue); 
     
     if (wr->executing == 0) {
       setupMemory(wr); 
       kernelSelect(wr); 
-      cudaEventRecord(wr->completionEvent, 0); 
+      // cudaEventRecord(wr->completionEvent, 0);
+      wr->executing = 1; 
       return; 
     }  
-    else if (cudaEventQuery(wr->completionEvent) == cudaSuccess ) {      
+    // else if (cudaEventQuery(wr->completionEvent) == cudaSuccess ) {      
+    else if (cudaStreamQuery(0) == cudaSuccess ) {      
       cleanupMemory(wr);
       dequeue(wrQueue);
       wr->callbackFn();
@@ -75,6 +93,9 @@ void gpuProgressFn() {
   }
 }
 
+/* exitHybridAPI
+   cleans up and deletes memory allocated for the queue
+*/
 void exitHybridAPI() {
-  delete_wrqueue(wrQueue); 
+  deleteWRqueue(wrQueue); 
 }
index bcc9f6c995cd8ed681d7d20a368829f744605b11..cb47688c18e86190a228895e2c7c68b077ed59d9 100644 (file)
 #ifndef __CUDA_HYBRID_API_H__
 #define __CUDA_HYBRID_API_H__
 
-//#include "wrqueue.h"
+#ifdef __cplusplus
+extern "C" {
+#endif
 
-/* initAPI
+/* initHybridAPI
    initializes the work request queue
 */
 void initHybridAPI(); 
@@ -29,22 +31,13 @@ void initHybridAPI();
    and invoke subsequent kernel */
 void gpuProgressFn();
 
-/* setupMemory
-   set up memory on the gpu for this kernel's execution */
-//void setupMemory(workRequest *wr); 
-
-/* cleanupMemory
-   free memory no longer needed on the gpu */ 
-//void cleanupMemory(workRequest *wr); 
-
-/* kernelSelect
-   a switch statement defined by user to allow the library to execute
-   the correct kernel */ 
-//void kernelSelect(workRequest *wr);
-
-/* exitAPI
+/* exitHybridAPI
    cleans up and deletes memory allocated for the queue
 */
 void exitHybridAPI(); 
 
+#ifdef __cplusplus
+}
+#endif
+
 #endif
diff --git a/src/arch/cuda/hybridAPI/wr.h b/src/arch/cuda/hybridAPI/wr.h
new file mode 100644 (file)
index 0000000..d9911ec
--- /dev/null
@@ -0,0 +1,115 @@
+/* 
+ * wr.h
+ *
+ * by Lukasz Wesolowski
+ * 06.02.2008
+ *
+ * header containing declarations needed by the user of hybridAPI
+ *
+ */
+
+#ifndef __WR_H__
+#define __WR_H__
+
+/* struct workRequest
+ * 
+ * purpose:  
+ * structure for holding information about work requests on the GPU
+ *
+ * usage model: 
+ * 1. declare a pointer to a workRequest 
+ * 2. allocate dynamic memory for the work request
+ * 3. call setupMemory to copy over the data to the GPU 
+ * 4. enqueue the work request by using addWorkRequest
+ */
+typedef struct workRequest {
+
+  /* parameters for kernel execution */
+
+  dim3 dimGrid; 
+  dim3 dimBlock; 
+  int smemSize;
+  
+  /* pointers to queues and their lengths on the device(gpu) and
+     host(cpu)  */
+
+  void *readWriteDevicePtr;
+  void *readWriteHostPtr; 
+  int readWriteLen; 
+
+  void *readOnlyDevicePtr; 
+  void *readOnlyHostPtr; 
+  int readOnlyLen; 
+
+  void *writeOnlyDevicePtr;
+  void *writeOnlyHostPtr; 
+  int writeOnlyLen; 
+
+  /* to be called after the kernel finishes executing on the GPU */ 
+
+  void (*callbackFn)(); 
+
+  /* id to select the correct kernel in kernelSelect */
+
+  int id; 
+
+  /* event which will be polled to check if kernel has finished
+     execution */
+
+  cudaEvent_t completionEvent;  
+
+  /* flags */
+
+  int executing; 
+
+} workRequest; 
+
+
+/* struct workRequestQueue 
+ *
+ * purpose: container/mechanism for GPU work requests 
+ *
+ * usage model: 
+ * 1. declare a workRequestQueue
+ * 2. call init to allocate memory for the queue and initialize
+ *    bookkeeping variables
+ * 3. call enqueue for each request which needs to be 
+ *    executed on the GPU
+ * 4. in the hybrid API gpuProgressFn will execute periodically to
+ *    handle the details of executing the work request on the GPU
+ *             
+ * implementation notes: 
+ * the queue is implemented using a circular array; if the array fills
+ * up, requests are transferred to a queue having additional
+ * QUEUE_EXPANSION_SIZE slots, and the memory for the old queue is freed
+ */
+typedef struct {
+
+  /* array of requests */
+  workRequest* requests; 
+
+  /* array index for the logically first item in the queue */
+  int head; 
+
+  /* array index for the last item in the queue */ 
+  int tail; 
+
+  /* number of work requests in the queue */
+  int size; 
+
+  /* size of the array of work requests */ 
+  int capacity; 
+
+} workRequestQueue; 
+
+/* enqueue
+ *
+ * add a work request to the queue to be later executed on the GPU
+ *
+ */
+void enqueue(workRequestQueue *q, workRequest *wr); 
+
+
+#endif
+
+
index 626d4e38e43c833551d9067d1599f118ae86615f..d35c19f9d7157eeddfaa12b971fded0bd34c63fb 100644 (file)
 
 #include "wrqueue.h"
 
-void init_wrqueue(workRequestQueue *q) {
+void initWRqueue(workRequestQueue **qptr) {
 
-  q = (workRequestQueue*) malloc(sizeof(workRequestQueue));  
+  (*qptr) = (workRequestQueue*) malloc(sizeof(workRequestQueue));  
 
-  q->head = -1
-  q->tail = -1;
-  q->size = 0; 
-  q->capacity = QUEUE_SIZE_INIT; 
+  (*qptr)->head = 0
+  (*qptr)->tail = -1;
+  (*qptr)->size = 0; 
+  (*qptr)->capacity = QUEUE_SIZE_INIT; 
 
-  q->requests = (workRequest *) malloc(QUEUE_SIZE_INIT * sizeof(workRequest)); 
+  (*qptr)->requests = (workRequest *) malloc(QUEUE_SIZE_INIT * sizeof(workRequest)); 
 
 }
 
@@ -64,7 +64,7 @@ void enqueue(workRequestQueue *q, workRequest *wr) {
     q->tail = 0; 
   }
 
-  memcpy(&q->requests[q->tail], wr, sizeof(workRequest));
+  memcpy(&q->requests[q->tail], &(*wr), sizeof(workRequest));
   free(wr); 
 
   q->size++; 
@@ -78,7 +78,7 @@ void dequeue(workRequestQueue *q) {
   q->size--; 
 }
 
-int delete_wrqueue(workRequestQueue *q) {
+int deleteWRqueue(workRequestQueue *q) {
   if (q->size != 0) {
     return -1; 
   }
index b85f5b4b731bd17f292e36bf766561a91dc77cbc..08700f681ebe6728c728917a2056b3dbb7f34bf5 100644 (file)
@@ -11,6 +11,8 @@
 #ifndef __WR_QUEUE_H__
 #define __WR_QUEUE_H__
 
+#include "wr.h"
+
 /* number of work requests the queue is initialized to handle */
 #define QUEUE_SIZE_INIT 10
 
    units */ 
 #define QUEUE_EXPANSION_SIZE 10
 
-/* struct workRequest
- * 
- * purpose:  
- * structure for holding information about work requests on the GPU
- *
- * usage model: 
- * 1. declare a pointer to a workRequest 
- * 2. allocate dynamic memory for the work request
- * 3. call setupMemory to copy over the data to the GPU 
- * 4. enqueue the work request by using addWorkRequest
- */
-
-typedef struct workRequest {
-
-  /* parameters for kernel execution */
-
-  dim3 dimGrid; 
-  dim3 dimBlock; 
-  int smemSize;
-  
-  /* pointers to queues and their lengths on the device(gpu) and
-     host(cpu)  */
-
-  void *readWriteDevicePtr;
-  void *readWriteHostPtr; 
-  int readWriteLen; 
-
-  void *readOnlyDevicePtr; 
-  void *readOnlyHostPtr; 
-  int readOnlyLen; 
-
-  void *writeOnlyDevicePtr;
-  void *writeOnlyHostPtr; 
-  int writeOnlyLen; 
-
-  /* to be called after the kernel finishes executing on the GPU */ 
-
-  void (*callbackFn)(); 
-
-  /* to select the correct kernel in the switch statement */
-
-  int switchNo; 
-
-  /* event which will be polled to check if kernel has finished
-     execution */
-
-  cudaEvent_t completionEvent;  
-
-  /* flags */
-
-  int executing; 
-
-} workRequest; 
-
-
-/* struct workRequestQueue 
- *
- * purpose: container/mechanism for GPU work requests 
- *
- * usage model: 
- * 1. declare a workRequestQueue
- * 2. call init to allocate memory for the queue and initialize
- *    bookkeeping variables
- * 3. call enqueue for each request which needs to be 
- *    executed on the GPU
- * 4. in the hybrid API gpuProgressFn will execute periodically to
- *    handle the details of executing the work request on the GPU
- *             
- * implementation notes: 
- * the queue is implemented using a circular array; if the array fills
- * up, requests are transferred to a queue having additional
- * QUEUE_EXPANSION_SIZE slots, and the memory for the old queue is freed
- */
-
-typedef struct {
-
-  /* array of requests */
-  workRequest* requests; 
-
-  /* array index for the logically first item in the queue */
-  int head; 
-
-  /* array index for the last item in the queue */ 
-  int tail; 
-
-  /* number of work requests in the queue */
-  int size; 
-
-  /* size of the array of work requests */ 
-  int capacity; 
-
-} workRequestQueue; 
-
-/* init_wrqueue
+/* initWRqueue
  *
  * allocate memory for the queue and initialize bookkeeping variables
  *
  */
-void init_wrqueue(workRequestQueue *q); 
-
-/* enqueue
- *
- * add a work request to the queue to be later executed on the GPU
- *
- */
-
-void enqueue(workRequestQueue *q, workRequest *wr); 
+void initWRqueue(workRequestQueue **qptr); 
 
 /* dequeue
  *
@@ -134,7 +35,7 @@ void enqueue(workRequestQueue *q, workRequest *wr);
  */
 void dequeue(workRequestQueue *q); 
 
-/* delete_wrqueue
+/* deleteWRqueue
  *
  * if queue is nonempty, return -1  
  * if queue is empty, delete the queue, freeing dynamically allocated 
@@ -142,15 +43,13 @@ void dequeue(workRequestQueue *q);
  *
  *
  */
-
-int delete_wrqueue(workRequestQueue *q); 
+int deleteWRqueue(workRequestQueue *q); 
 
 /* head
  * 
  * returns the first element in the queue 
  *
  */
-
 workRequest * head(workRequestQueue *q);
 
 /*
@@ -160,7 +59,6 @@ workRequest * head(workRequestQueue *q);
  * 1 if queue has no pending requests stored
  * 0 otherwise
  */
-
 int isEmpty(workRequestQueue *q);
 
 #endif
index 8d58a09855f886da4b118eee35a6b05f0c196d08..074ffbb54e62f2414bc8d4944b0b10f1e05982e0 100644 (file)
@@ -11,7 +11,8 @@ CMK_LD="$CMK_CC $CMK_AMD64 "
 CMK_LDXX="$CMK_CXX $CMK_AMD64 "
 CMK_LD_SHARED="-shared"
 CMK_LD_LIBRARY_PATH="-Wl,-rpath,$CHARMLIBSO/"
-CMK_LIBS="-lckqt -lcuda -lcudasrt -lGL -lGLU -lcutil"
+CMK_LIBDIR="-L/usr/local/cuda/lib -L$HOME/NVIDIA_CUDA_SDK/lib"
+CMK_LIBS="-lckqt -lcuda -lcudart -lGL -lGLU -lcutil -lcudahybridapi"
 CMK_RANLIB="ranlib"
 
 # native compiler for compiling charmxi, etc
index 8bc93c54f8a9c36b9bc848972ed182ec3a6da811..4b6660d4ec2ea4cfb28553ec6ef0b34c412f14de 100755 (executable)
@@ -5,7 +5,7 @@ echo "---------------------- special.sh for cuda executing ----------------"
 ./gathertree ../../src/arch/cuda .
 
 # make links
-test ! -f "../include/cuda-hybrid-api.h" && ./system_ln "../tmp/hybridAPI/cuda-hybrid-api.h" ../include && test ! -f "../include/wrqueue.h" && ./system_ln "../tmp/hybridAPI/wrqueue.h" ../include
+test ! -f "../include/cuda-hybrid-api.h" && ./system_ln "../tmp/hybridAPI/cuda-hybrid-api.h" ../include && test ! -f "../include/wr.h" && ./system_ln "../tmp/hybridAPI/wr.h" ../include && test ! -f "../include/wrqueue.h" && ./system_ln "../tmp/hybridAPI/wrqueue.h" ../include
 
 #make library
 export CHARMINC=../include
index 3693f0d78c13e5bdffe27b66f614abec297b7f89..6f1cb8bf1879360eb6f63f9bcf354faeeb3a98c5 100644 (file)
@@ -107,6 +107,10 @@ extern void CmiInitCPUAffinity(char **);
 #include <sys/timeb.h>
 #endif
 
+#ifdef CMK_CUDA
+#include "cuda-hybrid-api.h"
+#endif
+
 #include "quiescence.h"
 
 int cur_restart_phase = 1;      /* checkpointing/restarting phase counter */
@@ -2837,7 +2841,7 @@ void ConverseCommonInit(char **argv)
 #endif
 
 #ifdef CMK_CUDA
-  initAPI(); 
+  initHybridAPI(); 
 #endif
 
   /* main thread is suspendable */
@@ -2872,7 +2876,7 @@ void ConverseCommonExit(void)
 #endif
 
 #if CMK_CUDA
-  exitAPI(); 
+  exitHybridAPI(); 
 #endif
 
 }
index 8812b8c78e4caaddde3baf68d9ec65bfc7e7bd4f..fc655b2103dec4f50ef6df745b7ef5c297ec5f41 100644 (file)
@@ -1588,9 +1588,6 @@ CpvExtern(char *,_validProcessors);
 #if CMK_CELL
 #include "cell-api.h"
 #endif
-#ifdef CMK_CUDA
-#include "cuda-hybrid-api.h"
-#endif
 #ifdef __ONESIDED_IMPL
 #include "conv-onesided.h"
 #endif
index ef22a894d81c47b874252b9be9dd654c853b7b57..4d30cadc36e237bead60cdb29fb189e039675af0 100644 (file)
@@ -334,8 +334,7 @@ THREADLIBS=$(L)/libthreads-default.o $(L)/libthreads-qt.o   \
 CVLIBS=$(L)/libconv-core.a \
        $(L)/libconv-cplus-y.a $(L)/libconv-cplus-n.a \
        $(L)/libconv-util.a $(L)/libconv-utilf.a \
-       $(CLBLIBS) $(TRACELIBS) $(MEMLIBS) $(THREADLIBS) \
-       $(L)/libccs-client.a
+       $(CLBLIBS) $(TRACELIBS) $(MEMLIBS) $(THREADLIBS)
 
 LIBCONV_CORE= convcore.o conv-conds.o queueing.o msgmgr.o \
        cpm.o cpthreads.o futures.o cldb.o topology.o random.o \
@@ -367,9 +366,9 @@ LIBCONV_UTILF=pup_f.o
 
 converse: basics
        $(MAKE) QuickThreads/libqt.a
-       $(MAKE) converse-target
+       $(MAKE) converse-target 
 
-converse-target:  $(CVLIBS) fmain-ok charmrun-target swapglobal-target conv-cpm
+converse-target:  $(CVLIBS) fmain-ok charmrun-target swapglobal-target conv-cpm 
 
 charmrun-target: sockRoutines.c sockRoutines.h ccs-server.c ccs-server.h sockRoutines-seq.o
        if [ -d charmrun ] ; then ( cd charmrun ; $(MAKE) OPTS='$(OPTS)' ) ; fi
@@ -401,7 +400,7 @@ CONVLIBS: converse
        cd libs; $(MAKE) convlibs OPTS='$(OPTS)'
 
 $(L)/libconv-core.a: $(LIBCONV_CORE)
-       $(CHARMC) -o $@ $(LIBCONV_CORE)
+       $(CHARMC) -fPIC -o $@ $(LIBCONV_CORE)
 
 $(L)/libconv-cplus-y.a: $(L)/libconv-cplus-n.a
 
@@ -624,7 +623,7 @@ LIBCK_CORE=trace-common.o tracef.o init.o register.o qd.o ck.o main.o  \
           NullLB.o LBSimulation.o $(COMLIB_CORE_OBJS) \
           charmProjections.o
 
-charm-target: converse $(L)/libck.a $(L)/libckf.a loadbalancers default_libs comlib_objs
+charm-target: converse $(L)/libck.a $(L)/libckf.a loadbalancers default_libs comlib_objs 
 
 CHARMLIBS: charm++
        cd libs; $(MAKE) charmlibs OPTS='$(OPTS)'
index a19b018f4afeb82f6987c12d41321aa71c8cc396..e5570e90489cbea713a026bdd39880c7aa732a1b 100644 (file)
@@ -6,7 +6,7 @@ OBJS = pingpong.o
 all:   cifiles pgm
 
 pgm: $(OBJS)
-       $(CHARMC) -language charm++ -o pgm $(OBJS)
+       $(CHARMC) -language charm++ -pg -o pgm $(OBJS)
 
 cifiles: pingpong.ci
        $(CHARMC)  pingpong.ci
@@ -15,7 +15,11 @@ clean:
        rm -f *.decl.h *.def.h conv-host *.o pgm charmrun
 
 pingpong.o: pingpong.C
+<<<<<<< Makefile
+       $(CHARMC) -pg -c pingpong.C
+=======
        $(CHARMC) -I$(SRC)/conv-core pingpong.C
+>>>>>>> 1.4
 
 test: all
        @echo "Intra-processor Pingpong.."