tricky way to find out Ugly Number(at LeetCode)

For the given question in LeetCode.com, https://leetcode.com/problems/ugly-number/

I coded kind of tricky logic to solve this question.

Ugly number is number which is only divided by three prime numbers; 2, 3, 5 (1 is considered as Ugly Number)

At first, I tried this and got time limit error.

    bool isUgly(int num) {
        if(num == 1) return true;
        do{
            if(num % 2 == 0) { num /= 2; continue; }
            if(num % 3 == 0) { num /= 3; continue; }
            if(num % 5 == 0) { num /= 5; continue; }
            if(num == 1) return true;
            return false;
        } while(1);
    }

Well, if very big number is given, there might be lots of loop.


So I added one line below and I got Accepted with 8 msec runtime.

if(num % 7 == 0) { return false; }

Fine, but unsatisfied because someone got faster result.

So I added more conditions and final code is this:

    bool isUgly(int num) {
        if(num == 1) return true;
        do{
            if(num % 7 == 0) { return false; }
            else if(num % 11 == 0) { return false; }
            else if(num % 13 == 0) { return false; }
            else if(num % 17 == 0) { return false; }
            else if(num % 19 == 0) { return false; }
            else if(num % 23 == 0) { return false; }

            if(num % 2 == 0) { num /= 2; continue; }
            if(num % 3 == 0) { num /= 3; continue; }
            if(num % 5 == 0) { num /= 5; continue; }
            
            if(num == 1) return true;
            return false;
        } while(1);
        
        
    }

Finally I got Accepted with 4 msec runtime.

Heh, lucky.

This code might be silly like a undergraduate student's job but shows good performance for the LeetCode examples.

I don't know how big numbers should be evaluated for the Ugly Number but most numbers might have small prime factors like 2, 3, 5, 7, 11, 13 blar blar..

I am not a mathematician and I cannot speculate the way of getting prime factors effectively.

However, If you just divide the given number with few prime numbers such as 7, 11, 13 ... 23, you can quickly make it.

According to Wikipedia, the number of prime numbers below a given number is approximately like this: 

Computer system's stability and reliability

We usually say A computer system is stable or reliable.
Both of them are used when we decorate system's non-functional quality.
Those are similar but different adjective. What's difference?

