diff --git a/backends/ispc.c b/backends/ispc.c index 0073ea3e..fb7a8c72 100644 --- a/backends/ispc.c +++ b/backends/ispc.c @@ -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, @@ -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); diff --git a/backends/opencl.c b/backends/opencl.c index f1f7ea19..aa83d5b0 100644 --- a/backends/opencl.c +++ b/backends/opencl.c @@ -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; }; @@ -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); @@ -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); \ @@ -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); \ } @@ -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) @@ -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; @@ -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; } diff --git a/backends/sycl.cpp b/backends/sycl.cpp index b985ee27..de6f2f22 100644 --- a/backends/sycl.cpp +++ b/backends/sycl.cpp @@ -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; @@ -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) { @@ -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; diff --git a/backends/unified-cuda-hip-impl.h b/backends/unified-cuda-hip-impl.h index 48165a1c..40a4dc44 100644 --- a/backends/unified-cuda-hip-impl.h +++ b/backends/unified-cuda-hip-impl.h @@ -57,7 +57,7 @@ #define backend_t TOKEN_PASTE(DRIVER, _backend_t) struct backend_t { - int device_id; + int device; backendDeviceProp_t prop; }; @@ -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, @@ -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) \ { \ @@ -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; @@ -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; } @@ -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 diff --git a/include/nomp-impl.h b/include/nomp-impl.h index 7b4601a2..cba72f00 100644 --- a/include/nomp-impl.h +++ b/include/nomp-impl.h @@ -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; @@ -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, @@ -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 @@ -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 @@ -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 @@ -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 @@ -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 } diff --git a/src/nomp.c b/src/nomp.c index ffd1dd37..ed2f13c7 100644 --- a/src/nomp.c +++ b/src/nomp.c @@ -25,9 +25,9 @@ static inline int check_cmd_line(struct nomp_backend_t *bnd, int argc, if (!strncmp("--nomp-backend", argv[i], NOMP_MAX_BUFSIZ)) { strncpy(bnd->backend, argv[i + 1], NOMP_MAX_BUFSIZ); } else if (!strncmp("--nomp-platform", argv[i], NOMP_MAX_BUFSIZ)) { - bnd->platform_id = nomp_str_toui(argv[i + 1], NOMP_MAX_BUFSIZ); + bnd->platform = nomp_str_toui(argv[i + 1], NOMP_MAX_BUFSIZ); } else if (!strncmp("--nomp-device", argv[i], NOMP_MAX_BUFSIZ)) { - bnd->device_id = nomp_str_toui(argv[i + 1], NOMP_MAX_BUFSIZ); + bnd->device = nomp_str_toui(argv[i + 1], NOMP_MAX_BUFSIZ); } else if (!strncmp("--nomp-verbose", argv[i], NOMP_MAX_BUFSIZ)) { bnd->verbose = nomp_str_toui(argv[i + 1], NOMP_MAX_BUFSIZ); } else if (!strncmp("--nomp-profile", argv[i], NOMP_MAX_BUFSIZ)) { @@ -49,10 +49,10 @@ static inline int check_cmd_line(struct nomp_backend_t *bnd, int argc, static inline int check_env_vars(struct nomp_backend_t *bnd) { char *tmp = getenv("NOMP_PLATFORM"); if (tmp) - bnd->platform_id = nomp_str_toui(tmp, NOMP_MAX_BUFSIZ); + bnd->platform = nomp_str_toui(tmp, NOMP_MAX_BUFSIZ); if ((tmp = getenv("NOMP_DEVICE"))) - bnd->device_id = nomp_str_toui(tmp, NOMP_MAX_BUFSIZ); + bnd->device = nomp_str_toui(tmp, NOMP_MAX_BUFSIZ); if ((tmp = getenv("NOMP_VERBOSE"))) bnd->verbose = nomp_str_toui(tmp, NOMP_MAX_BUFSIZ); @@ -76,7 +76,7 @@ static inline int init_configs(int argc, const char **argv, struct nomp_backend_t *bnd) { // verbose, profile, device and platform id are all initialized to zero. // Everything else has to be set by user explicitly. - bnd->verbose = bnd->profile = bnd->device_id = bnd->platform_id = 0; + bnd->verbose = bnd->profile = bnd->device = bnd->platform = 0; strcpy(bnd->backend, ""), strcpy(bnd->install_dir, ""); nomp_check(check_cmd_line(bnd, argc, argv)); @@ -94,8 +94,8 @@ static inline int init_configs(int argc, const char **argv, check_if_valid(bnd->verbose < 0, "--nomp-verbose", "NOMP_VERBOSE"); check_if_valid(bnd->profile < 0, "--nomp-profile", "NOMP_PROFILE"); - check_if_valid(bnd->platform_id < 0, "--nomp-platform", "NOMP_PLATFORM"); - check_if_valid(bnd->device_id < 0, "--nomp-device", "NOMP_DEVICE"); + check_if_valid(bnd->platform < 0, "--nomp-platform", "NOMP_PLATFORM"); + check_if_valid(bnd->device < 0, "--nomp-device", "NOMP_DEVICE"); check_if_valid(strlen(bnd->backend) == 0, "--nomp-backend", "NOMP_BACKEND"); check_if_valid(strlen(bnd->install_dir) == 0, "--nomp-install-dir", "NOMP_INSTALL_DIR"); @@ -105,7 +105,7 @@ static inline int init_configs(int argc, const char **argv, // Append nomp python directory to sys.path. char abs_dir[PATH_MAX + 32]; strncpy(abs_dir, bnd->install_dir, PATH_MAX); - strncat(abs_dir, "/python", 32); + strncat(abs_dir, "/python", 16); nomp_check(nomp_py_append_to_sys_path(abs_dir)); return 0; } @@ -137,23 +137,23 @@ static inline int init_backend(struct nomp_backend_t *bnd) { if (strncmp(bnd->backend, "opencl", NOMP_MAX_BUFSIZ) == 0) { #if defined(OPENCL_ENABLED) - nomp_check(opencl_init(&nomp, bnd->platform_id, bnd->device_id)); + nomp_check(opencl_init(&nomp, bnd->platform, bnd->device)); #endif } else if (strncmp(bnd->backend, "cuda", NOMP_MAX_BUFSIZ) == 0) { #if defined(CUDA_ENABLED) - nomp_check(cuda_init(&nomp, bnd->platform_id, bnd->device_id)); + nomp_check(cuda_init(&nomp, bnd->platform, bnd->device)); #endif } else if (strncmp(bnd->backend, "hip", NOMP_MAX_BUFSIZ) == 0) { #if defined(HIP_ENABLED) - nomp_check(hip_init(&nomp, bnd->platform_id, bnd->device_id)); + nomp_check(hip_init(&nomp, bnd->platform, bnd->device)); #endif } else if (strncmp(bnd->backend, "sycl", NOMP_MAX_BUFSIZ) == 0) { #if defined(SYCL_ENABLED) - nomp_check(sycl_init(&nomp, bnd->platform_id, bnd->device_id)); + nomp_check(sycl_init(&nomp, bnd->platform, bnd->device)); #endif } else if (strncmp(bnd->backend, "ispc", NOMP_MAX_BUFSIZ) == 0) { #if defined(ISPC_ENABLED) - nomp_check(ispc_init(&nomp, bnd->platform_id, bnd->device_id)); + nomp_check(ispc_init(&nomp, bnd->platform, bnd->device)); #endif } else { return nomp_log(NOMP_USER_INPUT_IS_INVALID, NOMP_ERROR, diff --git a/tests/nomp-api-020.c b/tests/nomp-api-020.c index 0510efb0..4f90863a 100644 --- a/tests/nomp-api-020.c +++ b/tests/nomp-api-020.c @@ -50,7 +50,7 @@ static int test_missing_argument(const char *backend) { } int main(int argc, const char *argv[]) { - char backend[NOMP_TEST_MAX_BUFSIZ]; + char backend[NOMP_TEST_MAX_BUFSIZ + 1]; for (unsigned i = 0; i < (unsigned)argc; i++) { if (strncmp(argv[i], "--nomp-backend", NOMP_TEST_MAX_BUFSIZ) == 0) { assert(i + 1 < (unsigned)argc); diff --git a/tests/nomp-api-021.c b/tests/nomp-api-021.c index 316566de..cc21929f 100644 --- a/tests/nomp-api-021.c +++ b/tests/nomp-api-021.c @@ -35,7 +35,7 @@ static int test_invalid_nomp_backend(int argc, const char **argv) { } // NOMP_PLATFORM environment variable with invalid value. -static int test_invalid_platform_id(int argc, const char **argv) { +static int test_invalid_platform(int argc, const char **argv) { char *platform = NULL; set_test_env(platform, "NOMP_PLATFORM", "invalid"); @@ -53,7 +53,7 @@ static int test_invalid_platform_id(int argc, const char **argv) { } // NOMP_DEVICE environment variable with invalid value. -static int test_invalid_device_id(int argc, const char **argv) { +static int test_invalid_device(int argc, const char **argv) { char *device = NULL; set_test_env(device, "NOMP_DEVICE", "invalid"); @@ -96,7 +96,7 @@ static int test_invalid_nomp_verbose(int argc, const char **argv) { } // Run with a valid NOMP_PLATFORM environment variable. -static int test_valid_platform_id(int argc, const char **argv) { +static int test_valid_platform(int argc, const char **argv) { char *platform = NULL; set_test_env(platform, "NOMP_PLATFORM", "0"); @@ -109,7 +109,7 @@ static int test_valid_platform_id(int argc, const char **argv) { } // Run with a valid NOMP_DEVICE environment variable. -static int test_valid_device_id(int argc, const char **argv) { +static int test_valid_device(int argc, const char **argv) { char *device = NULL; set_test_env(device, "NOMP_DEVICE", "0"); @@ -124,8 +124,8 @@ static int test_valid_device_id(int argc, const char **argv) { int main(int argc, const char *argv[]) { int err = 0; err |= SUBTEST(test_invalid_nomp_backend, argc, argv); - err |= SUBTEST(test_invalid_platform_id, argc, argv); - err |= SUBTEST(test_invalid_device_id, argc, argv); + err |= SUBTEST(test_invalid_platform, argc, argv); + err |= SUBTEST(test_invalid_device, argc, argv); err |= SUBTEST(test_invalid_nomp_verbose, argc, argv); nomp_test_assert(argc <= 64); @@ -137,7 +137,7 @@ int main(int argc, const char *argv[]) { if (strncmp(argv[i], "--nomp-platform", NOMP_TEST_MAX_BUFSIZ)) argvn[argcn] = strndup(argv[i], NOMP_TEST_MAX_BUFSIZ), argcn++; } - err |= SUBTEST(test_valid_platform_id, argcn, (const char **)argvn); + err |= SUBTEST(test_valid_platform, argcn, (const char **)argvn); for (unsigned i = 0; i < argcn; i++) nomp_free(&argvn[i]); @@ -147,7 +147,7 @@ int main(int argc, const char *argv[]) { if (strncmp(argv[i], "--nomp-device", NOMP_TEST_MAX_BUFSIZ)) argvn[argcn] = strndup(argv[i], NOMP_TEST_MAX_BUFSIZ), argcn++; } - err |= SUBTEST(test_valid_device_id, argcn, (const char **)argvn); + err |= SUBTEST(test_valid_device, argcn, (const char **)argvn); for (unsigned i = 0; i < argcn; i++) nomp_free(&argvn[i]); diff --git a/tests/nomp-api-200-impl.h b/tests/nomp-api-200-impl.h index d1c484d2..3a85cdec 100644 --- a/tests/nomp-api-200-impl.h +++ b/tests/nomp-api-200-impl.h @@ -1,5 +1,58 @@ #include "nomp-test.h" +#define nomp_api_200_aux_ui TOKEN_PASTE(nomp_api_200_aux_ui, TEST_SUFFIX) +static int nomp_api_200_aux_ui(const char *fmt, TEST_TYPE *a, TEST_TYPE *b, + unsigned n) { + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_TO)); + nomp_test_chk(nomp_update(b, 0, n, sizeof(TEST_TYPE), NOMP_TO)); + + int id = -1; + const char *clauses[4] = {"transform", "nomp-api-200", "transform", 0}; + char *knl = generate_knl(fmt, 2, TOSTRING(TEST_TYPE), TOSTRING(TEST_TYPE)); + nomp_test_chk(nomp_jit(&id, knl, clauses, 3, "a", sizeof(TEST_TYPE), NOMP_PTR, + "b", sizeof(TEST_TYPE), NOMP_PTR, "N", + sizeof(unsigned), NOMP_UINT)); + nomp_free(&knl); + + nomp_test_chk(nomp_run(id, a, b, &n)); + + nomp_test_chk(nomp_sync()); + + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_FROM)); + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_FREE)); + nomp_test_chk(nomp_update(b, 0, n, sizeof(TEST_TYPE), NOMP_FREE)); + + return 0; +} + +#define nomp_api_200_add_ui TOKEN_PASTE(nomp_api_200_add_ui, TEST_SUFFIX) +static int nomp_api_200_add_ui(unsigned n) { + nomp_test_assert(n <= TEST_MAX_SIZE); + + TEST_TYPE a[TEST_MAX_SIZE], b[TEST_MAX_SIZE]; + for (unsigned i = 0; i < n; i++) + a[i] = n - i, b[i] = i; + + const char *knl_fmt = + "void foo(%s *a, %s *b, int N) { \n" + " for (int i = 0; i < N; i++) \n" + " a[i] += b[i]; \n" + "} \n"; + nomp_api_200_aux_ui(knl_fmt, a, b, n); + +#if defined(TEST_TOL) + for (unsigned i = 0; i < n; i++) + nomp_test_assert(fabs(a[i] - n) < TEST_TOL); +#else + for (unsigned i = 0; i < n; i++) + nomp_test_assert(a[i] == (TEST_TYPE)n); +#endif + + return 0; +} +#undef nomp_api_200_add_ui +#undef nomp_api_200_aux_ui + #define nomp_api_200_aux TOKEN_PASTE(nomp_api_200_aux, TEST_SUFFIX) static int nomp_api_200_aux(const char *fmt, TEST_TYPE *a, TEST_TYPE *b, int n) { @@ -161,8 +214,8 @@ static int nomp_api_200_square(unsigned n) { } #undef nomp_api_200_square -#define nomp_api_200_linear TOKEN_PASTE(nomp_api_200_linear, TEST_SUFFIX) -static int nomp_api_200_linear(unsigned n) { +#define nomp_api_200_assign TOKEN_PASTE(nomp_api_200_assign, TEST_SUFFIX) +static int nomp_api_200_assign(unsigned n) { nomp_test_assert(n <= TEST_MAX_SIZE); TEST_TYPE a[TEST_MAX_SIZE] = {0}, b[TEST_MAX_SIZE] = {1, 2, 3, 4, 5}; @@ -184,5 +237,5 @@ static int nomp_api_200_linear(unsigned n) { return 0; } -#undef nomp_api_200_linear +#undef nomp_api_200_assign #undef nomp_api_200_aux diff --git a/tests/nomp-api-200.c b/tests/nomp-api-200.c index a4580344..d2d4db9c 100644 --- a/tests/nomp-api-200.c +++ b/tests/nomp-api-200.c @@ -11,6 +11,13 @@ static int test_vector_addition(void) { return err; } +static int test_vector_addition_unsigned_bounds(void) { + int err = 0; + TEST_BUILTIN_TYPES(200_add_ui, 10) + TEST_BUILTIN_TYPES(200_add_ui, 50) + return err; +} + static int test_vector_subtraction(void) { int err = 0; TEST_BUILTIN_TYPES(200_sub, 10) @@ -39,10 +46,10 @@ static int test_vector_square_sum(void) { return err; } -static int test_vector_linear(void) { +static int test_vector_assign(void) { int err = 0; - TEST_BUILTIN_TYPES(200_linear, 10) - TEST_BUILTIN_TYPES(200_linear, 50) + TEST_BUILTIN_TYPES(200_assign, 10) + TEST_BUILTIN_TYPES(200_assign, 50) return err; } @@ -51,11 +58,12 @@ int main(int argc, const char *argv[]) { nomp_check(err); err |= SUBTEST(test_vector_addition); + err |= SUBTEST(test_vector_addition_unsigned_bounds); err |= SUBTEST(test_vector_subtraction); err |= SUBTEST(test_vector_multiplication_sum); err |= SUBTEST(test_vector_multiplication); err |= SUBTEST(test_vector_square_sum); - err |= SUBTEST(test_vector_linear); + err |= SUBTEST(test_vector_assign); err |= nomp_finalize(); nomp_check(err); diff --git a/tests/nomp-api-300-impl.h b/tests/nomp-api-300-impl.h index 0a908570..32d45c16 100644 --- a/tests/nomp-api-300-impl.h +++ b/tests/nomp-api-300-impl.h @@ -2,6 +2,62 @@ #define TEST_MAX_SIZE2 (TEST_MAX_SIZE * TEST_MAX_SIZE) +#define nomp_api_300_aux_ui TOKEN_PASTE(nomp_api_300_aux_ui, TEST_SUFFIX) +static int nomp_api_300_aux_ui(const char *fmt, TEST_TYPE *a, TEST_TYPE *b, + unsigned rows, unsigned cols, int n) { + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_TO)); + nomp_test_chk(nomp_update(b, 0, n, sizeof(TEST_TYPE), NOMP_TO)); + + int id = -1; + const char *clauses[4] = {"transform", "nomp-api-300", "transform", 0}; + char *knl = generate_knl(fmt, 2, TOSTRING(TEST_TYPE), TOSTRING(TEST_TYPE)); + nomp_test_chk(nomp_jit(&id, knl, clauses, 4, "a", sizeof(TEST_TYPE), NOMP_PTR, + "b", sizeof(TEST_TYPE), NOMP_PTR, "rows", + sizeof(unsigned), NOMP_UINT, "cols", sizeof(unsigned), + NOMP_UINT)); + nomp_free(&knl); + + nomp_test_chk(nomp_run(id, a, b, &rows, &cols)); + + nomp_test_chk(nomp_sync()); + + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_FROM)); + nomp_test_chk(nomp_update(a, 0, n, sizeof(TEST_TYPE), NOMP_FREE)); + nomp_test_chk(nomp_update(b, 0, n, sizeof(TEST_TYPE), NOMP_FREE)); + + return 0; +} + +#define nomp_api_300_add_ui TOKEN_PASTE(nomp_api_300_add_ui, TEST_SUFFIX) +static int nomp_api_300_add_ui(unsigned rows, unsigned cols) { + const unsigned n = rows * cols; + nomp_test_assert(n <= TEST_MAX_SIZE2); + + TEST_TYPE a[TEST_MAX_SIZE2], b[TEST_MAX_SIZE2]; + for (unsigned i = 0; i < n; i++) + a[i] = 2 * n - i, b[i] = i; + + const char *knl_fmt = + "void foo(%s *a, %s *b, unsigned rows, unsigned cols) { \n" + " for (int e = 0; e < rows; e++) \n" + " for (int i = 0; i < cols; i++) \n" + " a[e * cols + i] = a[e * cols + i] + b[e * cols + i]; \n" + "} \n"; + nomp_api_300_aux_ui(knl_fmt, a, b, rows, cols, n); + +#if defined(TEST_TOL) + for (unsigned i = 0; i < n; i++) + nomp_test_assert(fabs(a[i] - 2 * n) < TEST_TOL); +#else + for (unsigned i = 0; i < n; i++) + nomp_test_assert(a[i] == (TEST_TYPE)(2 * n)); +#endif + + return 0; +} +#undef nomp_api_300_add_ui +#undef nomp_api_300_aux_ui + #define nomp_api_300_aux TOKEN_PASTE(nomp_api_300_aux, TEST_SUFFIX) static int nomp_api_300_aux(const char *fmt, TEST_TYPE *a, TEST_TYPE *b, int rows, int cols, int n) { diff --git a/tests/nomp-api-300.c b/tests/nomp-api-300.c index 9e02e02c..25c56140 100644 --- a/tests/nomp-api-300.c +++ b/tests/nomp-api-300.c @@ -11,6 +11,13 @@ static int test_matrix_addition(void) { return err; } +static int test_matrix_addition_unsigned_bounds(void) { + int err = 0; + TEST_BUILTIN_TYPES(300_add_ui, 40, 5) + TEST_BUILTIN_TYPES(300_add_ui, 16, 16) + return err; +} + static int test_matrix_transpose(void) { int err = 0; TEST_BUILTIN_TYPES(300_transpose, 40, 5) @@ -37,6 +44,7 @@ int main(int argc, const char *argv[]) { nomp_check(err); err |= SUBTEST(test_matrix_addition); + err |= SUBTEST(test_matrix_addition_unsigned_bounds); err |= SUBTEST(test_matrix_transpose); err |= SUBTEST(test_matrix_vector_multiplication); err |= SUBTEST(test_matrix_matrix_multiplication); diff --git a/tests/nomp-api-500-impl.h b/tests/nomp-api-500-impl.h index c4c7b563..f047a8da 100644 --- a/tests/nomp-api-500-impl.h +++ b/tests/nomp-api-500-impl.h @@ -188,7 +188,7 @@ static int nomp_api_500_dot(unsigned N) { #if defined(TEST_TOL) nomp_test_assert(fabs(total - N * (2 * N - 1) * (N - 1) / 6) < TEST_TOL); #else - nomp_test_assert(total == N * (2 * N - 1) * (N - 1) / 6); + nomp_test_assert(total == (TEST_TYPE)(N * (2 * N - 1) * (N - 1) / 6)); #endif return 0; @@ -229,7 +229,7 @@ static int nomp_api_500_multiple_reductions(unsigned N, unsigned iterations) { "} \n"; TEST_TYPE a[TEST_MAX_SIZE]; - for (unsigned i = 1; i < iterations; ++i) { + for (unsigned i = 0; i < iterations; ++i) { for (unsigned j = 0; j < N; j++) a[j] = i * j; @@ -239,7 +239,7 @@ static int nomp_api_500_multiple_reductions(unsigned N, unsigned iterations) { #if defined(TEST_TOL) nomp_test_assert(fabs(total - (N - 1) * N * i / 2) < TEST_TOL); #else - nomp_test_assert(total == (N - 1) * N * i / 2); + nomp_test_assert(total == (TEST_TYPE)((N - 1) * N * i / 2)); #endif }