simple implementation of all to all communication using send/recv @ 12/23/2010

programming & research
I was writing a Pregel framework for SCC when I encountered the problem of all to all communication. I used RCCE communication library as the underlying communication. However, RCCE provides very limited functionality, not much aside from blocking send/recv actually. Pregel requires all to all communication to exchange the messages sent from vertexes. That's why I started searching for a way to implement all to all communication using send/recv.

After googling for a while, I find that most of the algorithms are just too complex for me. So I dug into the RCCE_Comm library, which provides a RCCE_Alltoall function to perform all to all communication on fixed size data structures. The code seems very straightforward, and I gladly used the same algorithm in my implementation.

Here is the code snippets:

    int np=get_num_workers();
    int me=get_worker_id();
    for(int i=0;i<np;i++){
        int partner=(i-me+np)%np;
        if( me!=partner ){
            if( me<partner ){
                send( send_buf, send_size, partner );
                recv( recv_buf, recv_size, partner );
            }
            else{
                recv( recv_buf, recv_size, partner ); 
                send( send_buf, send_size, partner );           
            }
        }
    }

发布于 12/23/2010 18:49:08 | 评论:0

我的面经 @ 12/15/2010

a box of chocolate
我找工作的一点经验总结,希望对后来者有点用。

简历
    简历直接决定你能不能得到面试机会,也决定了面试官会问你什么问题,以及面试官对你的期望。
    简历要简洁。第一轮挑简历的是HR,他们会面对大量的简历。如果你的简历太长,他们可能没有那么多耐心去看完。但是,你的简历一定要跟别人的区别开来,能让HR一眼看到你的强项。如果你发了很多牛文章,那就把文章放在显眼的位置,适当的加粗。如果从事的项目比较多,那就重点强调项目。
    尽量不要写“精通XXX”这一类的话,除非你真的很精通。如果你的简历经过了第一轮筛选,得到面试机会,那么你的简历会被HR传给面试官。面试官对你的第一印象就来自你的简历。面试官往往跟你是同行,所以这种很主观的,甚至有点自夸性质的描述不要写。要表现自己的优点,经历往往比这种主观的词汇更有说服力。比如你精通C++,那可以说你“在xx开源C++项目中贡献过xxx行代码”,或者“作为主力完成xxx项目的实现”。
    适当的在项目或者文章列表部分加一些简单的描述。比如项目的性质,规模,你在里面所起的作用,等等。这样面试官才能找得到话题,而不会凭空的去想一些面试题来考你。显然,在项目相关的问题上,你应该比面试官熟悉。而如果是面试官找问题来考你,那面试官在该问题上一般都会有比较深入的理解。所以,最好还是用简历来引导面试官的问题,而不是让他来主导。这叫以己之长,击彼之短。

面试前准备
    面试要不要准备?这个分情况,但大多数情况下还是要的。
    IT从业人员的面试,一般分三种情况,一种是考算法跟编程,一种是问项目或者文章相关的问题,最后一种是所谓“闻味”。
    “闻味”的面试一般就不需要准备了,反正是问一些性格方面的问题,照实答就是了。如果你真的跟公司的文化冲突很大,那去了也难受。
    如果是问项目或者文章相关的问题,一般还是需要稍微准备一下的。毕竟大多数的项目和文章都完成很久了,有些细节还是会忘的。可以在面试之前稍微回忆一下。
    而算法跟编程,就需要准备了。微软跟Google都喜欢问算法题。网上有很多经典题,可以看一看。面试的时候不一定会碰到你看过的题,但提前让自己的脑子动一动还是有好处的。我推荐看一下《算法导论》中的分治,partition算法和动态规划。其实面试中考的最多的还是分治法,具体问题不一样,但是分而治之的思想都一样。partition算法和动态规划很少考,但考的时候至少要知道“这个问题可以用partition/动态规划来解决”。编程的话,我只被考过C++,也没有太多变态的东西。当然,如果对自己没信心,事先看一些面试宝典之类的东西也还是有用的。
    最后,面试官还可能会问一些常识性的问题,比如线程进程的区别啦,如何同步啦,这一类。这种问题我觉得就没什么必要准备了,知道的就是知道,不知道的真是一时半会补不回来的。

