summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/3d')
-rw-r--r--cuda/3d/cone_bp.cu62
-rw-r--r--cuda/3d/cone_fp.cu66
-rw-r--r--cuda/3d/par3d_bp.cu53
-rw-r--r--cuda/3d/par3d_fp.cu68
-rw-r--r--cuda/3d/util3d.cu22
5 files changed, 120 insertions, 151 deletions
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index e265304..41d781c 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,12 @@ struct DevConeParams {
__constant__ DevConeParams gC_C[g_MaxAngles];
-bool bindProjDataTexture(const cudaArray* array)
-{
- 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;
-}
-
-
//__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 +113,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 +134,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 +200,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 +293,9 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
const SDimensions3D& dims, const SConeProjection* angles,
const SProjectorParams3D& params)
{
- bindProjDataTexture(D_projArray);
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(D_projArray, D_texObj))
+ return false;
float fOutputScale;
if (params.bFDKWeighting) {
@@ -323,14 +305,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 +327,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,
diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu
index 2c3d1f6..2ef58ee 100644
--- a/cuda/3d/cone_fp.cu
+++ b/cuda/3d/cone_fp.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_coneVolumeTexture;
-
namespace astraCUDA3d {
static const unsigned int g_anglesPerBlock = 4;
@@ -63,24 +59,6 @@ __constant__ float gC_DetVY[g_MaxAngles];
__constant__ float gC_DetVZ[g_MaxAngles];
-bool bindVolumeDataTexture(const cudaArray* array)
-{
- cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
-
- gT_coneVolumeTexture.addressMode[0] = cudaAddressModeBorder;
- gT_coneVolumeTexture.addressMode[1] = cudaAddressModeBorder;
- gT_coneVolumeTexture.addressMode[2] = cudaAddressModeBorder;
- gT_coneVolumeTexture.filterMode = cudaFilterModeLinear;
- gT_coneVolumeTexture.normalized = false;
-
- cudaBindTextureToArray(gT_coneVolumeTexture, array, channelDesc);
-
- // TODO: error value?
-
- return true;
-}
-
-
// x=0, y=1, z=2
struct DIR_X {
__device__ float nSlices(const SDimensions3D& dims) const { return dims.iVolX; }
@@ -89,7 +67,7 @@ struct DIR_X {
__device__ float c0(float x, float y, float z) const { return x; }
__device__ float c1(float x, float y, float z) const { return y; }
__device__ float c2(float x, float y, float z) const { return z; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f0, f1, f2); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f0, f1, f2); }
__device__ float x(float f0, float f1, float f2) const { return f0; }
__device__ float y(float f0, float f1, float f2) const { return f1; }
__device__ float z(float f0, float f1, float f2) const { return f2; }
@@ -103,7 +81,7 @@ struct DIR_Y {
__device__ float c0(float x, float y, float z) const { return y; }
__device__ float c1(float x, float y, float z) const { return x; }
__device__ float c2(float x, float y, float z) const { return z; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f1, f0, f2); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f1, f0, f2); }
__device__ float x(float f0, float f1, float f2) const { return f1; }
__device__ float y(float f0, float f1, float f2) const { return f0; }
__device__ float z(float f0, float f1, float f2) const { return f2; }
@@ -117,7 +95,7 @@ struct DIR_Z {
__device__ float c0(float x, float y, float z) const { return z; }
__device__ float c1(float x, float y, float z) const { return x; }
__device__ float c2(float x, float y, float z) const { return y; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f1, f2, f0); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f1, f2, f0); }
__device__ float x(float f0, float f1, float f2) const { return f1; }
__device__ float y(float f0, float f1, float f2) const { return f2; }
__device__ float z(float f0, float f1, float f2) const { return f0; }
@@ -144,6 +122,7 @@ struct SCALE_NONCUBE {
template<class COORD, class SCALE>
__global__ void cone_FP_t(float* D_projData, unsigned int projPitch,
+ cudaTextureObject_t tex,
unsigned int startSlice,
unsigned int startAngle, unsigned int endAngle,
const SDimensions3D dims,
@@ -208,7 +187,7 @@ __global__ void cone_FP_t(float* D_projData, unsigned int projPitch,
for (int s = startSlice; s < endSlice; ++s)
{
- fVal += c.tex(f0, f1, f2);
+ fVal += c.tex(tex, f0, f1, f2);
f0 += 1.0f;
f1 += a1;
f2 += a2;
@@ -222,6 +201,7 @@ __global__ void cone_FP_t(float* D_projData, unsigned int projPitch,
template<class COORD>
__global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch,
+ cudaTextureObject_t tex,
unsigned int startSlice,
unsigned int startAngle, unsigned int endAngle,
const SDimensions3D dims, int iRaysPerDetDim,
@@ -295,7 +275,7 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch,
for (int s = startSlice; s < endSlice; ++s)
{
- fVal += c.tex(f0, f1, f2);
+ fVal += c.tex(tex, f0, f1, f2);
f0 += 1.0f;
f1 += a1;
f2 += a2;
@@ -313,7 +293,9 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch,
bool ConeFP_Array_internal(cudaPitchedPtr D_projData,
- const SDimensions3D& dims, unsigned int angleCount, const SConeProjection* angles,
+ cudaTextureObject_t D_texObj,
+ const SDimensions3D& dims,
+ unsigned int angleCount, const SConeProjection* angles,
const SProjectorParams3D& params)
{
// transfer angles to constant memory
@@ -419,29 +401,29 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData,
for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube);
else
- cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX);
+ cone_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeX);
else
- cone_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX);
+ cone_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX);
} else if (blockDirection == 1) {
for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube);
else
- cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY);
+ cone_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeY);
else
- cone_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY);
+ cone_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY);
} else if (blockDirection == 2) {
for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube);
else
- cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ);
+ cone_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeZ);
else
- cone_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ);
+ cone_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ);
}
}
@@ -470,10 +452,14 @@ bool ConeFP(cudaPitchedPtr D_volumeData,
const SProjectorParams3D& params)
{
// transfer volume to array
-
cudaArray* cuArray = allocateVolumeArray(dims);
transferVolumeToArray(D_volumeData, cuArray, dims);
- bindVolumeDataTexture(cuArray);
+
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(cuArray, D_texObj)) {
+ cudaFreeArray(cuArray);
+ return false;
+ }
bool ret;
@@ -485,7 +471,7 @@ bool ConeFP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_subprojData = D_projData;
D_subprojData.ptr = (char*)D_projData.ptr + iAngle * D_projData.pitch;
- ret = ConeFP_Array_internal(D_subprojData,
+ ret = ConeFP_Array_internal(D_subprojData, D_texObj,
dims, iEndAngle - iAngle, angles + iAngle,
params);
if (!ret)
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu
index 1dc75ce..27d95fe 100644
--- a/cuda/3d/par3d_bp.cu
+++ b/cuda/3d/par3d_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_par3DProjTexture;
-
namespace astraCUDA3d {
static const unsigned int g_volBlockZ = 6;
@@ -58,26 +54,8 @@ __constant__ DevPar3DParams gC_C[g_MaxAngles];
__constant__ float gC_scale[g_MaxAngles];
-static bool bindProjDataTexture(const cudaArray* array)
-{
- cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
-
- gT_par3DProjTexture.addressMode[0] = cudaAddressModeBorder;
- gT_par3DProjTexture.addressMode[1] = cudaAddressModeBorder;
- gT_par3DProjTexture.addressMode[2] = cudaAddressModeBorder;
- gT_par3DProjTexture.filterMode = cudaFilterModeLinear;
- gT_par3DProjTexture.normalized = false;
-
- cudaBindTextureToArray(gT_par3DProjTexture, array, channelDesc);
-
- // TODO: error value?
-
- return true;
-}
-
-
template<unsigned int ZSIZE>
-__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
+__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale)
{
float* volData = (float*)D_volData;
@@ -125,7 +103,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn
for (int idx = 0; idx < ZSIZE; ++idx) {
- float fVal = tex3D(gT_par3DProjTexture, fU, fAngle, fV);
+ float fVal = tex3D<float>(tex, fU, fAngle, fV);
Z[idx] += fVal * fS;
fU += fCu.z;
@@ -144,7 +122,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn
}
// supersampling version
-__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale)
+__global__ void dev_par3D_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;
@@ -206,7 +184,7 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star
const float fU = fCu.w + fXs * fCu.x + fYs * fCu.y + fZs * fCu.z;
const float fV = fCv.w + fXs * fCv.x + fYs * fCv.y + fZs * fCv.z;
- fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV) * fS;
+ fVal += tex3D<float>(tex, fU, fAngle, fV) * fS;
fZs += fSubStep;
}
fYs += fSubStep;
@@ -259,18 +237,22 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
const SDimensions3D& dims, const SPar3DProjection* angles,
const SProjectorParams3D& params)
{
- bindProjDataTexture(D_projArray);
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(D_projArray, D_texObj))
+ return false;
float 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);
@@ -283,23 +265,26 @@ bool Par3DBP_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.iRaysPerVoxelDim == 1) {
if (dims.iVolZ == 1) {
- dev_par3D_BP<1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+ dev_par3D_BP<1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
} else {
- dev_par3D_BP<g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+ dev_par3D_BP<g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale);
}
} else
- dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale);
+ dev_par3D_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;
}
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu
index 5daddc1..b2178ec 100644
--- a/cuda/3d/par3d_fp.cu
+++ b/cuda/3d/par3d_fp.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_par3DVolumeTexture;
-
namespace astraCUDA3d {
static const unsigned int g_anglesPerBlock = 4;
@@ -63,24 +59,6 @@ __constant__ float gC_DetVY[g_MaxAngles];
__constant__ float gC_DetVZ[g_MaxAngles];
-static bool bindVolumeDataTexture(const cudaArray* array)
-{
- cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
-
- gT_par3DVolumeTexture.addressMode[0] = cudaAddressModeBorder;
- gT_par3DVolumeTexture.addressMode[1] = cudaAddressModeBorder;
- gT_par3DVolumeTexture.addressMode[2] = cudaAddressModeBorder;
- gT_par3DVolumeTexture.filterMode = cudaFilterModeLinear;
- gT_par3DVolumeTexture.normalized = false;
-
- cudaBindTextureToArray(gT_par3DVolumeTexture, array, channelDesc);
-
- // TODO: error value?
-
- return true;
-}
-
-
// x=0, y=1, z=2
struct DIR_X {
__device__ float nSlices(const SDimensions3D& dims) const { return dims.iVolX; }
@@ -89,7 +67,7 @@ struct DIR_X {
__device__ float c0(float x, float y, float z) const { return x; }
__device__ float c1(float x, float y, float z) const { return y; }
__device__ float c2(float x, float y, float z) const { return z; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_par3DVolumeTexture, f0, f1, f2); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f0, f1, f2); }
__device__ float x(float f0, float f1, float f2) const { return f0; }
__device__ float y(float f0, float f1, float f2) const { return f1; }
__device__ float z(float f0, float f1, float f2) const { return f2; }
@@ -103,7 +81,7 @@ struct DIR_Y {
__device__ float c0(float x, float y, float z) const { return y; }
__device__ float c1(float x, float y, float z) const { return x; }
__device__ float c2(float x, float y, float z) const { return z; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_par3DVolumeTexture, f1, f0, f2); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f1, f0, f2); }
__device__ float x(float f0, float f1, float f2) const { return f1; }
__device__ float y(float f0, float f1, float f2) const { return f0; }
__device__ float z(float f0, float f1, float f2) const { return f2; }
@@ -117,7 +95,7 @@ struct DIR_Z {
__device__ float c0(float x, float y, float z) const { return z; }
__device__ float c1(float x, float y, float z) const { return x; }
__device__ float c2(float x, float y, float z) const { return y; }
- __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_par3DVolumeTexture, f1, f2, f0); }
+ __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D<float>(tex, f1, f2, f0); }
__device__ float x(float f0, float f1, float f2) const { return f1; }
__device__ float y(float f0, float f1, float f2) const { return f2; }
__device__ float z(float f0, float f1, float f2) const { return f0; }
@@ -145,6 +123,7 @@ struct SCALE_NONCUBE {
template<class COORD, class SCALE>
__global__ void par3D_FP_t(float* D_projData, unsigned int projPitch,
+ cudaTextureObject_t tex,
unsigned int startSlice,
unsigned int startAngle, unsigned int endAngle,
const SDimensions3D dims,
@@ -210,7 +189,7 @@ __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch,
for (int s = startSlice; s < endSlice; ++s)
{
- fVal += c.tex(f0, f1, f2);
+ fVal += c.tex(tex, f0, f1, f2);
f0 += 1.0f;
f1 += a1;
f2 += a2;
@@ -225,6 +204,7 @@ __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch,
// Supersampling version
template<class COORD>
__global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch,
+ cudaTextureObject_t tex,
unsigned int startSlice,
unsigned int startAngle, unsigned int endAngle,
const SDimensions3D dims, int iRaysPerDetDim,
@@ -301,7 +281,7 @@ __global__ void par3D_FP_SS_t(float* D_projData, unsigned int projPitch,
for (int s = startSlice; s < endSlice; ++s)
{
- fVal += c.tex(f0, f1, f2);
+ fVal += c.tex(tex, f0, f1, f2);
f0 += 1.0f;
f1 += a1;
f2 += a2;
@@ -415,7 +395,9 @@ __global__ void par3D_FP_SumSqW_t(float* D_projData, unsigned int projPitch,
bool Par3DFP_Array_internal(cudaPitchedPtr D_projData,
- const SDimensions3D& dims, unsigned int angleCount, const SPar3DProjection* angles,
+ cudaTextureObject_t D_texObj,
+ const SDimensions3D& dims,
+ unsigned int angleCount, const SPar3DProjection* angles,
const SProjectorParams3D& params)
{
// transfer angles to constant memory
@@ -520,29 +502,29 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData,
for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube);
else
- par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX);
+ par3D_FP_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, snoncubeX);
else
- par3D_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX);
+ par3D_FP_SS_t<DIR_X><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX);
} else if (blockDirection == 1) {
for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, scube);
else
- par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY);
+ par3D_FP_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, snoncubeY);
else
- par3D_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY);
+ par3D_FP_SS_t<DIR_Y><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY);
} else if (blockDirection == 2) {
for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices)
if (params.iRaysPerDetDim == 1)
if (cube)
- par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube);
+ par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, scube);
else
- par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ);
+ par3D_FP_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, snoncubeZ);
else
- par3D_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ);
+ par3D_FP_SS_t<DIR_Z><<<dimGrid, dimBlock, 0, stream>>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj,i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ);
}
}
@@ -569,10 +551,16 @@ bool Par3DFP(cudaPitchedPtr D_volumeData,
const SDimensions3D& dims, const SPar3DProjection* angles,
const SProjectorParams3D& params)
{
+
// transfer volume to array
cudaArray* cuArray = allocateVolumeArray(dims);
transferVolumeToArray(D_volumeData, cuArray, dims);
- bindVolumeDataTexture(cuArray);
+
+ cudaTextureObject_t D_texObj;
+ if (!createTextureObject3D(cuArray, D_texObj)) {
+ cudaFreeArray(cuArray);
+ return false;
+ }
bool ret;
@@ -584,7 +572,7 @@ bool Par3DFP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_subprojData = D_projData;
D_subprojData.ptr = (char*)D_projData.ptr + iAngle * D_projData.pitch;
- ret = Par3DFP_Array_internal(D_subprojData,
+ ret = Par3DFP_Array_internal(D_subprojData, D_texObj,
dims, iEndAngle - iAngle, angles + iAngle,
params);
if (!ret)
@@ -593,6 +581,8 @@ bool Par3DFP(cudaPitchedPtr D_volumeData,
cudaFreeArray(cuArray);
+ cudaDestroyTextureObject(D_texObj);
+
return 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,
generated by cgit v1.2.3 (git 2.25.1) at 2024-11-14 08:15:25 +0000