'stable' means literally system is not easy to die. In computer system, critical SW errors such as memory access violation and dead lock causes system shutdown or halting.
'stable' system has strong exception handling to not cause those unstable situation.
(This is totally different with logical error which results unexpected output against program's purpose. We call this 'bug', not 'unstable' system).

stable system is actually used in control system engineering. stable control system is a strong controller which can accept noise or sharply changed input data and make the system stable.

'reliable' is much more computer-engineering idiom. It means there are looooooooooots of efforts to keep the system away from unstable situation.
For example, even if unstable situation comes, there might be some backup solution such as redundancy design or system can recover and roll back its status. And data loss prevention is also important axis for evaluating system's reliability.

The key point is that both stability and reliability do not have specific standardization.
There is just an evaluation basis that should be satisfied. Because there is not system satisfying 100% of quality.
One more thing, developers commonly should make a loooooooooooooot of efforts.






컴퓨터 시스템의 stability VS reliability.

흔히들 컴퓨터 시스템이 stable and reliable 하다 라고 한다.
일단 둘다 비기능적 품질을 형용할 때 자주 사용되는 말이다.
비슷하면서도 뭔가 다른 이 두 단어는 어떨 때 쓸까?

stable(안정적) 하다는 말은 한마디로 시스템이 안죽는다는 뜻이다. 컴퓨터 시스템 구조상 잘못된 메모리 접근 이나 데드락 같은 심각한 소프트웨어 오류들은 시스템을 죽게 만들거나 먹통으로 만들어 버린다.
stable 한 시스템은 이러한 문제가 거의 발생하지 않도록 예외 상황을 꼼꼼히 처리해 두었다는 뜻이다. (원래 의도와 잘못된 동작을 하는 로직 오류와는 전혀 다르다. 이건 bug가 있다라고 하지 unstable 하다 라고 하지 않는다)
stable system은 원래 제어 공학에서 control system의 input data에 noise(간섭) 또는 급격한 데이터 변화(큰 error)가 들어 와도 이를 잘 상쇄 시켜서 제어를 안정적으로 하는 시스템을 표현할 때 쓴다.
SW의 로직이나 알고리즘(제어기)에서 error(noise) 가 발생하더라도 예외 상황을 잘 처리한다(상쇄) 라고 비유하면 이해가 쉽다. (전산학의 많은 용어들이 전통 공학들에서 사용하는 용어들을 많이 가져왔다)


reliable(신뢰할만한) 하다는 말은 조금더 매우 컴퓨터 공학적인 용어인데, 시스템이 unstable 한 상황이 발생하지 않도록 노오오오오오력을 많이 했다는 이야기다. unstable한 상황이 발생하더라도 그것을 대체할 수 있는 차선책이 있다던가(redundancy design 등), stable 한 상태로 온전하게 복구를 할 수 있다(roll back)는 뜻이다. 또 스스로 복구가 되었을 때 예전 상태를 온전하게 유지할 수 있는지(data loss prevention) 여부도 reliability를 결정하는데 중요한 요소이다.


중요한 것은, stability와 reliability는 둘 다 딱히 정확한 기준이라고 하는 것이 없다.
그냥 평가 기준을 만족 하느냐 아니냐만 있을 뿐이다. 왜냐하면 100% 완벽한 시스템은 없기 때문이다.
또 하나, 개발자들이 노오오오오오오오력을 많이 해야 한다는 공통점이 있다.

Heterogeneous compute case study: image convolution filtering(by Doug Watt)

Heterogeneous compute case study: image convolution filtering

by Doug Watt on September 4, 2015 in Multimedia, PowerVR Graphics

I translated this article, originally written by Dough Watt, into Korean.

refer to link below
http://blog.imgtec.com/powervr/heterogeneous-compute-case-study-image-convolution-filtering
  

이기종 컴퓨팅 적용 사례 : 이미지 컨볼루션 필터링

      

In a previously published article, I offered a quick guide to writing OpenCL kernels for PowerVR Rogue GPUs; this sets the scene for what follows next: a practical case study that analyzes image convolution kernels written using OpenCL.

Many image processing tasks such as blurring, sharpening and edge detection can be implemented by means of a convolution between an image and a matrix of numbers (or kernel). The figure below illustrates a 3×3 kernel that implements a smoothing filter, which replaces each pixel value in an image with the mean average value of its neighbours including itself.

예전에 나는 PowerVR Rogue GPU를 위한  a quick guide to writing OpenCL kernels 라는 글을 기고했었다. 여기에서 언급할 내용은 OpenCL로 작성된 이미지 컨볼루션 커널을 분석하는 실제 적용 사례에 대한 것이다.
Blurring, sharpening, edge detection 등과 같이 많은 이미지 처리 연산들은 하나의 이미지와 여러 개의 matrix 사이의 컨볼루션으로 정의되고 구현된다.(또는 커널 이라고도 한다). 아래 그림을 보면 3x3 커널이 보이고 여기서는 smoothing filter연산을 구현한 것이다. 이미지 내의 각각의 픽셀은 그 주위의 픽셀 값들의 평균 값으로 대체된다.

Kernel convolution usually requires values from pixels outside of the image boundaries. A variety of methods can be used to handle image edges, for example by extending the nearest border pixels to provide values for the convolutions (as shown above) or cropping pixels in the output image that would require values beyond the edge of an input image, which reduces the output image size.

커널 컨볼루션은 주로 이미지경계 밖의 외곽선의 픽셀 값을 필요로 한다. 이미지 경계선을 처리하기 위해 다양한 방법들이 사용되는데, 예를 들면 (위 그림처럼) 컨볼루션 연산을 위한 값을 가장 가까운 경계 픽셀로부터 단순히 확장해서 채우는 방법 또는 입력 이미지의 경계 부분의 값을 그대로 사용하고 결과 이미지의 외곽선 픽셀을 잘라내는 방법이 있을 수 있다. 후자의 경우 결과 이미지의 크기는 줄어든다.

The table below shows the algorithmic pseudo code for this filter on the left along with a C implementation on the right. In the C program, it is assumed that each pixel is represented by a 32-bit integer comprising four 8-bit values for R, G, B and A; the macro MUL4 therefore performs four separate multiplications, one for each of these 8-bit values.

아래 테이블은 위 예제를 코드로 표현한 것인데, 왼쪽은 간단한 알고리즘의 슈도 코드를, 오른쪽은 그에 대응하는 C 코드를 보여준다. C 프로그램에서 각각의 픽셀은 32비트 정수로 표현되며 각각 R, G, B, A(알파) 채널에 해당하는 색상을 8비트로 표현한다. 코드의 “MUL4” 매크로는 네 개의 분리된 곱 연산을 수행하며, 그 연산들은 각각 8비트 값에 대한 연산이다.

Convolution filter pseudo code	Convolution filter in C
	void blur(int src[Gx][],int dst[Gx][], char *weight) { int x, y, pixel, acc;
for each image row in input image:	for (y=1; y<Gx-1; y++) {
for each pixel in image row:	for (x=1; x<Gy-1; x++) {
set accumulator to zero	acc = 0;
for each kernel row in kernel:	for (j=-1; j<=1; j++) {
for each element in kernel row:	for (i=-1; i<=1; i++) {
multiply element value corresponding to pixel value	pixel = MUL4(src[y+j][x+i], weight[j+1][i+1], 16);
add result to accumulator	acc += pixel;        } }
set output image pixel to accumulator	dst[y][x] = acc;      }    }  }

By using the same approach as introduced in my previous article to extract the inner compute kernel from the serial control flow (the outer two nested loops), and applying the programming guidelines found here, the following OpenCL kernel is produced.

예전 글에서 소개한 것과 같은 방식으로 순차적인 코드의 control flow를 아래와 같은 OpenCL 커널로 작성할 수 있다..

__attribute__((reqd_work_group_size(8, 4, 1))) __kernel void blur(image2d_t src, image2d_t dst, sampler_t s, float *weight) {   int x = get_global_id(0);   int y = get_global_id(1);   float4 pixel = 0.0f;   for (j=-1; j<=1; j++) {     for (i=-1; i<=1; i++)       pixel += read_imagef(src, (int2)(x+i,y+j), s) * weight[j+1][i+1];   }   write_imagef(dst, (int2)(x,y), pixel/9.f); }

__attribute__((reqd_work_group_size(8, 4, 1)))

The statement sets the workgroup size at compile-time to 32 (8x4). This restricts the host program to enqueuing this kernel as a 2D range with an 8×4 configuration, but improves the performance of the kernel when compiled.

이 명령줄에서 workgroup 크기는 컴파일 시에 32로 정해진다(8x4). 이는 호스트 프로그램이 8x4의 설정 값으로 커널을 2D range로 enqueue 하도록 제한한다. 하지만 일단 컴파일이 되면 커널의 성능은 증가한다.

__kernel void blur(image2d_t src, image2d_t dst, sampler_t s, …

The function declaration specifies OpenCL image parameters and a sampler for accessing image data from system memory. To implement the border pixel behaviour in the above example, the host should configure the sampler as clamp-to-nearest-border (not shown).

여기에 선언된 커널 함수는 OpenCL 이미지와 시스템 메모리로부터 이미지 데이터를 접근하기 위한 sampler를 파라미터로 선언한다. 경계 픽셀 영역을 고려하기 위해 호스트는 sampler를 clamp-to-nearest-border로 설정해야 한다.(여기선 안보이는데, 커널 함수 호출 시 파라미터에 넣는 sampler 이미지를 선언할 때 sampler state를 설정하는 부분이 있다. 다음 링크를 참고하라. 역자 주 https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/sampler_t.html  )

float4 pixel = 0.0f;

The statement defines a vector of four 32-bit floating point values. The threads use floating-point arithmetic to perform the convolution, which offers higher throughput compared to integer or character data types.

이 명령줄은 4개의 32비트 floating 포인트 값에 대한 벡터를 정의한다.(4차원 벡터이다) 스레드는 컨볼루션을 수행하기 위해 floating-point 산술 연산을 사용하는데, floating-point를 쓰는 것이 정수나 문자열 데이터 타입에 비해 훨씬 빠른 연산 속도를 제공한다.

read_imagef(src, s, (int2)(x+i,y+j))

The statement causes the TPU to sample a pixel from system memory into private memory, converting the constituent R, G, B and A values into four 32-bit floating-point values and placing these into a four-wide vector. This conversion is performed efficiently by the hardware, requiring the multiprocessor to issue just a single instruction.

이 명령줄은 시스템 메모리로부터 픽셀 데이터를 읽어서 private memory 영역으로 복사를 하도록 TPU(Texture Processing Unit)를 사용한다. 이 때 R,G,B,A 값들을 4개의 32비트 floating-point 값으로 변환을 하고 4차원 벡터에 넣는다. 이 컨볼루션은 하드웨어에 의해 매우 빠르게 수행되며, 단 하나의 멀티프로세서 instruction 으로 수행된다.

write_imagef(dst, (int2)(x,y), pixel/9.f);

The statement writes a (normalized) output pixel value back to system memory.

이 명령줄은 결과 픽셀 값들을 시스템 메모리에 쓰는 역할을 한다.

Caching frequently-used data in the common store

자주 사용되는 데이터를 common store에 캐쉬하기

In the example in the previous section, all work-items operate independently of one another, each work-item independently sampling nine input pixels to calculate one output pixel. Overall the kernel has a fairly low arithmetic intensity (i.e. a low ratio of multiply-and-accumulate operations to memory sampling operations), which can lead to low performance.

이전에 보여준 예제에서 모든 work-item들은 독립적으로 수행하고, 각각의 work-item들은 하나의 결과 픽셀을 내놓기 위해 9개의 입력 픽셀을 독립적으로 샘플링한다. 이는 커널 전반적으로 동등하게 낮은 산술 집적도를 갖는다. (메모리 샘플링 연산에 비해 multiply-and-accumulate 연산의 비율이 낮다.) 당연한 이야기겠지만 낮은 산술 집적도는 낮은 성능을 야기시킨다.

For a workgroup size of 32, each workgroup performs a total of 288 (9×32) sampling operations. However, as show below, adjacent work-items use six of the same overlapping pixels from the input image.

workgroup 사이즈를 32로 정할 때, 각각의 workgroup은 총 288(9x32)번의 샘플링 연산을 수행한다. 하지만 아래에서 보여지듯, 두 개의 인접한 work-item들 사이에 동일한 입력 이미지 안에서 6개의 픽셀 영역이 겹쳐진다.

The common store is a fast on-chip memory that you can use to optimize access to frequently-used data, and also to share data between work-items in a workgroup, in this case enabling reduction of the number of sampling operations performed by a workgroup. A typical programming pattern is to stage data from system memory into common store by having each work-item in a workgroup:

Common store는 빠른 on-chip 메모리로서 자주 사용되는 데이터에 대한 접근을 최적화 할 수 있게 한다. 또한 workgroup 안의 work-item 들 간의 데이터 공유 또한 가능하게 한다. 이 예시에서는 workgroup에 의해 수행되는 sampling 연산의 수를 줄인다. 일반적인 프로그래밍 패턴은 하나의 workgroup 안에 각각의 work-item을 갖도록 만들어서 시스템 메모리로부터 common store로 데이터를 위치시킨다.
아래에 그 방법을 순서대로 설명하였다.

·         Load data from global memory to local memory.

먼저 Global memory로부터 local memory로 데이터를 읽어 들인다.

·         Synchronize with all other work-items in the workgroup, ensuring all work-items block until all memory reads have completed.

·         모든 work-item들이 모든 메모리 접근이 끝날 때까지 block되도록 Workgroup 내 모든 work-item들을 동기화 한다.

·         Process the data in local memory.

Local memory에 있는 데이터를 처리한다.

·         Synchronize with all other work-items in the workgroup, ensuring all work-items finish writing their results to local memory.

모든 work-item들이 결과를 local memory에 쓰도록 Workgroup 내의 모든 work-item들을 동기화 한다.

·         Write the results back to global memory.

결과를 global memory로 돌려준다.

The example program below is a refinement of the previous program, rewritten to use local memory to reduce the number of sampling operations to system memory.

아래의 프로그램 예제는 앞의 예제 프로그램에서 시스템 메모리에 대한 sampling 연산의 수를 줄이도록 local memory를 사용하도록 수정한 것이다. 

__attribute__((reqd_work_group_size(8, 4, 1))) __kernel void blur (image2d_t src, image2d_t dst, sampler_t s, float *weight) {   int2 gid = (int2)(get_group(id(0)*8, get_group_id(1)*4);   int2 lid = (int2)(get_local_id(0),   get_local_id(1));   float4 pixel = 0.0f;   __local float4 rgb[10][6];   prefetch_texture_samples_8x4(src, sampler, rgb, gid, lid);   for (j=-1; j<=1; j++) {     for (i=-1; i<=1; i++)       pixel += rgb[lid.x+1+i][lid.y+1+i]) * weight[j+1][i+1]);   }   write_imagef(dst, (int2)(x, y), pixel/9.f); } void prefetch_texture_samples_8x4(image2d_t src, sampler_t s, __local float4 rgb [10][6], int2 gid, int2 lid) {   if (lid.x == 0) {     // work-item 1 fetches all 60 rgb samples     for (int i=-1; i<9; i++) {       for (int j=-1; j<5; j++)         rgb[i+1][j+1] = read_imagef(src, s, gid+(int2)(i, j));     }   }   barrier(CLK_LOCAL_MEM_FENCE); }

__local float4 rgb[10][6];

The statement declares a local array, which is allocated in the common store.

이 명령줄은 local 메모리 영역에 배열을 선언한다. Local 메모리에 선언된 배열은 common store에 할당된다.

The kernel first calls the function

커널은 먼저 아래 함수를 호출한다.

void prefetch_texture_samples_8x4( …

In this function, all work-items in a work-group first test their local ID together, and work-item 0 samples data from memory into the common store; all work-items then synchronize on a barrier. This synchronization operation is necessary to prevent the other work-items from attempting to read uninitialized data from the common-store. In the main kernel, calls to read_imagef are replaced by reads from the local memory array rgb.

이 함수에서 workgroup 내의 모든 work-item들은 제일 먼저 local ID를 테스트 한다. 그리고 work-item 0은 데이터를 메모리로부터 common store로 샘플링한다. 그 다음 모든 work-item 들은 barrier 연산에서 동기화가 된다. 이 동기화 연산은 common-store에 초기화 되지 않은 데이터를 읽으려는 다른 work-item들의 작업을 막는다. 메인 커널에서는 read_imagef()를 호출하던 것이 local memory 배열인 ‘rgb’로부터 데이터를 읽는 것으로 대체되었다.

In this optimized program each work-group performs a total of 60 sample operations, all during initialization, compared to the 288 in-line sampling operations performed in the original program. This reduction in memory bandwidth can significantly improve performance.

이 최적화된 프로그램은, 기존 프로그램에서 모든 초기화 과정에서 288 개의 in-line 샘플링 연산을 수행하던 것을 각각의 workgroup이 단지 총 60개의 샘플 연상을 수행하도록 수정된 것이다.

The prefetch function can be further improved so that instead of a single work-item fetching 60 samples in sequence, 30 work-items each fetch two samples in sequence. The following example shows one way in which this can be implemented.

prefetch 함수를 통해, 60개의 단일 work-item 에 대한 샘플링이 순차적으로 실행되는 대신에 30개의 work-itme들이 각각 두 개의 샘플 연산을 순서대로 처리함으로써 성능을 개선할 수 있다. 아래의 예제는 이를 구현한 것을 보여준다.

inline void prefetch_8x4_optimized(image2d_t src, sampler_t s, __local float4 rgb[10][6]) {   // Coord of wi0 in NRDange   int2 wi0Coord = (int2)(get_group_id(0)*8, get_group_id(1)*4);   // 2D to 1D address (from 8x4 to 32x1)   int flatLocal = get_local_id(1)*8 + get_local_id(0);   // Only first 30 work-items load, each loads 2 values in sequence   if (flatLocal < 30)   {     /* Convert from flatLocal 1D id to 2D, 10x3 */     int i = flatLocal % 10; // Width     int j = flatLocal / 10; // Height         /* 30 work items reads 10x3 values,      * values 0-9, 10-19, 20-29 from 10x6 - top half      */     rgb[j][i] = read_imagef(src, s, (int2)(wi0Coord.x + i - 1, wi0Coord.y + j - 1));         /* 30 work iteams reads 10x3 values,      * values 30-39, 40-49, 50-59 from 10x6 - bottom half      */     rgb[j + 3][i] = read_imagef(src, s, (int2)(wi0Coord.x + i - 1, wi0Coord.y + j + 3 - 1));   }   barrier(CLK_LOCAL_MEM_FENCE); }

In the best case, work-items can fetch data from the common store in a single cycle. In practice, however, a number of conditions must be met to achieve this efficiency.

여기에서 work-item들은 단일 사이클 내에서 common store 로부터 데이터를 fetch 한다. 그러나 이를 달성하기 위해서는 실제로는 수많은 조건들이 따른다.

Computer vision is what we’ll be focusing on for the next section of our heterogeneous compute series; stay tuned for an overview of how you build a computer vision platform for mobile and embedded devices.

다음에는 이기종 컴퓨팅 시리즈 섹션에서 컴퓨터 비전에 대한 내용에 초점을 맞춰 다루도록 하겠다. 모바일과 임베디드 디바이스에서 어떻게 컴퓨터 비전 플랫폼이 구성되고 최적화를 하는지 설명할 것이다.