Skip to content
Snippets Groups Projects
Commit 8057fa89 authored by Stephan Seitz's avatar Stephan Seitz
Browse files

Fix Memcpy2D

parent 16b089f2
No related branches found
No related tags found
No related merge requests found
...@@ -187,7 +187,7 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + ( ...@@ -187,7 +187,7 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + (
copy_params = f'cpy_{texture_name}_params' copy_params = f'cpy_{texture_name}_params'
return f"""cudaMemcpy3DParms {copy_params}{{}}; return f"""cudaMemcpy3DParms {copy_params}{{}};
{copy_params}.srcPtr = {{{self._device_ptr}, {copy_params}.srcPtr = {{{self._device_ptr},
{self._texture.field.strides[-1] * self._texture.field.shape[-2]}, {self._texture.field.strides[-2] * self._texture.field.dtype.numpy_dtype.itemsize},
{self._texture.field.shape[-1]}, {self._texture.field.shape[-1]},
{self._texture.field.shape[-2]}}}; {self._texture.field.shape[-2]}}};
{copy_params}.dstArray = {array}; {copy_params}.dstArray = {array};
...@@ -196,13 +196,12 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + ( ...@@ -196,13 +196,12 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + (
cudaMemcpy3D(&{copy_params});""" # noqa cudaMemcpy3D(&{copy_params});""" # noqa
elif self._texture.field.ndim == 2: elif self._texture.field.ndim == 2:
# noqa: cudaMemcpy2DToArray(cudaArray_t dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind); # noqa: cudaMemcpy2DToArray(cudaArray_t dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);
return f"""cudaMemcpy2DToArray({array}, return f"""cudaMemcpy2DToArray({array},
0u, 0u,
0u, 0u,
{self._device_ptr}, {self._device_ptr},
{self._texture.field.strides[-1]}, {self._texture.field.strides[-2] * self._texture.field.dtype.numpy_dtype.itemsize},
{self._texture.field.shape[-1]}, {self._texture.field.shape[-1] * self._texture.field.dtype.numpy_dtype.itemsize}, // Dafaq, this has to be in bytes, but only columns and only in memcpy2D
{self._texture.field.shape[-2]}, {self._texture.field.shape[-2]},
cudaMemcpyDeviceToDevice); cudaMemcpyDeviceToDevice);
""" # noqa """ # noqa
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment