Skip to content

Add tests with unsigned loop bounds (kernel arguments) #227

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

Closed
thilinarmtb wants to merge 2 commits into from
Closed

Add tests with unsigned loop bounds (kernel arguments) #227

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
dd7982f
Add tests with unsigned loop bounds/kernel args
thilinarmtb Jul 5, 2023
19b9957
Fix warnings when compiling with gcc
thilinarmtb Jul 5, 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
8 changes: 4 additions & 4 deletions backends/ispc.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -188,7 +188,7 @@ static int ispc_chk_env(struct ispc_backend *ispc) {
static int ispc_sync(struct nomp_backend_t *bnd) { return 0; }

int ispc_init(struct nomp_backend_t *bnd, const int platform_type,
const int device_id) {
const int device) {
ispcrtSetErrorFunc(ispcrt_error);
if (platform_type < 0 | platform_type >= 2) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
Expand All @@ -198,11 +198,11 @@ int ispc_init(struct nomp_backend_t *bnd, const int platform_type,
uint32_t num_devices =
ispcrtGetDeviceCount(nomp_to_ispc_device[platform_type]);
chk_ispcrt("get device count", rt_error);
if (device_id < 0 || device_id >= num_devices) {
if (device < 0 || device >= num_devices) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
ERR_STR_USER_DEVICE_IS_INVALID, device_id);
ERR_STR_USER_DEVICE_IS_INVALID, device);
}
ISPCRTDevice device = ispcrtGetDevice(platform_type, device_id);
ISPCRTDevice device = ispcrtGetDevice(platform_type, device);
chk_ispcrt("device get", rt_error);

struct ispc_backend *ispc = bnd->bptr = nomp_calloc(struct ispc_backend, 1);
Expand Down
61 changes: 31 additions & 30 deletions backends/opencl.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ static const char *ERR_STR_OPENCL_FAILURE = "%s failed with error code: %d.";
}

