Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

split and add preamble #36

Open
wants to merge 2 commits into
base: clesperanto_kernels
Choose a base branch
from
Open

split and add preamble #36

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
b37382d
split and add preamble
StRigaud Sep 25, 2023
cc1ac82
add define guards
StRigaud Sep 26, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
115 changes: 115 additions & 0 deletions preamble/preamble_def.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,115 @@
#ifndef PREAMBLE_DEFINE
#define PREAMBLE_DEFINE

#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
// #pragma OPENCL EXTENSION cl_amd_printf : enable

#ifndef M_PI
#define M_PI 3.14159265358979323846f /* pi */
#endif

#ifndef M_LOG2E
#define M_LOG2E 1.4426950408889634074f /* log_2 e */
#endif

#ifndef M_LOG10E
#define M_LOG10E 0.43429448190325182765f /* log_10 e */
#endif

#ifndef M_LN2
#define M_LN2 0.69314718055994530942f /* log_e 2 */
#endif

#ifndef M_LN10
#define M_LN10 2.30258509299404568402f /* log_e 10 */
#endif

inline uchar clij_convert_uchar_sat(float value) {
if (value > 255) {
return 255;
}
if (value < 0) {
return 0;
}
return (uchar)value;
}

inline char clij_convert_char_sat(float value) {
if (value > 127) {
return 127;
}
if (value < -128) {
return -128;
}
return (char)value;
}

inline ushort clij_convert_ushort_sat(float value) {
if (value > 65535) {
return 65535;
}
if (value < 0) {
return 0;
}
return (ushort)value;
}

inline short clij_convert_short_sat(float value) {
if (value > 32767) {
return 32767;
}
if (value < -32768) {
return -32768;
}
return (short)value;
}

inline uint clij_convert_uint_sat(float value) {
if (value > 4294967295) {
return 4294967295;
}
if (value < 0) {
return 0;
}
return (uint)value;
}

inline int clij_convert_int_sat(float value) {
if (value > 2147483647) {
return 2147483647;
}
if (value < -2147483648) {
return -2147483648;
}
return (int)value;
}

inline uint clij_convert_ulong_sat(float value) {
if (value > 18446744073709551615) {
return 18446744073709551615;
}
if (value < 0) {
return 0;
}
return (ulong)value;
}

inline int clij_convert_long_sat(float value) {
if (value > 9223372036854775807) {
return 9223372036854775807;
}
if (value < -9223372036854775808 ) {
return -9223372036854775808 ;
}
return (long)value;
}

inline float clij_convert_float_sat(float value) {
return value;
}

#define READ_IMAGE(a,b,c) READ_ ## a ## _IMAGE(a,b,c)
#define WRITE_IMAGE(a,b,c) WRITE_ ## a ## _IMAGE(a,b,c)

#endif // PREAMBLE_DEFINE
171 changes: 171 additions & 0 deletions preamble/preamble_def.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,171 @@
#ifndef PREAMBLE_DEFINE
#define PREAMBLE_DEFINE

#define sampler_t int

#define FLT_MIN 1.19209e-07
#define FLT_MAX 1e+37
#define MAX_ARRAY_SIZE 1000

#define uchar unsigned char
#define ushort unsigned short
#define uint unsigned int
#define ulong unsigned long

__device__ inline int2 operator+(int2 a, int2 b)
{
return make_int2(a.x + b.x, a.y + b.y);
}

__device__ inline int4 operator+(int4 a, int4 b)
{
return make_int4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
}

__device__ inline int2 operator*(int b, int2 a)
{
return make_int2(b * a.x, b * a.y);
}

__device__ inline int4 operator*(int b, int4 a)
{
return make_int4(b * a.x, b * a.y, b * a.z, b * a.w);
}

__device__ inline float pow(float x, int y) {
return pow(float(x), float(y));
}

__device__ inline float2 sqrt(float2 a) {
return make_float2(sqrt(a.x), sqrt(a.y));
}

__device__ inline float4 cross(float4 a, float4 b)
{
return make_float4(a.y * b.z - a.z * b.y, a.z * b.x - a.x * b.z, a.x * b.y - a.y * b.x, 0);
}

__device__ inline float dot(float4 a, float4 b)
{
return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w;
}

__device__ inline float length(float4 v)
{
return sqrtf(dot(v, v));
}

__device__ inline unsigned int atomic_add(unsigned int* address, unsigned int value) {
return atomicAdd(address, value);
}

__device__ inline uchar clij_convert_uchar_sat(float value) {
if (value > 255) {
return 255;
}
if (value < 0) {
return 0;
}
return (uchar)value;
}

__device__ inline char clij_convert_char_sat(float value) {
if (value > 127) {
return 127;
}
if (value < -128) {
return -128;
}
return (char)value;
}

__device__ inline ushort clij_convert_ushort_sat(float value) {
if (value > 65535) {
return 65535;
}
if (value < 0) {
return 0;
}
return (ushort)value;
}

__device__ inline short clij_convert_short_sat(float value) {
if (value > 32767) {
return 32767;
}
if (value < -32768) {
return -32768;
}
return (short)value;
}

__device__ inline uint clij_convert_uint_sat(float value) {
if (value > 4294967295) {
return 4294967295;
}
if (value < 0) {
return 0;
}
return (uint)value;
}

__device__ inline uint convert_uint_sat(float value) {
if (value > 4294967295) {
return 4294967295;
}
if (value < 0) {
return 0;
}
return (uint)value;
}

__device__ inline int clij_convert_int_sat(float value) {
if (value > 2147483647) {
return 2147483647;
}
if (value < -2147483648) {
return -2147483648;
}
return (int)value;
}

__device__ inline uint clij_convert_ulong_sat(float value) {
if (value > 18446744073709551615) {
return 18446744073709551615;
}
if (value < 0) {
return 0;
}
return (ulong)value;
}

__device__ inline int clij_convert_long_sat(float value) {
if (value > 9223372036854775807) {
return 9223372036854775807;
}
if (value < -9223372036854775808 ) {
return -9223372036854775808 ;
}
return (long)value;
}

__device__ inline float clij_convert_float_sat(float value) {
return value;
}

#define READ_IMAGE(a,b,c) READ_ ## a ## _IMAGE(a,b,c)
#define WRITE_IMAGE(a,b,c) WRITE_ ## a ## _IMAGE(a,b,c)

#define GET_IMAGE_WIDTH(image_key) IMAGE_SIZE_ ## image_key ## _WIDTH
#define GET_IMAGE_HEIGHT(image_key) IMAGE_SIZE_ ## image_key ## _HEIGHT
#define GET_IMAGE_DEPTH(image_key) IMAGE_SIZE_ ## image_key ## _DEPTH

#define CLK_NORMALIZED_COORDS_FALSE 1
#define CLK_ADDRESS_CLAMP_TO_EDGE 2
#define CLK_FILTER_NEAREST 4
#define CLK_NORMALIZED_COORDS_TRUE 8
#define CLK_ADDRESS_CLAMP 16
#define CLK_FILTER_LINEAR 32
#define CLK_ADDRESS_NONE 64

#endif // PREAMBLE_DEFINE
80 changes: 80 additions & 0 deletions preamble/preamble_io.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,80 @@
#ifdef BUFFER_READ_WRITE
#define BUFFER_READ_WRITE
#define MINMAX_TYPE int

inline {pixel_type}2 read_buffer3d{short_pixel_type}(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global {pixel_type} * buffer_var, sampler_t sampler, int4 position )
{
int4 pos = (int4){position.x, position.y, position.z, 0};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
}
int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
return ({pixel_type}2){0, 0};
}
return ({pixel_type}2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer3d{short_pixel_type}(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global {pixel_type} * buffer_var, int4 pos, short value )
{
int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
return;
}
buffer_var[pos_in_buffer] = value;
}

inline {pixel_type}2 read_buffer2d{short_pixel_type}(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global {pixel_type} * buffer_var, sampler_t sampler, int2 position )
{
int2 pos = (int2){position.x, position.y};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
}
int pos_in_buffer = pos.x + pos.y * read_buffer_width;
if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
return ({pixel_type}2){0, 0};
}
return ({pixel_type}2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer2d{short_pixel_type}(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global {pixel_type} * buffer_var, int2 pos, ushort value )
{
int pos_in_buffer = pos.x + pos.y * write_buffer_width;
if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
return;
}
buffer_var[pos_in_buffer] = value;
}

inline {pixel_type}2 read_buffer1d{short_pixel_type}(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global {pixel_type} * buffer_var, sampler_t sampler, int position )
{
int pos = (int){position};
if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
}
int pos_in_buffer = pos;
if (pos < 0 || pos >= read_buffer_width) {
return ({pixel_type}2){0,0};
}
return ({pixel_type}2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer1d{short_pixel_type}(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global {pixel_type} * buffer_var, int pos, short value )
{
int pos_in_buffer = pos;
if (pos < 0 || pos >= write_buffer_width) {
return;
}
buffer_var[pos_in_buffer] = value;
}

#endif // BUFFER_READ_WRITE
Loading