/* StarPU --- Runtime system for heterogeneous multicore architectures. * * Copyright (C) 2009-2016 Université de Bordeaux * Copyright (C) 2010, 2011, 2012, 2013, 2014, 2015, 2016 CNRS * Copyright (C) 2013 Corentin Salingue * * StarPU is free software; you can redistribute it and/or modify * it under the terms of the GNU Lesser General Public License as published by * the Free Software Foundation; either version 2.1 of the License, or (at * your option) any later version. * * StarPU is distributed in the hope that it will be useful, but * WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. * * See the GNU Lesser General Public License in COPYING.LGPL for more details. */ #ifdef STARPU_USE_CUDA #ifndef _GNU_SOURCE #define _GNU_SOURCE #endif #include #endif #include #include #include #include #include #include #ifdef HAVE_UNISTD_H #include #endif #include #include #include #include #include #ifdef STARPU_USE_OPENCL #include #endif #ifdef STARPU_HAVE_WINDOWS #include #endif #if defined(HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX) && HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX #include #endif #define SIZE (32*1024*1024*sizeof(char)) #define NITER 32 #ifndef STARPU_SIMGRID static void _starpu_bus_force_sampling(void); #endif /* timing is in µs per byte (i.e. slowness, inverse of bandwidth) */ struct dev_timing { int cpu_id; double timing_htod; double latency_htod; double timing_dtoh; double latency_dtoh; }; /* TODO: measure latency */ static double bandwidth_matrix[STARPU_MAXNODES][STARPU_MAXNODES]; static double latency_matrix[STARPU_MAXNODES][STARPU_MAXNODES]; static unsigned was_benchmarked = 0; #ifndef STARPU_SIMGRID static unsigned ncpus = 0; #endif static unsigned ncuda = 0; static unsigned nopencl = 0; static unsigned nmic = 0; static unsigned nmpi_ms = 0; /* Benchmarking the performance of the bus */ #ifndef STARPU_SIMGRID static uint64_t cuda_size[STARPU_MAXCUDADEVS]; #endif #ifdef STARPU_USE_CUDA /* preference order of cores (logical indexes) */ static int cuda_affinity_matrix[STARPU_MAXCUDADEVS][STARPU_MAXCPUS]; #ifndef STARPU_SIMGRID static double cudadev_timing_htod[STARPU_MAXNODES] = {0.0}; static double cudadev_latency_htod[STARPU_MAXNODES] = {0.0}; static double cudadev_timing_dtoh[STARPU_MAXNODES] = {0.0}; static double cudadev_latency_dtoh[STARPU_MAXNODES] = {0.0}; #ifdef HAVE_CUDA_MEMCPY_PEER static double cudadev_timing_dtod[STARPU_MAXNODES][STARPU_MAXNODES] = {{0.0}}; static double cudadev_latency_dtod[STARPU_MAXNODES][STARPU_MAXNODES] = {{0.0}}; #endif #endif static struct dev_timing cudadev_timing_per_cpu[STARPU_MAXNODES*STARPU_MAXCPUS]; #endif #ifndef STARPU_SIMGRID static uint64_t opencl_size[STARPU_MAXCUDADEVS]; #endif #ifdef STARPU_USE_OPENCL /* preference order of cores (logical indexes) */ static int opencl_affinity_matrix[STARPU_MAXOPENCLDEVS][STARPU_MAXCPUS]; #ifndef STARPU_SIMGRID static double opencldev_timing_htod[STARPU_MAXNODES] = {0.0}; static double opencldev_latency_htod[STARPU_MAXNODES] = {0.0}; static double opencldev_timing_dtoh[STARPU_MAXNODES] = {0.0}; static double opencldev_latency_dtoh[STARPU_MAXNODES] = {0.0}; #endif static struct dev_timing opencldev_timing_per_cpu[STARPU_MAXNODES*STARPU_MAXCPUS]; #endif #ifdef STARPU_USE_MIC static double mic_time_host_to_device[STARPU_MAXNODES] = {0.0}; static double mic_time_device_to_host[STARPU_MAXNODES] = {0.0}; #endif /* STARPU_USE_MIC */ #ifdef STARPU_USE_MPI_MASTER_SLAVE static double mpi_time_host_to_device[STARPU_MAXNODES] = {0.0}; static double mpi_time_device_to_host[STARPU_MAXNODES] = {0.0}; static double mpi_latency_host_to_device[STARPU_MAXNODES] = {0.0}; static double mpi_latency_device_to_host[STARPU_MAXNODES] = {0.0}; #endif #ifdef STARPU_HAVE_HWLOC static hwloc_topology_t hwtopology; #endif #if (defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL)) && !defined(STARPU_SIMGRID) #ifdef STARPU_USE_CUDA static void measure_bandwidth_between_host_and_dev_on_cpu_with_cuda(int dev, int cpu, struct dev_timing *dev_timing_per_cpu) { struct _starpu_machine_config *config = _starpu_get_machine_config(); _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); size_t size = SIZE; /* Initialize CUDA context on the device */ /* We do not need to enable OpenGL interoperability at this point, * since we cleanly shutdown CUDA before returning. */ cudaSetDevice(dev); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* hack to force the initialization */ cudaFree(0); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Get the maximum size which can be allocated on the device */ struct cudaDeviceProp prop; cudaError_t cures; cures = cudaGetDeviceProperties(&prop, dev); if (STARPU_UNLIKELY(cures)) STARPU_CUDA_REPORT_ERROR(cures); cuda_size[dev] = prop.totalGlobalMem; if (size > prop.totalGlobalMem/4) size = prop.totalGlobalMem/4; /* Allocate a buffer on the device */ unsigned char *d_buffer; cures = cudaMalloc((void **)&d_buffer, size); STARPU_ASSERT(cures == cudaSuccess); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Allocate a buffer on the host */ unsigned char *h_buffer; cures = cudaHostAlloc((void **)&h_buffer, size, 0); STARPU_ASSERT(cures == cudaSuccess); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Fill them */ memset(h_buffer, 0, size); cudaMemset(d_buffer, 0, size); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); unsigned iter; double timing; double start; double end; /* Measure upload bandwidth */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpy(d_buffer, h_buffer, size, cudaMemcpyHostToDevice); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_htod = timing/NITER/size; /* Measure download bandwidth */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpy(h_buffer, d_buffer, size, cudaMemcpyDeviceToHost); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_dtoh = timing/NITER/size; /* Measure upload latency */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpy(d_buffer, h_buffer, 1, cudaMemcpyHostToDevice); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_htod = timing/NITER; /* Measure download latency */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpy(h_buffer, d_buffer, 1, cudaMemcpyDeviceToHost); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_dtoh = timing/NITER; /* Free buffers */ cudaFreeHost(h_buffer); cudaFree(d_buffer); cudaThreadExit(); } #ifdef HAVE_CUDA_MEMCPY_PEER static void measure_bandwidth_between_dev_and_dev_cuda(int src, int dst) { size_t size = SIZE; int can; /* Get the maximum size which can be allocated on the device */ struct cudaDeviceProp prop; cudaError_t cures; cures = cudaGetDeviceProperties(&prop, src); if (STARPU_UNLIKELY(cures)) STARPU_CUDA_REPORT_ERROR(cures); if (size > prop.totalGlobalMem/4) size = prop.totalGlobalMem/4; cures = cudaGetDeviceProperties(&prop, dst); if (STARPU_UNLIKELY(cures)) STARPU_CUDA_REPORT_ERROR(cures); if (size > prop.totalGlobalMem/4) size = prop.totalGlobalMem/4; /* Initialize CUDA context on the source */ /* We do not need to enable OpenGL interoperability at this point, * since we cleanly shutdown CUDA before returning. */ cudaSetDevice(src); if (starpu_get_env_number("STARPU_ENABLE_CUDA_GPU_GPU_DIRECT") != 0) { cures = cudaDeviceCanAccessPeer(&can, src, dst); if (!cures && can) { cures = cudaDeviceEnablePeerAccess(dst, 0); if (!cures) _STARPU_DISP("GPU-Direct %d -> %d\n", dst, src); } } /* Allocate a buffer on the device */ unsigned char *s_buffer; cures = cudaMalloc((void **)&s_buffer, size); STARPU_ASSERT(cures == cudaSuccess); cudaMemset(s_buffer, 0, size); /* Initialize CUDA context on the destination */ /* We do not need to enable OpenGL interoperability at this point, * since we cleanly shutdown CUDA before returning. */ cudaSetDevice(dst); if (starpu_get_env_number("STARPU_ENABLE_CUDA_GPU_GPU_DIRECT") != 0) { cures = cudaDeviceCanAccessPeer(&can, dst, src); if (!cures && can) { cures = cudaDeviceEnablePeerAccess(src, 0); if (!cures) _STARPU_DISP("GPU-Direct %d -> %d\n", src, dst); } } /* Allocate a buffer on the device */ unsigned char *d_buffer; cures = cudaMalloc((void **)&d_buffer, size); STARPU_ASSERT(cures == cudaSuccess); cudaMemset(d_buffer, 0, size); unsigned iter; double timing; double start; double end; /* Measure upload bandwidth */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpyPeer(d_buffer, dst, s_buffer, src, size); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; cudadev_timing_dtod[src+1][dst+1] = timing/NITER/size; /* Measure upload latency */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { cudaMemcpyPeer(d_buffer, dst, s_buffer, src, 1); cudaThreadSynchronize(); } end = starpu_timing_now(); timing = end - start; cudadev_latency_dtod[src+1][dst+1] = timing/NITER; /* Free buffers */ cudaFree(d_buffer); cudaSetDevice(src); cudaFree(s_buffer); cudaThreadExit(); } #endif #endif #ifdef STARPU_USE_OPENCL static void measure_bandwidth_between_host_and_dev_on_cpu_with_opencl(int dev, int cpu, struct dev_timing *dev_timing_per_cpu) { cl_context context; cl_command_queue queue; cl_int err=0; size_t size = SIZE; int not_initialized; struct _starpu_machine_config *config = _starpu_get_machine_config(); _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Is the context already initialised ? */ starpu_opencl_get_context(dev, &context); not_initialized = (context == NULL); if (not_initialized == 1) _starpu_opencl_init_context(dev); /* Get context and queue */ starpu_opencl_get_context(dev, &context); starpu_opencl_get_queue(dev, &queue); /* Get the maximum size which can be allocated on the device */ cl_device_id device; cl_ulong maxMemAllocSize, totalGlobalMem; starpu_opencl_get_device(dev, &device); err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxMemAllocSize), &maxMemAllocSize, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); if (size > (size_t)maxMemAllocSize/4) size = maxMemAllocSize/4; err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE , sizeof(totalGlobalMem), &totalGlobalMem, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); opencl_size[dev] = totalGlobalMem; if (_starpu_opencl_get_device_type(dev) == CL_DEVICE_TYPE_CPU) { /* Let's not use too much RAM when running OpenCL on a CPU: it * would make the OS swap like crazy. */ size /= 2; } /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Allocate a buffer on the device */ cl_mem d_buffer; d_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &err); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Allocate a buffer on the host */ unsigned char *h_buffer; h_buffer = (unsigned char *)malloc(size); STARPU_ASSERT(h_buffer); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); /* Fill them */ memset(h_buffer, 0, size); err = clEnqueueWriteBuffer(queue, d_buffer, CL_TRUE, 0, size, h_buffer, 0, NULL, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); clFinish(queue); /* hack to avoid third party libs to rebind threads */ _starpu_bind_thread_on_cpu(config, cpu, STARPU_NOWORKERID); unsigned iter; double timing; double start; double end; /* Measure upload bandwidth */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { err = clEnqueueWriteBuffer(queue, d_buffer, CL_TRUE, 0, size, h_buffer, 0, NULL, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); clFinish(queue); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_htod = timing/NITER/size; /* Measure download bandwidth */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { err = clEnqueueReadBuffer(queue, d_buffer, CL_TRUE, 0, size, h_buffer, 0, NULL, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); clFinish(queue); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_dtoh = timing/NITER/size; /* Measure upload latency */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { err = clEnqueueWriteBuffer(queue, d_buffer, CL_TRUE, 0, 1, h_buffer, 0, NULL, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); clFinish(queue); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_htod = timing/NITER; /* Measure download latency */ start = starpu_timing_now(); for (iter = 0; iter < NITER; iter++) { err = clEnqueueReadBuffer(queue, d_buffer, CL_TRUE, 0, 1, h_buffer, 0, NULL, NULL); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); clFinish(queue); } end = starpu_timing_now(); timing = end - start; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_dtoh = timing/NITER; /* Free buffers */ err = clReleaseMemObject(d_buffer); if (STARPU_UNLIKELY(err != CL_SUCCESS)) STARPU_OPENCL_REPORT_ERROR(err); free(h_buffer); /* Uninitiliaze OpenCL context on the device */ if (not_initialized == 1) _starpu_opencl_deinit_context(dev); } #endif /* NB: we want to sort the bandwidth by DECREASING order */ static int compar_dev_timing(const void *left_dev_timing, const void *right_dev_timing) { const struct dev_timing *left = (const struct dev_timing *)left_dev_timing; const struct dev_timing *right = (const struct dev_timing *)right_dev_timing; double left_dtoh = left->timing_dtoh; double left_htod = left->timing_htod; double right_dtoh = right->timing_dtoh; double right_htod = right->timing_htod; double timing_sum2_left = left_dtoh*left_dtoh + left_htod*left_htod; double timing_sum2_right = right_dtoh*right_dtoh + right_htod*right_htod; /* it's for a decreasing sorting */ return (timing_sum2_left > timing_sum2_right); } #ifdef STARPU_HAVE_HWLOC static int find_numa_node(hwloc_obj_t obj) { STARPU_ASSERT(obj); hwloc_obj_t current = obj; while (current->depth != HWLOC_OBJ_NODE) { current = current->parent; /* If we don't find a "node" obj before the root, this means * hwloc does not know whether there are numa nodes or not, so * we should not use a per-node sampling in that case. */ STARPU_ASSERT(current); } STARPU_ASSERT(current->depth == HWLOC_OBJ_NODE); return current->logical_index; } #endif static void measure_bandwidth_between_cpus_and_dev(int dev, struct dev_timing *dev_timing_per_cpu, char *type) { /* Either we have hwloc and we measure the bandwith between each GPU * and each NUMA node, or we don't have such NUMA information and we * measure the bandwith for each pair of (CPU, GPU), which is slower. * */ #ifdef STARPU_HAVE_HWLOC int cpu_depth = hwloc_get_type_depth(hwtopology, HWLOC_OBJ_PU); int nnuma_nodes = hwloc_get_nbobjs_by_depth(hwtopology, HWLOC_OBJ_NODE); /* If no NUMA node was found, we assume that we have a single memory * bank. */ const unsigned no_node_obj_was_found = (nnuma_nodes == 0); unsigned *is_available_per_numa_node = NULL; double *dev_timing_htod_per_numa_node = NULL; double *dev_latency_htod_per_numa_node = NULL; double *dev_timing_dtoh_per_numa_node = NULL; double *dev_latency_dtoh_per_numa_node = NULL; if (!no_node_obj_was_found) { is_available_per_numa_node = (unsigned *)malloc(nnuma_nodes * sizeof(unsigned)); STARPU_ASSERT(is_available_per_numa_node); dev_timing_htod_per_numa_node = (double *)malloc(nnuma_nodes * sizeof(double)); STARPU_ASSERT(dev_timing_htod_per_numa_node); dev_latency_htod_per_numa_node = (double *)malloc(nnuma_nodes * sizeof(double)); STARPU_ASSERT(dev_latency_htod_per_numa_node); dev_timing_dtoh_per_numa_node = (double *)malloc(nnuma_nodes * sizeof(double)); STARPU_ASSERT(dev_timing_dtoh_per_numa_node); dev_latency_dtoh_per_numa_node = (double *)malloc(nnuma_nodes * sizeof(double)); STARPU_ASSERT(dev_latency_dtoh_per_numa_node); memset(is_available_per_numa_node, 0, nnuma_nodes*sizeof(unsigned)); } #endif unsigned cpu; for (cpu = 0; cpu < ncpus; cpu++) { dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].cpu_id = cpu; #ifdef STARPU_HAVE_HWLOC int numa_id = 0; if (!no_node_obj_was_found) { hwloc_obj_t obj = hwloc_get_obj_by_depth(hwtopology, cpu_depth, cpu); numa_id = find_numa_node(obj); if (is_available_per_numa_node[numa_id]) { /* We reuse the previous numbers for that NUMA node */ dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_htod = dev_timing_htod_per_numa_node[numa_id]; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_htod = dev_latency_htod_per_numa_node[numa_id]; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_dtoh = dev_timing_dtoh_per_numa_node[numa_id]; dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_dtoh = dev_latency_dtoh_per_numa_node[numa_id]; continue; } } #endif #ifdef STARPU_USE_CUDA if (strncmp(type, "CUDA", 4) == 0) measure_bandwidth_between_host_and_dev_on_cpu_with_cuda(dev, cpu, dev_timing_per_cpu); #endif #ifdef STARPU_USE_OPENCL if (strncmp(type, "OpenCL", 6) == 0) measure_bandwidth_between_host_and_dev_on_cpu_with_opencl(dev, cpu, dev_timing_per_cpu); #endif #ifdef STARPU_HAVE_HWLOC if (!no_node_obj_was_found && !is_available_per_numa_node[numa_id]) { /* Save the results for that NUMA node */ dev_timing_htod_per_numa_node[numa_id] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_htod; dev_latency_htod_per_numa_node[numa_id] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_htod; dev_timing_dtoh_per_numa_node[numa_id] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_dtoh; dev_latency_dtoh_per_numa_node[numa_id] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].latency_dtoh; is_available_per_numa_node[numa_id] = 1; } #endif } #ifdef STARPU_HAVE_HWLOC if (!no_node_obj_was_found) { free(is_available_per_numa_node); free(dev_timing_htod_per_numa_node); free(dev_latency_htod_per_numa_node); free(dev_timing_dtoh_per_numa_node); free(dev_latency_dtoh_per_numa_node); } #endif /* STARPU_HAVE_HWLOC */ } static void measure_bandwidth_between_host_and_dev(int dev, double *dev_timing_htod, double *dev_latency_htod, double *dev_timing_dtoh, double *dev_latency_dtoh, struct dev_timing *dev_timing_per_cpu, char *type) { measure_bandwidth_between_cpus_and_dev(dev, dev_timing_per_cpu, type); /* sort the results */ qsort(&(dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS]), ncpus, sizeof(struct dev_timing), compar_dev_timing); #ifdef STARPU_VERBOSE unsigned cpu; for (cpu = 0; cpu < ncpus; cpu++) { unsigned current_cpu = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].cpu_id; double bandwidth_dtoh = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_dtoh; double bandwidth_htod = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+cpu].timing_htod; double bandwidth_sum2 = bandwidth_dtoh*bandwidth_dtoh + bandwidth_htod*bandwidth_htod; _STARPU_DISP("(%10s) BANDWIDTH GPU %d CPU %u - htod %f - dtoh %f - %f\n", type, dev, current_cpu, bandwidth_htod, bandwidth_dtoh, sqrt(bandwidth_sum2)); } unsigned best_cpu = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+0].cpu_id; _STARPU_DISP("(%10s) BANDWIDTH GPU %d BEST CPU %u\n", type, dev, best_cpu); #endif /* The results are sorted in a decreasing order, so that the best * measurement is currently the first entry. */ dev_timing_dtoh[dev+1] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+0].timing_dtoh; dev_latency_dtoh[dev+1] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+0].latency_dtoh; dev_timing_htod[dev+1] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+0].timing_htod; dev_latency_htod[dev+1] = dev_timing_per_cpu[(dev+1)*STARPU_MAXCPUS+0].latency_htod; } #endif /* defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) */ static void benchmark_all_gpu_devices(void) { #ifdef STARPU_SIMGRID _STARPU_DISP("can not measure bus in simgrid mode, please run starpu_calibrate_bus in non-simgrid mode to make sure the bus performance model was calibrated\n"); STARPU_ABORT(); #else /* !SIMGRID */ #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) || defined(STARPU_USE_MIC) || defined(STARPU_USE_MPI_MASTER_SLAVE) unsigned i; #endif #ifdef HAVE_CUDA_MEMCPY_PEER unsigned j; #endif _STARPU_DEBUG("Benchmarking the speed of the bus\n"); #ifdef STARPU_HAVE_HWLOC hwloc_topology_init(&hwtopology); hwloc_topology_load(hwtopology); #endif #ifdef STARPU_HAVE_HWLOC hwloc_bitmap_t former_cpuset = hwloc_bitmap_alloc(); hwloc_get_cpubind(hwtopology, former_cpuset, HWLOC_CPUBIND_THREAD); #elif __linux__ /* Save the current cpu binding */ cpu_set_t former_process_affinity; int ret; ret = sched_getaffinity(0, sizeof(former_process_affinity), &former_process_affinity); if (ret) { perror("sched_getaffinity"); STARPU_ABORT(); } #else #warning Missing binding support, StarPU will not be able to properly benchmark NUMA topology #endif struct _starpu_machine_config *config = _starpu_get_machine_config(); ncpus = _starpu_topology_get_nhwcpu(config); #ifdef STARPU_USE_CUDA ncuda = _starpu_get_cuda_device_count(); for (i = 0; i < ncuda; i++) { _STARPU_DISP("CUDA %d...\n", i); /* measure bandwidth between Host and Device i */ measure_bandwidth_between_host_and_dev(i, cudadev_timing_htod, cudadev_latency_htod, cudadev_timing_dtoh, cudadev_latency_dtoh, cudadev_timing_per_cpu, "CUDA"); } #ifdef HAVE_CUDA_MEMCPY_PEER for (i = 0; i < ncuda; i++) for (j = 0; j < ncuda; j++) if (i != j) { _STARPU_DISP("CUDA %d -> %d...\n", i, j); /* measure bandwidth between Host and Device i */ measure_bandwidth_between_dev_and_dev_cuda(i, j); } #endif #endif #ifdef STARPU_USE_OPENCL nopencl = _starpu_opencl_get_device_count(); for (i = 0; i < nopencl; i++) { _STARPU_DISP("OpenCL %d...\n", i); /* measure bandwith between Host and Device i */ measure_bandwidth_between_host_and_dev(i, opencldev_timing_htod, opencldev_latency_htod, opencldev_timing_dtoh, opencldev_latency_dtoh, opencldev_timing_per_cpu, "OpenCL"); } #endif #ifdef STARPU_USE_MIC /* TODO: implement real calibration ! For now we only put an arbitrary * value for each device during at the declaration as a bug fix, else * we get problems on heft scheduler */ nmic = _starpu_mic_src_get_device_count(); for (i = 0; i < STARPU_MAXNODES; i++) { mic_time_host_to_device[i] = 0.1; mic_time_device_to_host[i] = 0.1; } #endif /* STARPU_USE_MIC */ #ifdef STARPU_USE_MPI_MASTER_SLAVE _starpu_mpi_common_measure_bandwidth_latency(mpi_time_host_to_device, mpi_time_device_to_host, mpi_latency_host_to_device, mpi_latency_device_to_host); #endif /* STARPU_USE_MPI_MASTER_SLAVE */ #ifdef STARPU_HAVE_HWLOC hwloc_set_cpubind(hwtopology, former_cpuset, HWLOC_CPUBIND_THREAD); hwloc_bitmap_free(former_cpuset); #elif __linux__ /* Restore the former affinity */ ret = sched_setaffinity(0, sizeof(former_process_affinity), &former_process_affinity); if (ret) { perror("sched_setaffinity"); STARPU_ABORT(); } #endif #ifdef STARPU_HAVE_HWLOC hwloc_topology_destroy(hwtopology); #endif _STARPU_DEBUG("Benchmarking the speed of the bus is done.\n"); was_benchmarked = 1; #endif /* !SIMGRID */ } static void get_bus_path(const char *type, char *path, size_t maxlen) { char hostname[65]; _starpu_gethostname(hostname, sizeof(hostname)); snprintf(path, maxlen, "%s%s.%s", _starpu_get_perf_model_dir_bus(), hostname, type); } /* * Affinity */ static void get_affinity_path(char *path, size_t maxlen) { get_bus_path("affinity", path, maxlen); } #ifndef STARPU_SIMGRID static void load_bus_affinity_file_content(void) { #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) FILE *f; char path[256]; get_affinity_path(path, sizeof(path)); _STARPU_DEBUG("loading affinities from %s\n", path); f = fopen(path, "r"); STARPU_ASSERT(f); _starpu_frdlock(f); struct _starpu_machine_config *config = _starpu_get_machine_config(); ncpus = _starpu_topology_get_nhwcpu(config); unsigned gpu; #ifdef STARPU_USE_CUDA ncuda = _starpu_get_cuda_device_count(); for (gpu = 0; gpu < ncuda; gpu++) { int ret; unsigned dummy; _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &dummy); STARPU_ASSERT(ret == 1); STARPU_ASSERT(dummy == gpu); unsigned cpu; for (cpu = 0; cpu < ncpus; cpu++) { ret = fscanf(f, "%d\t", &cuda_affinity_matrix[gpu][cpu]); STARPU_ASSERT(ret == 1); } ret = fscanf(f, "\n"); STARPU_ASSERT(ret == 0); } #endif /* !STARPU_USE_CUDA */ #ifdef STARPU_USE_OPENCL nopencl = _starpu_opencl_get_device_count(); for (gpu = 0; gpu < nopencl; gpu++) { int ret; unsigned dummy; _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &dummy); STARPU_ASSERT(ret == 1); STARPU_ASSERT(dummy == gpu); unsigned cpu; for (cpu = 0; cpu < ncpus; cpu++) { ret = fscanf(f, "%d\t", &opencl_affinity_matrix[gpu][cpu]); STARPU_ASSERT(ret == 1); } ret = fscanf(f, "\n"); STARPU_ASSERT(ret == 0); } #endif /* !STARPU_USE_OPENCL */ _starpu_frdunlock(f); fclose(f); #endif /* !(STARPU_USE_CUDA_ || STARPU_USE_OPENCL */ } #ifndef STARPU_SIMGRID static void write_bus_affinity_file_content(void) { STARPU_ASSERT(was_benchmarked); #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) FILE *f; char path[256]; get_affinity_path(path, sizeof(path)); _STARPU_DEBUG("writing affinities to %s\n", path); f = fopen(path, "w+"); if (!f) { perror("fopen write_buf_affinity_file_content"); _STARPU_DISP("path '%s'\n", path); fflush(stderr); STARPU_ABORT(); } _starpu_frdlock(f); unsigned cpu; unsigned gpu; fprintf(f, "# GPU\t"); for (cpu = 0; cpu < ncpus; cpu++) fprintf(f, "CPU%u\t", cpu); fprintf(f, "\n"); #ifdef STARPU_USE_CUDA for (gpu = 0; gpu < ncuda; gpu++) { fprintf(f, "%d\t", gpu); for (cpu = 0; cpu < ncpus; cpu++) { fprintf(f, "%d\t", cudadev_timing_per_cpu[(gpu+1)*STARPU_MAXCPUS+cpu].cpu_id); } fprintf(f, "\n"); } #endif #ifdef STARPU_USE_OPENCL for (gpu = 0; gpu < nopencl; gpu++) { fprintf(f, "%d\t", gpu); for (cpu = 0; cpu < ncpus; cpu++) { fprintf(f, "%d\t", opencldev_timing_per_cpu[(gpu+1)*STARPU_MAXCPUS+cpu].cpu_id); } fprintf(f, "\n"); } #endif _starpu_frdunlock(f); fclose(f); #endif } #endif /* STARPU_SIMGRID */ static void generate_bus_affinity_file(void) { if (!was_benchmarked) benchmark_all_gpu_devices(); #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Slaves don't write files */ if (!_starpu_mpi_common_is_src_node()) return; #endif write_bus_affinity_file_content(); } static void load_bus_affinity_file(void) { int res; char path[256]; get_affinity_path(path, sizeof(path)); res = access(path, F_OK); if (res) { /* File does not exist yet */ generate_bus_affinity_file(); } load_bus_affinity_file_content(); } #ifdef STARPU_USE_CUDA int *_starpu_get_cuda_affinity_vector(unsigned gpuid) { return cuda_affinity_matrix[gpuid]; } #endif /* STARPU_USE_CUDA */ #ifdef STARPU_USE_OPENCL int *_starpu_get_opencl_affinity_vector(unsigned gpuid) { return opencl_affinity_matrix[gpuid]; } #endif /* STARPU_USE_OPENCL */ void starpu_bus_print_affinity(FILE *f) { #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) unsigned cpu; unsigned gpu; #endif fprintf(f, "# GPU\tCPU in preference order (logical index)\n"); #ifdef STARPU_USE_CUDA fprintf(f, "# CUDA\n"); for(gpu = 0 ; gpu. Expected a number\n", path); fclose(f); return 0; } n = getc(f); if (n == '\n') break; if (n != '\t') { _STARPU_DISP("bogus character '%c' (%d) in latency file %s\n", n, n, path); fclose(f); return 0; } latency_matrix[src][dst] = latency; /* Look out for \t\n */ n = getc(f); if (n == '\n') break; ungetc(n, f); n = '\t'; } /* No more values, take NAN */ for ( ; dst < STARPU_MAXNODES; dst++) latency_matrix[src][dst] = NAN; while (n == '\t') { /* Look out for \t\n */ n = getc(f); if (n == '\n') break; ungetc(n, f); n = _starpu_read_double(f, "%le", &latency); if (n && !isnan(latency)) { _STARPU_DISP("Too many nodes in latency file %s for this configuration (%d)\n", path, STARPU_MAXNODES); fclose(f); return 0; } n = getc(f); } if (n != '\n') { _STARPU_DISP("Bogus character '%c' (%d) in latency file %s\n", n, n, path); fclose(f); return 0; } /* Look out for EOF */ n = getc(f); if (n == EOF) break; ungetc(n, f); } _starpu_frdunlock(f); fclose(f); /* No more values, take NAN */ for ( ; src < STARPU_MAXNODES; src++) for (dst = 0; dst < STARPU_MAXNODES; dst++) latency_matrix[src][dst] = NAN; return 1; } #ifndef STARPU_SIMGRID static void write_bus_latency_file_content(void) { unsigned src, dst, maxnode; FILE *f; STARPU_ASSERT(was_benchmarked); char path[256]; get_latency_path(path, sizeof(path)); _STARPU_DEBUG("writing latencies to %s\n", path); f = fopen(path, "w+"); if (!f) { perror("fopen write_bus_latency_file_content"); _STARPU_DISP("path '%s'\n", path); fflush(stderr); STARPU_ABORT(); } _starpu_fwrlock(f); _starpu_fftruncate(f, 0); fprintf(f, "# "); for (dst = 0; dst < STARPU_MAXNODES; dst++) fprintf(f, "to %d\t\t", dst); fprintf(f, "\n"); maxnode = ncuda; #ifdef STARPU_USE_OPENCL maxnode += nopencl; #endif #ifdef STARPU_USE_MIC maxnode += nmic; #endif #ifdef STARPU_USE_MPI_MASTER_SLAVE maxnode += nmpi_ms; #endif for (src = 0; src < STARPU_MAXNODES; src++) { for (dst = 0; dst < STARPU_MAXNODES; dst++) { double latency = 0.0; if ((src > maxnode) || (dst > maxnode)) { /* convention */ latency = NAN; } else if (src == dst) { latency = 0.0; } else { /* µs */ #ifdef STARPU_USE_CUDA #ifdef HAVE_CUDA_MEMCPY_PEER if (src && src < ncuda && dst && dst <= ncuda) latency = cudadev_latency_dtod[src][dst]; else #endif { if (src && src <= ncuda) latency += cudadev_latency_dtoh[src]; if (dst && dst <= ncuda) latency += cudadev_latency_htod[dst]; } #endif #ifdef STARPU_USE_OPENCL if (src > ncuda && src <= ncuda + nopencl) latency += opencldev_latency_dtoh[src-ncuda]; if (dst > ncuda && dst <= ncuda + nopencl) latency += opencldev_latency_htod[dst-ncuda]; #endif /* TODO Latency MIC */ #ifdef STARPU_USE_MPI_MASTER_SLAVE if (src > ncuda + nopencl + nmic && src <= ncuda + nopencl + nmic + nmpi_ms) latency += mpi_latency_device_to_host[src - (ncuda + nopencl + nmic) - 1]; if (dst > ncuda + nopencl + nmic && dst <= ncuda + nopencl + nmic + nmpi_ms) latency += mpi_latency_host_to_device[dst - (ncuda + nopencl + nmic) - 1]; #endif } if (dst) fputc('\t', f); _starpu_write_double(f, "%e", latency); } fprintf(f, "\n"); } _starpu_fwrunlock(f); fclose(f); } #endif static void generate_bus_latency_file(void) { if (!was_benchmarked) benchmark_all_gpu_devices(); #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Slaves don't write files */ if (!_starpu_mpi_common_is_src_node()) return; #endif #ifndef STARPU_SIMGRID write_bus_latency_file_content(); #endif } static void load_bus_latency_file(void) { int res; char path[256]; get_latency_path(path, sizeof(path)); res = access(path, F_OK); if (res || !load_bus_latency_file_content()) { /* File does not exist yet or is bogus */ generate_bus_latency_file(); } } /* * Bandwidth */ static void get_bandwidth_path(char *path, size_t maxlen) { get_bus_path("bandwidth", path, maxlen); } static int load_bus_bandwidth_file_content(void) { int n; unsigned src, dst; FILE *f; double bandwidth; char path[256]; get_bandwidth_path(path, sizeof(path)); _STARPU_DEBUG("loading bandwidth from %s\n", path); f = fopen(path, "r"); if (!f) { perror("fopen load_bus_bandwidth_file_content"); _STARPU_DISP("path '%s'\n", path); fflush(stderr); STARPU_ABORT(); } _starpu_frdlock(f); for (src = 0; src < STARPU_MAXNODES; src++) { _starpu_drop_comments(f); for (dst = 0; dst < STARPU_MAXNODES; dst++) { n = _starpu_read_double(f, "%le", &bandwidth); if (n != 1) { _STARPU_DISP("Error while reading bandwidth file <%s>. Expected a number\n", path); fclose(f); return 0; } n = getc(f); if (n == '\n') break; if (n != '\t') { _STARPU_DISP("bogus character '%c' (%d) in bandwidth file %s\n", n, n, path); fclose(f); return 0; } bandwidth_matrix[src][dst] = bandwidth; /* Look out for \t\n */ n = getc(f); if (n == '\n') break; ungetc(n, f); n = '\t'; } /* No more values, take NAN */ for ( ; dst < STARPU_MAXNODES; dst++) bandwidth_matrix[src][dst] = NAN; while (n == '\t') { /* Look out for \t\n */ n = getc(f); if (n == '\n') break; ungetc(n, f); n = _starpu_read_double(f, "%le", &bandwidth); if (n && !isnan(bandwidth)) { _STARPU_DISP("Too many nodes in bandwidth file %s for this configuration (%d)\n", path, STARPU_MAXNODES); fclose(f); return 0; } n = getc(f); } if (n != '\n') { _STARPU_DISP("Bogus character '%c' (%d) in bandwidth file %s\n", n, n, path); fclose(f); return 0; } /* Look out for EOF */ n = getc(f); if (n == EOF) break; ungetc(n, f); } _starpu_frdunlock(f); fclose(f); /* No more values, take NAN */ for ( ; src < STARPU_MAXNODES; src++) for (dst = 0; dst < STARPU_MAXNODES; dst++) latency_matrix[src][dst] = NAN; return 1; } #ifndef STARPU_SIMGRID static void write_bus_bandwidth_file_content(void) { unsigned src, dst, maxnode; FILE *f; STARPU_ASSERT(was_benchmarked); char path[256]; get_bandwidth_path(path, sizeof(path)); _STARPU_DEBUG("writing bandwidth to %s\n", path); f = fopen(path, "w+"); STARPU_ASSERT(f); _starpu_fwrlock(f); _starpu_fftruncate(f, 0); fprintf(f, "# "); for (dst = 0; dst < STARPU_MAXNODES; dst++) fprintf(f, "to %d\t\t", dst); fprintf(f, "\n"); maxnode = ncuda; #ifdef STARPU_USE_OPENCL maxnode += nopencl; #endif #ifdef STARPU_USE_MIC maxnode += nmic; #endif #ifdef STARPU_USE_MPI_MASTER_SLAVE maxnode += nmpi_ms; #endif for (src = 0; src < STARPU_MAXNODES; src++) { for (dst = 0; dst < STARPU_MAXNODES; dst++) { double bandwidth; if ((src > maxnode) || (dst > maxnode)) { bandwidth = NAN; } #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) || defined(STARPU_USE_MIC) || defined(STARPU_USE_MPI_MASTER_SLAVE) else if (src != dst) { double slowness = 0.0; /* Total bandwidth is the harmonic mean of bandwidths */ #ifdef STARPU_USE_CUDA #ifdef HAVE_CUDA_MEMCPY_PEER if (src && src <= ncuda && dst && dst <= ncuda) /* Direct GPU-GPU transfert */ slowness = cudadev_timing_dtod[src][dst]; else #endif { if (src && src <= ncuda) slowness += cudadev_timing_dtoh[src]; if (dst && dst <= ncuda) slowness += cudadev_timing_htod[dst]; } #endif /* TODO: generalize computation */ #ifdef STARPU_USE_OPENCL if (src > ncuda && src <= ncuda + nopencl) slowness += opencldev_timing_dtoh[src-ncuda]; if (dst > ncuda && dst <= ncuda + nopencl) slowness += opencldev_timing_htod[dst-ncuda]; #endif #ifdef STARPU_USE_MIC if (src > ncuda + nopencl && src <= ncuda + nopencl + nmic) slowness += mic_time_device_to_host[src - (ncuda + nopencl)]; if (dst > ncuda + nopencl && dst <= ncuda + nopencl + nmic) slowness += mic_time_host_to_device[dst - (ncuda + nopencl)]; #endif #ifdef STARPU_USE_MPI_MASTER_SLAVE /* here we have bandwidth */ if (src > ncuda + nopencl + nmic && src <= ncuda + nopencl + nmic + nmpi_ms) slowness += 1.0/mpi_time_device_to_host[src - (ncuda + nopencl + nmic) - 1]; if (dst > ncuda + nopencl + nmic && dst <= ncuda + nopencl + nmic +nmpi_ms) slowness += 1.0/mpi_time_host_to_device[dst - (ncuda + nopencl + nmic) - 1]; #endif bandwidth = 1.0/slowness; } #endif else { /* convention */ bandwidth = 0.0; } if (dst) fputc('\t', f); _starpu_write_double(f, "%e", bandwidth); } fprintf(f, "\n"); } _starpu_fwrunlock(f); fclose(f); } #endif /* STARPU_SIMGRID */ void starpu_bus_print_filenames(FILE *output) { char bandwidth_path[256]; char affinity_path[256]; char latency_path[256]; get_bandwidth_path(bandwidth_path, sizeof(bandwidth_path)); get_affinity_path(affinity_path, sizeof(affinity_path)); get_latency_path(latency_path, sizeof(latency_path)); fprintf(output, "bandwidth: <%s>\n", bandwidth_path); fprintf(output, " affinity: <%s>\n", affinity_path); fprintf(output, " latency: <%s>\n", latency_path); } void starpu_bus_print_bandwidth(FILE *f) { unsigned src, dst, maxnode; maxnode = ncuda; #ifdef STARPU_USE_OPENCL maxnode += nopencl; #endif #ifdef STARPU_USE_MIC maxnode += nmic; #endif #ifdef STARPU_USE_MPI_MASTER_SLAVE maxnode += nmpi_ms; #endif fprintf(f, "from/to\t"); fprintf(f, "RAM\t"); for (dst = 0; dst < ncuda; dst++) fprintf(f, "CUDA %d\t", dst); for (dst = 0; dst < nopencl; dst++) fprintf(f, "OpenCL%d\t", dst); for (dst = 0; dst < nmic; dst++) fprintf(f, "MIC%d\t", dst); for (dst = 0; dst < nmpi_ms; dst++) fprintf(f, "MPI_MS%d\t", dst); fprintf(f, "\n"); for (src = 0; src <= maxnode; src++) { if (!src) fprintf(f, "RAM\t"); else if (src <= ncuda) fprintf(f, "CUDA %d\t", src-1); else if (src <= ncuda + nopencl) fprintf(f, "OpenCL%d\t", src-ncuda-1); else if (src <= ncuda + nopencl + nmic) fprintf(f, "MIC%d\t", src-ncuda-nopencl-1); else fprintf(f, "MPI_MS%d\t", src-ncuda-nopencl-nmic-1); for (dst = 0; dst <= maxnode; dst++) fprintf(f, "%.0f\t", bandwidth_matrix[src][dst]); fprintf(f, "\n"); } fprintf(f, "\n"); for (src = 0; src <= maxnode; src++) { if (!src) fprintf(f, "RAM\t"); else if (src <= ncuda) fprintf(f, "CUDA %d\t", src-1); else if (src <= ncuda + nopencl) fprintf(f, "OpenCL%d\t", src-ncuda-1); else if (src <= ncuda + nopencl + nmic) fprintf(f, "MIC%d\t", src-ncuda-nopencl-1); else fprintf(f, "MPI_MS%d\t", src-ncuda-nopencl-nmic-1); for (dst = 0; dst <= maxnode; dst++) fprintf(f, "%.0f\t", latency_matrix[src][dst]); fprintf(f, "\n"); } #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_OPENCL) if (ncuda != 0 || nopencl != 0) fprintf(f, "\nGPU\tCPU in preference order (logical index), host-to-device, device-to-host\n"); for (src = 1; src <= ncuda + nopencl; src++) { struct dev_timing *timing; struct _starpu_machine_config *config = _starpu_get_machine_config(); unsigned config_ncpus = _starpu_topology_get_nhwcpu(config); unsigned cpu; #ifdef STARPU_USE_CUDA if (src <= ncuda) { fprintf(f, "CUDA %d\t", src-1); for (cpu = 0; cpu < config_ncpus; cpu++) { timing = &cudadev_timing_per_cpu[src*STARPU_MAXCPUS+cpu]; if (timing->timing_htod) fprintf(f, "%2d %.0f %.0f\t", timing->cpu_id, 1/timing->timing_htod, 1/timing->timing_dtoh); else fprintf(f, "%2d\t", cuda_affinity_matrix[src-1][cpu]); } } #ifdef STARPU_USE_OPENCL else #endif #endif #ifdef STARPU_USE_OPENCL { fprintf(f, "OpenCL%d\t", src-ncuda-1); for (cpu = 0; cpu < config_ncpus; cpu++) { timing = &opencldev_timing_per_cpu[(src-ncuda)*STARPU_MAXCPUS+cpu]; if (timing->timing_htod) fprintf(f, "%2d %.0f %.0f\t", timing->cpu_id, 1/timing->timing_htod, 1/timing->timing_dtoh); else fprintf(f, "%2d\t", opencl_affinity_matrix[src-1][cpu]); } } #endif fprintf(f, "\n"); } #endif } static void generate_bus_bandwidth_file(void) { if (!was_benchmarked) benchmark_all_gpu_devices(); #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Slaves don't write files */ if (!_starpu_mpi_common_is_src_node()) return; #endif #ifndef STARPU_SIMGRID write_bus_bandwidth_file_content(); #endif } static void load_bus_bandwidth_file(void) { int res; char path[256]; get_bandwidth_path(path, sizeof(path)); res = access(path, F_OK); if (res || !load_bus_bandwidth_file_content()) { /* File does not exist yet or is bogus */ generate_bus_bandwidth_file(); } } #ifndef STARPU_SIMGRID /* * Config */ static void get_config_path(char *path, size_t maxlen) { get_bus_path("config", path, maxlen); } static void compare_value_and_recalibrate(char * msg, unsigned val_file, unsigned val_detected) { if (val_file != val_detected) { #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Only the master prints the message */ if (_starpu_mpi_common_is_src_node()) #endif _STARPU_DISP("Current configuration does not match the bus performance model (%s: (stored) %d != (current) %d), recalibrating...\n", msg, val_file, val_detected); _starpu_bus_force_sampling(); #ifdef STARPU_USE_MPI_MASTER_SLAVE if (_starpu_mpi_common_is_src_node()) #endif _STARPU_DISP("... done\n"); } } static void check_bus_config_file(void) { int res; char path[256]; struct _starpu_machine_config *config = _starpu_get_machine_config(); get_config_path(path, sizeof(path)); res = access(path, F_OK); if (res || config->conf.bus_calibrate > 0) { if (res) _STARPU_DISP("No performance model for the bus, calibrating...\n"); _starpu_bus_force_sampling(); if (res) _STARPU_DISP("... done\n"); } else { FILE *f; int ret; unsigned read_cuda = -1, read_opencl = -1, read_mic = -1, read_mpi_ms = -1; unsigned read_cpus = -1; // Loading configuration from file f = fopen(path, "r"); STARPU_ASSERT(f); _starpu_frdlock(f); _starpu_drop_comments(f); ret = fscanf(f, "%u\t", &read_cpus); STARPU_ASSERT(ret == 1); _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &read_cuda); STARPU_ASSERT(ret == 1); _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &read_opencl); STARPU_ASSERT(ret == 1); _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &read_mic); if (ret == 0) read_mic = 0; _starpu_drop_comments(f); ret = fscanf(f, "%d\t", &read_mpi_ms); if (ret == 0) read_mpi_ms = 0; _starpu_drop_comments(f); _starpu_frdunlock(f); fclose(f); // Loading current configuration ncpus = _starpu_topology_get_nhwcpu(config); #ifdef STARPU_USE_CUDA ncuda = _starpu_get_cuda_device_count(); #endif #ifdef STARPU_USE_OPENCL nopencl = _starpu_opencl_get_device_count(); #endif #ifdef STARPU_USE_MIC nmic = _starpu_mic_src_get_device_count(); #endif /* STARPU_USE_MIC */ #ifdef STARPU_USE_MPI_MASTER_SLAVE nmpi_ms = _starpu_mpi_src_get_device_count(); #endif /* STARPU_USE_MPI_MASTER_SLAVE */ // Checking if both configurations match compare_value_and_recalibrate("CPUS", read_cpus, ncpus); compare_value_and_recalibrate("CUDA", read_cuda, ncuda); compare_value_and_recalibrate("OpenCL", read_opencl, nopencl); compare_value_and_recalibrate("MIC", read_mic, nmic); compare_value_and_recalibrate("MPI Master-Slave", read_mpi_ms, nmpi_ms); } } static void write_bus_config_file_content(void) { FILE *f; char path[256]; STARPU_ASSERT(was_benchmarked); get_config_path(path, sizeof(path)); _STARPU_DEBUG("writing config to %s\n", path); f = fopen(path, "w+"); STARPU_ASSERT(f); _starpu_fwrlock(f); _starpu_fftruncate(f, 0); fprintf(f, "# Current configuration\n"); fprintf(f, "%u # Number of CPUs\n", ncpus); fprintf(f, "%d # Number of CUDA devices\n", ncuda); fprintf(f, "%d # Number of OpenCL devices\n", nopencl); fprintf(f, "%d # Number of MIC devices\n", nmic); fprintf(f, "%d # Number of MPI devices\n", nmpi_ms); _starpu_fwrunlock(f); fclose(f); } static void generate_bus_config_file(void) { if (!was_benchmarked) benchmark_all_gpu_devices(); #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Slaves don't write files */ if (!_starpu_mpi_common_is_src_node()) return; #endif write_bus_config_file_content(); } #endif /* !SIMGRID */ void _starpu_simgrid_get_platform_path(int version, char *path, size_t maxlen) { if (version == 3) get_bus_path("platform.xml", path, maxlen); else get_bus_path("platform.v4.xml", path, maxlen); } #ifndef STARPU_SIMGRID /* * Compute the precise PCI tree bandwidth and link shares * * We only have measurements from one leaf to another. We assume that the * available bandwidth is greater at lower levels, and thus measurements from * increasingly far GPUs provide the PCI bridges bandwidths at each level. * * The bandwidth of a PCI bridge is thus computed as the maximum of the speed * of the various transfers that we have achieved through it. We thus browse * the PCI tree three times: * * - first through all CUDA-CUDA possible transfers to compute the maximum * measured bandwidth on each PCI link and hub used for that. * - then through the whole tree to emit links for each PCI link and hub. * - then through all CUDA-CUDA possible transfers again to emit routes. */ #if defined(STARPU_USE_CUDA) && defined(HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX) && HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX && defined(HAVE_CUDA_MEMCPY_PEER) /* Records, for each PCI link and hub, the maximum bandwidth seen through it */ struct pci_userdata { /* Uplink max measurement */ double bw_up; double bw_down; /* Hub max measurement */ double bw; }; /* Allocate a pci_userdata structure for the given object */ static void allocate_userdata(hwloc_obj_t obj) { struct pci_userdata *data; if (obj->userdata) return; data = obj->userdata = malloc(sizeof(*data)); data->bw_up = 0.0; data->bw_down = 0.0; data->bw = 0.0; } /* Update the maximum bandwidth seen going to upstream */ static void update_bandwidth_up(hwloc_obj_t obj, double bandwidth) { struct pci_userdata *data; if (obj->type != HWLOC_OBJ_BRIDGE && obj->type != HWLOC_OBJ_PCI_DEVICE) return; allocate_userdata(obj); data = obj->userdata; if (data->bw_up < bandwidth) data->bw_up = bandwidth; } /* Update the maximum bandwidth seen going from upstream */ static void update_bandwidth_down(hwloc_obj_t obj, double bandwidth) { struct pci_userdata *data; if (obj->type != HWLOC_OBJ_BRIDGE && obj->type != HWLOC_OBJ_PCI_DEVICE) return; allocate_userdata(obj); data = obj->userdata; if (data->bw_down < bandwidth) data->bw_down = bandwidth; } /* Update the maximum bandwidth seen going through this Hub */ static void update_bandwidth_through(hwloc_obj_t obj, double bandwidth) { struct pci_userdata *data; allocate_userdata(obj); data = obj->userdata; if (data->bw < bandwidth) data->bw = bandwidth; } /* find_* functions perform the first step: computing maximum bandwidths */ /* Our trafic had to go through the host, go back from target up to the host, * updating uplink downstream bandwidth along the way */ static void find_platform_backward_path(hwloc_obj_t obj, double bandwidth) { if (!obj) /* Oops, we should have seen a host bridge. Well, too bad. */ return; /* Update uplink bandwidth of PCI Hub */ update_bandwidth_down(obj, bandwidth); /* Update internal bandwidth of PCI Hub */ update_bandwidth_through(obj, bandwidth); if (obj->type == HWLOC_OBJ_BRIDGE && obj->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) /* Finished */ return; /* Continue up */ find_platform_backward_path(obj->parent, bandwidth); } /* Same, but update uplink upstream bandwidth */ static void find_platform_forward_path(hwloc_obj_t obj, double bandwidth) { if (!obj) /* Oops, we should have seen a host bridge. Well, too bad. */ return; /* Update uplink bandwidth of PCI Hub */ update_bandwidth_up(obj, bandwidth); /* Update internal bandwidth of PCI Hub */ update_bandwidth_through(obj, bandwidth); if (obj->type == HWLOC_OBJ_BRIDGE && obj->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) /* Finished */ return; /* Continue up */ find_platform_forward_path(obj->parent, bandwidth); } /* Find the path from obj1 through parent down to obj2 (without ever going up), * and update the maximum bandwidth along the path */ static int find_platform_path_down(hwloc_obj_t parent, hwloc_obj_t obj1, hwloc_obj_t obj2, double bandwidth) { unsigned i; /* Base case, path is empty */ if (parent == obj2) return 1; /* Try to go down from parent */ for (i = 0; i < parent->arity; i++) if (parent->children[i] != obj1 && find_platform_path_down(parent->children[i], NULL, obj2, bandwidth)) { /* Found it down there, update bandwidth of parent */ update_bandwidth_down(parent->children[i], bandwidth); update_bandwidth_through(parent, bandwidth); return 1; } return 0; } /* Find the path from obj1 to obj2, and update the maximum bandwidth along the * path */ static int find_platform_path_up(hwloc_obj_t obj1, hwloc_obj_t obj2, double bandwidth) { int ret; hwloc_obj_t parent = obj1->parent; if (!parent) { /* Oops, we should have seen a host bridge. Act as if we had seen it. */ find_platform_backward_path(obj2, bandwidth); return 1; } if (find_platform_path_down(parent, obj1, obj2, bandwidth)) /* obj2 was a mere (sub)child of our parent */ return 1; /* obj2 is not a (sub)child of our parent, we have to go up through the parent */ if (parent->type == HWLOC_OBJ_BRIDGE && parent->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) { /* We have to go up to the Host, so obj2 is not in the same PCI * tree, so we're for for obj1 to Host, and just find the path * from obj2 to Host too. */ find_platform_backward_path(obj2, bandwidth); update_bandwidth_up(parent, bandwidth); update_bandwidth_through(parent, bandwidth); return 1; } /* Not at host yet, just go up */ ret = find_platform_path_up(parent, obj2, bandwidth); update_bandwidth_up(parent, bandwidth); update_bandwidth_through(parent, bandwidth); return ret; } /* find the path between cuda i and cuda j, and update the maximum bandwidth along the path */ static int find_platform_cuda_path(hwloc_topology_t topology, unsigned i, unsigned j, double bandwidth) { hwloc_obj_t cudai, cudaj; cudai = hwloc_cuda_get_device_osdev_by_index(topology, i); cudaj = hwloc_cuda_get_device_osdev_by_index(topology, j); if (!cudai || !cudaj) return 0; return find_platform_path_up(cudai, cudaj, bandwidth); } /* emit_topology_bandwidths performs the second step: emitting link names */ /* Emit the link name of the object */ static void emit_pci_hub(FILE *f, hwloc_obj_t obj) { STARPU_ASSERT(obj->type == HWLOC_OBJ_BRIDGE); fprintf(f, "PCI:%04x:[%02x-%02x]", obj->attr->bridge.downstream.pci.domain, obj->attr->bridge.downstream.pci.secondary_bus, obj->attr->bridge.downstream.pci.subordinate_bus); } static void emit_pci_dev(FILE *f, struct hwloc_pcidev_attr_s *pcidev) { fprintf(f, "PCI:%04x:%02x:%02x.%1x", pcidev->domain, pcidev->bus, pcidev->dev, pcidev->func); } /* Emit the links of the object */ static void emit_topology_bandwidths(FILE *f, hwloc_obj_t obj, const char *Bps, const char *s) { unsigned i; if (obj->userdata) { struct pci_userdata *data = obj->userdata; if (obj->type == HWLOC_OBJ_BRIDGE) { /* Uplink */ fprintf(f, " \n", data->bw_up, Bps, s); fprintf(f, " \n", data->bw_down, Bps, s); /* PCI Switches are assumed to have infinite internal bandwidth */ if (!obj->name || !strstr(obj->name, "Switch")) { /* We assume that PCI Hubs have double bandwidth in * order to support full duplex but not more */ fprintf(f, " \n", data->bw * 2, Bps, s); } } else if (obj->type == HWLOC_OBJ_PCI_DEVICE) { fprintf(f, " attr->pcidev); fprintf(f, " up\" bandwidth=\"%f%s\" latency=\"0.000000%s\"/>\n", data->bw_up, Bps, s); fprintf(f, " attr->pcidev); fprintf(f, " down\" bandwidth=\"%f%s\" latency=\"0.000000%s\"/>\n", data->bw_down, Bps, s); } } for (i = 0; i < obj->arity; i++) emit_topology_bandwidths(f, obj->children[i], Bps, s); } /* emit_pci_link_* functions perform the third step: emitting the routes */ static void emit_pci_link(FILE *f, hwloc_obj_t obj, const char *suffix) { if (obj->type == HWLOC_OBJ_BRIDGE) { fprintf(f, " \n", suffix); } else if (obj->type == HWLOC_OBJ_PCI_DEVICE) { fprintf(f, " attr->pcidev); fprintf(f, " %s\"/>\n", suffix); } } /* Go to upstream */ static void emit_pci_link_up(FILE *f, hwloc_obj_t obj) { emit_pci_link(f, obj, "up"); } /* Go from upstream */ static void emit_pci_link_down(FILE *f, hwloc_obj_t obj) { emit_pci_link(f, obj, "down"); } /* Go through PCI hub */ static void emit_pci_link_through(FILE *f, hwloc_obj_t obj) { /* We don't care about trafic going through PCI switches */ if (obj->type == HWLOC_OBJ_BRIDGE) { if (!obj->name || !strstr(obj->name, "Switch")) emit_pci_link(f, obj, "through"); else { fprintf(f, " \n"); } } } /* Our trafic has to go through the host, go back from target up to the host, * using uplink downstream along the way */ static void emit_platform_backward_path(FILE *f, hwloc_obj_t obj) { if (!obj) /* Oops, we should have seen a host bridge. Well, too bad. */ return; /* Go through PCI Hub */ emit_pci_link_through(f, obj); /* Go through uplink */ emit_pci_link_down(f, obj); if (obj->type == HWLOC_OBJ_BRIDGE && obj->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) { /* Finished, go through host */ fprintf(f, " \n"); return; } /* Continue up */ emit_platform_backward_path(f, obj->parent); } /* Same, but use upstream link */ static void emit_platform_forward_path(FILE *f, hwloc_obj_t obj) { if (!obj) /* Oops, we should have seen a host bridge. Well, too bad. */ return; /* Go through PCI Hub */ emit_pci_link_through(f, obj); /* Go through uplink */ emit_pci_link_up(f, obj); if (obj->type == HWLOC_OBJ_BRIDGE && obj->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) { /* Finished, go through host */ fprintf(f, " \n"); return; } /* Continue up */ emit_platform_forward_path(f, obj->parent); } /* Find the path from obj1 through parent down to obj2 (without ever going up), * and use the links along the path */ static int emit_platform_path_down(FILE *f, hwloc_obj_t parent, hwloc_obj_t obj1, hwloc_obj_t obj2) { unsigned i; /* Base case, path is empty */ if (parent == obj2) return 1; /* Try to go down from parent */ for (i = 0; i < parent->arity; i++) if (parent->children[i] != obj1 && emit_platform_path_down(f, parent->children[i], NULL, obj2)) { /* Found it down there, path goes through this hub */ emit_pci_link_down(f, parent->children[i]); emit_pci_link_through(f, parent); return 1; } return 0; } /* Find the path from obj1 to obj2, and use the links along the path */ static int emit_platform_path_up(FILE *f, hwloc_obj_t obj1, hwloc_obj_t obj2) { int ret; hwloc_obj_t parent = obj1->parent; if (!parent) { /* Oops, we should have seen a host bridge. Act as if we had seen it. */ emit_platform_backward_path(f, obj2); return 1; } if (emit_platform_path_down(f, parent, obj1, obj2)) /* obj2 was a mere (sub)child of our parent */ return 1; /* obj2 is not a (sub)child of our parent, we have to go up through the parent */ if (parent->type == HWLOC_OBJ_BRIDGE && parent->attr->bridge.upstream_type == HWLOC_OBJ_BRIDGE_HOST) { /* We have to go up to the Host, so obj2 is not in the same PCI * tree, so we're for for obj1 to Host, and just find the path * from obj2 to Host too. */ emit_platform_backward_path(f, obj2); fprintf(f, " \n"); emit_pci_link_up(f, parent); emit_pci_link_through(f, parent); return 1; } /* Not at host yet, just go up */ ret = emit_platform_path_up(f, parent, obj2); emit_pci_link_up(f, parent); emit_pci_link_through(f, parent); return ret; } /* Clean our mess in the topology before destroying it */ static void clean_topology(hwloc_obj_t obj) { unsigned i; if (obj->userdata) free(obj->userdata); for (i = 0; i < obj->arity; i++) clean_topology(obj->children[i]); } #endif static void write_bus_platform_file_content(int version) { FILE *f; char path[256]; unsigned i; const char *speed, *flops, *Bps, *s; char dash; if (version == 3) { speed = "power"; flops = ""; Bps = ""; s = ""; dash = '_'; } else { speed = "speed"; flops = "f"; Bps = "Bps"; s = "s"; dash = '-'; } STARPU_ASSERT(was_benchmarked); _starpu_simgrid_get_platform_path(version, path, sizeof(path)); _STARPU_DEBUG("writing platform to %s\n", path); f = fopen(path, "w+"); if (!f) { perror("fopen write_bus_platform_file_content"); _STARPU_DISP("path '%s'\n", path); fflush(stderr); STARPU_ABORT(); } _starpu_fwrlock(f); _starpu_fftruncate(f, 0); fprintf(f, "\n" "\n" " \n" " \n" " \n" " \n" " \n" " \n" " \n" " \n", version == 3 ? "http://simgrid.gforge.inria.fr/simgrid.dtd" : "http://simgrid.gforge.inria.fr/simgrid/simgrid.dtd", version, dash, dash, dash, speed, flops); for (i = 0; i < ncpus; i++) /* TODO: host memory for out-of-core simulation */ fprintf(f, " \n", i, speed, flops); for (i = 0; i < ncuda; i++) { fprintf(f, " \n", i, speed, flops); fprintf(f, " \n", (unsigned long long) cuda_size[i]); #ifdef HAVE_CUDA_MEMCPY_PEER fprintf(f, " \n"); #endif fprintf(f, " \n"); } for (i = 0; i < nopencl; i++) { fprintf(f, " \n", i, speed, flops); fprintf(f, " \n", (unsigned long long) opencl_size[i]); fprintf(f, " \n"); } fprintf(f, "\n \n", speed, flops); /* * Compute maximum bandwidth, taken as host bandwidth */ double max_bandwidth = 0; #ifdef STARPU_USE_CUDA for (i = 0; i < ncuda; i++) { double down_bw = 1.0 / cudadev_timing_dtoh[1+i]; double up_bw = 1.0 / cudadev_timing_htod[1+i]; if (max_bandwidth < down_bw) max_bandwidth = down_bw; if (max_bandwidth < up_bw) max_bandwidth = up_bw; } #endif #ifdef STARPU_USE_OPENCL for (i = 0; i < nopencl; i++) { double down_bw = 1.0 / opencldev_timing_dtoh[1+i]; double up_bw = 1.0 / opencldev_timing_htod[1+i]; if (max_bandwidth < down_bw) max_bandwidth = down_bw; if (max_bandwidth < up_bw) max_bandwidth = up_bw; } #endif fprintf(f, "\n \n\n", max_bandwidth*1000000, Bps, s); /* * OpenCL links */ #ifdef STARPU_USE_OPENCL for (i = 0; i < nopencl; i++) { char i_name[16]; snprintf(i_name, sizeof(i_name), "OpenCL%d", i); fprintf(f, " \n", i_name, 1000000 / opencldev_timing_htod[1+i], Bps, opencldev_latency_htod[1+i]/1000000., s); fprintf(f, " \n", i_name, 1000000 / opencldev_timing_dtoh[1+i], Bps, opencldev_latency_dtoh[1+i]/1000000., s); } fprintf(f, "\n"); #endif /* * CUDA links and routes */ #ifdef STARPU_USE_CUDA /* Write RAM/CUDA bandwidths and latencies */ for (i = 0; i < ncuda; i++) { char i_name[16]; snprintf(i_name, sizeof(i_name), "CUDA%d", i); fprintf(f, " \n", i_name, 1000000. / cudadev_timing_htod[1+i], Bps, cudadev_latency_htod[1+i]/1000000., s); fprintf(f, " \n", i_name, 1000000. / cudadev_timing_dtoh[1+i], Bps, cudadev_latency_dtoh[1+i]/1000000., s); } fprintf(f, "\n"); #ifdef HAVE_CUDA_MEMCPY_PEER /* Write CUDA/CUDA bandwidths and latencies */ for (i = 0; i < ncuda; i++) { unsigned j; char i_name[16]; snprintf(i_name, sizeof(i_name), "CUDA%d", i); for (j = 0; j < ncuda; j++) { char j_name[16]; if (j == i) continue; snprintf(j_name, sizeof(j_name), "CUDA%d", j); fprintf(f, " \n", i_name, j_name, 1000000. / cudadev_timing_dtod[1+i][1+j], Bps, cudadev_latency_dtod[1+i][1+j]/1000000., s); } } #endif #if defined(HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX) && HAVE_DECL_HWLOC_CUDA_GET_DEVICE_OSDEV_BY_INDEX && defined(HAVE_CUDA_MEMCPY_PEER) /* If we have enough hwloc information, write PCI bandwidths and routes */ if (!starpu_get_env_number_default("STARPU_PCI_FLAT", 0)) { hwloc_topology_t topology; hwloc_topology_init(&topology); hwloc_topology_set_flags(topology, HWLOC_TOPOLOGY_FLAG_IO_DEVICES | HWLOC_TOPOLOGY_FLAG_IO_BRIDGES); hwloc_topology_load(topology); /* First find paths and record measured bandwidth along the path */ for (i = 0; i < ncuda; i++) { unsigned j; for (j = 0; j < ncuda; j++) if (i != j) if (!find_platform_cuda_path(topology, i, j, 1000000. / cudadev_timing_dtod[1+i][1+j])) { clean_topology(hwloc_get_root_obj(topology)); hwloc_topology_destroy(topology); goto flat_cuda; } /* Record RAM/CUDA bandwidths */ find_platform_forward_path(hwloc_cuda_get_device_osdev_by_index(topology, i), 1000000. / cudadev_timing_dtoh[1+i]); find_platform_backward_path(hwloc_cuda_get_device_osdev_by_index(topology, i), 1000000. / cudadev_timing_htod[1+i]); } /* Ok, found path in all cases, can emit advanced platform routes */ fprintf(f, "\n"); emit_topology_bandwidths(f, hwloc_get_root_obj(topology), Bps, s); fprintf(f, "\n"); for (i = 0; i < ncuda; i++) { unsigned j; for (j = 0; j < ncuda; j++) if (i != j) { fprintf(f, " \n", i, j); fprintf(f, " \n", i, j); emit_platform_path_up(f, hwloc_cuda_get_device_osdev_by_index(topology, i), hwloc_cuda_get_device_osdev_by_index(topology, j)); fprintf(f, " \n"); } fprintf(f, " \n", i); fprintf(f, " \n", i); emit_platform_forward_path(f, hwloc_cuda_get_device_osdev_by_index(topology, i)); fprintf(f, " \n"); fprintf(f, " \n", i); fprintf(f, " \n", i); emit_platform_backward_path(f, hwloc_cuda_get_device_osdev_by_index(topology, i)); fprintf(f, " \n"); } clean_topology(hwloc_get_root_obj(topology)); hwloc_topology_destroy(topology); } else { flat_cuda: #else { #endif /* If we don't have enough hwloc information, write trivial routes always through host */ for (i = 0; i < ncuda; i++) { char i_name[16]; snprintf(i_name, sizeof(i_name), "CUDA%d", i); fprintf(f, " \n", i_name, i_name); fprintf(f, " \n", i_name, i_name); } #ifdef HAVE_CUDA_MEMCPY_PEER for (i = 0; i < ncuda; i++) { unsigned j; char i_name[16]; snprintf(i_name, sizeof(i_name), "CUDA%d", i); for (j = 0; j < ncuda; j++) { char j_name[16]; if (j == i) continue; snprintf(j_name, sizeof(j_name), "CUDA%d", j); fprintf(f, " \n", i_name, j_name, i_name, j_name); } } #endif } /* defined(STARPU_HAVE_HWLOC) && defined(HAVE_CUDA_MEMCPY_PEER) */ fprintf(f, "\n"); #endif /* STARPU_USE_CUDA */ /* * OpenCL routes */ #ifdef STARPU_USE_OPENCL for (i = 0; i < nopencl; i++) { char i_name[16]; snprintf(i_name, sizeof(i_name), "OpenCL%d", i); fprintf(f, " \n", i_name, i_name); fprintf(f, " \n", i_name, i_name); } #endif fprintf(f, " \n" " \n" ); _starpu_fwrunlock(f); fclose(f); } static void generate_bus_platform_file(void) { if (!was_benchmarked) benchmark_all_gpu_devices(); #ifdef STARPU_USE_MPI_MASTER_SLAVE /* Slaves don't write files */ if (!_starpu_mpi_common_is_src_node()) return; #endif write_bus_platform_file_content(3); write_bus_platform_file_content(4); } static void check_bus_platform_file(void) { int res; char path[256]; _starpu_simgrid_get_platform_path(4, path, sizeof(path)); res = access(path, F_OK); if (!res) { _starpu_simgrid_get_platform_path(3, path, sizeof(path)); res = access(path, F_OK); } if (res) { /* File does not exist yet */ generate_bus_platform_file(); } } /* * Generic */ static void _starpu_bus_force_sampling(void) { _STARPU_DEBUG("Force bus sampling ...\n"); _starpu_create_sampling_directory_if_needed(); generate_bus_affinity_file(); generate_bus_latency_file(); generate_bus_bandwidth_file(); generate_bus_config_file(); generate_bus_platform_file(); } #endif /* !SIMGRID */ void _starpu_load_bus_performance_files(void) { _starpu_create_sampling_directory_if_needed(); #if defined(STARPU_USE_CUDA) || defined(STARPU_USE_SIMGRID) ncuda = _starpu_get_cuda_device_count(); #endif #if defined(STARPU_USE_OPENCL) || defined(STARPU_USE_SIMGRID) nopencl = _starpu_opencl_get_device_count(); #endif #if defined(STARPU_USE_MPI_MASTER_SLAVE) || defined(STARPU_USE_SIMGRID) nmpi_ms = _starpu_mpi_src_get_device_count(); #endif #ifndef STARPU_SIMGRID check_bus_config_file(); #endif #ifdef STARPU_USE_MPI_MASTER_SLAVE /* be sure that master wrote the perf files */ _starpu_mpi_common_barrier(); #endif #ifndef STARPU_SIMGRID load_bus_affinity_file(); #endif load_bus_latency_file(); load_bus_bandwidth_file(); #ifndef STARPU_SIMGRID check_bus_platform_file(); #endif } /* (in MB/s) */ double starpu_transfer_bandwidth(unsigned src_node, unsigned dst_node) { return bandwidth_matrix[src_node][dst_node]; } /* (in µs) */ double starpu_transfer_latency(unsigned src_node, unsigned dst_node) { return latency_matrix[src_node][dst_node]; } /* (in µs) */ double starpu_transfer_predict(unsigned src_node, unsigned dst_node, size_t size) { double bandwidth = bandwidth_matrix[src_node][dst_node]; double latency = latency_matrix[src_node][dst_node]; struct _starpu_machine_topology *topology = &_starpu_get_machine_config()->topology; return latency + (size/bandwidth)*2*(topology->ncudagpus+topology->nopenclgpus); } /* calculate save bandwidth and latency */ /* bandwidth in MB/s - latency in µs */ void _starpu_save_bandwidth_and_latency_disk(double bandwidth_write, double bandwidth_read, double latency_write, double latency_read, unsigned node) { unsigned int i, j; double slowness_disk_between_main_ram, slowness_main_ram_between_node; /* save bandwith */ for(i = 0; i < STARPU_MAXNODES; ++i) { for(j = 0; j < STARPU_MAXNODES; ++j) { if (i == j && j == node) /* source == destination == node */ { bandwidth_matrix[i][j] = 0; } else if (i == node) /* source == disk */ { /* convert in slowness */ if(bandwidth_read != 0) slowness_disk_between_main_ram = 1/bandwidth_read; else slowness_disk_between_main_ram = 0; if(bandwidth_matrix[STARPU_MAIN_RAM][j] != 0) slowness_main_ram_between_node = 1/bandwidth_matrix[STARPU_MAIN_RAM][j]; else slowness_main_ram_between_node = 0; bandwidth_matrix[i][j] = 1/(slowness_disk_between_main_ram+slowness_main_ram_between_node); } else if (j == node) /* destination == disk */ { /* convert in slowness */ if(bandwidth_write != 0) slowness_disk_between_main_ram = 1/bandwidth_write; else slowness_disk_between_main_ram = 0; if(bandwidth_matrix[i][STARPU_MAIN_RAM] != 0) slowness_main_ram_between_node = 1/bandwidth_matrix[i][STARPU_MAIN_RAM]; else slowness_main_ram_between_node = 0; bandwidth_matrix[i][j] = 1/(slowness_disk_between_main_ram+slowness_main_ram_between_node); } else if (j > node || i > node) /* not affected by the node */ { bandwidth_matrix[i][j] = NAN; } } } /* save latency */ for(i = 0; i < STARPU_MAXNODES; ++i) { for(j = 0; j < STARPU_MAXNODES; ++j) { if (i == j && j == node) /* source == destination == node */ { latency_matrix[i][j] = 0; } else if (i == node) /* source == disk */ { latency_matrix[i][j] = (latency_write+latency_matrix[STARPU_MAIN_RAM][j]); } else if (j == node) /* destination == disk */ { latency_matrix[i][j] = (latency_read+latency_matrix[i][STARPU_MAIN_RAM]); } else if (j > node || i > node) /* not affected by the node */ { latency_matrix[i][j] = NAN; } } } }