面试中注意事项
    面试官不一定比你强,也不一定比你弱
    有人怕面试官,有人BS面试官。其实我觉得,面试官们都可能是我们以后的同事,所以不需要特别害怕或者BS他们。他们有可能比你强,也有可能比你弱。所以,如果你碰到不会的问题,没必要沮丧,可能面试官只是在这方面比你懂的多。反之,如果面试官问一些弱智题,或者他自己都不懂,那也没必要BS他,因为他可能只是这方面不行,或者他刚好是公司里最弱的一个。。。
    要抱着平常心态去面试,大多数时候面试官的水平应该与你差不多。

    面试官想要什么样的答案
    有些面试题可能有很多种答案,或者有很多种答法。面试官问问题的时候往往心里会有一个自己的答案,他期待你答出来。那么,怎么去猜测面试官想要的是哪个答案?
    举个例子,那个“进程和线程的区别”的面试题,回答可以是多种多样的。你可以简单的说“进程是有独立地址空间的,线程没有”,也可以说“Linux下进程跟线程是一个东西,只不过线程是跟其它线程共享一个地址映射表的”。面试官想要哪个答案,取决于他自己的知识水平,以及他对你的期待。所以这个时候就需要你去问他:你要的是哪个答案?当然,你不能开口这么问。所以你可以先说“进程是有独立地址空间的,线程没有”,然后问他还需不需要描述更多细节。
    面试是一个交流的过程。面试官试图在有限的时间内,对你做更多的了解。而你则是要更多的表现你的优点。大家的目的其实是一样的。多多交流和互动,对大家都有好处。

    碰到不会的题怎么办
    所谓“不会”的题,可能你没理解面试官想问什么,可能是某些知识你根本就不知道,或者是某些算法一时想不出来。
    如果你没理解面试官想问什么,你可以请求他给你一些提示。不要试图去猜测,这样太浪费时间,也没有必要。面试官肯定希望你能尽快的答完问题,这样大家都能休息了。比如说,面试官可能会问一些很抽象的问题,比如“有一个数据,更改他的时间很长,但是我又需要经常读,怎么办?”他的脑子里其实是另一个问题,可能是一个很具体的应用场景,但是他把问题抽象出来,于是变成一个很难懂的问题。这时候你可以要求他把问题具体化到一个场景中来。于是他可能会把问题变成“我有一个倒排表,更新倒排表的过程很慢,但是更新过程中我需要保证搜索引擎依然能访问它”。这个问题就简单多了。
    对于根本就不知道的知识,不需要装懂。对企业来说,你不懂什么不重要,重要的是你懂什么。每个人的知识面都是很窄的,碰到不懂的题很正常。你可以说,这个我的确不懂,不过我知道另一些相关的知识,比如xxx。有问题不懂正常,但是你得让面试官知道,你不是什么都不懂。
    对于算法题,其实考的是解题思路。比如算法,可能用分治法可以在NlogN时间内解出来的,但是在面试的时间内,很难把整个算法做出来。其实包括面试官本身可能也在短时间内解不出来的,所以没有关系,你只需要让他知道,给你比较充分的时间,你能解得出来。所以,碰到这种规定了复杂度的算法题一般是这样的:(1)先提出一个高复杂度的算法,以证明你至少能解决这个问题;(2)告诉面试官说,这道题可以用分治法,partition或者动态规划来解,复杂度会是多少;(3)试图在有限时间内把问题用高效的方法解出来。这就把一个“会做-不会做”的问题分成了三部分,我相信大多数人至少能搞定(1)和(2),而大多数应该都完不成第三步。但是不要紧,做完(1)和(2)的话,你至少已经得到了三分之二的分数。

    其实还有很重要的一点是,不要冷场。面试的时间很宝贵,不管是对你还是对面试官来说。所以你要在面试的时间内,尽量多的展示你的优点。冷场很显然不符合大家的利益。冷场一般发生在你答不上来的时候,这时候,不要再冥思苦想,冷场的压力太大了,大多数人都没办法集中精神的。其实你可能有一点思路,但是需要时间,你可以一边思考一边把你的思路讲给面试官听。这样一方面不至于冷场,另一方面在你的思路有偏差或者你的回答不符合面试官的预定答案的话,他也有机会提醒你。面试是面试官和面试者的交流过程,你是闲不下来的,也不要让面试官闲下来。充分的利用面试的时间,这才是王道。
发布于 12/15/2010 16:46:10 | 评论:3

拔个草 @ 12/13/2010

a box of chocolate
又好久没写过blog了,过来拔个草

最近在Intel实习,做SCC相关的东东。上星期去无锡给他们做了个presentation,算是有个交待了。

然后在找工作,明年七月毕业,面了IBM, M$, Google,华为。都给offer了,目前基本上定了去MSRA,不过今天又被人种草,说去阿里巴巴分股票。算了,就先去MSRA混两年再说吧。除这些外,HP给了个offer,没面试;Intel随便聊了一次,但没有正式面试,也不准备去了...不折腾了

现在得开始准备毕业了,可能还得再弄篇论文,毕业的时候会容易一些。以后准备做research,得开始养成看paper的习惯了。计划每天至少看一篇,要是能拿看网络小说的热情来看paper就好了...
发布于 12/13/2010 17:23:36 | 评论:2

HTC HD2刷Android的方法 @ 8/23/2010

未分类
HTC HD2是我见过的屏幕最大的手机了。拿到HD2的第一感觉是,大。第二感觉是"shit, windows!"

不是我歧视windows,windows作为桌面OS是非常强大的,但是windows mobile...对话框的关闭键在右上角,对话框能拖动,窗口滚动得靠拖动滚动条。。。算了我还是用Android吧

于是,在折腾了几天后,终于给HD2装了个Android,现在使用起来基本正常。

HD2上装Android跟其它手机不一样,现在并没有原生的Android ROM for HD2。也就是说,不能通过把Android刷到手机ROM里实现。目前为止(至2010年8月23日),还只能把Android文件放到SD卡里,然后在Windows里加载。所以,在HD2上装Android需要以下步骤:
1. 刷一个支持加载Android的Windows Mobile ROM
2. 刷一个新的Radio版本
3. 把Android文件拷到SD卡上
4. 加载Android

1. 刷Windows
有些windows ROM在加载Android时会出错(比如我原来的港行ROM),所以需要刷新的windows ROM。当然,因为刷了Android之后,我主要以Android为主,windows没啥用,所以我就找了一个轻量级的windows ROM:
Chucky ROM

2. 刷Radio
之所以要刷Radio,一个是2.07以下版本的Radio在Android下不工作,另一个是,某些ROM会要求特殊的Radio版本。我刷的是Radio_2_12_50_02_2。刷Radio的过程跟刷windows一样的。
Radio thread

3. 把Android文件拷到SD卡上
首先,确保SD卡格式是FAT32,其它格式不行。然后,去下载一个Android。XDA-developer论坛上有一堆不同的Android。最漂亮的还是Energy的版本,可惜非常费电,新机几乎只能撑一天。所以后来还是用了Darkstone HD2Froyo V2.1,这个版本跟原版的Android几乎一样,不过不知道为什么,没有Google Maps,得自己装,囧。。。
下载下来的是压缩包,解压完,把Android目录拷到SD卡的要目录下。
Energy
Darkstone HD2Froyo V2.1

4. 加载Android
手机开机到windows下,打开文件管理器,进入 sd card -> Android,运行里面的CLRCAD.exe,再运行HARET.exe,就开始加载Android了。这时候应该出现一只企鹅,然后是字符界面。第一次加载需要创建文件,所以比较慢,以后就快了。


