Hang in there!

加油!

Forza!
¡Ánimo!
Bon courage!
Hang in there!
Zlom vaz!
בהצלחה!
geluk!
Viel Erfolg!
بالتوفيق
İyi şanslar!

from Tokyo

Happy Canada Day!

Hello Vancouver

Good bye Montreal
Hello Vancouver

New tab homepage and firefox 59

It looks firefox 59 has a nice feature to kindly disable newtab homepage add-on in favor of their cool bookmark page.
If it's not working check the add-on manager.

VILKU VILKA Crafts

A friend of mine has opened an online shop.

Vancouver based Latvian Lifestyle Shop
Latvian Zakka* + Original Handcrafts

Check out their crafts here!
http://vilkuvilka.ca/onlineshop/

Canada 150

Happy Canada day!
Joyeuse Fête du Canada!

Fracture

Real example here.

Kd-tree implementation

 I have implemented my own kd-tree library based on Accelerating kd-tree searches for all k -nearest neighbours (pdf), compared the result with ANN to confirm that it outputs the correct result. It was surprising my implementation was several hundred times faster (below graph, for the same queries, ANN gets 14728 samples while my code gets 15). I still doubt if there's something wrong with my code or profiling because it's too fast even after considering the fact that my kd-tree is specialized in finding nearest *one* point and uses Eigen::Vector4f for SSE optimisation.

 I tested several sampling profiler: perf, gperftools, oprofile, codeXL.  It took some time for me to get the correct call graph. Many profilers output a call graph where malloc/new has no parent, probably because libstdc++ is not compiled with no-omit-frame-pointer and cannot find the caller function. Newer profilers such as recent perf and gperftools can handle it correctly. I couldn't though find a way to make perf report's "-G" option work when I took a profile with

perf record -call-graph dwarf ...

so I used gperftools for it (above graph). It still took a while until I found pprof's 'callgrind' option (kcachegrind) could not output the result correctly. 'web' worked fine.

Amazing

from Wikipedia:

Andromeda Software Development (or simply ASD) is a Greek demogroup that was formed in 1992. They produced a number of small intros and demos in the mid-1990s for the PC, most notably CounterFactual (winner of the first Greek demo party ever, The Gardening 1995) and Beyond (Placed 4th in The Gardening 1996). ASD was quiet for the following years until 2001, when they presented Cadence & Cascade - their first accelerated demo - and won the Digital Nexus demoparty, held in Athens, Greece.

Testing unified memory

I just tested CUDA6's unified memory. This code worked as I expected.

#include <iostream>
#include <string.h>

#include <cuda_runtime.h>

__global__ void set(char* buf, unsigned int num)
{
 int i = blockDim.x * blockIdx.x + threadIdx.x;
 if (i < num)
 {
  buf[i] = 1;
 }
}

int main(void)
{
 unsigned int num = 16384;
 char *buf;
 cudaMallocManaged(&buf, num);
 memset(buf, 0, num);                         //##1## D->H(16384bytes). page fault.
 set<<<4, 128>>>(buf, num);                   //##2## H->D(16384bytes).
 cudaDeviceSynchronize();
 std::cout << (int)buf[0] << std::endl;       //##3## D->H(16384bytes). page fault.
 std::cout << (int)buf[10000] << std::endl;   //No data transfer.
 set<<<4, 128>>>(buf, num);
 cudaDeviceSynchronize();
 buf[0] = 5;                                  //##4## D->H(16384bytes). page fault.
 set<<<4, 128>>>(buf, num);                   //##5## H->D(4096bytes).
 cudaFree(buf);
 cudaDeviceReset();
   return 0;
}

