futhark 0.12.2 → 0.12.3
raw patch · 51 files changed
+1057/−398 lines, 51 files
Files
- docs/index.rst +1/−0
- futhark.cabal +4/−3
- futlib/math.fut +23/−0
- package.yaml +1/−1
- rts/c/cuda.h +26/−43
- rts/c/free_list.h +7/−7
- rts/c/opencl.h +88/−13
- rts/c/values.h +34/−126
- rts/csharp/scalar.cs +76/−0
- rts/python/scalar.py +20/−0
- src/Futhark/Analysis/CallGraph.hs +7/−8
- src/Futhark/Analysis/Range.hs +1/−0
- src/Futhark/CLI/Dev.hs +2/−0
- src/Futhark/CodeGen/Backends/CCUDA.hs +0/−1
- src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs +16/−1
- src/Futhark/CodeGen/Backends/COpenCL.hs +32/−13
- src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs +47/−17
- src/Futhark/CodeGen/Backends/CSOpenCL.hs +0/−2
- src/Futhark/CodeGen/Backends/GenericC.hs +29/−16
- src/Futhark/CodeGen/Backends/GenericCSharp.hs +2/−2
- src/Futhark/CodeGen/Backends/GenericPython.hs +1/−1
- src/Futhark/CodeGen/Backends/PyOpenCL.hs +0/−2
- src/Futhark/CodeGen/Backends/SequentialC.hs +13/−0
- src/Futhark/CodeGen/Backends/SimpleRepresentation.hs +132/−16
- src/Futhark/CodeGen/ImpCode.hs +2/−2
- src/Futhark/CodeGen/ImpCode/OpenCL.hs +0/−1
- src/Futhark/CodeGen/ImpGen.hs +15/−1
- src/Futhark/CodeGen/ImpGen/Kernels.hs +4/−4
- src/Futhark/CodeGen/ImpGen/Kernels/Base.hs +64/−5
- src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs +18/−34
- src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs +1/−1
- src/Futhark/Compiler.hs +4/−2
- src/Futhark/Internalise.hs +14/−5
- src/Futhark/Internalise/Monomorphise.hs +4/−2
- src/Futhark/Optimise/InliningDeadFun.hs +58/−27
- src/Futhark/Optimise/Simplify/Rules.hs +4/−4
- src/Futhark/Optimise/Sink.hs +186/−0
- src/Futhark/Passes.hs +3/−0
- src/Futhark/Pipeline.hs +3/−1
- src/Futhark/Representation/Primitive.hs +31/−1
- src/Futhark/Representation/Ranges.hs +14/−0
- src/Futhark/Representation/SOACS/Simplify.hs +21/−10
- src/Futhark/Transform/CopyPropagate.hs +10/−2
- src/Futhark/TypeCheck.hs +19/−7
- src/Language/Futhark/Interpreter.hs +1/−1
- src/Language/Futhark/Parser/Parser.y +4/−3
- src/Language/Futhark/Pretty.hs +8/−6
- src/Language/Futhark/TypeChecker/Monad.hs +2/−3
- src/Language/Futhark/TypeChecker/Terms.hs +2/−1
- src/Language/Futhark/TypeChecker/Types.hs +2/−2
- src/Language/Futhark/TypeChecker/Unify.hs +1/−1
docs/index.rst view
@@ -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
futhark.cabal view
@@ -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
futlib/math.fut view
@@ -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..<n let replicate 'v (n: i8) (x: v) = map (const x) (iota n)@@ -323,6 +332,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.popc16+ let clz = intrinsics.clz16 let iota (n: i16) = 0i16..1i16..<n let replicate 'v (n: i16) (x: v) = map (const x) (iota n)@@ -396,6 +407,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.popc32+ let clz = intrinsics.clz32 let iota (n: i32) = 0..1..<n let replicate 'v (n: i32) (x: v) = map (const x) (iota n)@@ -469,6 +482,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))) | intrinsics.zext_i32_i64 (b intrinsics.<< bit))+ let popc = intrinsics.popc64+ let clz = intrinsics.clz64 let iota (n: i64) = 0i64..1i64..<n let replicate 'v (n: i64) (x: v) = map (const x) (iota n)@@ -542,6 +557,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 x = intrinsics.popc8 (sign x)+ let clz x = intrinsics.clz8 (sign x) let iota (n: u8) = 0u8..1u8..<n let replicate 'v (n: u8) (x: v) = map (const x) (iota n)@@ -615,6 +632,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 x = intrinsics.popc16 (sign x)+ let clz x = intrinsics.clz16 (sign x) let iota (n: u16) = 0u16..1u16..<n let replicate 'v (n: u16) (x: v) = map (const x) (iota n)@@ -688,6 +707,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 x = intrinsics.popc32 (sign x)+ let clz x = intrinsics.clz32 (sign x) let iota (n: u32) = 0u32..1u32..<n let replicate 'v (n: u32) (x: v) = map (const x) (iota n)@@ -761,6 +782,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 x = intrinsics.popc64 (sign x)+ let clz x = intrinsics.clz64 (sign x) let iota (n: u64) = 0u64..1u64..<n let replicate 'v (n: u64) (x: v) = map (const x) (iota n)
package.yaml view
@@ -1,5 +1,5 @@ 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
rts/c/cuda.h view
@@ -4,8 +4,7 @@ #define NVRTC_SUCCEED(x) nvrtc_api_succeed(x, #x, __FILE__, __LINE__) static inline void cuda_api_succeed(CUresult res, const char *call,- const char *file, int line)-{+ const char *file, int line) { if (res != CUDA_SUCCESS) { const char *err_str; cuGetErrorString(res, &err_str);@@ -16,8 +15,7 @@ } static inline void nvrtc_api_succeed(nvrtcResult res, const char *call,- const char *file, int line)-{+ const char *file, int line) { if (res != NVRTC_SUCCESS) { const char *err_str = nvrtcGetErrorString(res); panic(-1, "%s:%d: NVRTC call\n %s\nfailed with error code %d (%s)\n",@@ -52,13 +50,12 @@ const char **size_classes; }; -void cuda_config_init(struct cuda_config *cfg,- int num_sizes,- const char *size_names[],- const char *size_vars[],- size_t *size_values,- const char *size_classes[])-{+static void cuda_config_init(struct cuda_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->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) {
rts/c/free_list.h view
@@ -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;
rts/c/opencl.h view
@@ -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) {
rts/c/values.h view
@@ -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<elem_size/2; j++) {+ unsigned char head = elem[j];+ int tail_index = elem_size-1-j;+ elem[j] = elem[tail_index];+ elem[tail_index] = head;+ }+}+ // On Windows we need to explicitly set the file mode to not mangle // newline characters. On *nix there is no difference. #ifdef _WIN32@@ -450,100 +459,11 @@ } #endif -// Reading little-endian byte sequences. On big-endian hosts, we flip-// the resulting bytes.- static int read_byte(void* dest) { int num_elems_read = fread(dest, 1, 1, stdin); return num_elems_read == 1 ? 0 : 1; } -static int read_le_2byte(void* dest) {- uint16_t x;- int num_elems_read = fread(&x, 2, 1, stdin);- if (IS_BIG_ENDIAN) {- x = (x>>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; i<dims; i++) { uint64_t bin_shape;- ret = read_le_8byte(&bin_shape);- if (ret != 0) { panic(1, "binary-input: Couldn't read size for dimension %i of array.\n", i); }+ ret = fread(&bin_shape, sizeof(bin_shape), 1, stdin);+ if (ret != 1) {+ panic(1, "binary-input: Couldn't read size for dimension %i of array.\n", i);+ }+ if (IS_BIG_ENDIAN) {+ flip_bytes(sizeof(bin_shape), (unsigned char*) &bin_shape);+ } elem_count *= bin_shape; shape[i] = (int64_t) bin_shape; }@@ -715,16 +627,7 @@ // If we're on big endian platform we must change all multibyte elements // from using little endian to big endian if (IS_BIG_ENDIAN && elem_size != 1) {- char* elems = (char*) *data;- for (uint64_t i=0; i<elem_count; i++) {- char* elem = elems+(i*elem_size);- for (unsigned int j=0; j<elem_size/2; j++) {- char head = elem[j];- int tail_index = elem_size-1-j;- elem[j] = elem[tail_index];- elem[tail_index] = head;- }- }+ flip_bytes(elem_size, (unsigned char*) *data); } return 0;@@ -824,7 +727,12 @@ return expected_type->read_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; } }
rts/csharp/scalar.cs view
@@ -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);}
rts/python/scalar.py view
@@ -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)
src/Futhark/Analysis/CallGraph.hs view
@@ -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
src/Futhark/Analysis/Range.hs view
@@ -3,6 +3,7 @@ ( rangeAnalysis , runRangeM , RangeM+ , analyseFun , analyseExp , analyseLambda , analyseBody
src/Futhark/CLI/Dev.hs view
@@ -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"
src/Futhark/CodeGen/Backends/CCUDA.hs view
@@ -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
src/Futhark/CodeGen/Backends/CCUDA/Boilerplate.hs view
@@ -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,
src/Futhark/CodeGen/Backends/COpenCL.hs view
@@ -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 <CL/cl.h>", "#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); }
src/Futhark/CodeGen/Backends/COpenCL/Boilerplate.hs view
@@ -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); } |]
src/Futhark/CodeGen/Backends/CSOpenCL.hs view
@@ -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
src/Futhark/CodeGen/Backends/GenericC.hs view
@@ -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 <stdint.h>") $esc:("#include <stddef.h>")@@ -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 } }|]
src/Futhark/CodeGen/Backends/GenericCSharp.hs view
@@ -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
src/Futhark/CodeGen/Backends/GenericPython.hs view
@@ -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]
src/Futhark/CodeGen/Backends/PyOpenCL.hs view
@@ -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
src/Futhark/CodeGen/Backends/SequentialC.hs view
@@ -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 =
src/Futhark/CodeGen/Backends/SimpleRepresentation.hs view
@@ -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]
src/Futhark/CodeGen/ImpCode.hs view
@@ -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] }
src/Futhark/CodeGen/ImpCode/OpenCL.hs view
@@ -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
src/Futhark/CodeGen/ImpGen.hs view
@@ -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
src/Futhark/CodeGen/ImpGen/Kernels.hs view
@@ -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
src/Futhark/CodeGen/ImpGen/Kernels/Base.hs view
@@ -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
src/Futhark/CodeGen/ImpGen/Kernels/SegHist.hs view
@@ -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
src/Futhark/CodeGen/ImpGen/Kernels/Transpose.hs view
@@ -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).
src/Futhark/Compiler.hs view
@@ -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) =>
src/Futhark/Internalise.hs view
@@ -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
src/Futhark/Internalise/Monomorphise.hs view
@@ -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
src/Futhark/Optimise/InliningDeadFun.hs view
@@ -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
src/Futhark/Optimise/Simplify/Rules.hs view
@@ -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
+ src/Futhark/Optimise/Sink.hs view
@@ -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
src/Futhark/Passes.hs view
@@ -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 ]
src/Futhark/Pipeline.hs view
@@ -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'
src/Futhark/Representation/Primitive.hs view
@@ -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
src/Futhark/Representation/Ranges.hs view
@@ -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
src/Futhark/Representation/SOACS/Simplify.hs view
@@ -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
src/Futhark/Transform/CopyPropagate.hs view
@@ -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
src/Futhark/TypeCheck.hs view
@@ -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 ()
src/Language/Futhark/Interpreter.hs view
@@ -115,7 +115,7 @@ ppr (ValueRecord m) = prettyRecord m ppr ValueFun{} = text "#<fun>"- 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.
src/Language/Futhark/Parser/Parser.y view
@@ -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
src/Language/Futhark/Pretty.hs view
@@ -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
src/Language/Futhark/TypeChecker/Monad.hs view
@@ -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
src/Language/Futhark/TypeChecker/Terms.hs view
@@ -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
src/Language/Futhark/TypeChecker/Types.hs view
@@ -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',
src/Language/Futhark/TypeChecker/Unify.hs view
@@ -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