关于省电
首先,得找个省电点的Android。Darkstone是我试过的几个中最省电的,有人说shubCRAFT也挺省电,不过实在太丑,我一看到那界面就败退了。
其次,装个SetCPU,设置CPU频率
第三,装个advanced task kill,在屏幕不亮时把进程都杀掉
最后,如果一开机手机下方SD卡附近就很热的话,那就重启机器,一般重启一次就好了
发布于 8/23/2010 16:12:52 | 评论:4

Sonny.Pacman @ 5/18/2010

life in IT industry
这周末抽空把原来的天气预报程序改了改,把代码整理了一下,放到google code做个开源项目。项目名字叫Sonny4Pacman。Sonny 是《I.Robot》里那个机器人。Pacman是我现在所在的研究小组的名称。

Sonny基本上是个Fetion+Gtalk接口的应答程序。用Python写的,用PyFetion来登陆Fetion服务器,用xmpppy登陆GTalk。另外还有个脚本从weather.com.cn读天气预报信息。

现在提供的功能是:
1. 定时发天气预报 (现在定在每天晚上19点)
2. 提供短信查询天气预报。只需要发送 "天气 [空格] 北京"到 1252013693549822 就能查询北京的天气。因为飞信的限制,目前只有好友能使用这个功能
3. 提供GTalk管理界面。这个主要是因为我自己不能用手机给自己的飞信帐号发短信,所以只能用GTalk管理。我的G1可以方便的上GTalk,所以就用GTalk做管理接口了

以下是google code项目地址,有兴趣的可以去下下来玩:
http://code.google.com/p/sonny4pacman/
发布于 5/18/2010 10:25:17 | 评论:0

突发奇想--有没有那种可以调节光线的玻璃 @ 3/2/2010

a box of chocolate
玻璃是一种很常见的东西,主要作用就是在透光的基础上提供挡风功能。但有时候我们又不需要那么强的光线,于是就有了窗帘这种东西。窗帘的作用就是,在人们需要的时候,为人提供一定的遮挡。
窗帘提供了方便的调节光线的功能,但同时它也是有代价的--一次性购买费用就不说了,它还占空间,而且需要经常清洗。那么,有没有一种玻璃,能够提供调节光线的作用,而不需要额外加窗帘呢?
这样的玻璃主要有两种用途:一种是在空间狭小的地方,比如车上;另一种是在高楼大厦的外墙玻璃,现在这种外墙玻璃一般是配合百叶窗使用的,如果使用可调节玻璃,显然更多好看一些。

我的思路是这样的:使用液晶材料做一个贴层,利用电流改变贴层的透光度来达到调节的作用。最低成本的做法就是,整个玻璃做成一大块,电流大小会改变整个玻璃的透光度,这样的成本应该可以很容易的控制。然后再高点,可以做成一定大小的阵列,可以独立改变每个阵列中的小块。最高档的就做成现在的TFT那种样子--还可以显示各种色彩,连壁画都省了。

当然,只是瞎想。具体这东西成本会有多高,用的时候会不会因为阳光照射而快速老化,我都不知道。。。
发布于 3/2/2010 10:33:05 | 评论:5

get thread id in pthread @ 12/19/2009

programming & research
Thread ID starting from 0 is useful when writing multi-threaded program. In OpenMP, you can simply use omp_get_thread_num() to get the current thread ID. But unfortunately, there is no such thing in PTHREAD. Though pthread_self() returns pthread_t, it does not return an ID starting from 0.

The following code is a simple and naive implementation of GetThreadID() in PTHREAD. By wrapping the standard PTHREAD calls, we are able to provide a mechanism to determine the current thread ID.  The code basically uses pthread_key, which is provided by PTHREAD to manipulate data on a per-thread basis, so that we can keep a thread ID for every thread, and easily retrieve the information within each thread. For more information, please refer to PTHREAD specification.

Note: the code actually has several flaws, one of them is that CreateThreadWithID() is not thread-safe. This is because it uses a global variable thread_barrier to synchronize with the child process. But if there are more than one thread trying to perform CreateThreadWithID(), there will be race conditions. I will try to come up with a thread-safe version later, but currently, I am satisfied with it. It meets my purpose, in which a master thread creates all the other slaves.