I used a profiler and added comments where data transfers have occurred. Whenever an un-transferred host memory is accessed for either read or write, Data transfer D->H occurs, not only before read but before write. In my understanding it is because unless it copies the memory D->H on write, once you wrote a value to a memory and then read a memory near the address where you have modified the value, data transfer D->H on read would overwrite your modification. It produces a bit of unnecessary data transfer (##1##).
You can also see when ##5##, only 4096 bytes of data transfer occurs, not full 16384 bytes. It is because it keep tracks of the page you have written (I have asked it to a NVIDIA guy).

I modified the main function a little bit and created another buffer buf2.

int main(void)
{
 unsigned int num = 16384;

 //Below code adds D->H(16384bytes), H->D(16384bytes). 1 page fault.
 char *buf2;
 cudaMallocManaged(&buf2, num);
 memset(buf2, 0, num);

 //Same as above.
 char *buf;
 cudaMallocManaged(&buf, num);
 memset(buf, 0, num);                         //D->H(16384bytes). page fault.
 set<<<4, 128>>>(buf, num);                   //H->D(16384bytes).
 cudaDeviceSynchronize();
 std::cout << (int)buf[0] << std::endl;       //D->H(16384bytes). page fault.
 std::cout << (int)buf[10000] << std::endl;   //No data migration.
 set<<<4, 128>>>(buf, num);
 cudaDeviceSynchronize();
 buf[0] = 5;                                  //D->H(16384bytes). page fault.
 set<<<4, 128>>>(buf, num);                   //H->D(4096bytes).
 cudaFree(buf);
 cudaFree(buf2);
 cudaDeviceReset();
   return 0;
}

In this code I allocated a unified memory to buf2 and called memset() but buf2 is not used elsewhere. It adds D->H(16384bytes) when memset() is called and it is natural, but it also adds H->D(16384bytes). We know the kernel doesn't use buf2 but it is too complicated for CUDA to know that. So before the first kernel call it also transfers buf2, which is another unnecessary data transfer. I have concluded unified memory makes the program simpler but not as efficient as I have expected. I heard Pascal has cleverer memory management. I will wait and see if is good enough.

CPU/GPU memory abstraction

When you write a program that uses GPU (either CUDA of OpenCL), you may want to implement both CPU code and GPU code, and use only one of them depending on the user's choice. For this reason I wanted to know how to treat different kinds of memories generically, i.e. the way  to abstract host and device memory.

 The easiest one is to follow the approach of (or just use) Thrust. It has host_vector and device_vector which interfaces are quite similar to std::vector. When you have one of them, you can copy it to another with assignment. Host/device copy will be done automatically under the hood.

 I wanted another way of abstraction to debug the program easily. What interested me was to have a mechanism where input data was either in host or device memory, and the the code does not have to know where it is. After several try and error I implemented one like this;

template < typename T >
class Buffer
{
public:
    T* get(MemoryType mType);
    const T* get(MemoryType mType) const;
    void setClean(MemoryType mType, bool isClean=true);
    void sync(MemoryType mType) const;
    void allocate(MemoryType mType) const;
    void free(MemoryType mType);

private:
    mutable MemoryType m_cleanState; //Clean state. Bitwise OR of HOST and DEVICE.
    mutable void* m_addrs[2]; //Host and device.
};

It can have both host and device memory, and knows if the data stored in the host/device memory is up to date or not. If the one in the host memory is up to date and one in the device memory is not, it copies data from the host memory to the device by calling sync(DEVICE). if the device memory is already up to date, sync() does nothing.
 You can use this class like this,

someCalculationCpu(const Buffer < float > * input, Buffer < float > * output)
{
    input.sync(HOST);
    float* ip = input.get(HOST);
    float* op = output.get(HOST);
    op[0] = ip[0]; /*Do some calculation with CPU*/
    output.setClean(HOST); //Tell the buffer that the data stored in the device memory is up to date.
}

someCalculationGpu(const Buffer < float > * input, Buffer < float > * output)
{
    input.sync(DEVICE);
    float* ip = input.get(DEVICE);
    float* op = output.get(DEVICE);
    op[0] = ip[0]; /*Do some calculation with GPU*/
    output.setClean(DEVICE); //Tell the buffer that the data stored in the device memory is up to date.
}

anotherCalculationCpu(const Buffer < float > * input, Buffer < float > * output)
{
    /*Same style as someCalculationCpu() with another calculation.*/
}

anotherCalculationGpu(const Buffer < float > * input, Buffer < float > * output)
{
    /*Same style as someCalculationGpu() with another calculation.*/
}

Now the these are all valid,

   someCalculationCpu(input, output);
   anotherCalculationCpu(input, output);

   someCalculationCpu(input, output);
   anotherCalculationGpu(input, output);

   someCalculationGpu(input, output);
   anotherCalculationCpu(input, output);

   someCalculationGpu(input, output);
   anotherCalculationGpu(input, output);

I've already implemented so I'll keep using it but just wonder if there is already a tool or a way with Thrust. Please leave a comment if you know.

Merry Christmas

Wish you have a Merry Christmas.

Looks I haven't updated this blog for a long time. I'm not dead ;)

The eyeballing game

The eyeballing game

My score

Better than the median, yaay.

One of the coolest thing I have ever seen

What am I doing now?

Now I'm working at a CG production, playing with joints and handles. I found Maya Python API 2 is quite handly.
In my free time I'm learning Nuke SDK (the manual is a little bit confusing but its concept is simple), reading ambient occlusion papers (ambient occlusion volumes and HBAO, crytek one seems to produce only low quality ao if my understanding is correct).

Programming experience changes the sense of beauty

I had a chat with a guy today who hates things that have no intrinsic meaning.
Say we have a directory structure like

episode - sequence - shot

and shot dependent files are under the 'shot' directory.
(I'm sure this structure is familiar with you if you are working in a vfx studio)

When we want a directory to put data that are common to all the shots in an episode. We can think of two directory structures for the purpose.

episode - sequence - shot
               - common

episode - sequence - shot
               - common   - common

The latter has two directories named "common". Outer "common" always has only one directory: inner "common". So essentially it's enough if we have only one common directory like the former example. Having two commons are redundant and it can make people feel intuitively ugly. But in real life creating two directory has a good meaning - if we have two commons, we can always assume there are three directories, one for episode, one for sequence, and one for shot. This assumption simplifies making tools a lot. It slows down directory access a little bit, slightly increases disk usage and traffic, but in most cases you can ignore them. It will also make you irritated if you access to common files manually, e.g. with windows explorer, but it is usually wrapped up by tools and it should be, so there's no drawback in that sense as well.

In general, it is better if there are smaller number of rules while making things well defined (i.e. not like, "we have everything in 'any data' direcotry".  In the above example "common" direcotry should have well defined sub-directory). As you start getting used to programming, you become finding simplicity more beautiful, and you starts understanding that removing simplicity just for removing redundancy has no sense because it means sacrificing functional beauty just for visual appearance.


Jun 30 added:

Good generalization and bad generalization. Good generalization makes you extend functionality easily,  bad generalization puzzles you with lack of information on what data it is about. If you have a fully generalized structure (directory structure, framework structure,  work flow structure etc.) and well defined layer on top of it based on your demand, you can prevent having 'any data' stuffs that all you know about is creation date, and the structure still has room for future extensibility.

Hand editing pouring water

Again from my friends project.
It shows how the tool lets you hand edit water pouring into a wine glass.
He says it was a 5 minutes work. Nice usage example.

http://www.facebook.com/sparta.project

You can see the result first, and the tool being used at 0:55.
The number of particles are still low and it looks blobby but you can imagine how the tool can be used and be evolved.

Wine Video from lyouta on Vimeo.


1) Simulate liquid by nParticle
2) Exports BIN files
3) Edit BIN files by Sparta
4) Import BIN files as Maya Particle
5) Connect Maya Particle to nParticle
6) Polygonize liquid
7) Render liquid


 By the way it succeeded in cloud funding today. Congrats.

Espresso Machine!

Finally I bought  one for myself. I'd been telling my friends to buy one since it's kind of big and expensive and it was best to have one at my friends apartment, but recent ones are smaller and cheaper than before.

Buy the way I started working at a CG production again. It's good to be in the CG world.

Espresso machine and cow tit cup

Annular Solar Eclipse

6.44 Still bright outside.

















7:03

Sparta Project (Sculpting particles update)

Updated info for the previous entry. Now it's under crowd funding.

It's crowd funding page
http://camp-fire.jp/projects/view/237
(It accepts funding from overseas: see Sparta Project Is Inviting Your Help for more details)

Project's facebook page
https://www.facebook.com/sparta.project

A web news by a Japanese CG magazine.
http://cgworld.jp/flashnews/software/sparta-project.html
(Click "Translate" for google translation)