diff --git a/docs/index.rst b/docs/index.rst --- a/docs/index.rst +++ b/docs/index.rst @@ -41,6 +41,7 @@ :maxdepth: 1 man/futhark.rst + man/futhark-autotune.rst man/futhark-bench.rst man/futhark-c.rst man/futhark-csharp.rst diff --git a/futhark.cabal b/futhark.cabal --- a/futhark.cabal +++ b/futhark.cabal @@ -1,13 +1,13 @@ cabal-version: 1.12 --- This file has been generated from package.yaml by hpack version 0.32.0. +-- This file has been generated from package.yaml by hpack version 0.33.0. -- -- see: https://github.com/sol/hpack -- --- hash: 6f8ebcf4ed2412437b0237651acb344581aeadbae4524bcbcb98189e63ce02f5 +-- hash: 9481b7958055b85d7ecbaed32f9b09d4af6d92255d1a11cf5161472dae048594 name: futhark -version: 0.12.2 +version: 0.12.3 synopsis: An optimising compiler for a functional, array-oriented language. description: Futhark is a small programming language designed to be compiled to efficient parallel code. It is a statically typed, data-parallel, @@ -194,6 +194,7 @@ Futhark.Optimise.Simplify.Lore Futhark.Optimise.Simplify.Rule Futhark.Optimise.Simplify.Rules + Futhark.Optimise.Sink Futhark.Optimise.TileLoops Futhark.Optimise.Unstream Futhark.Pass diff --git a/futlib/math.fut b/futlib/math.fut --- a/futlib/math.fut +++ b/futlib/math.fut @@ -92,6 +92,13 @@ val num_bits: i32 val get_bit: i32 -> t -> i32 val set_bit: i32 -> t -> i32 -> t + + -- | Count number of one bits. + val popc: t -> i32 + + -- | Count number of zero bits preceding the most significant set + -- bit. + val clz: t -> i32 } -- | An extension of `size`@mtype that further includes facilities for @@ -253,6 +260,8 @@ let get_bit (bit: i32) (x: t) = to_i32 ((x >> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc = intrinsics.popc8 + let clz = intrinsics.clz8 let iota (n: i8) = 0i8..1i8..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc = intrinsics.popc16 + let clz = intrinsics.clz16 let iota (n: i16) = 0i16..1i16..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc = intrinsics.popc32 + let clz = intrinsics.clz32 let iota (n: i32) = 0..1..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | intrinsics.zext_i32_i64 (b intrinsics.<< bit)) + let popc = intrinsics.popc64 + let clz = intrinsics.clz64 let iota (n: i64) = 0i64..1i64..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc x = intrinsics.popc8 (sign x) + let clz x = intrinsics.clz8 (sign x) let iota (n: u8) = 0u8..1u8..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc x = intrinsics.popc16 (sign x) + let clz x = intrinsics.clz16 (sign x) let iota (n: u16) = 0u16..1u16..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc x = intrinsics.popc32 (sign x) + let clz x = intrinsics.clz32 (sign x) let iota (n: u32) = 0u32..1u32..> i32 bit) & i32 1) let set_bit (bit: i32) (x: t) (b: i32) = ((x & i32 (intrinsics.!(1 intrinsics.<< bit))) | i32 (b intrinsics.<< bit)) + let popc x = intrinsics.popc64 (sign x) + let clz x = intrinsics.clz64 (sign x) let iota (n: u64) = 0u64..1u64..debugging = 0; cfg->logging = 0; cfg->preferred_device = ""; @@ -105,8 +102,7 @@ #define CU_DEV_ATTR(x) (CU_DEVICE_ATTRIBUTE_##x) #define device_query(dev,attrib) _device_query(dev, CU_DEV_ATTR(attrib)) -static int _device_query(CUdevice dev, CUdevice_attribute attrib) -{ +static int _device_query(CUdevice dev, CUdevice_attribute attrib) { int val; CUDA_SUCCEED(cuDeviceGetAttribute(&val, attrib, dev)); return val; @@ -114,20 +110,17 @@ #define CU_FUN_ATTR(x) (CU_FUNC_ATTRIBUTE_##x) #define function_query(fn,attrib) _function_query(dev, CU_FUN_ATTR(attrib)) -static int _function_query(CUfunction dev, CUfunction_attribute attrib) -{ +static int _function_query(CUfunction dev, CUfunction_attribute attrib) { int val; CUDA_SUCCEED(cuFuncGetAttribute(&val, attrib, dev)); return val; } -void set_preferred_device(struct cuda_config *cfg, const char *s) -{ +static void set_preferred_device(struct cuda_config *cfg, const char *s) { cfg->preferred_device = s; } -static int cuda_device_setup(struct cuda_context *ctx) -{ +static int cuda_device_setup(struct cuda_context *ctx) { char name[256]; int count, chosen = -1, best_cc = -1; int cc_major_best, cc_minor_best; @@ -185,8 +178,7 @@ return 0; } -static char *concat_fragments(const char *src_fragments[]) -{ +static char *concat_fragments(const char *src_fragments[]) { size_t src_len = 0; const char **p; @@ -204,8 +196,7 @@ return src; } -static const char *cuda_nvrtc_get_arch(CUdevice dev) -{ +static const char *cuda_nvrtc_get_arch(CUdevice dev) { struct { int major; int minor; @@ -244,8 +235,7 @@ } static char *cuda_nvrtc_build(struct cuda_context *ctx, const char *src, - const char *extra_opts[]) -{ + const char *extra_opts[]) { nvrtcProgram prog; NVRTC_SUCCEED(nvrtcCreateProgram(&prog, src, "futhark-cuda", 0, NULL, NULL)); int arch_set = 0, num_extra_opts; @@ -387,16 +377,14 @@ } } -static void dump_string_to_file(const char *file, const char *buf) -{ +static void dump_string_to_file(const char *file, const char *buf) { FILE *f = fopen(file, "w"); assert(f != NULL); assert(fputs(buf, f) != EOF); assert(fclose(f) == 0); } -static void load_string_from_file(const char *file, char **obuf, size_t *olen) -{ +static void load_string_from_file(const char *file, char **obuf, size_t *olen) { char *buf; size_t len; FILE *f = fopen(file, "r"); @@ -419,8 +407,7 @@ static void cuda_module_setup(struct cuda_context *ctx, const char *src_fragments[], - const char *extra_opts[]) -{ + const char *extra_opts[]) { char *ptx = NULL, *src = NULL; if (ctx->cfg.load_ptx_from == NULL && ctx->cfg.load_program_from == NULL) { @@ -457,8 +444,7 @@ } } -void cuda_setup(struct cuda_context *ctx, const char *src_fragments[], const char *extra_opts[]) -{ +static void cuda_setup(struct cuda_context *ctx, const char *src_fragments[], const char *extra_opts[]) { CUDA_SUCCEED(cuInit(0)); if (cuda_device_setup(ctx) != 0) { @@ -479,18 +465,16 @@ cuda_module_setup(ctx, src_fragments, extra_opts); } -CUresult cuda_free_all(struct cuda_context *ctx); +static CUresult cuda_free_all(struct cuda_context *ctx); -void cuda_cleanup(struct cuda_context *ctx) -{ +static void cuda_cleanup(struct cuda_context *ctx) { CUDA_SUCCEED(cuda_free_all(ctx)); CUDA_SUCCEED(cuModuleUnload(ctx->module)); CUDA_SUCCEED(cuCtxDestroy(ctx->cu_ctx)); } -CUresult cuda_alloc(struct cuda_context *ctx, size_t min_size, - const char *tag, CUdeviceptr *mem_out) -{ +static CUresult cuda_alloc(struct cuda_context *ctx, size_t min_size, + const char *tag, CUdeviceptr *mem_out) { if (min_size < sizeof(int)) { min_size = sizeof(int); } @@ -524,9 +508,8 @@ return res; } -CUresult cuda_free(struct cuda_context *ctx, CUdeviceptr mem, - const char *tag) -{ +static CUresult cuda_free(struct cuda_context *ctx, CUdeviceptr mem, + const char *tag) { size_t size; CUdeviceptr existing_mem; @@ -546,7 +529,7 @@ return res; } -CUresult cuda_free_all(struct cuda_context *ctx) { +static CUresult cuda_free_all(struct cuda_context *ctx) { CUdeviceptr mem; free_list_pack(&ctx->free_list); while (free_list_first(&ctx->free_list, &mem) == 0) { diff --git a/rts/c/free_list.h b/rts/c/free_list.h --- a/rts/c/free_list.h +++ b/rts/c/free_list.h @@ -16,7 +16,7 @@ int used; // Number of valid entries. }; -void free_list_init(struct free_list *l) { +static void free_list_init(struct free_list *l) { l->capacity = 30; // Picked arbitrarily. l->used = 0; l->entries = (struct free_list_entry*) malloc(sizeof(struct free_list_entry) * l->capacity); @@ -26,7 +26,7 @@ } /* Remove invalid entries from the free list. */ -void free_list_pack(struct free_list *l) { +static void free_list_pack(struct free_list *l) { int p = 0; for (int i = 0; i < l->capacity; i++) { if (l->entries[i].valid) { @@ -39,12 +39,12 @@ l->capacity = l->used; } -void free_list_destroy(struct free_list *l) { +static void free_list_destroy(struct free_list *l) { assert(l->used == 0); free(l->entries); } -int free_list_find_invalid(struct free_list *l) { +static int free_list_find_invalid(struct free_list *l) { int i; for (i = 0; i < l->capacity; i++) { if (!l->entries[i].valid) { @@ -54,7 +54,7 @@ return i; } -void free_list_insert(struct free_list *l, size_t size, fl_mem_t mem, const char *tag) { +static void free_list_insert(struct free_list *l, size_t size, fl_mem_t mem, const char *tag) { int i = free_list_find_invalid(l); if (i == l->capacity) { @@ -78,7 +78,7 @@ /* Find and remove a memory block of at least the desired size and tag. Returns 0 on success. */ -int free_list_find(struct free_list *l, const char *tag, size_t *size_out, fl_mem_t *mem_out) { +static int free_list_find(struct free_list *l, const char *tag, size_t *size_out, fl_mem_t *mem_out) { int i; for (i = 0; i < l->capacity; i++) { if (l->entries[i].valid && l->entries[i].tag == tag) { @@ -95,7 +95,7 @@ /* Remove the first block in the free list. Returns 0 if a block was removed, and nonzero if the free list was already empty. */ -int free_list_first(struct free_list *l, fl_mem_t *mem_out) { +static int free_list_first(struct free_list *l, fl_mem_t *mem_out) { for (int i = 0; i < l->capacity; i++) { if (l->entries[i].valid) { l->entries[i].valid = 0; diff --git a/rts/c/opencl.h b/rts/c/opencl.h --- a/rts/c/opencl.h +++ b/rts/c/opencl.h @@ -23,6 +23,7 @@ struct opencl_config { int debugging; + int profiling; int logging; int preferred_device_num; const char *preferred_platform; @@ -49,14 +50,15 @@ const char **size_classes; }; -void opencl_config_init(struct opencl_config *cfg, - int num_sizes, - const char *size_names[], - const char *size_vars[], - size_t *size_values, - const char *size_classes[]) { +static void opencl_config_init(struct opencl_config *cfg, + int num_sizes, + const char *size_names[], + const char *size_vars[], + size_t *size_values, + const char *size_classes[]) { cfg->debugging = 0; cfg->logging = 0; + cfg->profiling = 0; cfg->preferred_device_num = 0; cfg->preferred_platform = ""; cfg->preferred_device = ""; @@ -84,6 +86,13 @@ cfg->size_classes = size_classes; } +// A record of something that happened. +struct profiling_record { + cl_event *event; + int *runs; + int64_t *runtime; +}; + struct opencl_context { cl_device_id device; cl_context ctx; @@ -100,6 +109,10 @@ size_t max_local_memory; size_t lockstep_width; + + struct profiling_record *profiling_records; + int profiling_records_capacity; + int profiling_records_used; }; struct opencl_device_option { @@ -226,12 +239,12 @@ } } -void set_preferred_platform(struct opencl_config *cfg, const char *s) { +static void set_preferred_platform(struct opencl_config *cfg, const char *s) { cfg->preferred_platform = s; cfg->ignore_blacklist = 1; } -void set_preferred_device(struct opencl_config *cfg, const char *s) { +static void set_preferred_device(struct opencl_config *cfg, const char *s) { int x = 0; if (*s == '#') { s++; @@ -734,19 +747,81 @@ OPENCL_SUCCEED_FATAL(clCreateContext_error); cl_int clCreateCommandQueue_error; - cl_command_queue queue = clCreateCommandQueue(ctx->ctx, device_option.device, 0, &clCreateCommandQueue_error); + cl_command_queue queue = + clCreateCommandQueue(ctx->ctx, + device_option.device, + ctx->cfg.profiling ? CL_QUEUE_PROFILING_ENABLE : 0, + &clCreateCommandQueue_error); OPENCL_SUCCEED_FATAL(clCreateCommandQueue_error); return setup_opencl_with_command_queue(ctx, queue, srcs, required_types, extra_build_opts); } +// Count up the runtime all the profiling_records that occured during execution. +// Also clears the buffer of profiling_records. +static cl_int opencl_tally_profiling_records(struct opencl_context *ctx) { + cl_int err; + for (int i = 0; i < ctx->profiling_records_used; i++) { + struct profiling_record record = ctx->profiling_records[i]; + + cl_ulong start_t, end_t; + + if ((err = clGetEventProfilingInfo(*record.event, + CL_PROFILING_COMMAND_START, + sizeof(start_t), + &start_t, + NULL)) != CL_SUCCESS) { + return err; + } + + if ((err = clGetEventProfilingInfo(*record.event, + CL_PROFILING_COMMAND_END, + sizeof(end_t), + &end_t, + NULL)) != CL_SUCCESS) { + return err; + } + + // OpenCL provides nanosecond resolution, but we want + // microseconds. + *record.runs += 1; + *record.runtime += (end_t - start_t)/1000; + + if ((err = clReleaseEvent(*record.event)) != CL_SUCCESS) { + return err; + } + free(record.event); + } + + ctx->profiling_records_used = 0; + + return CL_SUCCESS; +} + +// If profiling, produce an event associated with a profiling record. +static cl_event* opencl_get_event(struct opencl_context *ctx, int *runs, int64_t *runtime) { + if (ctx->profiling_records_used == ctx->profiling_records_capacity) { + ctx->profiling_records_capacity *= 2; + ctx->profiling_records = + realloc(ctx->profiling_records, + ctx->profiling_records_capacity * + sizeof(struct profiling_record)); + } + cl_event *event = malloc(sizeof(cl_event)); + ctx->profiling_records[ctx->profiling_records_used].event = event; + ctx->profiling_records[ctx->profiling_records_used].runs = runs; + ctx->profiling_records[ctx->profiling_records_used].runtime = runtime; + ctx->profiling_records_used++; + return event; +} + // Allocate memory from driver. The problem is that OpenCL may perform // lazy allocation, so we cannot know whether an allocation succeeded // until the first time we try to use it. Hence we immediately // perform a write to see if the allocation succeeded. This is slow, // but the assumption is that this operation will be rare (most things // will go through the free list). -int opencl_alloc_actual(struct opencl_context *ctx, size_t size, cl_mem *mem_out) { +static int opencl_alloc_actual(struct opencl_context *ctx, size_t size, cl_mem *mem_out) { int error; *mem_out = clCreateBuffer(ctx->ctx, CL_MEM_READ_WRITE, size, NULL, &error); @@ -765,7 +840,7 @@ return error; } -int opencl_alloc(struct opencl_context *ctx, size_t min_size, const char *tag, cl_mem *mem_out) { +static int opencl_alloc(struct opencl_context *ctx, size_t min_size, const char *tag, cl_mem *mem_out) { if (min_size < sizeof(int)) { min_size = sizeof(int); } @@ -826,7 +901,7 @@ return error; } -int opencl_free(struct opencl_context *ctx, cl_mem mem, const char *tag) { +static int opencl_free(struct opencl_context *ctx, cl_mem mem, const char *tag) { size_t size; cl_mem existing_mem; @@ -847,7 +922,7 @@ return error; } -int opencl_free_all(struct opencl_context *ctx) { +static int opencl_free_all(struct opencl_context *ctx) { cl_mem mem; free_list_pack(&ctx->free_list); while (free_list_first(&ctx->free_list, &mem) == 0) { diff --git a/rts/c/values.h b/rts/c/values.h --- a/rts/c/values.h +++ b/rts/c/values.h @@ -436,6 +436,15 @@ #define BINARY_FORMAT_VERSION 2 #define IS_BIG_ENDIAN (!*(unsigned char *)&(uint16_t){1}) +static void flip_bytes(int elem_size, unsigned char *elem) { + for (int j=0; j>8) | (x<<8); - } - *(uint16_t*)dest = x; - return num_elems_read == 1 ? 0 : 1; -} - -static int read_le_4byte(void* dest) { - uint32_t x; - int num_elems_read = fread(&x, 4, 1, stdin); - if (IS_BIG_ENDIAN) { - x = - ((x>>24)&0xFF) | - ((x>>8) &0xFF00) | - ((x<<8) &0xFF0000) | - ((x<<24)&0xFF000000); - } - *(uint32_t*)dest = x; - return num_elems_read == 1 ? 0 : 1; -} - -static int read_le_8byte(void* dest) { - uint64_t x; - int num_elems_read = fread(&x, 8, 1, stdin); - if (IS_BIG_ENDIAN) { - x = - ((x>>56)&0xFFull) | - ((x>>40)&0xFF00ull) | - ((x>>24)&0xFF0000ull) | - ((x>>8) &0xFF000000ull) | - ((x<<8) &0xFF00000000ull) | - ((x<<24)&0xFF0000000000ull) | - ((x<<40)&0xFF000000000000ull) | - ((x<<56)&0xFF00000000000000ull); - } - *(uint64_t*)dest = x; - return num_elems_read == 1 ? 0 : 1; -} - -static int write_byte(void* dest) { - int num_elems_written = fwrite(dest, 1, 1, stdin); - return num_elems_written == 1 ? 0 : 1; -} - -static int write_le_2byte(void* dest) { - uint16_t x = *(uint16_t*)dest; - if (IS_BIG_ENDIAN) { - x = (x>>8) | (x<<8); - } - int num_elems_written = fwrite(&x, 2, 1, stdin); - return num_elems_written == 1 ? 0 : 1; -} - -static int write_le_4byte(void* dest) { - uint32_t x = *(uint32_t*)dest; - if (IS_BIG_ENDIAN) { - x = - ((x>>24)&0xFF) | - ((x>>8) &0xFF00) | - ((x<<8) &0xFF0000) | - ((x<<24)&0xFF000000); - } - int num_elems_written = fwrite(&x, 4, 1, stdin); - return num_elems_written == 1 ? 0 : 1; -} - -static int write_le_8byte(void* dest) { - uint64_t x = *(uint64_t*)dest; - if (IS_BIG_ENDIAN) { - x = - ((x>>56)&0xFFull) | - ((x>>40)&0xFF00ull) | - ((x>>24)&0xFF0000ull) | - ((x>>8) &0xFF000000ull) | - ((x<<8) &0xFF00000000ull) | - ((x<<24)&0xFF0000000000ull) | - ((x<<40)&0xFF000000000000ull) | - ((x<<56)&0xFF00000000000000ull); - } - int num_elems_written = fwrite(&x, 8, 1, stdin); - return num_elems_written == 1 ? 0 : 1; -} - //// Types struct primtype_info_t { @@ -552,54 +472,41 @@ const int size; // in bytes const writer write_str; // Write in text format. const str_reader read_str; // Read in text format. - const writer write_bin; // Write in binary format. - const bin_reader read_bin; // Read in binary format. }; static const struct primtype_info_t i8_info = {.binname = " i8", .type_name = "i8", .size = 1, - .write_str = (writer)write_str_i8, .read_str = (str_reader)read_str_i8, - .write_bin = (writer)write_byte, .read_bin = (bin_reader)read_byte}; + .write_str = (writer)write_str_i8, .read_str = (str_reader)read_str_i8}; static const struct primtype_info_t i16_info = {.binname = " i16", .type_name = "i16", .size = 2, - .write_str = (writer)write_str_i16, .read_str = (str_reader)read_str_i16, - .write_bin = (writer)write_le_2byte, .read_bin = (bin_reader)read_le_2byte}; + .write_str = (writer)write_str_i16, .read_str = (str_reader)read_str_i16}; static const struct primtype_info_t i32_info = {.binname = " i32", .type_name = "i32", .size = 4, - .write_str = (writer)write_str_i32, .read_str = (str_reader)read_str_i32, - .write_bin = (writer)write_le_4byte, .read_bin = (bin_reader)read_le_4byte}; + .write_str = (writer)write_str_i32, .read_str = (str_reader)read_str_i32}; static const struct primtype_info_t i64_info = {.binname = " i64", .type_name = "i64", .size = 8, - .write_str = (writer)write_str_i64, .read_str = (str_reader)read_str_i64, - .write_bin = (writer)write_le_8byte, .read_bin = (bin_reader)read_le_8byte}; + .write_str = (writer)write_str_i64, .read_str = (str_reader)read_str_i64}; static const struct primtype_info_t u8_info = {.binname = " u8", .type_name = "u8", .size = 1, - .write_str = (writer)write_str_u8, .read_str = (str_reader)read_str_u8, - .write_bin = (writer)write_byte, .read_bin = (bin_reader)read_byte}; + .write_str = (writer)write_str_u8, .read_str = (str_reader)read_str_u8}; static const struct primtype_info_t u16_info = {.binname = " u16", .type_name = "u16", .size = 2, - .write_str = (writer)write_str_u16, .read_str = (str_reader)read_str_u16, - .write_bin = (writer)write_le_2byte, .read_bin = (bin_reader)read_le_2byte}; + .write_str = (writer)write_str_u16, .read_str = (str_reader)read_str_u16}; static const struct primtype_info_t u32_info = {.binname = " u32", .type_name = "u32", .size = 4, - .write_str = (writer)write_str_u32, .read_str = (str_reader)read_str_u32, - .write_bin = (writer)write_le_4byte, .read_bin = (bin_reader)read_le_4byte}; + .write_str = (writer)write_str_u32, .read_str = (str_reader)read_str_u32}; static const struct primtype_info_t u64_info = {.binname = " u64", .type_name = "u64", .size = 8, - .write_str = (writer)write_str_u64, .read_str = (str_reader)read_str_u64, - .write_bin = (writer)write_le_8byte, .read_bin = (bin_reader)read_le_8byte}; + .write_str = (writer)write_str_u64, .read_str = (str_reader)read_str_u64}; static const struct primtype_info_t f32_info = {.binname = " f32", .type_name = "f32", .size = 4, - .write_str = (writer)write_str_f32, .read_str = (str_reader)read_str_f32, - .write_bin = (writer)write_le_4byte, .read_bin = (bin_reader)read_le_4byte}; + .write_str = (writer)write_str_f32, .read_str = (str_reader)read_str_f32}; static const struct primtype_info_t f64_info = {.binname = " f64", .type_name = "f64", .size = 8, - .write_str = (writer)write_str_f64, .read_str = (str_reader)read_str_f64, - .write_bin = (writer)write_le_8byte, .read_bin = (bin_reader)read_le_8byte}; + .write_str = (writer)write_str_f64, .read_str = (str_reader)read_str_f64}; static const struct primtype_info_t bool_info = {.binname = "bool", .type_name = "bool", .size = 1, - .write_str = (writer)write_str_bool, .read_str = (str_reader)read_str_bool, - .write_bin = (writer)write_byte, .read_bin = (bin_reader)read_byte}; + .write_str = (writer)write_str_bool, .read_str = (str_reader)read_str_bool}; static const struct primtype_info_t* primtypes[] = { &i8_info, &i16_info, &i32_info, &i64_info, @@ -692,8 +599,13 @@ uint64_t elem_count = 1; for (int i=0; iread_str(buf, dest); } else { read_bin_ensure_scalar(expected_type); - return expected_type->read_bin(dest); + size_t elem_size = expected_type->size; + int num_elems_read = fread(dest, elem_size, 1, stdin); + if (IS_BIG_ENDIAN) { + flip_bytes(elem_size, (unsigned char*) dest); + } + return num_elems_read == 1 ? 0 : 1; } } diff --git a/rts/csharp/scalar.cs b/rts/csharp/scalar.cs --- a/rts/csharp/scalar.cs +++ b/rts/csharp/scalar.cs @@ -318,6 +318,82 @@ private static float futhark_lerp32(float v0, float v1, float t){return v0 + (v1-v0)*t;} private static double futhark_lerp64(double v0, double v1, double t){return v0 + (v1-v0)*t;} +int futhark_popc8 (sbyte x) { + int c = 0; + for (; x != 0; ++c) { + x &= (sbyte)(x - 1); + } + return c; + } + +int futhark_popc16 (short x) { + int c = 0; + for (; x != 0; ++c) { + x &= (short)(x - 1); + } + return c; +} + +int futhark_popc32 (int x) { + int c = 0; + for (; x != 0; ++c) { + x &= x - 1; + } + return c; +} + +int futhark_popc64 (long x) { + int c = 0; + for (; x != 0; ++c) { + x &= x - 1; + } + return c; +} + +int futhark_clzz8 (sbyte x) { + int n = 0; + int bits = sizeof(sbyte) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; +} + +int futhark_clzz16 (short x) { + int n = 0; + int bits = sizeof(short) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; +} + +int futhark_clzz32 (int x) { + int n = 0; + int bits = sizeof(int) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; +} + +int futhark_clzz64 (long x) { + int n = 0; + int bits = sizeof(long) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; +} + private static bool llt (bool x, bool y){return (!x && y);} private static bool lle (bool x, bool y){return (!x || y);} diff --git a/rts/python/scalar.py b/rts/python/scalar.py --- a/rts/python/scalar.py +++ b/rts/python/scalar.py @@ -197,6 +197,26 @@ sext_i8_i64 = sext_i16_i64 = sext_i32_i64 = sext_i64_i64 = sext_T_i64 itob_i8_bool = itob_i16_bool = itob_i32_bool = itob_i64_bool = itob_T_bool +def clz_T(x): + n = np.int32(0) + bits = x.itemsize * 8 + for i in range(bits): + if x < 0: + break + n += 1 + x <<= np.int8(1) + return n + +def popc_T(x): + c = np.int32(0) + while x != 0: + x &= x - np.int8(1) + c += np.int8(1) + return c + +futhark_popc8 = futhark_popc16 = futhark_popc32 = futhark_popc64 = popc_T +futhark_clzz8 = futhark_clzz16 = futhark_clzz32 = futhark_clzz64 = clz_T + def ssignum(x): return np.sign(x) diff --git a/src/Futhark/Analysis/CallGraph.hs b/src/Futhark/Analysis/CallGraph.hs --- a/src/Futhark/Analysis/CallGraph.hs +++ b/src/Futhark/Analysis/CallGraph.hs @@ -21,20 +21,19 @@ where expand ftab f = M.insert (funDefName f) f ftab -- | The call graph is just a mapping from a function name, i.e., the --- caller, to a list of the names of functions called by the function. --- The order of this list is not significant. +-- caller, to a set of the names of functions called *directly* (not +-- transitively!) by the function. type CallGraph = M.Map Name (S.Set Name) --- | @buildCallGraph prog@ build the program's Call Graph. The representation --- is a hashtable that maps function names to a list of callee names. +-- | @buildCallGraph prog@ build the program's call graph. buildCallGraph :: Prog SOACS -> CallGraph buildCallGraph prog = foldl' (buildCGfun ftable) M.empty entry_points - where entry_points = map funDefName $ filter (isJust . funDefEntryPoint) $ progFunctions prog + where entry_points = map funDefName $ filter (isJust . funDefEntryPoint) $ + progFunctions prog ftable = buildFunctionTable prog --- | @buildCallGraph ftable cg fname@ updates Call Graph @cg@ with the --- contributions of function @fname@, and recursively, with the --- contributions of the callees of @fname@. +-- | @buildCallGraph ftable cg fname@ updates @cg@ with the +-- contributions of function @fname@. buildCGfun :: FunctionTable -> CallGraph -> Name -> CallGraph buildCGfun ftable cg fname = -- Check if function is a non-builtin that we have not already diff --git a/src/Futhark/Analysis/Range.hs b/src/Futhark/Analysis/Range.hs --- a/src/Futhark/Analysis/Range.hs +++ b/src/Futhark/Analysis/Range.hs @@ -3,6 +3,7 @@ ( rangeAnalysis , runRangeM , RangeM + , analyseFun , analyseExp , analyseLambda , analyseBody diff --git a/src/Futhark/CLI/Dev.hs b/src/Futhark/CLI/Dev.hs --- a/src/Futhark/CLI/Dev.hs +++ b/src/Futhark/CLI/Dev.hs @@ -39,6 +39,7 @@ import Futhark.Pass.Simplify import Futhark.Optimise.InPlaceLowering import Futhark.Optimise.DoubleBuffer +import Futhark.Optimise.Sink import Futhark.Optimise.TileLoops import Futhark.Optimise.Unstream import Futhark.Pass.KernelBabysitting @@ -298,6 +299,7 @@ , kernelsPassOption babysitKernels [] , kernelsPassOption tileLoops [] , kernelsPassOption unstream [] + , kernelsPassOption sink [] , typedPassOption soacsProg Kernels extractKernels [] , typedPassOption kernelsProg ExplicitMemory explicitAllocations "a" diff --git a/src/Futhark/CodeGen/Backends/CCUDA.hs b/src/Futhark/CodeGen/Backends/CCUDA.hs --- a/src/Futhark/CodeGen/Backends/CCUDA.hs +++ b/src/Futhark/CodeGen/Backends/CCUDA.hs @@ -209,7 +209,6 @@ error $ "CUDA backend does not support '" ++ space ++ "' memory space." callKernel :: GC.OpCompiler OpenCL () -callKernel (HostCode c) = GC.compileCode c callKernel (GetSize v key) = GC.stm [C.cstm|$id:v = ctx->sizes.$id:key;|] callKernel (CmpSizeLe v key x) = do diff --git a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs --- a/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +++ b/src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs @@ -234,6 +234,7 @@ [C.cedecl|struct $id:s { int detail_memory; int debugging; + int profiling; typename lock_t lock; char *error; $sdecls:fields @@ -252,7 +253,7 @@ if (ctx == NULL) { return NULL; } - ctx->debugging = ctx->detail_memory = cfg->cu_cfg.debugging; + ctx->profiling = ctx->debugging = ctx->detail_memory = cfg->cu_cfg.debugging; ctx->cuda.cfg = cfg->cu_cfg; create_lock(&ctx->lock); @@ -286,6 +287,20 @@ [C.cedecl|char* $id:s(struct $id:ctx* ctx) { return ctx->error; }|]) + + + GC.publicDef_ "context_pause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + (void)ctx; + }|]) + + GC.publicDef_ "context_unpause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + (void)ctx; + }|]) + where loadKernelByName name = [C.cstm|CUDA_SUCCEED(cuModuleGetFunction(&ctx->$id:name, diff --git a/src/Futhark/CodeGen/Backends/COpenCL.hs b/src/Futhark/CodeGen/Backends/COpenCL.hs --- a/src/Futhark/CodeGen/Backends/COpenCL.hs +++ b/src/Futhark/CodeGen/Backends/COpenCL.hs @@ -26,9 +26,14 @@ res <- ImpGen.compileProg prog case res of Left err -> return $ Left err - Right (Program opencl_code opencl_prelude kernel_names types sizes prog') -> + Right (Program opencl_code opencl_prelude kernel_names types sizes prog') -> do + let cost_centres = + [copyDevToDev, copyDevToHost, copyHostToDev, + copyScalarToDev, copyScalarFromDev] + ++ kernel_names Right <$> GC.compileProg operations - (generateBoilerplate opencl_code opencl_prelude kernel_names types sizes) + (generateBoilerplate opencl_code opencl_prelude + cost_centres kernel_names types sizes) include_opencl_h [Space "device", DefaultSpace] cliOptions prog' where operations :: GC.Operations OpenCL () @@ -51,6 +56,20 @@ "#include ", "#endif"] +copyDevToDev, copyDevToHost, copyHostToDev, copyScalarToDev, copyScalarFromDev :: String +copyDevToDev = "copy_dev_to_dev" +copyDevToHost = "copy_dev_to_host" +copyHostToDev = "copy_host_to_dev" +copyScalarToDev = "copy_scalar_to_dev" +copyScalarFromDev = "copy_scalar_from_dev" + +profilingEvent :: String -> C.Exp +profilingEvent name = + [C.cexp|ctx->profiling_paused ? NULL + : opencl_get_event(&ctx->opencl, + &ctx->$id:(kernelRuns name), + &ctx->$id:(kernelRuntime name))|] + cliOptions :: [Option] cliOptions = [ Option { optionLongName = "platform" , optionShortName = Just 'p' @@ -148,6 +167,11 @@ panic(1, "When loading tuning from '%s': %s\n", optarg, ret); }}|] } + , Option { optionLongName = "profile" + , optionShortName = Just 'P' + , optionArgument = NoArgument + , optionAction = [C.cstm|futhark_context_config_set_profiling(cfg, 1);|] + } ] -- We detect the special case of writing a constant and turn it into a @@ -168,7 +192,7 @@ clEnqueueWriteBuffer(ctx->opencl.queue, $exp:mem, $exp:blocking, $exp:i * sizeof($ty:t), sizeof($ty:t), &$id:val', - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent copyScalarToDev))); }|] writeOpenCLScalar _ _ _ space _ _ = error $ "Cannot write to '" ++ space ++ "' memory space." @@ -181,7 +205,7 @@ clEnqueueReadBuffer(ctx->opencl.queue, $exp:mem, CL_TRUE, $exp:i * sizeof($ty:t), sizeof($ty:t), &$id:val, - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent copyScalarFromDev))); |] return [C.cexp|$id:val|] readOpenCLScalar _ _ _ space _ = @@ -199,7 +223,6 @@ deallocateOpenCLBuffer _ _ space = error $ "Cannot deallocate in '" ++ space ++ "' space" - copyOpenCLMemory :: GC.Copy OpenCL () -- The read/write/copy-buffer functions fail if the given offset is -- out of bounds, even if asked to read zero bytes. We protect with a @@ -211,7 +234,7 @@ clEnqueueReadBuffer(ctx->opencl.queue, $exp:srcmem, CL_TRUE, $exp:srcidx, $exp:nbytes, $exp:destmem + $exp:destidx, - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent copyHostToDev))); } |] copyOpenCLMemory destmem destidx (Space "device") srcmem srcidx DefaultSpace nbytes = @@ -221,7 +244,7 @@ clEnqueueWriteBuffer(ctx->opencl.queue, $exp:destmem, CL_TRUE, $exp:destidx, $exp:nbytes, $exp:srcmem + $exp:srcidx, - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent copyDevToHost))); } |] copyOpenCLMemory destmem destidx (Space "device") srcmem srcidx (Space "device") nbytes = @@ -234,7 +257,7 @@ $exp:srcmem, $exp:destmem, $exp:srcidx, $exp:destidx, $exp:nbytes, - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent copyDevToDev))); if (ctx->debugging) { OPENCL_SUCCEED_FATAL(clFinish(ctx->opencl.queue)); } @@ -299,8 +322,6 @@ callKernel (GetSizeMax v size_class) = let field = "max_" ++ pretty size_class in GC.stm [C.cstm|$id:v = ctx->opencl.$id:field;|] -callKernel (HostCode c) = - GC.compileCode c callKernel (LaunchKernel name args num_workgroups workgroup_size) = do zipWithM_ setKernelArg [(0::Int)..] args @@ -356,13 +377,11 @@ OPENCL_SUCCEED_OR_RETURN( clEnqueueNDRangeKernel(ctx->opencl.queue, ctx->$id:kernel_name, $int:kernel_rank, NULL, $id:global_work_size, $id:local_work_size, - 0, NULL, NULL)); + 0, NULL, $exp:(profilingEvent kernel_name))); if (ctx->debugging) { OPENCL_SUCCEED_FATAL(clFinish(ctx->opencl.queue)); $id:time_end = get_wall_time(); long int $id:time_diff = $id:time_end - $id:time_start; - ctx->$id:(kernelRuntime kernel_name) += $id:time_diff; - ctx->$id:(kernelRuns kernel_name)++; fprintf(stderr, "kernel %s runtime: %ldus\n", $string:kernel_name, $id:time_diff); } diff --git a/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs b/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs --- a/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs +++ b/src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs @@ -17,14 +17,14 @@ import Futhark.CodeGen.OpenCL.Heuristics import Futhark.Util (chunk, zEncodeString) -generateBoilerplate :: String -> String -> [String] -> [PrimType] +generateBoilerplate :: String -> String -> [String] -> [String] -> [PrimType] -> M.Map Name SizeClass -> GC.CompilerM OpenCL () () -generateBoilerplate opencl_code opencl_prelude kernel_names types sizes = do +generateBoilerplate opencl_code opencl_prelude profiling_centres kernel_names types sizes = do final_inits <- GC.contextFinalInits let (ctx_opencl_fields, ctx_opencl_inits, top_decls, later_top_decls) = - openClDecls kernel_names opencl_code opencl_prelude + openClDecls profiling_centres kernel_names opencl_code opencl_prelude GC.earlyDecls top_decls @@ -103,9 +103,15 @@ GC.publicDef_ "context_config_set_debugging" GC.InitDecl $ \s -> ([C.cedecl|void $id:s(struct $id:cfg* cfg, int flag);|], [C.cedecl|void $id:s(struct $id:cfg* cfg, int flag) { - cfg->opencl.logging = cfg->opencl.debugging = flag; + cfg->opencl.profiling = cfg->opencl.logging = cfg->opencl.debugging = flag; }|]) + GC.publicDef_ "context_config_set_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:cfg* cfg, int flag);|], + [C.cedecl|void $id:s(struct $id:cfg* cfg, int flag) { + cfg->opencl.profiling = flag; + }|]) + GC.publicDef_ "context_config_set_logging" GC.InitDecl $ \s -> ([C.cedecl|void $id:s(struct $id:cfg* cfg, int flag);|], [C.cedecl|void $id:s(struct $id:cfg* cfg, int flag) { @@ -221,6 +227,8 @@ [C.cedecl|struct $id:s { int detail_memory; int debugging; + int profiling; + int profiling_paused; int logging; typename lock_t lock; char *error; @@ -240,8 +248,15 @@ ctx->opencl.cfg = cfg->opencl; ctx->detail_memory = cfg->opencl.debugging; ctx->debugging = cfg->opencl.debugging; + ctx->profiling = cfg->opencl.profiling; + ctx->profiling_paused = 0; ctx->logging = cfg->opencl.logging; ctx->error = NULL; + ctx->opencl.profiling_records_capacity = 200; + ctx->opencl.profiling_records_used = 0; + ctx->opencl.profiling_records = + malloc(ctx->opencl.profiling_records_capacity * + sizeof(struct profiling_record)); create_lock(&ctx->lock); $stms:init_fields @@ -298,6 +313,8 @@ ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], [C.cedecl|void $id:s(struct $id:ctx* ctx) { free_lock(&ctx->lock); + opencl_tally_profiling_records(&ctx->opencl); + free(ctx->opencl.profiling_records); free(ctx); }|]) @@ -316,6 +333,18 @@ return error; }|]) + GC.publicDef_ "context_pause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + ctx->profiling_paused = 1; + }|]) + + GC.publicDef_ "context_unpause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + ctx->profiling_paused = 0; + }|]) + GC.publicDef_ "context_clear_caches" GC.InitDecl $ \s -> ([C.cedecl|int $id:s(struct $id:ctx* ctx);|], [C.cedecl|int $id:s(struct $id:ctx* ctx) { @@ -329,11 +358,12 @@ return ctx->opencl.queue; }|]) - mapM_ GC.debugReport $ openClReport kernel_names + GC.profileReport [C.citem|OPENCL_SUCCEED_FATAL(opencl_tally_profiling_records(&ctx->opencl));|] + mapM_ GC.profileReport $ openClReport profiling_centres -openClDecls :: [String] -> String -> String +openClDecls :: [String] -> [String] -> String -> String -> ([C.FieldGroup], [C.Stm], [C.Definition], [C.Definition]) -openClDecls kernel_names opencl_program opencl_prelude = +openClDecls profiling_centres kernel_names opencl_program opencl_prelude = (ctx_fields, ctx_inits, openCL_boilerplate, openCL_load) where opencl_program_fragments = -- Some C compilers limit the size of literal strings, so @@ -344,12 +374,13 @@ ctx_fields = [ [C.csdecl|int total_runs;|], [C.csdecl|long int total_runtime;|] ] ++ + [ [C.csdecl|typename cl_kernel $id:name;|] + | name <- kernel_names ] ++ concat - [ [ [C.csdecl|typename cl_kernel $id:name;|] - , [C.csdecl|int $id:(kernelRuntime name);|] + [ [ [C.csdecl|typename int64_t $id:(kernelRuntime name);|] , [C.csdecl|int $id:(kernelRuns name);|] ] - | name <- kernel_names ] + | name <- profiling_centres ] ctx_inits = [ [C.cstm|ctx->total_runs = 0;|], @@ -358,7 +389,7 @@ [ [ [C.cstm|ctx->$id:(kernelRuntime name) = 0;|] , [C.cstm|ctx->$id:(kernelRuns name) = 0;|] ] - | name <- kernel_names ] + | name <- profiling_centres ] openCL_load = [ [C.cedecl| @@ -382,7 +413,7 @@ $esc:("typedef cl_mem fl_mem_t;") $esc:free_list_h $esc:openCL_h - const char *opencl_program[] = {$inits:program_fragments};|] + static const char *opencl_program[] = {$inits:program_fragments};|] loadKernelByName :: String -> C.Stm loadKernelByName name = [C.cstm|{ @@ -405,9 +436,8 @@ report_kernels = concatMap reportKernel names format_string name = let padding = replicate (longest_name - length name) ' ' - in unwords ["Kernel", - name ++ padding, - "executed %6d times, with average runtime: %6ldus\tand total runtime: %6ldus\n"] + in unwords [name ++ padding, + "ran %5d times; avg: %8ldus; total: %8ldus\n"] reportKernel name = let runs = kernelRuns name total_runtime = kernelRuntime name @@ -422,8 +452,8 @@ [C.citem|ctx->total_runs += ctx->$id:runs;|]] report_total = [C.citem| - if (ctx->debugging) { - fprintf(stderr, "Ran %d kernels with cumulative runtime: %6ldus\n", + if (ctx->profiling) { + fprintf(stderr, "%d operations with cumulative runtime: %6ldus\n", ctx->total_runs, ctx->total_runtime); } |] diff --git a/src/Futhark/CodeGen/Backends/CSOpenCL.hs b/src/Futhark/CodeGen/Backends/CSOpenCL.hs --- a/src/Futhark/CodeGen/Backends/CSOpenCL.hs +++ b/src/Futhark/CodeGen/Backends/CSOpenCL.hs @@ -134,8 +134,6 @@ Imp.SizeThreshold{} -> "MaxThreshold" Imp.SizeLocalMemory -> "MaxLocalMemory" -callKernel (Imp.HostCode c) = CS.compileCode c - callKernel (Imp.LaunchKernel name args num_workgroups workgroup_size) = do num_workgroups' <- mapM CS.compileExp num_workgroups workgroup_size' <- mapM CS.compileExp workgroup_size diff --git a/src/Futhark/CodeGen/Backends/GenericC.hs b/src/Futhark/CodeGen/Backends/GenericC.hs --- a/src/Futhark/CodeGen/Backends/GenericC.hs +++ b/src/Futhark/CodeGen/Backends/GenericC.hs @@ -51,7 +51,7 @@ , headerDecl , publicDef , publicDef_ - , debugReport + , profileReport , HeaderSection(..) , libDecl , earlyDecls @@ -101,7 +101,7 @@ , compHeaderDecls :: M.Map HeaderSection (DL.DList C.Definition) , compLibDecls :: DL.DList C.Definition , compCtxFields :: DL.DList (String, C.Type, Maybe C.Exp) - , compDebugItems :: DL.DList C.BlockItem + , compProfileItems :: DL.DList C.BlockItem , compDeclaredMem :: [(VName,Space)] } @@ -116,7 +116,7 @@ , compHeaderDecls = mempty , compLibDecls = mempty , compCtxFields = mempty - , compDebugItems = mempty + , compProfileItems = mempty , compDeclaredMem = mempty } @@ -408,9 +408,9 @@ contextField name ty initial = modify $ \s -> s { compCtxFields = compCtxFields s <> DL.singleton (name,ty,initial) } -debugReport :: C.BlockItem -> CompilerM op s () -debugReport x = modify $ \s -> - s { compDebugItems = compDebugItems s <> DL.singleton x } +profileReport :: C.BlockItem -> CompilerM op s () +profileReport x = modify $ \s -> + s { compProfileItems = compProfileItems s <> DL.singleton x } stm :: C.Stm -> CompilerM op s () stm (C.Block items _) = mapM_ item items @@ -656,7 +656,7 @@ paramsTypes = map paramType -- Let's hope we don't need the size for anything, because we are -- just making something up. - where paramType (MemParam _ space) = Mem (ConstSize 0) space + where paramType (MemParam _ space) = Mem space paramType (ScalarParam _ t) = Scalar t --- Entry points. @@ -1083,7 +1083,7 @@ ty <- valueDescToCType vd item [C.citem|$ty:ty *$id:dest;|] - let t' = primTypeToCType t + let t' = signedPrimTypeToCType ept t rank = length dims name = arrayName t ept rank dims_exps = [ [C.cexp|$id:shape[$int:j]|] | j <- [0..rank-1] ] @@ -1160,11 +1160,18 @@ cli_entry_point_function_name = "futrts_cli_entry_" ++ entry_point_name entry_point_function_name <- publicName $ "entry_" ++ entry_point_name + pause_profiling <- publicName "context_pause_profiling" + unpause_profiling <- publicName "context_unpause_profiling" + let run_it = [C.citems| int r; /* Run the program once. */ $stms:pack_input assert($id:sync_ctx(ctx) == 0); + // Only profile last run. + if (profile_run) { + $id:unpause_profiling(ctx); + } t_start = get_wall_time(); r = $id:entry_point_function_name(ctx, $args:(map addrOf output_vals), @@ -1173,6 +1180,9 @@ panic(1, "%s", $id:error_ctx(ctx)); } assert($id:sync_ctx(ctx) == 0); + if (profile_run) { + $id:pause_profiling(ctx); + } t_end = get_wall_time(); long int elapsed_usec = t_end - t_start; if (time_runs && runtime_file != NULL) { @@ -1183,8 +1193,11 @@ return ([C.cedecl|static void $id:cli_entry_point_function_name($ty:ctx_ty *ctx) { typename int64_t t_start, t_end; - int time_runs; + int time_runs = 0, profile_run = 0; + // We do not want to profile all the initialisation. + $id:pause_profiling(ctx); + /* Declare and read input. */ set_binary_mode(stdin); $items:input_items @@ -1192,13 +1205,14 @@ /* Warmup run */ if (perform_warmup) { - time_runs = 0; $items:run_it $stms:free_outputs } time_runs = 1; /* Proper run. */ for (int run = 0; run < num_runs; run++) { + // Only profile last run. + profile_run = run == num_runs -1; $items:run_it if (run < num_runs-1) { $stms:free_outputs @@ -1282,7 +1296,7 @@ -- | Produce header and implementation files. asLibrary :: CParts -> (String, String) -asLibrary parts = (cHeader parts, cUtils parts <> cLib parts) +asLibrary parts = ("#pragma once\n\n" <> cHeader parts, cUtils parts <> cLib parts) -- | As executable with command-line interface. asExecutable :: CParts -> String @@ -1307,7 +1321,6 @@ option_parser = generateOptionParser "parse_options" $ benchmarkOptions++options let headerdefs = [C.cunit| -$esc:("#pragma once\n") $esc:("/*\n * Headers\n*/\n") $esc:("#include ") $esc:("#include ") @@ -1474,16 +1487,16 @@ entry_points <- mapM (uncurry onEntryPoint) $ filter (functionEntry . snd) funs extra mapM_ libDecl $ concat memfuns - debugreport <- gets $ DL.toList . compDebugItems + profilereport <- gets $ DL.toList . compProfileItems ctx_ty <- contextType headerDecl MiscDecl [C.cedecl|void futhark_debugging_report($ty:ctx_ty *ctx);|] libDecl [C.cedecl|void futhark_debugging_report($ty:ctx_ty *ctx) { - if (ctx->detail_memory) { + if (ctx->detail_memory || ctx->profiling) { $items:memreport } - if (ctx->debugging) { - $items:debugreport + if (ctx->profiling) { + $items:profilereport } }|] diff --git a/src/Futhark/CodeGen/Backends/GenericCSharp.hs b/src/Futhark/CodeGen/Backends/GenericCSharp.hs --- a/src/Futhark/CodeGen/Backends/GenericCSharp.hs +++ b/src/Futhark/CodeGen/Backends/GenericCSharp.hs @@ -325,7 +325,7 @@ paramsTypes = map paramType paramType :: Imp.Param -> Imp.Type -paramType (Imp.MemParam _ space) = Imp.Mem (Imp.ConstSize 0) space +paramType (Imp.MemParam _ space) = Imp.Mem space paramType (Imp.ScalarParam _ t) = Imp.Scalar t compileOutput :: Imp.Param -> (CSExp, CSType) @@ -584,7 +584,7 @@ compileType :: Imp.Type -> CSType compileType (Imp.Scalar p) = compilePrimTypeToAST p -compileType (Imp.Mem _ space) = rawMemCSType space +compileType (Imp.Mem space) = rawMemCSType space compilePrimTypeToAST :: PrimType -> CSType compilePrimTypeToAST (IntType Int8) = Primitive $ CSInt Int8T diff --git a/src/Futhark/CodeGen/Backends/GenericPython.hs b/src/Futhark/CodeGen/Backends/GenericPython.hs --- a/src/Futhark/CodeGen/Backends/GenericPython.hs +++ b/src/Futhark/CodeGen/Backends/GenericPython.hs @@ -230,7 +230,7 @@ paramsTypes :: [Imp.Param] -> [Imp.Type] paramsTypes = map paramType - where paramType (Imp.MemParam _ space) = Imp.Mem (Imp.ConstSize 0) space + where paramType (Imp.MemParam _ space) = Imp.Mem space paramType (Imp.ScalarParam _ t) = Imp.Scalar t compileOutput :: [Imp.Param] -> [PyExp] diff --git a/src/Futhark/CodeGen/Backends/PyOpenCL.hs b/src/Futhark/CodeGen/Backends/PyOpenCL.hs --- a/src/Futhark/CodeGen/Backends/PyOpenCL.hs +++ b/src/Futhark/CodeGen/Backends/PyOpenCL.hs @@ -138,8 +138,6 @@ callKernel (Imp.GetSizeMax v size_class) = Py.stm $ Assign (Var (Py.compileName v)) $ Var $ "self.max_" ++ pretty size_class -callKernel (Imp.HostCode c) = - Py.compileCode c callKernel (Imp.LaunchKernel name args num_workgroups workgroup_size) = do num_workgroups' <- mapM (fmap asLong . Py.compileExp) num_workgroups diff --git a/src/Futhark/CodeGen/Backends/SequentialC.hs b/src/Futhark/CodeGen/Backends/SequentialC.hs --- a/src/Futhark/CodeGen/Backends/SequentialC.hs +++ b/src/Futhark/CodeGen/Backends/SequentialC.hs @@ -72,6 +72,7 @@ [C.cedecl|struct $id:s { int detail_memory; int debugging; + int profiling; typename lock_t lock; char *error; $sdecls:fields @@ -112,6 +113,18 @@ ctx->error = NULL; return error; }|]) + + GC.publicDef_ "context_pause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + (void)ctx; + }|]) + + GC.publicDef_ "context_unpause_profiling" GC.InitDecl $ \s -> + ([C.cedecl|void $id:s(struct $id:ctx* ctx);|], + [C.cedecl|void $id:s(struct $id:ctx* ctx) { + (void)ctx; + }|]) copySequentialMemory :: GC.Copy Imp.Sequential () copySequentialMemory destmem destidx DefaultSpace srcmem srcidx DefaultSpace nbytes = diff --git a/src/Futhark/CodeGen/Backends/SimpleRepresentation.hs b/src/Futhark/CodeGen/Backends/SimpleRepresentation.hs --- a/src/Futhark/CodeGen/Backends/SimpleRepresentation.hs +++ b/src/Futhark/CodeGen/Backends/SimpleRepresentation.hs @@ -1,8 +1,7 @@ {-# LANGUAGE QuasiQuotes #-} -- | Simple C runtime representation. module Futhark.CodeGen.Backends.SimpleRepresentation - ( sameRepresentation - , tupleField + ( tupleField , tupleFieldExp , funName , defaultMemBlockType @@ -60,20 +59,6 @@ signedPrimTypeToCType TypeDirect (IntType t) = intTypeToCType t signedPrimTypeToCType _ t = primTypeToCType t --- | True if both types map to the same runtime representation. This --- is the case if they are identical modulo uniqueness. -sameRepresentation :: [Type] -> [Type] -> Bool -sameRepresentation ets1 ets2 - | length ets1 == length ets2 = - and $ zipWith sameRepresentation' ets1 ets2 - | otherwise = False - -sameRepresentation' :: Type -> Type -> Bool -sameRepresentation' (Scalar t1) (Scalar t2) = - t1 == t2 -sameRepresentation' (Mem _ space1) (Mem _ space2) = space1 == space2 -sameRepresentation' _ _ = False - -- | @tupleField i@ is the name of field number @i@ in a tuple. tupleField :: Int -> String tupleField i = "v" ++ show i @@ -98,6 +83,7 @@ cIntOps :: [C.Definition] cIntOps = concatMap (`map` [minBound..maxBound]) ops + ++ cIntPrimFuns where ops = [mkAdd, mkSub, mkMul, mkUDiv, mkUMod, mkSDiv, mkSMod, @@ -213,6 +199,136 @@ uintCmpOp s e t = [C.cedecl|static inline char $id:(taggedI s t)($ty:ct x, $ty:ct y) { return $exp:e; }|] where ct = uintTypeToCType t + +cIntPrimFuns :: [C.Definition] +cIntPrimFuns = + [C.cunit| +$esc:("#ifdef __OPENCL_VERSION__") + typename int32_t $id:(funName' "popc8") (typename int8_t x) { + return popcount(x); + } + typename int32_t $id:(funName' "popc16") (typename int16_t x) { + return popcount(x); + } + typename int32_t $id:(funName' "popc32") (typename int32_t x) { + return popcount(x); + } + typename int32_t $id:(funName' "popc64") (typename int64_t x) { + return popcount(x); + } +$esc:("#elif __CUDA_ARCH__") + typename int32_t $id:(funName' "popc8") (typename int8_t x) { + return __popc(zext_i8_i32(x)); + } + typename int32_t $id:(funName' "popc16") (typename int16_t x) { + return __popc(zext_i16_i32(x)); + } + typename int32_t $id:(funName' "popc32") (typename int32_t x) { + return __popc(x); + } + typename int32_t $id:(funName' "popc64") (typename int64_t x) { + return __popcll(x); + } +$esc:("#else") + typename int32_t $id:(funName' "popc8") (typename int8_t x) { + int c = 0; + for (; x; ++c) { + x &= x - 1; + } + return c; + } + typename int32_t $id:(funName' "popc16") (typename int16_t x) { + int c = 0; + for (; x; ++c) { + x &= x - 1; + } + return c; + } + typename int32_t $id:(funName' "popc32") (typename int32_t x) { + int c = 0; + for (; x; ++c) { + x &= x - 1; + } + return c; + } + typename int32_t $id:(funName' "popc64") (typename int64_t x) { + int c = 0; + for (; x; ++c) { + x &= x - 1; + } + return c; + } +$esc:("#endif") + +$esc:("#ifdef __OPENCL_VERSION__") + typename int32_t $id:(funName' "clz8") (typename int8_t x) { + return clz(x); + } + typename int32_t $id:(funName' "clz16") (typename int16_t x) { + return clz(x); + } + typename int32_t $id:(funName' "clz32") (typename int32_t x) { + return clz(x); + } + typename int32_t $id:(funName' "clz64") (typename int64_t x) { + return clz(x); + } +$esc:("#elif __CUDA_ARCH__") + typename int32_t $id:(funName' "clz8") (typename int8_t x) { + return __clz(zext_i8_i32(x))-24; + } + typename int32_t $id:(funName' "clz16") (typename int16_t x) { + return __clz(zext_i16_i32(x))-16; + } + typename int32_t $id:(funName' "clz32") (typename int32_t x) { + return __clz(x); + } + typename int32_t $id:(funName' "clz64") (typename int64_t x) { + return __clzll(x); + } +$esc:("#else") + typename int32_t $id:(funName' "clz8") (typename int8_t x) { + int n = 0; + int bits = sizeof(x) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; + } + typename int32_t $id:(funName' "clz16") (typename int16_t x) { + int n = 0; + int bits = sizeof(x) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; + } + typename int32_t $id:(funName' "clz32") (typename int32_t x) { + int n = 0; + int bits = sizeof(x) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; + } + typename int32_t $id:(funName' "clz64") (typename int64_t x) { + int n = 0; + int bits = sizeof(x) * 8; + for (int i = 0; i < bits; i++) { + if (x < 0) break; + n++; + x <<= 1; + } + return n; + } +$esc:("#endif") + |] cFloat32Ops :: [C.Definition] cFloat64Ops :: [C.Definition] diff --git a/src/Futhark/CodeGen/ImpCode.hs b/src/Futhark/CodeGen/ImpCode.hs --- a/src/Futhark/CodeGen/ImpCode.hs +++ b/src/Futhark/CodeGen/ImpCode.hs @@ -78,7 +78,7 @@ type MemSize = Size type DimSize = Size -data Type = Scalar PrimType | Mem MemSize Space +data Type = Scalar PrimType | Mem Space data Param = MemParam VName Space | ScalarParam VName PrimType @@ -126,7 +126,7 @@ data FunctionT a = Function { functionEntry :: Bool , functionOutput :: [Param] , functionInput :: [Param] - , functionbBody :: Code a + , functionBody :: Code a , functionResult :: [ExternalValue] , functionArgs :: [ExternalValue] } diff --git a/src/Futhark/CodeGen/ImpCode/OpenCL.hs b/src/Futhark/CodeGen/ImpCode/OpenCL.hs --- a/src/Futhark/CodeGen/ImpCode/OpenCL.hs +++ b/src/Futhark/CodeGen/ImpCode/OpenCL.hs @@ -60,7 +60,6 @@ -- | Host-level OpenCL operation. data OpenCL = LaunchKernel KernelName [KernelArg] [Exp] [Exp] - | HostCode Code | GetSize VName Name | CmpSizeLe VName Name Exp | GetSizeMax VName SizeClass diff --git a/src/Futhark/CodeGen/ImpGen.hs b/src/Futhark/CodeGen/ImpGen.hs --- a/src/Futhark/CodeGen/ImpGen.hs +++ b/src/Futhark/CodeGen/ImpGen.hs @@ -82,6 +82,8 @@ , sWrite, sUpdate , sLoopNest , (<--) + + , function ) where @@ -743,7 +745,7 @@ , entryArrayElemType = bt } --- | Like 'daringFParams', but does not create new declarations. +-- | Like 'dFParams', but does not create new declarations. -- Note: a hack to be used only for functions. addFParams :: ExplicitMemorish lore => [FParam lore] -> ImpM lore op () addFParams = mapM_ addFParam @@ -1303,3 +1305,15 @@ (<--) :: VName -> Imp.Exp -> ImpM lore op () x <-- e = emit $ Imp.SetScalar x e infixl 3 <-- + +-- | Constructing a non-entry point function. +function :: [Imp.Param] -> [Imp.Param] -> ImpM lore op () -> ImpM lore op (Imp.Function op) +function outputs inputs m = do + body <- collect $ do + mapM_ addParam $ outputs ++ inputs + m + return $ Imp.Function False outputs inputs body [] [] + where addParam (Imp.MemParam name space) = + addVar name $ MemVar Nothing $ MemEntry space + addParam (Imp.ScalarParam name bt) = + addVar name $ ScalarVar Nothing $ ScalarEntry bt diff --git a/src/Futhark/CodeGen/ImpGen/Kernels.hs b/src/Futhark/CodeGen/ImpGen/Kernels.hs --- a/src/Futhark/CodeGen/ImpGen/Kernels.hs +++ b/src/Futhark/CodeGen/ImpGen/Kernels.hs @@ -121,8 +121,8 @@ sIota (patElemName pe) n' x' s' et -expCompiler (Pattern _ [pe]) (BasicOp (Replicate shape se)) = - sReplicate (patElemName pe) shape se +expCompiler (Pattern _ [pe]) (BasicOp (Replicate _ se)) = + sReplicate (patElemName pe) se -- Allocation in the "local" space is just a placeholder. expCompiler _ (Op (Alloc _ (Space "local"))) = @@ -166,9 +166,9 @@ | otherwise = sCopy bt destloc srcloc n -mapTransposeForType :: PrimType -> ImpM ExplicitMemory Imp.HostOp Name +mapTransposeForType :: PrimType -> CallKernelGen Name mapTransposeForType bt = do - -- XXX: The leading underscore is to avoid clashes with a + -- FIXME: The leading underscore is to avoid clashes with a -- programmer-defined function of the same name (this is a bad -- solution...). let fname = nameFromString $ "_" <> mapTransposeName bt diff --git a/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs b/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs --- a/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs +++ b/src/Futhark/CodeGen/ImpGen/Kernels/Base.hs @@ -51,7 +51,7 @@ import Futhark.CodeGen.ImpCode.Kernels (elements) import Futhark.CodeGen.ImpGen import Futhark.Util.IntegralExp (quotRoundingUp, quot, rem) -import Futhark.Util (chunks, maybeNth, mapAccumLM, takeLast) +import Futhark.Util (chunks, maybeNth, mapAccumLM, takeLast, dropLast) type CallKernelGen = ImpM ExplicitMemory Imp.HostOp type InKernelGen = ImpM ExplicitMemory Imp.KernelOp @@ -548,8 +548,12 @@ -- } while(assumed != old); assumed <- dPrim "assumed" t run_loop <- dPrimV "run_loop" 1 - copyDWIM old [] (Var arr) bucket + -- XXX: CUDA may generate really bad code if this is not a volatile + -- read. Unclear why. The later reads are volatile, so maybe + -- that's it. + everythingVolatile $ copyDWIM old [] (Var arr) bucket + (arr', _a_space, bucket_offset) <- fullyIndexArray arr bucket -- While-loop: Try to insert your value @@ -1110,10 +1114,10 @@ } -- | Perform a Replicate with a kernel. -sReplicate :: VName -> Shape -> SubExp - -> CallKernelGen () -sReplicate arr (Shape ds) se = do +sReplicateKernel :: VName -> SubExp -> CallKernelGen () +sReplicateKernel arr se = do t <- subExpType se + ds <- dropLast (arrayRank t) . arrayDims <$> lookupType arr dims <- mapM toExp $ ds ++ arrayDims t (constants, set_constants) <- @@ -1127,6 +1131,61 @@ set_constants sWhen (kernelThreadActive constants) $ copyDWIM arr is' se $ drop (length ds) is' + +replicateFunction :: PrimType -> CallKernelGen Imp.Function +replicateFunction bt = do + mem <- newVName "mem" + num_elems <- newVName "num_elems" + val <- newVName "val" + + let params = [Imp.MemParam mem (Space "device"), + Imp.ScalarParam num_elems int32, + Imp.ScalarParam val bt] + shape = Shape [Var num_elems] + function [] params $ do + arr <- sArray "arr" bt shape $ ArrayIn mem $ IxFun.iota $ + map (primExpFromSubExp int32) $ shapeDims shape + sReplicateKernel arr $ Var val + +replicateName :: PrimType -> String +replicateName bt = "replicate_" ++ pretty bt + +replicateForType :: PrimType -> CallKernelGen Name +replicateForType bt = do + -- FIXME: The leading underscore is to avoid clashes with a + -- programmer-defined function of the same name (this is a bad + -- solution...). + let fname = nameFromString $ "_" <> replicateName bt + + exists <- hasFunction fname + unless exists $ emitFunction fname =<< replicateFunction bt + + return fname + +replicateIsFill :: VName -> SubExp -> CallKernelGen (Maybe (CallKernelGen ())) +replicateIsFill arr v = do + ArrayEntry (MemLocation arr_mem arr_shape arr_ixfun) _ <- lookupArray arr + v_t <- subExpType v + case v_t of + Prim v_t' + | IxFun.isLinear arr_ixfun -> return $ Just $ do + fname <- replicateForType v_t' + emit $ Imp.Call [] fname + [Imp.MemArg arr_mem, + Imp.ExpArg $ unCount $ product $ map dimSizeToExp arr_shape, + Imp.ExpArg $ toExp' v_t' v] + _ -> return Nothing + +-- | Perform a Replicate with a kernel. +sReplicate :: VName -> SubExp -> CallKernelGen () +sReplicate arr se = do + -- If the replicate is of a particularly common and simple form + -- (morally a memset()/fill), then we use a common function. + is_fill <- replicateIsFill arr se + + case is_fill of + Just m -> m + Nothing -> sReplicateKernel arr se -- | Perform an Iota with a kernel. sIota :: VName -> Imp.Exp -> Imp.Exp -> Imp.Exp -> IntType diff --git a/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs b/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs --- a/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs +++ b/src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs @@ -126,9 +126,7 @@ Imp.unCount (Imp.elements num_elems `Imp.withElemType` elemType dest_t) sAlloc_ subhistos_mem subhistos_mem_size $ Space "device" - sReplicate subhistos (Shape (map snd segment_dims ++ - [Var num_subhistos, histWidth op]) <> - histShape op) ne + sReplicate subhistos ne subhistos_t <- lookupType subhistos let slice = fullSliceNum (map (toExp' int32) $ arrayDims subhistos_t) $ map (unitSlice 0 . toExp' int32 . snd) segment_dims ++ @@ -182,9 +180,7 @@ prepareIntermediateArraysGlobal :: Imp.Exp -> Imp.Exp -> [SegHistSlug] -> CallKernelGen - [(VName, - [VName], - [Imp.Exp] -> InKernelGen ())] + [[Imp.Exp] -> InKernelGen ()] prepareIntermediateArraysGlobal hist_T hist_N = fmap snd . mapAccumLM onOp Nothing where @@ -286,7 +282,7 @@ (l', do_op') <- prepareAtomicUpdateGlobal l dests slug - return (l', (hist_M, dests, do_op')) + return (l', do_op') histKernelGlobal :: [PatElem ExplicitMemory] -> Count NumGroups SubExp -> Count GroupSize SubExp @@ -308,42 +304,30 @@ prepareIntermediateArraysGlobal num_threads (toExp' int32 $ last space_sizes) slugs - elems_per_thread_64 <- dPrimVE "elems_per_thread_64" $ - total_w_64 `quotRoundingUp` - ConvOpExp (SExt Int32 Int64) num_threads - sKernelThread "seghist_global" num_groups' group_size' (segFlat space) $ \constants -> do -- Compute subhistogram index for each thread, per histogram. - subhisto_inds <- forM histograms $ \(num_histograms, _, _) -> + subhisto_inds <- forM slugs $ \slug -> dPrimVE "subhisto_ind" $ kernelGlobalThreadId constants `quot` - (kernelNumThreads constants `quotRoundingUp` Imp.var num_histograms int32) - - sFor "flat_idx" elems_per_thread_64 $ \flat_idx -> do - -- Compute the offset into the input and output. To this a - -- thread can add its local ID to figure out which element it is - -- responsible for. The calculation is done with 64-bit - -- integers to avoid overflow, but the final segment indexes are - -- 32 bit. - offset <- dPrimVE "offset" $ - (i32Toi64 (kernelGroupId constants) * - (elems_per_thread_64 * - i32Toi64 (kernelGroupSize constants))) - + (flat_idx * i32Toi64 (kernelGroupSize constants)) + (kernelNumThreads constants `quotRoundingUp` Imp.vi32 (slugNumSubhistos slug)) - j <- dPrimVE "j" $ offset + i32Toi64 (kernelLocalThreadId constants) + -- Loop over flat offsets into the input and output. The + -- calculation is done with 64-bit integers to avoid overflow, + -- but the final unflattened segment indexes are 32 bit. + let gtid = i32Toi64 $ kernelGlobalThreadId constants + kernelLoop gtid (i32Toi64 num_threads) total_w_64 $ \offset -> do -- Construct segment indices. let setIndex v e = do dPrim_ v int32 v <-- e zipWithM_ setIndex space_is $ - map (ConvOpExp (SExt Int64 Int32)) $ unflattenIndex space_sizes_64 j + map (ConvOpExp (SExt Int64 Int32)) $ unflattenIndex space_sizes_64 offset -- We execute the bucket function once and update each histogram serially. -- We apply the bucket function if j=offset+ltid is less than -- num_elements. This also involves writing to the mapout -- arrays. - let input_in_bounds = j .<. total_w_64 + let input_in_bounds = offset .<. total_w_64 sWhen input_in_bounds $ compileStms mempty (kernelBodyStms kbody) $ do let (red_res, map_res) = splitFromEnd (length map_pes) $ kernelBodyResult kbody @@ -360,7 +344,7 @@ sComment "perform atomic updates" $ forM_ (zip5 (map slugOp slugs) histograms buckets (perOp vs) subhisto_inds) $ \(HistOp dest_w _ _ _ shape lam, - (_, _, do_op), bucket, vs', subhisto_ind) -> do + do_op, bucket, vs', subhisto_ind) -> do let bucket' = toExp' int32 $ kernelResultSubExp bucket dest_w' = toExp' int32 dest_w @@ -768,11 +752,11 @@ -- well as collapsing the subhistograms produced (which are always in -- global memory, but their number may vary). compileSegHist :: Pattern ExplicitMemory - -> Count NumGroups SubExp -> Count GroupSize SubExp - -> SegSpace - -> [HistOp ExplicitMemory] - -> KernelBody ExplicitMemory - -> CallKernelGen () + -> Count NumGroups SubExp -> Count GroupSize SubExp + -> SegSpace + -> [HistOp ExplicitMemory] + -> KernelBody ExplicitMemory + -> CallKernelGen () compileSegHist (Pattern _ pes) num_groups group_size space ops kbody = do group_size' <- traverse toExp group_size diff --git a/src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs b/src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs --- a/src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs +++ b/src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs @@ -145,7 +145,7 @@ -- Be extremely careful when editing this list to ensure that -- the names match up. Also, be careful that the tags on - -- these names do not conflicts with the tags of the + -- these names do not conflict with the tags of the -- surrounding code. We accomplish the latter by using very -- low tags (normal variables start at least in the low -- hundreds). diff --git a/src/Futhark/Compiler.hs b/src/Futhark/Compiler.hs --- a/src/Futhark/Compiler.hs +++ b/src/Futhark/Compiler.hs @@ -22,6 +22,7 @@ import System.IO import qualified Data.Text.IO as T +import qualified Futhark.Analysis.Alias as Alias import Futhark.Internalise import Futhark.Pipeline import Futhark.MonadFreshNames @@ -121,9 +122,10 @@ typeCheckInternalProgram :: I.Prog I.SOACS -> FutharkM () typeCheckInternalProgram prog = - case I.checkProg prog of - Left err -> internalErrorS ("After internalisation:\n" ++ show err) (Just prog) + case I.checkProg prog' of + Left err -> internalErrorS ("After internalisation:\n" ++ show err) (Just prog') Right () -> return () + where prog' = Alias.aliasAnalysis prog -- | Read and type-check a Futhark program, including all imports. readProgram :: (MonadError CompilerError m, MonadIO m) => diff --git a/src/Futhark/Internalise.hs b/src/Futhark/Internalise.hs --- a/src/Futhark/Internalise.hs +++ b/src/Futhark/Internalise.hs @@ -86,10 +86,10 @@ body' <- localScope constscope $ do msg <- case retdecl of - Just dt -> ErrorMsg . + Just dt -> errorMsg . ("Function return value does not match shape of type ":) <$> typeExpForError rcm dt - Nothing -> return $ ErrorMsg ["Function return value does not match shape of declared return type."] + Nothing -> return $ errorMsg ["Function return value does not match shape of declared return type."] internaliseBody body >>= ensureResultExtShape asserting msg loc (map I.fromDecl rettype') @@ -433,7 +433,7 @@ let parts = ["Value of (core language) shape ("] ++ intersperse ", " (map ErrorInt32 dims) ++ [") cannot match shape of type `"] ++ dt' ++ ["`."] - ensureExtShape asserting (ErrorMsg parts) loc (I.fromDecl t') desc e' + ensureExtShape asserting (errorMsg parts) loc (I.fromDecl t') desc e' internaliseExp desc (E.Negate e _) = do e' <- internaliseExp1 "negate_arg" e @@ -658,7 +658,7 @@ internaliseExp desc (E.Assert e1 e2 (Info check) loc) = do e1' <- internaliseExp1 "assert_cond" e1 c <- assertingOne $ letExp "assert_c" $ - I.BasicOp $ I.Assert e1' (ErrorMsg [ErrorString check]) (loc, mempty) + I.BasicOp $ I.Assert e1' (errorMsg [ErrorString check]) (loc, mempty) -- Make sure there are some bindings to certify. certifying c $ mapM rebind =<< internaliseExp desc e2 where rebind v = do @@ -867,7 +867,7 @@ (idxs', oks, parts) <- unzip3 <$> zipWithM internaliseDimIndex dims idxs c <- assertingOne $ do ok <- letSubExp "index_ok" =<< foldBinOp I.LogAnd (constant True) oks - let msg = ErrorMsg $ ["Index ["] ++ intercalate [", "] parts ++ + let msg = errorMsg $ ["Index ["] ++ intercalate [", "] parts ++ ["] out of bounds for array of shape ["] ++ intersperse "][" (map ErrorInt32 $ take (length idxs) dims) ++ ["]."] letExp "index_certs" $ I.BasicOp $ I.Assert ok msg (loc, mempty) @@ -1796,3 +1796,12 @@ dimDeclForError _ (ConstDim d) = return $ ErrorString $ pretty d dimDeclForError _ AnyDim = return "" + +-- A smart constructor that compacts neighbouring literals for easier +-- reading in the IR. +errorMsg :: [ErrorMsgPart a] -> ErrorMsg a +errorMsg = ErrorMsg . compact + where compact [] = [] + compact (ErrorString x : ErrorString y : parts) = + compact (ErrorString (x++y) : parts) + compact (x:y) = x : compact y diff --git a/src/Futhark/Internalise/Monomorphise.hs b/src/Futhark/Internalise/Monomorphise.hs --- a/src/Futhark/Internalise/Monomorphise.hs +++ b/src/Futhark/Internalise/Monomorphise.hs @@ -205,8 +205,10 @@ t' <- transformType t return $ Var (QualName qs fname') (Info t') loc -transformExp (Ascript e tp t loc) = - Ascript <$> transformExp e <*> pure tp <*> pure t <*> pure loc +transformExp (Ascript e tp t loc) = do + noticeDims $ unInfo t + Ascript <$> transformExp e <*> pure tp <*> + traverse transformType t <*> pure loc transformExp (LetPat pat e1 e2 (Info t) loc) = do (pat', rr) <- transformPattern pat diff --git a/src/Futhark/Optimise/InliningDeadFun.hs b/src/Futhark/Optimise/InliningDeadFun.hs --- a/src/Futhark/Optimise/InliningDeadFun.hs +++ b/src/Futhark/Optimise/InliningDeadFun.hs @@ -15,36 +15,58 @@ import qualified Data.Set as S import Futhark.Representation.SOACS -import Futhark.Representation.SOACS.Simplify (simplifyFun) +import Futhark.Representation.SOACS.Simplify (simpleSOACS, simplifyFun) +import Futhark.Transform.CopyPropagate (copyPropagateInFun) import Futhark.Transform.Rename import Futhark.Analysis.CallGraph import Futhark.Binder import Futhark.Pass aggInlining :: MonadFreshNames m => CallGraph -> [FunDef SOACS] -> m [FunDef SOACS] -aggInlining cg = fmap (filter keep) . recurse - where noInterestingCalls :: S.Set Name -> FunDef SOACS -> Bool - noInterestingCalls interesting fundec = +aggInlining cg = fmap (filter keep) . + recurse 0 . + filter isFunInCallGraph + where isFunInCallGraph fundec = + isJust $ M.lookup (funDefName fundec) cg + + noCallsTo :: (Name -> Bool) -> FunDef SOACS -> Bool + noCallsTo interesting fundec = case M.lookup (funDefName fundec) cg of - Just calls | not $ any (`elem` interesting') calls -> True - _ -> False - where interesting' = funDefName fundec `S.insert` interesting + Just calls | not $ any interesting calls -> True + _ -> False + -- The inverse rate at which we perform full simplification + -- after inlining. For the other steps we just do copy + -- propagation. The rate here has been determined + -- heuristically and is probably not optimal for any given + -- program. + simplifyRate :: Int + simplifyRate = 4 + -- We apply simplification after every round of inlining, -- because it is more efficient to shrink the program as soon -- as possible, rather than wait until it has balooned after -- full inlining. - recurse funs = do - let interesting = S.fromList $ map funDefName funs - (to_be_inlined, to_inline_in) = - partition (noInterestingCalls interesting) funs + recurse i funs = do + let remaining = S.fromList $ map funDefName funs + (to_be_inlined, maybe_inline_in) = + partition (noCallsTo (`S.member` remaining)) funs + (not_to_inline_in, to_inline_in) = + partition (noCallsTo + (`elem` map funDefName to_be_inlined)) + maybe_inline_in inlined_but_entry_points = filter (isJust . funDefEntryPoint) to_be_inlined if null to_be_inlined then return funs - else do let onFun = simplifyFun <=< renameFun . + else do let simplify + | i `rem` simplifyRate == 0 = simplifyFun + | otherwise = copyPropagateInFun simpleSOACS + + let onFun = simplify <=< renameFun . (`doInlineInCaller` to_be_inlined) - to_inline_in' <- recurse =<< mapM onFun to_inline_in + to_inline_in' <- recurse (i+1) . (not_to_inline_in++) =<< + mapM onFun to_inline_in return $ inlined_but_entry_points ++ to_inline_in' keep fundec = isJust (funDefEntryPoint fundec) || callsRecursive fundec @@ -68,26 +90,35 @@ in FunDef entry name rtp args body' inlineInBody :: [FunDef SOACS] -> Body -> Body -inlineInBody inlcallees (Body attr stms res) = Body attr stms' res - where stms' = stmsFromList (concatMap inline $ stmsToList stms) - - inline (Let pat aux (Apply fname args _ (safety,loc,locs))) +inlineInBody inlcallees (Body attr stms res) = + Body attr (stmsFromList $ inline (stmsToList stms)) res + where inline (Let pat aux (Apply fname args _ (safety,loc,locs)) : rest) | fun:_ <- filter ((== fname) . funDefName) inlcallees = - let param_stms = zipWith reshapeIfNecessary (map paramIdent $ funDefParams fun) (map fst args) - body_stms = stmsToList $ addLocations safety - (filter notNoLoc (loc:locs)) $ bodyStms $ funDefBody fun - res_stms = map (certify $ stmAuxCerts aux) $ - zipWith reshapeIfNecessary (patternIdents pat) $ - bodyResult $ funDefBody fun - in param_stms ++ body_stms ++ res_stms - inline stm = [inlineInStm inlcallees stm] + let param_names = + map paramName $ funDefParams fun + param_stms = + zipWith (reshapeIfNecessary param_names) + (map paramIdent $ funDefParams fun) (map fst args) + body_stms = + stmsToList $ + addLocations safety (filter notNoLoc (loc:locs)) $ + bodyStms $ funDefBody fun + res_stms = + certify (stmAuxCerts aux) <$> + zipWith (reshapeIfNecessary (patternNames pat)) + (patternIdents pat) (bodyResult $ funDefBody fun) + in param_stms <> body_stms <> res_stms <> inline rest + inline (stm : rest) = + inlineInStm inlcallees stm : inline rest + inline [] = mempty - reshapeIfNecessary ident se + reshapeIfNecessary dim_names ident se | t@Array{} <- identType ident, + any (`elem` dim_names) $ subExpVars $ arrayDims t, Var v <- se = mkLet [] [ident] $ shapeCoerce (arrayDims t) v | otherwise = - mkLet [] [ident] $ BasicOp $ SubExp se + mkLet [] [ident] $ BasicOp $ SubExp se notNoLoc :: SrcLoc -> Bool notNoLoc = (/=NoLoc) . locOf diff --git a/src/Futhark/Optimise/Simplify/Rules.hs b/src/Futhark/Optimise/Simplify/Rules.hs --- a/src/Futhark/Optimise/Simplify/Rules.hs +++ b/src/Futhark/Optimise/Simplify/Rules.hs @@ -802,12 +802,12 @@ | Just (vs, vcs) <- unzip <$> mapM isArrayLit (x:xs) = Simplify $ do rt <- rowType <$> lookupType x certifying (cs <> mconcat vcs) $ - letBind_ pat $ BasicOp $ ArrayLit vs rt + letBind_ pat $ BasicOp $ ArrayLit (concat vs) rt where isArrayLit v | Just (Replicate shape se, vcs) <- ST.lookupBasicOp v vtable, - unitShape shape = Just (se, vcs) - | Just (ArrayLit [se] _, vcs) <- ST.lookupBasicOp v vtable = - Just (se, vcs) + unitShape shape = Just ([se], vcs) + | Just (ArrayLit ses _, vcs) <- ST.lookupBasicOp v vtable = + Just (ses, vcs) | otherwise = Nothing diff --git a/src/Futhark/Optimise/Sink.hs b/src/Futhark/Optimise/Sink.hs new file mode 100644 --- /dev/null +++ b/src/Futhark/Optimise/Sink.hs @@ -0,0 +1,186 @@ +{-# LANGUAGE FlexibleContexts #-} +{-# LANGUAGE TypeFamilies #-} +-- | "Sinking" is conceptually the opposite of hoisting. The idea is +-- to take code that looks like this: +-- +-- @ +-- x = xs[i] +-- y = ys[i] +-- if x != 0 then { +-- y +-- } else { +-- 0 +-- } +-- @ +-- +-- and turn it into +-- +-- @ +-- x = xs[i] +-- if x != 0 then { +-- y = ys[i] +-- y +-- } else { +-- 0 +-- } +-- @ +-- +-- The idea is to delay loads from memory until (if) they are actually +-- needed. Code patterns like the above is particularly common in +-- code that makes use of pattern matching on sum types. +-- +-- We are currently quite conservative about when we do this. In +-- particular, if any consumption is going on in a body, we don't do +-- anything. This is far too conservative. Also, we are careful +-- never to duplicate work. +-- +-- This pass redundantly computes free-variable information a lot. If +-- you ever see this pass as being a compilation speed bottleneck, +-- start by caching that a bit. +-- +-- This pass is defined on the Kernels representation. This is not +-- because we do anything kernel-specific here, but simply because +-- more explicit indexing is going on after SOACs are gone. + +module Futhark.Optimise.Sink (sink) where + +import Control.Monad.State +import Data.List (foldl') +import qualified Data.Map as M +import qualified Data.Set as S + +import qualified Futhark.Analysis.Alias as Alias +import qualified Futhark.Analysis.Range as Range +import qualified Futhark.Analysis.SymbolTable as ST +import Futhark.MonadFreshNames +import Futhark.Representation.Aliases +import Futhark.Representation.Ranges +import Futhark.Representation.Kernels +import Futhark.Pass + +-- We do not care about ranges, but in order to use ST.SymbolTable +-- (which is a convenient way to handle aliases), we need range information. +type SinkLore = Ranges (Aliases Kernels) +type SymbolTable = ST.SymbolTable SinkLore +type Sinking = M.Map VName (Stm SinkLore) +type Sunk = S.Set VName + +-- | Given a statement, compute how often each of its free variables +-- are used. Not accurate: what we care about are only 1, and greater +-- than 1. +multiplicity :: Stm SinkLore -> M.Map VName Int +multiplicity stm = + case stmExp stm of + If cond tbranch fbranch _ -> + free cond 1 <> M.unionWith (+) (free tbranch 1) (free fbranch 1) + Op{} -> free stm 2 + DoLoop{} -> free stm 2 + _ -> free stm 1 + where free x k = M.fromList $ zip (namesToList $ freeIn x) $ repeat k + +optimiseBranch :: SymbolTable -> Sinking -> Body SinkLore + -> (Body SinkLore, Sunk) +optimiseBranch vtable sinking (Body attr stms res) = + let (stms', stms_sunk) = optimiseStms vtable sinking' stms + in (Body attr (sunk_stms <> stms') res, + sunk <> stms_sunk) + where free_in_stms = freeIn stms + (sinking_here, sinking') = M.partitionWithKey sunkHere sinking + sunk_stms = stmsFromList $ M.elems sinking_here + sunkHere v stm = + v `nameIn` free_in_stms && + all (`ST.available` vtable) (namesToList (freeIn stm)) + sunk = S.fromList $ concatMap (patternNames . stmPattern) sunk_stms + +optimiseStms :: SymbolTable -> Sinking -> Stms SinkLore + -> (Stms SinkLore, Sunk) +optimiseStms init_vtable init_sinking all_stms = + let (all_stms', sunk) = + optimiseStms' init_vtable init_sinking $ stmsToList all_stms + in (stmsFromList all_stms', sunk) + where + multiplicities = foldl' (M.unionWith (+)) mempty $ + map multiplicity $ stmsToList all_stms + + optimiseStms' _ _ [] = ([], mempty) + + optimiseStms' vtable sinking (stm : stms) + | BasicOp Index{} <- stmExp stm, + [pe] <- patternElements (stmPattern stm), + primType $ patElemType pe, + maybe True (==1) $ M.lookup (patElemName pe) multiplicities = + let (stms', sunk) = + optimiseStms' vtable' (M.insert (patElemName pe) stm sinking) stms + in if patElemName pe `S.member` sunk + then (stms', sunk) + else (stm : stms', sunk) + + | If cond tbranch fbranch ret <- stmExp stm = + let (tbranch', tsunk) = optimiseBranch vtable sinking tbranch + (fbranch', fsunk) = optimiseBranch vtable sinking fbranch + (stms', sunk) = optimiseStms' vtable' sinking stms + in (stm { stmExp = If cond tbranch' fbranch' ret } : stms', + tsunk <> fsunk <> sunk) + + | Op (SegOp op) <- stmExp stm = + let scope = scopeOfSegSpace $ segSpace op + (stms', stms_sunk) = optimiseStms' vtable' sinking stms + (op', op_sunk) = runState (mapSegOpM (opMapper scope) op) mempty + in (stm { stmExp = Op (SegOp op') } : stms', + stms_sunk <> op_sunk) + + | otherwise = + let (stms', stms_sunk) = optimiseStms' vtable' sinking stms + (e', stm_sunk) = runState (mapExpM mapper (stmExp stm)) mempty + in (stm { stmExp = e' } : stms', + stm_sunk <> stms_sunk) + + where vtable' = ST.insertStm stm vtable + mapper = + identityMapper + { mapOnBody = \scope body -> do + let (body', sunk) = + optimiseBody (ST.fromScope scope <> vtable) sinking body + modify (<>sunk) + return body' + } + + opMapper scope = + identitySegOpMapper + { mapOnSegOpLambda = \lam -> do + let (body, sunk) = + optimiseBody op_vtable sinking $ + lambdaBody lam + modify (<>sunk) + return lam { lambdaBody = body } + + , mapOnSegOpBody = \body -> do + let (body', sunk) = + optimiseKernelBody op_vtable sinking body + modify (<>sunk) + return body' + } + where op_vtable = ST.fromScope scope <> vtable + +optimiseBody :: SymbolTable -> Sinking -> Body SinkLore + -> (Body SinkLore, Sunk) +optimiseBody vtable sinking (Body attr stms res) = + let (stms', sunk) = optimiseStms vtable sinking stms + in (Body attr stms' res, sunk) + +optimiseKernelBody :: SymbolTable -> Sinking -> KernelBody SinkLore + -> (KernelBody SinkLore, Sunk) +optimiseKernelBody vtable sinking (KernelBody attr stms res) = + let (stms', sunk) = optimiseStms vtable sinking stms + in (KernelBody attr stms' res, sunk) + +optimiseFunDef :: MonadFreshNames m => FunDef Kernels -> m (FunDef Kernels) +optimiseFunDef fundef = do + let fundef' = Range.analyseFun $ Alias.analyseFun fundef + vtable = ST.insertFParams (funDefParams fundef') mempty + (body, _) = optimiseBody vtable mempty $ funDefBody fundef' + return fundef { funDefBody = removeBodyAliases $ removeBodyRanges body } + +sink :: Pass Kernels Kernels +sink = Pass "sink" "move memory loads closer to their uses" $ + intraproceduralTransformation optimiseFunDef diff --git a/src/Futhark/Passes.hs b/src/Futhark/Passes.hs --- a/src/Futhark/Passes.hs +++ b/src/Futhark/Passes.hs @@ -15,6 +15,7 @@ import Futhark.Optimise.Fusion import Futhark.Optimise.InPlaceLowering import Futhark.Optimise.InliningDeadFun +import Futhark.Optimise.Sink import Futhark.Optimise.TileLoops import Futhark.Optimise.DoubleBuffer import Futhark.Optimise.Unstream @@ -57,6 +58,7 @@ , unstream , performCSE True , simplifyKernels + , sink , inPlaceLowering ] @@ -65,6 +67,7 @@ standardPipeline >>> onePass firstOrderTransform >>> passes [ simplifyKernels + , sink , inPlaceLowering ] diff --git a/src/Futhark/Pipeline.hs b/src/Futhark/Pipeline.hs --- a/src/Futhark/Pipeline.hs +++ b/src/Futhark/Pipeline.hs @@ -33,6 +33,7 @@ import Prelude hiding (id, (.)) +import qualified Futhark.Analysis.Alias as Alias import Futhark.Error import Futhark.Representation.AST (Prog, PrettyLore) import Futhark.TypeCheck @@ -126,8 +127,9 @@ when (pipelineVerbose cfg) $ logMsg $ "Running pass " <> T.pack (passName pass) prog' <- runPass pass prog + let prog'' = Alias.aliasAnalysis prog' when (pipelineValidate cfg) $ - case checkProg prog' of + case checkProg prog'' of Left err -> validationError pass prog' $ show err Right () -> return () return prog' diff --git a/src/Futhark/Representation/Primitive.hs b/src/Futhark/Representation/Primitive.hs --- a/src/Futhark/Representation/Primitive.hs +++ b/src/Futhark/Representation/Primitive.hs @@ -845,6 +845,16 @@ , f32 "gamma32" tgammaf, f64 "gamma64" tgamma , f32 "lgamma32" lgammaf, f64 "lgamma64" lgamma + , i8 "clz8" $ IntValue . Int32Value . fromIntegral . countLeadingZeros + , i16 "clz16" $ IntValue . Int32Value . fromIntegral . countLeadingZeros + , i32 "clz32" $ IntValue . Int32Value . fromIntegral . countLeadingZeros + , i64 "clz64" $ IntValue . Int32Value . fromIntegral . countLeadingZeros + + , i8 "popc8" $ IntValue . Int32Value . fromIntegral . popCount + , i16 "popc16" $ IntValue . Int32Value . fromIntegral . popCount + , i32 "popc32" $ IntValue . Int32Value . fromIntegral . popCount + , i64 "popc64" $ IntValue . Int32Value . fromIntegral . popCount + , ("atan2_32", ([FloatType Float32, FloatType Float32], FloatType Float32, \case @@ -925,8 +935,28 @@ v0 + (v1-v0)*max 0 (min 1 t) _ -> Nothing)) ] - where f32 s f = (s, ([FloatType Float32], FloatType Float32, f32PrimFun f)) + where i8 s f = (s, ([IntType Int8], IntType Int32, i8PrimFun f)) + i16 s f = (s, ([IntType Int16], IntType Int32, i16PrimFun f)) + i32 s f = (s, ([IntType Int32], IntType Int32, i32PrimFun f)) + i64 s f = (s, ([IntType Int64], IntType Int32, i64PrimFun f)) + f32 s f = (s, ([FloatType Float32], FloatType Float32, f32PrimFun f)) f64 s f = (s, ([FloatType Float64], FloatType Float64, f64PrimFun f)) + + i8PrimFun f [IntValue (Int8Value x)] = + Just $ f x + i8PrimFun _ _ = Nothing + + i16PrimFun f [IntValue (Int16Value x)] = + Just $ f x + i16PrimFun _ _ = Nothing + + i32PrimFun f [IntValue (Int32Value x)] = + Just $ f x + i32PrimFun _ _ = Nothing + + i64PrimFun f [IntValue (Int64Value x)] = + Just $ f x + i64PrimFun _ _ = Nothing f32PrimFun f [FloatValue (Float32Value x)] = Just $ FloatValue $ Float32Value $ f x diff --git a/src/Futhark/Representation/Ranges.hs b/src/Futhark/Representation/Ranges.hs --- a/src/Futhark/Representation/Ranges.hs +++ b/src/Futhark/Representation/Ranges.hs @@ -38,6 +38,7 @@ import Futhark.Representation.AST.Syntax import Futhark.Representation.AST.Attributes +import Futhark.Representation.AST.Attributes.Aliases import Futhark.Representation.AST.Attributes.Ranges import Futhark.Representation.AST.Traversals import Futhark.Representation.AST.Pretty @@ -182,3 +183,16 @@ -> Stm (Ranges lore) mkRangedLetStm pat explore e = Let (addRangesToPattern pat e) (StmAux mempty explore) e + +-- It is convenient for a wrapped aliased lore to also be aliased. + +instance AliasesOf attr => AliasesOf ([Range], attr) where + aliasesOf = aliasesOf . snd + +instance AliasesOf attr => AliasesOf (Range, attr) where + aliasesOf = aliasesOf . snd + +instance (Aliased lore, CanBeRanged (Op lore), + AliasedOp (OpWithRanges (Op lore))) => Aliased (Ranges lore) where + bodyAliases = bodyAliases . removeBodyRanges + consumedInBody = consumedInBody . removeBodyRanges diff --git a/src/Futhark/Representation/SOACS/Simplify.hs b/src/Futhark/Representation/SOACS/Simplify.hs --- a/src/Futhark/Representation/SOACS/Simplify.hs +++ b/src/Futhark/Representation/SOACS/Simplify.hs @@ -8,6 +8,7 @@ ( simplifySOACS , simplifyLambda , simplifyFun + , simplifyFun' , simplifyStms , simpleSOACS @@ -66,6 +67,10 @@ simplifyFun = Simplify.simplifyFun simpleSOACS soacRules Engine.noExtraHoistBlockers +simplifyFun' :: MonadFreshNames m => FunDef SOACS -> m (FunDef SOACS) +simplifyFun' = + Simplify.simplifyFun simpleSOACS mempty Engine.noExtraHoistBlockers + simplifyLambda :: (HasScope SOACS m, MonadFreshNames m) => Lambda -> [Maybe VName] -> m Lambda simplifyLambda = @@ -600,11 +605,11 @@ fromArrayOp (ArrayRotate cs arr rots) = (cs, BasicOp $ Rotate rots arr) fromArrayOp (ArrayVar cs arr) = (cs, BasicOp $ SubExp $ Var arr) -arrayOps :: AST.Body (Wise SOACS) -> S.Set ArrayOp +arrayOps :: AST.Body (Wise SOACS) -> S.Set (AST.Pattern (Wise SOACS), ArrayOp) arrayOps = mconcat . map onStm . stmsToList . bodyStms - where onStm (Let _ aux e) = + where onStm (Let pat aux e) = case isArrayOp (stmAuxCerts aux) e of - Just op -> S.singleton op + Just op -> S.singleton (pat, op) Nothing -> execState (walkExpM walker e) mempty onOp = execWriter . mapSOACM identitySOACMapper { mapOnSOACLambda = onLambda } onLambda lam = do tell $ arrayOps $ lambdaBody lam @@ -647,7 +652,7 @@ simplifyMapIota :: TopDownRuleOp (Wise SOACS) simplifyMapIota vtable pat _ (Screma w (ScremaForm scan reduce map_lam) arrs) | Just (p, _) <- find isIota (zip (lambdaParams map_lam) arrs), - indexings <- filter (indexesWith (paramName p)) $ S.toList $ + indexings <- filter (indexesWith (paramName p)) $ map snd $ S.toList $ arrayOps $ lambdaBody map_lam, not $ null indexings = Simplify $ do -- For each indexing with iota, add the corresponding array to @@ -693,7 +698,7 @@ -- full array. moveTransformToInput :: TopDownRuleOp (Wise SOACS) moveTransformToInput vtable pat _ (Screma w (ScremaForm scan reduce map_lam) arrs) - | ops <- filter arrayIsMapParam $ S.toList $ arrayOps $ lambdaBody map_lam, + | ops <- map snd $ filter arrayIsMapParam $ S.toList $ arrayOps $ lambdaBody map_lam, not $ null ops = Simplify $ do (more_arrs, more_params, replacements) <- unzip3 . catMaybes <$> mapM mapOverArr ops @@ -709,21 +714,27 @@ letBind_ pat $ Op $ Screma w (ScremaForm scan reduce map_lam') (arrs <> more_arrs) where map_param_names = map paramName (lambdaParams map_lam) + topLevelPattern = (`elem` fmap stmPattern (bodyStms (lambdaBody map_lam))) + onlyUsedOnce arr = + case filter ((arr `nameIn`) . freeIn) $ stmsToList $ bodyStms $ lambdaBody map_lam of + _ : _ : _ -> False + _ -> True -- It's not just about whether the array is a parameter; -- everything else must be map-invariant. - arrayIsMapParam (ArrayIndexing cs arr slice) = + arrayIsMapParam (pat', ArrayIndexing cs arr slice) = arr `elem` map_param_names && all (`ST.elem` vtable) (namesToList $ freeIn cs <> freeIn slice) && - not (null slice) && not (null $ sliceDims slice) - arrayIsMapParam (ArrayRearrange cs arr perm) = + not (null slice) && + (not (null $ sliceDims slice) || (topLevelPattern pat' && onlyUsedOnce arr)) + arrayIsMapParam (_, ArrayRearrange cs arr perm) = arr `elem` map_param_names && all (`ST.elem` vtable) (namesToList $ freeIn cs) && not (null perm) - arrayIsMapParam (ArrayRotate cs arr rots) = + arrayIsMapParam (_, ArrayRotate cs arr rots) = arr `elem` map_param_names && all (`ST.elem` vtable) (namesToList $ freeIn cs <> freeIn rots) - arrayIsMapParam ArrayVar{} = + arrayIsMapParam (_, ArrayVar{}) = False mapOverArr op diff --git a/src/Futhark/Transform/CopyPropagate.hs b/src/Futhark/Transform/CopyPropagate.hs --- a/src/Futhark/Transform/CopyPropagate.hs +++ b/src/Futhark/Transform/CopyPropagate.hs @@ -4,16 +4,24 @@ -- simplifier with no rules, so hoisting and dead-code elimination may -- also take place. module Futhark.Transform.CopyPropagate - (copyPropagateInStms) + ( copyPropagateInStms + , copyPropagateInFun) where import Futhark.MonadFreshNames import Futhark.Representation.AST import Futhark.Optimise.Simplify --- | Run copy propagation. +-- | Run copy propagation on some statements. copyPropagateInStms :: (MonadFreshNames m, SimplifiableLore lore, HasScope lore m) => SimpleOps lore -> Stms lore -> m (Stms lore) copyPropagateInStms simpl = simplifyStms simpl mempty noExtraHoistBlockers + +-- | Run copy propagation on a function. +copyPropagateInFun :: (MonadFreshNames m, SimplifiableLore lore) => + SimpleOps lore + -> FunDef lore + -> m (FunDef lore) +copyPropagateInFun simpl = simplifyFun simpl mempty noExtraHoistBlockers diff --git a/src/Futhark/TypeCheck.hs b/src/Futhark/TypeCheck.hs --- a/src/Futhark/TypeCheck.hs +++ b/src/Futhark/TypeCheck.hs @@ -63,7 +63,6 @@ import Futhark.Analysis.PrimExp import Futhark.Construct (instantiateShapes) import Futhark.Representation.Aliases -import Futhark.Analysis.Alias import Futhark.Util import Futhark.Util.Pretty (Pretty, prettyDoc, indent, ppr, text, (<+>), align) @@ -456,7 +455,7 @@ -- yielding either a type error or a program with complete type -- information. checkProg :: Checkable lore => - Prog lore -> Either (TypeError lore) () + Prog (Aliases lore) -> Either (TypeError lore) () checkProg prog = do let typeenv = Env { envVtable = M.empty , envFtable = mempty @@ -468,11 +467,10 @@ local (\env -> env { envFtable = ftable }) $ checkFun fun (ftable, _) <- runTypeM typeenv buildFtable - sequence_ $ parMap rpar (onFunction ftable) $ progFunctions prog' + sequence_ $ parMap rpar (onFunction ftable) $ progFunctions prog where - prog' = aliasAnalysis prog - buildFtable = do table <- initialFtable prog' - foldM expand table $ progFunctions prog' + buildFtable = do table <- initialFtable prog + foldM expand table $ progFunctions prog expand ftable (FunDef _ name ret params _) | M.member name ftable = bad $ DupDefinitionError name @@ -915,9 +913,23 @@ context "When checking expression annotation" $ checkExpLore attr context ("When matching\n" ++ message " " pat ++ "\nwith\n" ++ message " " e) $ matchPattern pat e - binding (scopeOf stm) $ do + binding (maybeWithoutAliases $ scopeOf stm) $ do mapM_ checkPatElem (patternElements $ removePatternAliases pat) m + where + -- FIXME: this is wrong. However, the core language type system + -- is not strong enough to fully capture the aliases we want (see + -- issue #803). Since we eventually inline everything anyway, and + -- our intra-procedural alias analysis is much simpler and + -- correct, I could not justify spending time on improving the + -- inter-procedural alias analysis. If we ever stop inlining + -- everything, probably we need to go back and refine this. + maybeWithoutAliases = + case stmExp stm of + Apply{} -> M.map withoutAliases + _ -> id + withoutAliases (LetInfo (_, lattr)) = LetInfo (mempty, lattr) + withoutAliases info = info matchExtPattern :: Checkable lore => Pattern (Aliases lore) -> [ExtType] -> TypeM lore () diff --git a/src/Language/Futhark/Interpreter.hs b/src/Language/Futhark/Interpreter.hs --- a/src/Language/Futhark/Interpreter.hs +++ b/src/Language/Futhark/Interpreter.hs @@ -115,7 +115,7 @@ ppr (ValueRecord m) = prettyRecord m ppr ValueFun{} = text "#" - ppr (ValueSum n vs) = text "#" <> ppr n <+> sep (map ppr vs) + ppr (ValueSum n vs) = text "#" <> sep (ppr n : map ppr vs) -- | Create an array value; failing if that would result in an -- irregular array. diff --git a/src/Language/Futhark/Parser/Parser.y b/src/Language/Futhark/Parser/Parser.y --- a/src/Language/Futhark/Parser/Parser.y +++ b/src/Language/Futhark/Parser/Parser.y @@ -611,6 +611,8 @@ Atom :: { UncheckedExp } Atom : PrimLit { Literal (fst $1) (snd $1) } | Constr { Constr (fst $1) [] NoInfo (snd $1) } + | charlit { let L loc (CHARLIT x) = $1 + in IntLit (toInteger (ord x)) NoInfo loc } | intlit { let L loc (INTLIT x) = $1 in IntLit x NoInfo loc } | floatlit { let L loc (FLOATLIT x) = $1 in FloatLit x NoInfo loc } | stringlit { let L loc (STRINGLIT s) = $1 in @@ -675,9 +677,6 @@ | f32lit { let L loc (F32LIT num) = $1 in (FloatValue $ Float32Value num, loc) } | f64lit { let L loc (F64LIT num) = $1 in (FloatValue $ Float64Value num, loc) } - | charlit { let L loc (CHARLIT char) = $1 - in (SignedValue $ Int32Value $ fromIntegral $ ord char, loc) } - Exps1 :: { (UncheckedExp, [UncheckedExp]) } : Exps1_ { case reverse (snd $1 : fst $1) of [] -> (snd $1, []) @@ -783,6 +782,8 @@ CaseLiteral :: { (UncheckedExp, SrcLoc) } : PrimLit { (Literal (fst $1) (snd $1), snd $1) } + | charlit { let L loc (CHARLIT x) = $1 + in (IntLit (toInteger (ord x)) NoInfo loc, loc) } | intlit { let L loc (INTLIT x) = $1 in (IntLit x NoInfo loc, loc) } | floatlit { let L loc (FLOATLIT x) = $1 in (FloatLit x NoInfo loc, loc) } | stringlit { let L loc (STRINGLIT s) = $1 in diff --git a/src/Language/Futhark/Pretty.hs b/src/Language/Futhark/Pretty.hs --- a/src/Language/Futhark/Pretty.hs +++ b/src/Language/Futhark/Pretty.hs @@ -113,7 +113,8 @@ instance Pretty (ShapeDecl dim) => Pretty (ScalarTypeBase dim as) where ppr = pprPrec 0 pprPrec _ (Prim et) = ppr et - pprPrec _ (TypeVar _ u et targs) = + pprPrec p (TypeVar _ u et targs) = + parensIf (not (null targs) && p > 0) $ ppr u <> ppr (qualNameFromTypeName et) <+> spread (map (pprPrec 1) targs) pprPrec _ (Record fs) | Just ts <- areTupleFields fs = @@ -124,10 +125,10 @@ where ppField (name, t) = text (nameToString name) <> colon <+> ppr t fs' = map ppField $ M.toList fs pprPrec p (Arrow _ (Named v) t1 t2) = - parensIf (p > 0) $ - parens (pprName v <> colon <+> ppr t1) <+> text "->" <+> ppr t2 + parensIf (p > 1) $ + parens (pprName v <> colon <+> ppr t1) <+> text "->" <+> pprPrec 1 t2 pprPrec p (Arrow _ Unnamed t1 t2) = - parensIf (p > 0) $ pprPrec 1 t1 <+> text "->" <+> ppr t2 + parensIf (p > 1) $ pprPrec 2 t1 <+> text "->" <+> pprPrec 1 t2 pprPrec p (Sum cs) = parensIf (p > 0) $ oneLine (mconcat $ punctuate (text " | ") cs') @@ -137,7 +138,7 @@ instance Pretty (ShapeDecl dim) => Pretty (TypeBase dim as) where ppr = pprPrec 0 - pprPrec _ (Array _ u at shape) = ppr u <> ppr shape <> ppr at + pprPrec _ (Array _ u at shape) = ppr u <> ppr shape <> pprPrec 1 at pprPrec p (Scalar t) = pprPrec p t instance Pretty (ShapeDecl dim) => Pretty (TypeArg dim) where @@ -201,7 +202,8 @@ pprPrec _ (Var name _ _) = ppr name pprPrec _ (Parens e _) = align $ parens $ ppr e pprPrec _ (QualParens v e _) = ppr v <> text "." <> align (parens $ ppr e) - pprPrec _ (Ascript e t _ _) = pprPrec 0 e <> colon <+> pprPrec 0 t + pprPrec p (Ascript e t _ _) = + parensIf (p /= -1) $ pprPrec 0 e <+> colon <+> pprPrec 0 t pprPrec _ (Literal v _) = ppr v pprPrec _ (IntLit v _ _) = ppr v pprPrec _ (FloatLit v _ _) = ppr v diff --git a/src/Language/Futhark/TypeChecker/Monad.hs b/src/Language/Futhark/TypeChecker/Monad.hs --- a/src/Language/Futhark/TypeChecker/Monad.hs +++ b/src/Language/Futhark/TypeChecker/Monad.hs @@ -390,9 +390,8 @@ intrinsicsNameMap :: NameMap intrinsicsNameMap = M.fromList $ map mapping $ M.toList intrinsics - where mapping (v, IntrinsicType{}) = ((Type, baseName v), QualName [mod] v) - mapping (v, _) = ((Term, baseName v), QualName [mod] v) - mod = VName (nameFromString "intrinsics") 0 + where mapping (v, IntrinsicType{}) = ((Type, baseName v), QualName [] v) + mapping (v, _) = ((Term, baseName v), QualName [] v) topLevelNameMap :: NameMap topLevelNameMap = M.filterWithKey (\k _ -> atTopLevel k) intrinsicsNameMap diff --git a/src/Language/Futhark/TypeChecker/Terms.hs b/src/Language/Futhark/TypeChecker/Terms.hs --- a/src/Language/Futhark/TypeChecker/Terms.hs +++ b/src/Language/Futhark/TypeChecker/Terms.hs @@ -894,7 +894,8 @@ typeError loc $ "Type " ++ quote (pretty t') ++ " is not a subtype of " ++ quote (pretty decl_t') ++ "." - return $ Ascript e' decl' (Info (combineTypeShapes t $ fromStruct decl_t)) loc + let t'' = flip unifyTypeAliases t' $ combineTypeShapes t' $ fromStruct decl_t' + return $ Ascript e' decl' (Info t'') loc checkExp (BinOp (op, oploc) NoInfo (e1,_) (e2,_) NoInfo loc) = do (op', ftype) <- lookupVar oploc op diff --git a/src/Language/Futhark/TypeChecker/Types.hs b/src/Language/Futhark/TypeChecker/Types.hs --- a/src/Language/Futhark/TypeChecker/Types.hs +++ b/src/Language/Futhark/TypeChecker/Types.hs @@ -181,8 +181,8 @@ (tname', ps, t, l) <- lookupType tloc tname if length ps /= length targs then throwError $ TypeError tloc $ - "Type constructor " ++ pretty tname ++ " requires " ++ show (length ps) ++ - " arguments, but application at " ++ locStr tloc ++ " provides " ++ show (length targs) + "Type constructor " ++ quote (pretty tname) ++ " requires " ++ show (length ps) ++ + " arguments, but provided " ++ show (length targs) ++ "." else do (targs', substs) <- unzip <$> zipWithM checkArgApply ps targs return (foldl (\x y -> TEApply x y tloc) (TEVar tname' tname_loc) targs', diff --git a/src/Language/Futhark/TypeChecker/Unify.hs b/src/Language/Futhark/TypeChecker/Unify.hs --- a/src/Language/Futhark/TypeChecker/Unify.hs +++ b/src/Language/Futhark/TypeChecker/Unify.hs @@ -127,7 +127,7 @@ unify usage orig_t1 orig_t2 = do orig_t1' <- normaliseType orig_t1 orig_t2' <- normaliseType orig_t2 - breadCrumb (MatchingTypes orig_t1' orig_t2') $ subunify orig_t1 orig_t2 + breadCrumb (MatchingTypes orig_t1' orig_t2') $ subunify orig_t1' orig_t2' where subunify t1 t2 = do constraints <- getConstraints