// MyThread.h
        struct ThreadParam{
                unsigned int id;
                void *(*func)(void *);
                void * arg;
        };

        static pthread_key_t thread_key;
        static pthread_barrier_t thread_barrier;
        static pthread_once_t thread_once=PTHREAD_ONCE_INIT;

        void thread_make_key();
        void * FakeEntry( void * param );
        pthread_t CreateThreadWithID(void *(*func)(void *), void * arg, unsigned int id, const pthread_attr_t * attr);
        int ThreadJoin(pthread_t);
        unsigned int GetThreadID();



// MyThread.cpp
        void thread_make_key(){
                pthread_key_create(&thread_key,NULL);
        }

        void * FakeEntry(void * param){
                ThreadParam args=*(ThreadParam *)param;
                pthread_barrier_wait(&thread_barrier);
                pthread_once(&thread_once, thread_make_key);
                pthread_setspecific(thread_key, &args.id);
                args.func(args.arg);
        };

        pthread_t CreateThreadWithID(void *(*func)(void *), void * arg, unsigned int id, const pthread_attr_t * attr){
                ThreadParam param;
                param.id=id;
                param.func=func;
                param.arg=arg;

                pthread_barrier_init(&thread_barrier, NULL, 2);

                pthread_t thread;
                pthread_create(&thread, attr, FakeEntry, &param);

                pthread_barrier_wait(&thread_barrier);
                pthread_barrier_destroy(&thread_barrier);

                return thread;
        }

        int ThreadJoin(pthread_t thread){
                return pthread_join(thread,NULL);
        }

        unsigned int GetThreadID(){
                return *(unsigned int *)pthread_getspecific(thread_key);
        }



// testThread.cpp
#include "MyThread.h"

#include <iostream>
using namespace std;
using namespace HMM_CPU;

const int THREADS=4;

void * foo(void * arg){
        *(int *)arg=GetThreadID();
        return 0;
}

int main(){
        // create threads
        pthread_t * threads=new pthread_t[THREADS];
        int * ids=new int[THREADS];
        for(int i=0;i<THREADS;i++){
                threads[i]=CreateThreadWithID(foo, &ids[i], i, NULL);
        }

        for(int i=0;i<THREADS;i++){
                ThreadJoin(threads[i]);
                cout<<ids[i]<<endl;
        }

        delete[] ids;
        delete[] threads;
        return 0;
}
发布于 12/19/2009 21:49:57 | 评论:0

three ways of saying "count lines of code of my project" @ 12/4/2009

programming & research
C++ code:

// C++  code, counts all *.java code in a directory
#include <stddef.h>
#include <stdio.h>
#include <sys/types.h>
#include <dirent.h>
#include <iostream>
#include <fstream>
#include <sys/time.h>
#include <string>
using namespace std;

long long count_file(const string & path){
        string str;
        long long loc=0;
        ifstream in(path.c_str());
        while(in.good()){
                getline(in,str);
                loc++;
        }
        return loc;
}

long long count_dir(const string & dir){
        DIR *dp;
        struct dirent *ep;
        long long loc=0;

        dp = opendir (dir.c_str());
        if (dp != NULL) {
                while (ep = readdir (dp)){
                        string filename=ep->d_name;
                        // skip . and ..
                        if(filename=="." || filename==".." || ep->d_type==DT_LNK)
                                continue;

                        string path=dir+"/"+filename;
                        if(ep->d_type==DT_DIR){
                                cout<<"*********** entering dir: "<<path<<endl;
                                loc+=count_dir(path);
                        }
                        else{
                                cout<<"checksum file: "<<path<<endl;
                                loc+=count_file(path);
                        }
                }
                closedir (dp);
        }
        else{
                cout<<"Couldn't open the directory "<<dir<<endl;
        }
        return loc;
}

int main (int argc, char * argv[]){
        if(argc!=2){
                cout<<"usage: "<<argv[0]<<" dir"<<endl;
                return 1;
        }

        long long loc=count_dir(argv[1]);
        cout<<"================="<<endl;
        cout<<"total loc: "<<loc<<endl;

        return 0;
}


python code:

#!/usr/bin/python

import os
import sys

total_loc=0

def file_loc(file_path):
        global total_loc
        print("counting %s\n"%file_path)
        loc=0
        f=open(file_path,'r')
        for line in f:
                loc+=1
        total_loc+=loc

def dir_loc(exts, dir_name, files):
        loc=0
        for file in files:
                file_wanted=0
                for ext in exts:
                        if os.path.splitext(file)[1]==ext:
                                file_wanted=1
                                break
                if file_wanted:
                        file_loc(dir_name+"/"+file)

def main():
        global total_loc
        if(len(sys.argv)<3):
                print("usage: %s dir file_exts ...\n" % sys.argv[0])
                return
        dir=sys.argv[1]
        exts=[]
        for ext in sys.argv[2:]:
                exts.append("."+ext)
        total_loc=0
        os.path.walk(dir, dir_loc, exts)
        print("===================\ntotal loc: %d \n"%total_loc)

main()


shell code:

find . -name *.py | xargs wc -l


These code are written by myself. I am a newbie to python and shell programming. If you have better method for attacking this problem, or if you have other options, such as perl, Ruby, please let me know. :-)
发布于 12/4/2009 10:18:21 | 评论:1

simple logging mechanism for CUDA device code @ 10/13/2009

life in IT industry
I have been trying to debug a piece of CUDA code, and found it extremely hard. Sometimes, the bug won't reproduce in -deviceemu, nor -G. And sometimes, when it does come up in -deviceemu, I am not able to set a breakpoint in the right function, because it has been inlined. So I composed this simple logging "library". I would like to share it with you guys, hope it would be helpful for someone.

/***********************************************************
*  DLog.h
*  Simple logging mechanism for CUDA device code
*
*  Usage:
*       1. include it in your .cu file
*       2. call DLog_Init(size) to reserve size bytes as output buffer
*           on device
*       3. call DLog<<data in your CUDA device code
*       4. call DLog_Dump() after kernel invocation to dump the output
*
*  This implementation is very simple, and inefficient. It supports only
*     char, int and const char *.
*
*  Author: Chuntao HONG (chuntao.hong@gmail.com)
*  Last modified: Oct. 13th, 2009
************************************************************/
#ifndef DLOG_H
#define DLOG_H

#include <cuda.h>
#include <iostream>
using namespace std;

const int MIN_ALIGN=4;

// align n to b bytes
template<class T>
__device__ T Align(T n, uint32_t b){
        return ((uint32_t)n&(b-1))==NULL ? n : n+b-((uint32_t)n&(b-1));
}

template<class T>
__device__ T MinAlign(T n){
        return ((uint32_t)n&(MIN_ALIGN-1))==NULL ? n : n+MIN_ALIGN-((uint32_t)n&(MIN_ALIGN-1));
}

// set some empty bytes, so that the logging data structures won't be polluted by out-of-bound writes
const int SAFTY_BYTES=1024*1024;

// data types supported, now only three
enum DataType{
        CHAR,
        INT,
        STRING
};

struct LogPack{
        int size;
        DataType type;
        union{
                char c;
                int i;
                char str[];
        }data;
};

__device__ int requiredLogSpace(const char*);
__device__ int requiredLogSpace(const int);
__device__ int requiredLogSpace(const unsigned int);
__device__ int requiredLogSpace(const char);

__device__ void assignLog(LogPack *, const char c);
__device__ void assignLog(LogPack *, const char * c);
__device__ void assignLog(LogPack *, const int c);
__device__ void assignLog(LogPack *, const unsigned int c);


