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

Fix texture uploading for the 3-d case

parent 5ebe72cd
Branches
Tags
No related merge requests found
Pipeline #19184 failed
...@@ -258,14 +258,15 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + ( ...@@ -258,14 +258,15 @@ cudaMalloc{ndim}Array(&{array}, &channel_desc_{texture_name}, """ + (
array = 'array_' + texture_name array = 'array_' + texture_name
if self._texture.field.ndim == 3: if self._texture.field.ndim == 3:
copy_params = f'cpy_{texture_name}_params' copy_params = f'cpy_{texture_name}_params'
return f"""cudaMemcpy3DParams {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[-1] * self._dtype.numpy_dtype.itemsize}, {self._texture.field.strides[-1] * self._texture.field.shape[-1] * self._dtype.numpy_dtype.itemsize},
{self._texture.field.shape[-1], self._texture.field.shape[-2]}}}; {self._texture.field.shape[-1]},
{copy_params}.dsrArray = {array}; {self._texture.field.shape[-2]}}};
copyParams.extent = {{{", ".join(reversed(self._shape))}}}; {copy_params}.dstArray = {array};
copyParams.kind = cudaMemcpyDeviceToDevice; {copy_params}.extent = {{{", ".join(str(s) for s in reversed(self._shape))}}};
cudaMemcpy3D(&{{copy_params}});""" # noqa {copy_params}.kind = cudaMemcpyDeviceToDevice;
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);
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment