From d4645875801d29402b1e3f2a42a5d9902a37a718 Mon Sep 17 00:00:00 2001
From: Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>
Date: Mon, 22 Nov 2021 14:14:55 +0100
Subject: De-duplicate 3D texture object creation
---
cuda/3d/cone_bp.cu | 26 ++------------------------
cuda/3d/cone_fp.cu | 32 ++++++--------------------------
cuda/3d/par3d_bp.cu | 26 ++------------------------
cuda/3d/par3d_fp.cu | 31 ++++++-------------------------
cuda/3d/util3d.cu | 22 ++++++++++++++++++++++
include/astra/cuda/3d/util3d.h | 2 ++
6 files changed, 40 insertions(+), 99 deletions(-)
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index 3092467..41d781c 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -53,29 +53,6 @@ struct DevConeParams {
__constant__ DevConeParams gC_C[g_MaxAngles];
-bool bindProjDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
-{
- cudaChannelFormatDesc channelDesc =
- cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
-
- cudaResourceDesc resDesc;
- memset(&resDesc, 0, sizeof(resDesc));
- resDesc.resType = cudaResourceTypeArray;
- resDesc.res.array.array = array;
-
- cudaTextureDesc texDesc;
- memset(&texDesc, 0, sizeof(texDesc));
- texDesc.addressMode[0] = cudaAddressModeBorder;
- texDesc.addressMode[1] = cudaAddressModeBorder;
- texDesc.addressMode[2] = cudaAddressModeBorder;
- texDesc.filterMode = cudaFilterModeLinear;
- texDesc.readMode = cudaReadModeElementType;
- texDesc.normalizedCoords = 0;
-
- return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "cone_bp texture");
-}
-
-
//__launch_bounds__(32*16, 4)
template<bool FDKWEIGHT, unsigned int ZSIZE>
__global__ void dev_cone_BP(void* D_volData, unsigned int volPitch,
@@ -317,7 +294,8 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
const SProjectorParams3D& params)
{
cudaTextureObject_t D_texObj;
- bindProjDataTexture(D_projArray, D_texObj);
+ if (!createTextureObject3D(D_projArray, D_texObj))
+ return false;
float fOutputScale;
if (params.bFDKWeighting) {
diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu
index e49ea24..2ef58ee 100644
--- a/cuda/3d/cone_fp.cu
+++ b/cuda/3d/cone_fp.cu
@@ -59,29 +59,6 @@ __constant__ float gC_DetVY[g_MaxAngles];
__constant__ float gC_DetVZ[g_MaxAngles];
-bool bindVolumeDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
-{
- cudaChannelFormatDesc channelDesc =
- cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
-
- cudaResourceDesc resDesc;
- memset(&resDesc, 0, sizeof(resDesc));
- resDesc.resType = cudaResourceTypeArray;
- resDesc.res.array.array = array;
-
- cudaTextureDesc texDesc;
- memset(&texDesc, 0, sizeof(texDesc));
- texDesc.addressMode[0] = cudaAddressModeBorder;
- texDesc.addressMode[1] = cudaAddressModeBorder;
- texDesc.addressMode[2] = cudaAddressModeBorder;
- texDesc.filterMode = cudaFilterModeLinear;
- texDesc.readMode = cudaReadModeElementType;
- texDesc.normalizedCoords = 0;
-
- return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "cone_fp texture");
-}
-
-
// x=0, y=1, z=2
struct DIR_X {
__device__ float nSlices(const SDimensions3D& dims) const { return dims.iVolX; }
@@ -474,12 +451,15 @@ bool ConeFP(cudaPitchedPtr D_volumeData,
const SDimensions3D& dims, const SConeProjection* angles,
const SProjectorParams3D& params)
{
- cudaTextureObject_t D_texObj;
-
// transfer volume to array
cudaArray* cuArray = allocateVolumeArray(dims);
transferVolumeToArray(D_volumeData, cuArray, dims);
- bindVolumeDataTexture(cuArray, D_texObj);
+
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(cuArray, D_texObj)) {
+ cudaFreeArray(cuArray);
+ return false;
+ }
bool ret;
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu
index 748086e..27d95fe 100644
--- a/cuda/3d/par3d_bp.cu
+++ b/cuda/3d/par3d_bp.cu
@@ -54,29 +54,6 @@ __constant__ DevPar3DParams gC_C[g_MaxAngles];
__constant__ float gC_scale[g_MaxAngles];
-static bool bindProjDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
-{
- cudaChannelFormatDesc channelDesc =
- cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
-
- cudaResourceDesc resDesc;
- memset(&resDesc, 0, sizeof(resDesc));
- resDesc.resType = cudaResourceTypeArray;
- resDesc.res.array.array = array;
-
- cudaTextureDesc texDesc;
- memset(&texDesc, 0, sizeof(texDesc));
- texDesc.addressMode[0] = cudaAddressModeBorder;
- texDesc.addressMode[1] = cudaAddressModeBorder;
- texDesc.addressMode[2] = cudaAddressModeBorder;
- texDesc.filterMode = cudaFilterModeLinear;
- texDesc.readMode = cudaReadModeElementType;
- texDesc.normalizedCoords = 0;
-
- return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "par3d_bp texture");
-}
-
-
template<unsigned int ZSIZE>
__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
{
@@ -261,7 +238,8 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
const SProjectorParams3D& params)
{
cudaTextureObject_t D_texObj;
- bindProjDataTexture(D_projArray, D_texObj);
+ if (!createTextureObject3D(D_projArray, D_texObj))
+ return false;
float fOutputScale = params.fOutputScale * params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ;
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu
index cae75f1..b2178ec 100644
--- a/cuda/3d/par3d_fp.cu
+++ b/cuda/3d/par3d_fp.cu
@@ -59,29 +59,6 @@ __constant__ float gC_DetVY[g_MaxAngles];
__constant__ float gC_DetVZ[g_MaxAngles];
-static bool bindVolumeDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
-{
- cudaChannelFormatDesc channelDesc =
- cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
-
- cudaResourceDesc resDesc;
- memset(&resDesc, 0, sizeof(resDesc));
- resDesc.resType = cudaResourceTypeArray;
- resDesc.res.array.array = array;
-
- cudaTextureDesc texDesc;
- memset(&texDesc, 0, sizeof(texDesc));
- texDesc.addressMode[0] = cudaAddressModeBorder;
- texDesc.addressMode[1] = cudaAddressModeBorder;
- texDesc.addressMode[2] = cudaAddressModeBorder;
- texDesc.filterMode = cudaFilterModeLinear;
- texDesc.readMode = cudaReadModeElementType;
- texDesc.normalizedCoords = 0;
-
- return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "par3d_fp texture");
-}
-
-
// x=0, y=1, z=2
struct DIR_X {
__device__ float nSlices(const SDimensions3D& dims) const { return dims.iVolX; }
@@ -574,12 +551,16 @@ bool Par3DFP(cudaPitchedPtr D_volumeData,
const SDimensions3D& dims, const SPar3DProjection* angles,
const SProjectorParams3D& params)
{
- cudaTextureObject_t D_texObj;
// transfer volume to array
cudaArray* cuArray = allocateVolumeArray(dims);
transferVolumeToArray(D_volumeData, cuArray, dims);
- bindVolumeDataTexture(cuArray, D_texObj);
+
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(cuArray, D_texObj)) {
+ cudaFreeArray(cuArray);
+ return false;
+ }
bool ret;
diff --git a/cuda/3d/util3d.cu b/cuda/3d/util3d.cu
index 71b5668..3dc915d 100644
--- a/cuda/3d/util3d.cu
+++ b/cuda/3d/util3d.cu
@@ -378,6 +378,28 @@ bool transferHostProjectionsToArray(const float *projData, cudaArray* array, con
return checkCuda(cudaMemcpy3D(&p), "transferHostProjectionsToArray 3D");
}
+bool createTextureObject3D(cudaArray* array, cudaTextureObject_t& texObj)
+{
+ cudaChannelFormatDesc channelDesc =
+ cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
+
+ cudaResourceDesc resDesc;
+ memset(&resDesc, 0, sizeof(resDesc));
+ resDesc.resType = cudaResourceTypeArray;
+ resDesc.res.array.array = array;
+
+ cudaTextureDesc texDesc;
+ memset(&texDesc, 0, sizeof(texDesc));
+ texDesc.addressMode[0] = cudaAddressModeBorder;
+ texDesc.addressMode[1] = cudaAddressModeBorder;
+ texDesc.addressMode[2] = cudaAddressModeBorder;
+ texDesc.filterMode = cudaFilterModeLinear;
+ texDesc.readMode = cudaReadModeElementType;
+ texDesc.normalizedCoords = 0;
+
+ return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "createTextureObject3D");
+}
+
float dotProduct3D(cudaPitchedPtr data, unsigned int x, unsigned int y,
diff --git a/include/astra/cuda/3d/util3d.h b/include/astra/cuda/3d/util3d.h
index 9fa254d..210d944 100644
--- a/include/astra/cuda/3d/util3d.h
+++ b/include/astra/cuda/3d/util3d.h
@@ -60,6 +60,8 @@ bool zeroVolumeArray(cudaArray* array, const SDimensions3D& dims);
cudaArray* allocateProjectionArray(const SDimensions3D& dims);
cudaArray* allocateVolumeArray(const SDimensions3D& dims);
+bool createTextureObject3D(cudaArray* array, cudaTextureObject_t& texObj);
+
float dotProduct3D(cudaPitchedPtr data, unsigned int x, unsigned int y, unsigned int z);
int calcNextPowerOfTwo(int _iValue);
--
cgit v1.2.3