opencl
커널 기본 사항
수색…
소개
이 주제는 opencl을위한 커널 작성의 기본 원리를 설명하기위한 것입니다.
그레이 스케일 커널
회색 음영 이미지를 생성하는 커널을 빌드 할 수 있습니다. 각 구성 요소 및 RGBA 순서로 uint를 사용하여 정의 된 이미지 데이터를 사용합니다.
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_LINEAR;
__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
int2 gid = (int2)(get_global_id(0), get_global_id(1));
int2 size = get_image_dim(input);
if(all(gid < size)){
uint4 pixel = read_imageui(input, sampler, gid);
float4 color = convert_float4(pixel) / 255;
color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
pixel = convert_uint4_rte(color * 255);
write_imageui(output, gid, pixel);
}
}
이제 코드를 단계별로 살펴 보겠습니다. 첫 번째 줄은 sampler_t 유형의 __constant 메모리 영역에 변수를 만듭니다. 이 샘플러는 이미지 데이터에 대한 액세스를 추가로 지정하는 데 사용됩니다. 전체 문서는 크로노스 문서를 참조하십시오.
우리는 커널을 호출하기 전에 입력을 read_only로 할당하고 write_only로 출력을 할당하므로 여기에 수정자를 추가합니다.
image2d와 image3d는 항상 전역 메모리에 할당되므로 여기서 __global 수정자를 생략 할 수 있습니다.
그런 다음 우리가 회색 음영으로 변환 할 픽셀을 정의하는 스레드 ID를 얻습니다. 또한 스레드가 할당되지 않은 메모리에 액세스하지 않도록 크기를 쿼리합니다. 잊어 버리면 커널을 충돌시킬 수 있습니다.
우리가 합법적 인 스레드라는 것을 확인한 후에는 입력 이미지에서 픽셀을 읽습니다. 그런 다음 소수점의 손실을 피하기 위해 부동 소수점으로 변환하고 계산을 수행 한 다음 다시 변환하여 출력에 씁니다.
커널 스켈레톤
가장 간단한 커널과 그 변형을 살펴볼 수 있습니다.
__kernel void myKernel() {
}
주 코드에서 시작할 수있는 커널은 __kernel 키워드로 식별됩니다. 커널 함수는 반환 유형 void 만 가질 수 있습니다.
__kernel void myKernel(float a, uint b, byte c) {
}
물론 커널로 노출되지 않는 더 많은 함수를 생성 할 수 있습니다. 이 경우 __kernel 한정자를 생략하면됩니다.
함수는 다른 C / C ++ 함수처럼 변수를 노출 할 수 있습니다. 유일한 차이점은 메모리를 참조하려는 경우입니다. 이것은 인수이든 코드에서 사용 되든간에 모든 포인터에 적용됩니다.
float* ptr;
실행중인 스레드 만 액세스 할 수있는 메모리 영역에 대한 포인터입니다. 사실 그것은
__private float* ptr;
사용할 수있는 메모리 영역 수정 자에는 네 가지가 있습니다. 커널 내부에서는 보통 걱정할 필요가 없지만 인수에 대해서는 필수적입니다.
- __global :이 수정자는 전역 메모리에 배치 된 포인터를 참조합니다.
- __constant : 상수 메모리 포인터를 참조합니다.
- __local : 공유 메모리 포인터를 참조합니다.
- __private : 로컬 메모리 포인터를 참조합니다.
또한 메모리에 액세스하는 방법을 정의 할 수 있습니다.
- 수정 자 없음 : 읽기 및 쓰기
- __read_only
- __write_only
이러한 플래그는 호스트에서 메모리 버퍼를 다시 할당하는 방식과 일치해야합니다.
커널 ID
각 스레드가 스레드 블록 / 전역 스레드 풀에서 해당 위치를 알아야하는 데이터로 올바르게 작업하십시오. 이것은 다음과 같이 성취 될 수있다.
get_local_id($dim);
get_global_id($dim);
이 두 함수는 스레드 블록 또는 모든 스레드를 기준으로 스레드의 위치를 반환합니다.
get_working_dim();
커널이 시작된 총 크기를 가져옵니다.
get_local_size($dim);
get_global_size($dim);
스레드 블록의 총 스레드 수 또는 지정된 차원의 총 스레드 수를 가져옵니다.
주의 사항 : 스레드가 데이터 크기를 초과하지 않는지 항상 확인하십시오. 이것은 일어날 가능성이 매우 높으며 항상 확인해야합니다.
OpenCL의 벡터
각 기본 opencl 유형에는 벡터 버전이 있습니다. 유형 뒤에 원하는 구성 요소의 수를 추가하여 벡터 유형을 사용할 수 있습니다. 지원되는 구성 요소의 수는 2,3,4,8 및 16입니다. OpenCL 1.0은 세 가지 구성 요소를 제공하지 않습니다.
다음 두 가지 방법으로 벡터를 초기화 할 수 있습니다.
- 단일 스칼라 제공
- 모든 구성 요소 만족
float4 a = (float4)(1); //a = (1, 1, 1, 1)
또는
float4 b = (float4)(1, 2, 3, 4);
float4 c = (float4)(1, (float3)(2));
또는 성분들의 수를 만족시키는 임의의 다른 벡터들의 조합을 포함 할 수있다. 벡터 요소에 액세스하려면 다른 방법을 사용할 수 있습니다. 색인 생성을 사용할 수도 있습니다.
a[0] = 2;
또는 리터럴을 사용하십시오. 리터럴의 장점은 원하는 방식으로 결합 할 수 있다는 것입니다. 다음과 함께 모든 벡터 컴포넌트에 액세스 할 수 있습니다.
a.s0 = 2; // same as a[0] = 2
당신은 또한 새로운 벡터로 여러 구성 요소를 결합 할 수 있습니다
a.s02 = (float2)(0, 0); // same as a[0] = 0; a[2] = 0; or even a.s20 = (float2)(0, 0)
원하는 방식으로 구성 요소의 순서를 변경할 수 있습니다.
a.s1423 = a.s4132; // flip the vector
하지만 너는 뭔가를 할 수 없다.
a.s11 = ... // twice the same component is not possible
벡터 구성 요소에 액세스하기위한 편리한 편리한 약도가 있습니다. 뒤에 오는 속기는 크기 2, 4, 8 및 16에서만 적용한다
a.hi //=a.s23 for vectors of size 4, a.4567 for size 8 and so on.
a.lo //=a.s01
a.even //=a.s02
a.odd //=a.13
벡터 크기 2,3 및 4에는 몇 가지 추가 약식이 있습니다.
a.x //=a.s0
a.y //=a.s1
a.z //=a.s2
a.w //=a.s3
감마 보정 커널
감마 보정 커널을 살펴 봅니다.
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_LINEAR;
__kernel void Gamma(__read_only image2d_t input, __write_only image2d_t output, __constant float gamma) {
int2 gid = (int2)(get_global_id(0), get_global_id(1));
int2 size = get_image_dim(input);
if(all(gid < size)){
uint4 pixel = read_imageui(input, sampler, gid);
float4 color = convert_float4(pixel) / 255;
color = pow(color, (float4)(gamma));
pixel = convert_uint4_rte(color * 255);
write_imageui(output, gid, pixel);
}
}
이제 코드를 단계별로 살펴 보겠습니다. 첫 번째 줄은 sampler_t 유형의 __constant 메모리 영역에 변수를 만듭니다. 이 샘플러는 이미지 데이터에 대한 액세스를 추가로 지정하는 데 사용됩니다. 전체 문서는 크로노스 문서를 참조하십시오.
우리는 커널을 호출하기 전에 입력을 read_only로 할당하고 write_only로 출력을 할당하므로 여기에 수정자를 추가합니다.
image2d와 image3d는 항상 전역 메모리에 할당되므로 여기서 __global 수정자를 생략 할 수 있습니다. 우리의 감마 값은 __constant 메모리에 저장되어 있으므로이를 지정합니다.
그런 다음 감마 보정 할 픽셀을 정의하는 스레드 ID를 얻습니다. 또한 스레드가 할당되지 않은 메모리에 액세스하지 않도록 크기를 쿼리합니다. 잊어 버리면 커널을 충돌시킬 수 있습니다.
우리가 합법적 인 스레드라는 것을 확인한 후에는 입력 이미지에서 픽셀을 읽습니다. 그런 다음 소수점의 손실을 피하기 위해 부동 소수점으로 변환하고 계산을 수행 한 다음 다시 변환하여 출력에 씁니다.