From 5a5c4c227f9f34427a2c30a04e39c4b8886e65e3 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 19 Nov 2024 09:42:47 +0100 Subject: [PATCH 1/2] Now in all `new[]` expressions the size is now explicitly casted to integer. The reason is if `pow` is used as size. --- dace/codegen/instrumentation/data/data_dump.py | 4 ++-- dace/codegen/targets/cpu.py | 4 ++-- dace/codegen/targets/snitch.py | 3 ++- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/dace/codegen/instrumentation/data/data_dump.py b/dace/codegen/instrumentation/data/data_dump.py index 5fc487f94d..ce85905a1f 100644 --- a/dace/codegen/instrumentation/data/data_dump.py +++ b/dace/codegen/instrumentation/data/data_dump.py @@ -45,7 +45,7 @@ def _generate_copy_to_host(self, node: nodes.AccessNode, desc: dt.Array, ptr: st # Emit synchronous memcpy preamble = f''' {{ - {new_desc.as_arg(name=new_ptr)} = new {desc.dtype.ctype}[{csize}]; + {new_desc.as_arg(name=new_ptr)} = new {desc.dtype.ctype}[std::size_t({csize})]; {self.backend}Memcpy({new_ptr}, {ptr}, sizeof({desc.dtype.ctype}) * ({csize}), {self.backend}MemcpyDeviceToHost); ''' @@ -65,7 +65,7 @@ def _generate_copy_to_device(self, node: nodes.AccessNode, desc: dt.Array, ptr: # Emit synchronous memcpy preamble = f''' {{ - {new_desc.as_arg(name=new_ptr)} = new {desc.dtype.ctype}[{csize}]; + {new_desc.as_arg(name=new_ptr)} = new {desc.dtype.ctype}[std::size_t({csize})]; ''' postamble = f''' diff --git a/dace/codegen/targets/cpu.py b/dace/codegen/targets/cpu.py index 9ba202757e..c57ab0a478 100644 --- a/dace/codegen/targets/cpu.py +++ b/dace/codegen/targets/cpu.py @@ -498,7 +498,7 @@ def allocate_array(self, sdfg: SDFG, cfg: ControlFlowRegion, dfg: StateSubgraphV if not declared: declaration_stream.write(f'{nodedesc.dtype.ctype} *{name};\n', cfg, state_id, node) allocation_stream.write( - "%s = new %s DACE_ALIGN(64)[%s];\n" % (alloc_name, nodedesc.dtype.ctype, cpp.sym2cpp(arrsize)), cfg, + "%s = new %s DACE_ALIGN(64)[std::size_t(%s)];\n" % (alloc_name, nodedesc.dtype.ctype, cpp.sym2cpp(arrsize)), cfg, state_id, node) define_var(name, DefinedType.Pointer, ctypedef) @@ -548,7 +548,7 @@ def allocate_array(self, sdfg: SDFG, cfg: ControlFlowRegion, dfg: StateSubgraphV """ #pragma omp parallel {{ - {name} = new {ctype} DACE_ALIGN(64)[{arrsize}];""".format(ctype=nodedesc.dtype.ctype, + {name} = new {ctype} DACE_ALIGN(64)[std::size_t({arrsize})];""".format(ctype=nodedesc.dtype.ctype, name=alloc_name, arrsize=cpp.sym2cpp(arrsize)), cfg, diff --git a/dace/codegen/targets/snitch.py b/dace/codegen/targets/snitch.py index 5a62ca2995..d6c6738653 100644 --- a/dace/codegen/targets/snitch.py +++ b/dace/codegen/targets/snitch.py @@ -408,7 +408,7 @@ def allocate_array(self, sdfg: SDFG, cfg: ControlFlowRegion, dfg: StateSubgraphV #pragma omp parallel {{ #error "malloc is not threadsafe" - {name} = new {ctype} [{arrsize}];""".format(ctype=nodedesc.dtype.ctype, + {name} = new {ctype} [std::size_t({arrsize})];""".format(ctype=nodedesc.dtype.ctype, name=alloc_name, arrsize=cpp.sym2cpp(arrsize)), cfg, @@ -1108,6 +1108,7 @@ def gen_code_snitch(sdfg): # change new/delete to malloc/free code._code = re.sub(r"new (.+) \[(\d*)\];", r"(\1*)malloc(\2*sizeof(\1));", code._code) + code._code = re.sub(r"new (.+) \[std::size_t\((\d*)\)\];", r"(\1*)malloc(\2*sizeof(\1));", code._code) code._code = re.sub(r"new ([a-zA-Z0-9 _]*);", r"(\1*)malloc(sizeof(\1));", code._code) code._code = re.sub(r"delete (.*);", r"free(\1);", code._code) code._code = re.sub(r"delete\[\] (.*);", r"free(\1);", code._code) From e972978519a322f0f57e57c9a79bc789e66c21e2 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 19 Nov 2024 09:48:52 +0100 Subject: [PATCH 2/2] Enabled the Stockham FFT tests. --- tests/npbench/misc/stockham_fft_test.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/tests/npbench/misc/stockham_fft_test.py b/tests/npbench/misc/stockham_fft_test.py index 5878cf621a..eeedaf2034 100644 --- a/tests/npbench/misc/stockham_fft_test.py +++ b/tests/npbench/misc/stockham_fft_test.py @@ -155,12 +155,10 @@ def run_stockham_fft(device_type: dace.dtypes.DeviceType): return sdfg -@pytest.mark.skip(reason="Assertion error in read_and_write_sets") def test_cpu(): run_stockham_fft(dace.dtypes.DeviceType.CPU) -@pytest.mark.skip(reason="Assertion error in read_and_write_sets") @pytest.mark.gpu def test_gpu(): run_stockham_fft(dace.dtypes.DeviceType.GPU)