struct MyLog{
public:
        char empty_bytes[SAFTY_BYTES];
        char * buf;
        unsigned int curPos;
        char empty_bytes2[SAFTY_BYTES];
public:
        template <class T>
        __device__ MyLog & operator<<(const T & t){
                int packSize=MinAlign(requiredLogSpace(t));
                int start=atomicAdd(&curPos, packSize);
                LogPack * pack=(LogPack *)(buf+start);
                assignLog(pack,t);
                pack->size=packSize;
                return *this;
        }
};

__device__ MyLog DLog;

__device__ int requiredLogSpace(const int i){
        return sizeof(i)+sizeof(DataType)+sizeof(int);
}
__device__ int requiredLogSpace(const char * str){
        int len=0;
        while(*str++){
                len++;
        }
        return len+1+sizeof(DataType)+sizeof(int);
}
__device__ static int requiredLogSpace(const unsigned int i){
        return sizeof(i)+sizeof(DataType)+sizeof(int);
}
__device__ static int requiredLogSpace(const char c){
        return sizeof(c)+sizeof(DataType)+sizeof(int);
}

__device__ void assignLog(LogPack * pack, const char c){
        pack->data.c=c;
        pack->type=CHAR;
}
__device__ void assignLog(LogPack * pack, const int i){
        pack->data.i=i;
        pack->type=INT;
}
__device__ void assignLog(LogPack * pack, const unsigned int i){
        pack->data.i=i;
        pack->type=INT;
}
__device__ void assignLog(LogPack * pack, const char * str){
        int i=0;
        while(*str){
                pack->data.str[i++]=*str;
                str++;
        }
        pack->type=STRING;
}


__host__ void DLog_Init(unsigned int size){
        MyLog localLog;
        localLog.curPos=SAFTY_BYTES;
        cudaMalloc((void**)&localLog.buf, size+SAFTY_BYTES);
        cudaMemset(localLog.buf, 0, size);
        cudaMemcpyToSymbol(DLog, &localLog, sizeof(MyLog), 0, cudaMemcpyHostToDevice);
}

__host__ void DLog_Dump(){
        MyLog localLog;
        cudaMemcpyFromSymbol(&localLog, DLog, sizeof(MyLog), 0, cudaMemcpyDeviceToHost);
        char * h_buf=new char[localLog.curPos+1];
        cudaMemcpy(h_buf, localLog.buf, localLog.curPos, cudaMemcpyDeviceToHost);
        cudaFree(localLog.buf);

        char * ptr=h_buf+SAFTY_BYTES;
        while(ptr<h_buf+localLog.curPos){
                LogPack * pack=(LogPack *)ptr;
                if(pack->type==CHAR){
                        cout<<pack->data.c<<endl;
                }
                else if(pack->type==INT){
                        cout<<"0x"<<hex<<pack->data.i<<endl;
                }
                else{
                        cout<<pack->data.str<<endl;
                }
                ptr+=pack->size;
        }

        delete[] h_buf;
}

#endif
发布于 10/13/2009 9:53:33 | 评论:0

CUDA-GDB error: undefined reference to '$gpu_registers' @ 4/28/2009

未分类
I encountered this problem when I was trying to compile MARS with -G option.

I first tried to copy the ptxvars.cu to the project directory, and compile it with the project files. But the compiler complained about conflicting declarations of threadIdx, etc.

I then tried to compile a simple .cu file with -G, and it worked all right. So I figured there must be something wrong with my Makefile. When I compiled the .cu file, I just used "nvcc test.cu -G", while the files in my projects was first compiled into .o files and then linked together. Then I compiled test.cu into test.o, and then tried to use "nvcc test.o" to generate the executable. Then the error message "undefined reference to ..." appeared. Suddenly I realized that the difference between the two approaches was that nvcc was not passed the -G option when linking. So I tried "nvcc test.o -G", and it worked perfectly. By adding the -G option to my link command in Makefile, I have successfully built my project.

So, remember, when you use -G to compile .cu files, also add -G to the link command line.
发布于 4/28/2009 16:48:07 | 评论:0
总是很忙,同时又很闲

QQ:  26765600
MSN:  hctiaa@21cn.com
categories
recent comments
archives
links
statistics
  • 网志数:200
  • 评论数:345