Replace texref by texobj in cone_bp

1 files changed, 42 insertions, 34 deletions

diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index e265304..3092467 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -35,10 +35,6 @@ along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>.
 
 #include <cuda.h>
 
-typedef texture<float, 3, cudaReadModeElementType> texture3D;
-
-static texture3D gT_coneProjTexture;
-
 namespace astraCUDA3d {
 
 static const unsigned int g_volBlockZ = 6;
@@ -57,28 +53,35 @@ struct DevConeParams {
 
 __constant__ DevConeParams gC_C[g_MaxAngles];
 
-bool bindProjDataTexture(const cudaArray* array)
+bool bindProjDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
 {
-	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
-
-	gT_coneProjTexture.addressMode[0] = cudaAddressModeBorder;
-	gT_coneProjTexture.addressMode[1] = cudaAddressModeBorder;
-	gT_coneProjTexture.addressMode[2] = cudaAddressModeBorder;
-	gT_coneProjTexture.filterMode = cudaFilterModeLinear;
-	gT_coneProjTexture.normalized = false;
-
-	cudaBindTextureToArray(gT_coneProjTexture, array, channelDesc);
-
-	// TODO: error value?
-
-	return true;
+	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, int startAngle,
-                            int angleOffset, const astraCUDA3d::SDimensions3D dims,
+__global__ void dev_cone_BP(void* D_volData, unsigned int volPitch,
+                            cudaTextureObject_t tex,
+                            int startAngle, int angleOffset,
+                            const astraCUDA3d::SDimensions3D dims,
                             float fOutputScale)
 {
 	float* volData = (float*)D_volData;
@@ -133,7 +136,7 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng
 				fr = __fdividef(1.0f, fDen);
 				fU = fUNum * fr;
 				fV = fVNum * fr;
-				float fVal = tex3D(gT_coneProjTexture, fU, fAngle, fV);
+				float fVal = tex3D<float>(tex, fU, fAngle, fV);
 				Z[idx] += fr*fr*fVal;
 
 				fUNum += fCu.z;
@@ -154,7 +157,7 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng
 
 
 // supersampling version
-__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale)
+__global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale)
 {
 	float* volData = (float*)D_volData;
 
@@ -220,7 +223,7 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start
 				const float fU = fUNum * fr;
 				const float fV = fVNum * fr;
 
-				fVal += tex3D(gT_coneProjTexture, fU, fAngle, fV) * fr * fr;
+				fVal += tex3D<float>(tex, fU, fAngle, fV) * fr * fr;
 
 				fZs += fSubStep;
 			}
@@ -313,7 +316,8 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
                   const SDimensions3D& dims, const SConeProjection* angles,
                   const SProjectorParams3D& params)
 {
-	bindProjDataTexture(D_projArray);
+	cudaTextureObject_t D_texObj;
+	bindProjDataTexture(D_projArray, D_texObj);
 
 	float fOutputScale;
 	if (params.bFDKWeighting) {
@@ -323,14 +327,16 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
 		fOutputScale = params.fOutputScale * (params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ);
 	}
 
+	bool ok = true;
+
 	for (unsigned int th = 0; th < dims.iProjAngles; th += g_MaxAngles) {
 		unsigned int angleCount = g_MaxAngles;
 		if (th + angleCount > dims.iProjAngles)
 			angleCount = dims.iProjAngles - th;
 
-		bool ok = transferConstants(angles, angleCount, params);
+		ok = transferConstants(angles, angleCount, params);
 		if (!ok)
-			return false;
+			break;
 
 		dim3 dimBlock(g_volBlockX, g_volBlockY);
 
@@ -343,31 +349,33 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
 		// printf("Calling BP: %d, %dx%d, %dx%d to %p\n", i, dimBlock.x, dimBlock.y, dimGrid.x, dimGrid.y, (void*)D_volumeData.ptr); 
 			if (params.bFDKWeighting) {
 				if (dims.iVolZ == 1) {
-					dev_cone_BP<true, 1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+					dev_cone_BP<true, 1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
 				} else {
-					dev_cone_BP<true, g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+					dev_cone_BP<true, g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
 				}
 			} else if (params.iRaysPerVoxelDim == 1) {
 				if (dims.iVolZ == 1) {
-					dev_cone_BP<false, 1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+					dev_cone_BP<false, 1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
 				} else {
-					dev_cone_BP<false, g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+					dev_cone_BP<false, g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
 				}
 			} else
-				dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale);
+				dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, params.iRaysPerVoxelDim, fOutputScale);
 		}
 
 		// TODO: Consider not synchronizing here, if possible.
-		if (!checkCuda(cudaThreadSynchronize(), "cone_bp"))
-			return false;
+		ok = checkCuda(cudaThreadSynchronize(), "cone_bp");
+		if (!ok)
+			break;
 
 		angles = angles + angleCount;
 		// printf("%f\n", toc(t));
 
 	}
 
+	cudaDestroyTextureObject(D_texObj);
 
-	return true;
+	return ok;
 }
 
 bool ConeBP(cudaPitchedPtr D_volumeData,