struct opencl_backend_t {
cl_device_id device_id;
cl_device_id device;
cl_command_queue queue;
cl_context ctx;
};
Expand Down Expand Up @@ -81,10 +81,10 @@ static int opencl_knl_build(struct nomp_backend_t *bnd, struct nomp_prog_t *prg,
ocl_prg->prg = NULL, ocl_prg->knl = NULL;

size_t log_size;
clGetProgramBuildInfo(ocl_prg->prg, ocl->device_id, CL_PROGRAM_BUILD_LOG, 0,
clGetProgramBuildInfo(ocl_prg->prg, ocl->device, CL_PROGRAM_BUILD_LOG, 0,
NULL, &log_size);
char *log = nomp_calloc(char, log_size);
clGetProgramBuildInfo(ocl_prg->prg, ocl->device_id, CL_PROGRAM_BUILD_LOG,
clGetProgramBuildInfo(ocl_prg->prg, ocl->device, CL_PROGRAM_BUILD_LOG,
log_size, log, NULL);
int err = nomp_log(NOMP_OPENCL_FAILURE, NOMP_ERROR,
"clBuildProgram failed with error:\n %s.", log);
Expand Down Expand Up @@ -145,7 +145,8 @@ static int opencl_finalize(struct nomp_backend_t *bnd) {
return 0;
}

static int opencl_device_query(struct nomp_backend_t *bnd, cl_device_id id) {
static int opencl_device_query(struct nomp_backend_t *bnd,
cl_device_id device) {
#define set_string_aux(KEY, VAL) \
{ \
PyObject *obj = PyUnicode_FromString(VAL); \
Expand All @@ -156,7 +157,7 @@ static int opencl_device_query(struct nomp_backend_t *bnd, cl_device_id id) {
#define set_string_info(PARAM, KEY) \
{ \
char string[BUFSIZ]; \
check(clGetDeviceInfo(id, PARAM, sizeof(string), string, NULL), \
check(clGetDeviceInfo(device, PARAM, sizeof(string), string, NULL), \
"clGetDeviceInfo"); \
set_string_aux(KEY, string); \
}
Expand All @@ -166,7 +167,7 @@ static int opencl_device_query(struct nomp_backend_t *bnd, cl_device_id id) {
set_string_info(CL_DRIVER_VERSION, "device::driver");

cl_device_type type;
check(clGetDeviceInfo(id, CL_DEVICE_TYPE, sizeof(type), &type, NULL),
check(clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(type), &type, NULL),
"clGetDeviceInfo");
PyObject *obj = NULL;
if (type & CL_DEVICE_TYPE_CPU)
Expand All @@ -186,45 +187,41 @@ static int opencl_device_query(struct nomp_backend_t *bnd, cl_device_id id) {
return 0;
}

int opencl_init(struct nomp_backend_t *bnd, const int platform_id,
const int device_id) {
cl_uint num_platforms;
check(clGetPlatformIDs(0, NULL, &num_platforms), "clGetPlatformIDs");
if (platform_id < 0 | platform_id >= (int)num_platforms) {
int opencl_init(struct nomp_backend_t *bnd, const int platform,
const int device) {
cl_uint n_platforms;
check(clGetPlatformIDs(0, NULL, &n_platforms), "clGetPlatformIDs");
if (platform < 0 || platform >= (int)n_platforms) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
"Platform id %d provided to libnomp is not valid.",
platform_id);
platform);
}

cl_platform_id *cl_platforms = nomp_calloc(cl_platform_id, num_platforms);
check(clGetPlatformIDs(num_platforms, cl_platforms, &num_platforms),
cl_platform_id *platforms = nomp_calloc(cl_platform_id, n_platforms);
check(clGetPlatformIDs(n_platforms, platforms, &n_platforms),
"clGetPlatformIDs");
cl_platform_id platform = cl_platforms[platform_id];
nomp_free(&cl_platforms);

cl_uint num_devices;
check(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices),
cl_uint n_devices;
check(clGetDeviceIDs(platforms[platform], CL_DEVICE_TYPE_ALL, 0, NULL,
&n_devices),
"clGetDeviceIDs");
if (device_id < 0 || device_id >= (int)num_devices) {
if (device < 0 || device >= (int)n_devices) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
ERR_STR_USER_DEVICE_IS_INVALID, device_id);
ERR_STR_USER_DEVICE_IS_INVALID, device);
}

cl_device_id *cl_devices = nomp_calloc(cl_device_id, num_devices);
check(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, cl_devices,
&num_devices),
cl_device_id *devices = nomp_calloc(cl_device_id, n_devices);
check(clGetDeviceIDs(platforms[platform], CL_DEVICE_TYPE_ALL, n_devices,
devices, &n_devices),
"clGetDeviceIDs");
cl_device_id device = cl_devices[device_id];
nomp_free(&cl_devices);

nomp_check(opencl_device_query(bnd, device));

struct opencl_backend_t *ocl = nomp_calloc(struct opencl_backend_t, 1);
ocl->device_id = device;
ocl->device = devices[device];
cl_int err;
ocl->ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
ocl->ctx = clCreateContext(NULL, 1, &ocl->device, NULL, NULL, &err);
check(err, "clCreateContext");
ocl->queue = clCreateCommandQueueWithProperties(ocl->ctx, device, 0, &err);
ocl->queue =
clCreateCommandQueueWithProperties(ocl->ctx, ocl->device, 0, &err);
check(err, "clCreateCommandQueueWithProperties");

bnd->bptr = (void *)ocl;
Expand All @@ -235,6 +232,10 @@ int opencl_init(struct nomp_backend_t *bnd, const int platform_id,
bnd->sync = opencl_sync;
bnd->finalize = opencl_finalize;

nomp_check(opencl_device_query(bnd, ocl->device));

nomp_free(&devices), nomp_free(&platforms);

return 0;
}

Expand Down
24 changes: 12 additions & 12 deletions backends/sycl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ static const char *ERR_STR_SYCL_FAILURE = "SYCL backend failed with error: %s.";
}

struct sycl_backend {
sycl::device device_id;
sycl::device device;
sycl::queue queue;
sycl::context ctx;
char *compiler, *compiler_flags;
Expand All @@ -30,7 +30,7 @@ static int sycl_update(struct nomp_backend_t *bnd, struct nomp_mem_t *m,

if (op & NOMP_ALLOC) {
chk_sycl(m->bptr = sycl::malloc_device(NOMP_MEM_BYTES(start, end, usize),
sycl->device_id, sycl->ctx););
sycl->device, sycl->ctx););
}

if (op & NOMP_TO) {
Expand Down Expand Up @@ -139,27 +139,27 @@ static int check_env(struct sycl_backend *sycl) {
return 0;
}

int sycl_init(struct nomp_backend_t *bnd, const int platform_id,
const int device_id) {
int sycl_init(struct nomp_backend_t *bnd, const int platform,
const int device) {
struct sycl_backend *sycl = nomp_calloc(struct sycl_backend, 1);
bnd->bptr = (void *)sycl;

auto sycl_platforms = sycl::platform().get_platforms();
if (platform_id < 0 | platform_id >= sycl_platforms.size()) {
if (platform < 0 | platform >= sycl_platforms.size()) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
"Platform id %d provided to libnomp is not valid.",
platform_id);
platform);
}

auto sycl_pdevices = sycl_platforms[platform_id].get_devices();
if (device_id < 0 || device_id >= sycl_pdevices.size()) {
auto sycl_pdevices = sycl_platforms[platform].get_devices();
if (device < 0 || device >= sycl_pdevices.size()) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
ERR_STR_USER_DEVICE_IS_INVALID, device_id);
ERR_STR_USER_DEVICE_IS_INVALID, device);
}

sycl->device_id = sycl_pdevices[device_id];
sycl->ctx = sycl::context(sycl->device_id);
sycl->queue = sycl::queue(sycl->ctx, sycl->device_id);
sycl->device = sycl_pdevices[device];
sycl->ctx = sycl::context(sycl->device);
sycl->queue = sycl::queue(sycl->ctx, sycl->device);
check_env(sycl);

bnd->update = sycl_update;
Expand Down
30 changes: 12 additions & 18 deletions backends/unified-cuda-hip-impl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,7 @@

#define backend_t TOKEN_PASTE(DRIVER, _backend_t)
struct backend_t {
int device_id;
int device;
backendDeviceProp_t prop;
};

Expand Down Expand Up @@ -104,11 +104,6 @@ static int backend_update(struct nomp_backend_t *bnd, struct nomp_mem_t *m,
return 0;
}

#define backend_update_ptr TOKEN_PASTE(DRIVER, _update_ptr)
static void backend_update_ptr(void **p, size_t *size, struct nomp_mem_t *m) {
*p = (void *)m->bptr, *size = sizeof(m->bptr);
}

#define backend_knl_build TOKEN_PASTE(DRIVER, _knl_build)
static int backend_knl_build(struct nomp_backend_t *bnd,
struct nomp_prog_t *prg, const char *source,
Expand Down Expand Up @@ -193,9 +188,9 @@ static int backend_finalize(struct nomp_backend_t *bnd) {
}

#define backend_device_query TOKEN_PASTE(DRIVER, _device_query)
static int backend_device_query(struct nomp_backend_t *bnd, int device_id) {
static int backend_device_query(struct nomp_backend_t *bnd, int device) {
backendDeviceProp_t prop;
check_driver(backendGetDeviceProperties(&prop, device_id));
check_driver(backendGetDeviceProperties(&prop, device));

#define set_string_aux(KEY, VAL) \
{ \
Expand Down Expand Up @@ -232,23 +227,21 @@ static int backend_device_query(struct nomp_backend_t *bnd, int device_id) {
}

#define backend_init TOKEN_PASTE(DRIVER, _init)
int backend_init(struct nomp_backend_t *bnd, const int platform_id,
const int device_id) {
int backend_init(struct nomp_backend_t *bnd, const int platform,
const int device) {
int num_devices;
check_driver(backendGetDeviceCount(&num_devices));
if (device_id < 0 || device_id >= num_devices) {
if (device < 0 || device >= num_devices) {
return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR,
ERR_STR_USER_DEVICE_IS_INVALID, device_id);
ERR_STR_USER_DEVICE_IS_INVALID, device);
}

check_driver(backendSetDevice(device_id));
check_driver(backendSetDevice(device));
check_driver(backendFree(0));

nomp_check(backend_device_query(bnd, device_id));

struct backend_t *backend = nomp_calloc(struct backend_t, 1);
backend->device_id = device_id;
check_driver(backendGetDeviceProperties(&backend->prop, device_id));
backend->device = device;
check_driver(backendGetDeviceProperties(&backend->prop, device));

bnd->bptr = (void *)backend;
bnd->update = backend_update;
Expand All @@ -258,6 +251,8 @@ int backend_init(struct nomp_backend_t *bnd, const int platform_id,
bnd->sync = backend_sync;
bnd->finalize = backend_finalize;

nomp_check(backend_device_query(bnd, device));

return 0;
}

Expand All @@ -267,7 +262,6 @@ int backend_init(struct nomp_backend_t *bnd, const int platform_id,
#undef backend_knl_free
#undef backend_knl_run
#undef backend_knl_build
#undef backend_update_ptr
#undef backend_update
#undef backend_compile

Expand Down
42 changes: 21 additions & 21 deletions include/nomp-impl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ struct nomp_mem_t {
#define NOMP_MEM_BYTES(start, end, usize) (((end) - (start)) * (usize))

struct nomp_arg_t {
char name[NOMP_MAX_BUFSIZ];
char name[NOMP_MAX_BUFSIZ + 1];
size_t size;
unsigned type;
void *ptr;
Expand Down Expand Up @@ -69,8 +69,8 @@ struct nomp_prog_t {

struct nomp_backend_t {
// User configurations of the backend.
int platform_id, device_id, verbose, profile;
char backend[NOMP_MAX_BUFSIZ], install_dir[PATH_MAX];
int platform, device, verbose, profile;
char backend[NOMP_MAX_BUFSIZ + 1], install_dir[PATH_MAX + 1];
// Pointers to backend functions used for backend dispatch.
int (*update)(struct nomp_backend_t *, struct nomp_mem_t *,
const nomp_map_direction_t op, size_t start, size_t end,
Expand Down Expand Up @@ -110,12 +110,12 @@ extern "C" {
* occurred during the initialization, otherwise returns 0.
*
* @param[in] backend Target backend for code generation.
* @param[in] platform_id Target platform id.
* @param[in] device_id Target device id.
* @param[in] platform Target platform id.
* @param[in] device Target device id.
* @return int
*/
int opencl_init(struct nomp_backend_t *backend, const int platform_id,
const int device_id);
int opencl_init(struct nomp_backend_t *backend, const int platform,
const int device);

/**
* @ingroup nomp_backend_init
Expand All @@ -127,12 +127,12 @@ int opencl_init(struct nomp_backend_t *backend, const int platform_id,
* occurred during the initialization, otherwise returns 0.
*
* @param[in] backend Target backend for code generation.
* @param[in] platform_id Target platform id.
* @param[in] device_id Target device id.
* @param[in] platform Target platform id.
* @param[in] device Target device id.
* @return int
*/
int sycl_init(struct nomp_backend_t *backend, const int platform_id,
const int device_id);
int sycl_init(struct nomp_backend_t *backend, const int platform,
const int device);

/**
* @ingroup nomp_backend_init
Expand All @@ -143,12 +143,12 @@ int sycl_init(struct nomp_backend_t *backend, const int platform_id,
* error occurred during the initialization, otherwise returns 0.
*
* @param[in] backend Target backend for code generation.
* @param[in] platform_id Target platform id.
* @param[in] device_id Target device id.
* @param[in] platform Target platform id.
* @param[in] device Target device id.
* @return int
*/
int cuda_init(struct nomp_backend_t *backend, const int platform_id,
const int device_id);
int cuda_init(struct nomp_backend_t *backend, const int platform,
const int device);

/**
* @ingroup nomp_backend_init
Expand All @@ -159,12 +159,12 @@ int cuda_init(struct nomp_backend_t *backend, const int platform_id,
* error occurred during the initialization, otherwise returns 0.
*
* @param[in] backend Target backend for code generation.
* @param[in] platform_id Target platform id.
* @param[in] device_id Target device id.
* @param[in] platform Target platform id.
* @param[in] device Target device id.
* @return int
*/
int hip_init(struct nomp_backend_t *backend, const int platform_id,
const int device_id);
int hip_init(struct nomp_backend_t *backend, const int platform,
const int device);

/**
* @ingroup nomp_backend_init
Expand All @@ -176,11 +176,11 @@ int hip_init(struct nomp_backend_t *backend, const int platform_id,
*
* @param[in] backend Target backend for code generation.
* @param[in] platform_type Target platform type.
* @param[in] device_id Target device id.
* @param[in] device Target device id.
* @return int
*/
int ispc_init(struct nomp_backend_t *backend, const int platform_type,
const int device_id);
const int device);

#ifdef __cplusplus
}
Expand Down
Loading