Skip to content

Commit f610f8c

Browse files
purepanipurepani
committed
Successful compilation of HIP
1 parent 29e4f6a commit f610f8c

11 files changed

+75
-70
lines changed

‎Common/CUDA/GD_AwTV.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -542,7 +542,7 @@ void aw_pocs_tv(float* img,float* dst,float alpha,const long* image_size, int ma
542542
size_t dimgridRed = (total_pixels + MAXTHREADS - 1) / MAXTHREADS;
543543

544544
hipStreamSynchronize(stream[dev*nStream_device+1]);
545-
reduceNorm2 << <dimgridRed, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>> >(d_norm2[dev], d_norm2aux[dev], total_pixels);
545+
reduceNorm2 <<<dimgridRed, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>>>(d_norm2[dev], d_norm2aux[dev], total_pixels);
546546

547547
}
548548
for (dev = 0; dev < deviceCount; dev++){
@@ -553,7 +553,7 @@ void aw_pocs_tv(float* img,float* dst,float alpha,const long* image_size, int ma
553553
size_t dimgridRed = (total_pixels + MAXTHREADS - 1) / MAXTHREADS;
554554

555555
if (dimgridRed > 1) {
556-
reduceSum << <1, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device] >> >(d_norm2aux[dev], d_norm2[dev], dimgridRed);
556+
reduceSum <<<1, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>>>(d_norm2aux[dev], d_norm2[dev], dimgridRed);
557557
hipStreamSynchronize(stream[dev*nStream_device]);
558558
hipMemcpyAsync(&sumnorm2[dev], d_norm2[dev], sizeof(float), hipMemcpyDeviceToHost,stream[dev*nStream_device+1]);
559559
}

‎Common/CUDA/GD_TV.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -526,7 +526,7 @@ do { \
526526
size_t dimgridRed = (total_pixels + MAXTHREADS - 1) / MAXTHREADS;
527527

528528
hipStreamSynchronize(stream[dev*nStream_device+1]);
529-
reduceNorm2 << <dimgridRed, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>> >(d_norm2[dev], d_norm2aux[dev], total_pixels);
529+
reduceNorm2 <<<dimgridRed, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>>>(d_norm2[dev], d_norm2aux[dev], total_pixels);
530530

531531
}
532532
for (dev = 0; dev < deviceCount; dev++){
@@ -537,7 +537,7 @@ do { \
537537
size_t dimgridRed = (total_pixels + MAXTHREADS - 1) / MAXTHREADS;
538538

539539
if (dimgridRed > 1) {
540-
reduceSum << <1, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device] >> >(d_norm2aux[dev], d_norm2[dev], dimgridRed);
540+
reduceSum <<<1, dimblockRed, MAXTHREADS*sizeof(float),stream[dev*nStream_device]>>>(d_norm2aux[dev], d_norm2[dev], dimgridRed);
541541
hipStreamSynchronize(stream[dev*nStream_device]);
542542
hipMemcpyAsync(&sumnorm2[dev], d_norm2[dev], sizeof(float), hipMemcpyDeviceToHost,stream[dev*nStream_device+1]);
543543
}

‎Common/CUDA/RandomNumberGenerator.cu

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,8 @@
4848
#include <stdlib.h>
4949
#include <hip/hip_runtime.h>
5050
#include <hiprand/hiprand_kernel.h>
51-
#include <hiprand.h>
51+
#include <hiprand/hiprand.h>
52+
#include <hiprand/hiprand.h>
5253

5354
#include "gpuUtils.hpp"
5455
#include "RandomNumberGenerator.hpp"

‎Common/CUDA/Siddon_projection.cu

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -230,16 +230,16 @@ __global__ void kernelPixelDetector( Geometry geo,
230230
float ac=am;
231231
//eq (28), unit anlges
232232
float axu,ayu,azu;
233-
axu=__frcp_rd(fabsf(ray.x));
234-
ayu=__frcp_rd(fabsf(ray.y));
235-
azu=__frcp_rd(fabsf(ray.z));
233+
axu=__frcp_rn(fabsf(ray.x));
234+
ayu=__frcp_rn(fabsf(ray.y));
235+
azu=__frcp_rn(fabsf(ray.z));
236236
// eq(29), direction of update
237237
float iu,ju,ku;
238238
iu=(source.x< pixel1D.x)? 1.0f : -1.0f;
239239
ju=(source.y< pixel1D.y)? 1.0f : -1.0f;
240240
ku=(source.z< pixel1D.z)? 1.0f : -1.0f;
241241

242-
float maxlength=__fsqrt_rd(ray.x*ray.x*geo.dVoxelX*geo.dVoxelX+ray.y*ray.y*geo.dVoxelY*geo.dVoxelY+ray.z*ray.z*geo.dVoxelZ*geo.dVoxelZ);
242+
float maxlength=__fsqrt_rn(ray.x*ray.x*geo.dVoxelX*geo.dVoxelX+ray.y*ray.y*geo.dVoxelY*geo.dVoxelY+ray.z*ray.z*geo.dVoxelZ*geo.dVoxelZ);
243243
float sum=0.0f;
244244
unsigned long Np=(imax-imin+1)+(jmax-jmin+1)+(kmax-kmin+1); // Number of intersections
245245
// Go iterating over the line, intersection by intersection. If double point, no worries, 0 will be computed
@@ -601,7 +601,7 @@ void CreateTexture(const GpuIds& gpuids,const float* imagedata,Geometry geo,hipA
601601
//hipArray Descriptor
602602
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
603603
//cuda Array
604-
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent);
604+
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent, 0);
605605
}
606606
}
607607
for (unsigned int dev = 0; dev < num_devices; dev++){

‎Common/CUDA/Siddon_projection_parallel.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -491,7 +491,7 @@ void CreateTextureParallel(float* image,Geometry geo,hipArray** d_cuArrTex, hipT
491491
//hipArray Descriptor
492492
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
493493
//cuda Array
494-
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent);
494+
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent, 0);
495495

496496

497497
hipMemcpy3DParms copyParams = {0};

‎Common/CUDA/ray_interpolated_projection.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -162,7 +162,7 @@ template<bool sphericalrotation>
162162
P.z=(uvOrigin.z+pixelU*deltaU.z+pixelV*deltaV.z);
163163

164164
// Length is the ray length in normalized space
165-
float length=__fsqrt_rd((source.x-P.x)*(source.x-P.x)+(source.y-P.y)*(source.y-P.y)+(source.z-P.z)*(source.z-P.z));
165+
float length=__fsqrt_rn((source.x-P.x)*(source.x-P.x)+(source.y-P.y)*(source.y-P.y)+(source.z-P.z)*(source.z-P.z));
166166
//now legth is an integer of Nsamples that are required on this line
167167
length=ceilf(__fdividef(length,geo.accuracy));//Divide the directional vector by an integer
168168
vectX=__fdividef(P.x -source.x,length);
@@ -561,7 +561,7 @@ void CreateTextureInterp(const GpuIds& gpuids,const float* imagedata,Geometry ge
561561

562562
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
563563
//cuda Array
564-
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent);
564+
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent, 0);
565565
cudaCheckErrors("Texture memory allocation fail");
566566
}
567567

‎Common/CUDA/ray_interpolated_projection_parallel.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -419,7 +419,7 @@ void CreateTextureParallelInterp(float* image,Geometry geo,hipArray** d_cuArrTex
419419
//hipArray Descriptor
420420
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
421421
//cuda Array
422-
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent);
422+
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent, 0);
423423

424424

425425
hipMemcpy3DParms copyParams = {0};
@@ -447,4 +447,4 @@ void CreateTextureParallelInterp(float* image,Geometry geo,hipArray** d_cuArrTex
447447
texDescr.readMode = hipReadModeElementType;
448448
hipCreateTextureObject(&texImage[0], &texRes, &texDescr, NULL);
449449

450-
}
450+
}

‎Common/CUDA/voxel_backprojection.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -247,7 +247,7 @@ __global__ void kernelPixelBackprojectionFDK(const Geometry geo, float* image,co
247247

248248
weight=__fdividef(DSO+realy*sinalpha-realx*cosalpha,DSO);
249249

250-
weight=__frcp_rd(weight*weight);
250+
weight=__frcp_rn(weight*weight);
251251

252252
// Get Value in the computed (U,V) and multiply by the corresponding weight.
253253
// indAlpha is the ABSOLUTE number of projection in the projection array (NOT the current number of projection set!)
@@ -680,7 +680,7 @@ void CreateTexture(const GpuIds& gpuids, float* projectiondata,Geometry geo,hipA
680680
//hipArray Descriptor
681681
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
682682
//cuda Array
683-
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent);
683+
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent, 0);
684684

685685
}
686686
}

‎Common/CUDA/voxel_backprojection2.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -272,7 +272,7 @@ __global__ void kernelPixelBackprojection(const Geometry geo, float* image,const
272272
realD.y=-realDaux.x*sinalpha + realDaux.y*cosalpha; //sin(-x)=-sin(x) , cos(-x)=cos(x)
273273
float L,lsq;
274274

275-
L = __fsqrt_rd( (realS.x-realD.x)*(realS.x-realD.x)+ (realS.y-realD.y)*(realS.y-realD.y)+ (realD.z)*(realD.z)); // Sz=0 always.
275+
L = __fsqrt_rn( (realS.x-realD.x)*(realS.x-realD.x)+ (realS.y-realD.y)*(realS.y-realD.y)+ (realD.z)*(realD.z)); // Sz=0 always.
276276
lsq = (realS.x-realvoxel.x)*(realS.x-realvoxel.x)
277277
+ (realS.y-realvoxel.y)*(realS.y-realvoxel.y)
278278
+ (realS.z-realvoxel.z)*(realS.z-realvoxel.z);
@@ -665,7 +665,7 @@ void CreateTexture2(const GpuIds& gpuids, float* projectiondata,Geometry geo,hip
665665
//hipArray Descriptor
666666
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
667667
//cuda Array
668-
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent);
668+
hipMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent, 0);
669669

670670
}
671671
}

‎Common/CUDA/voxel_backprojection_parallel.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -595,7 +595,7 @@ void CreateTextureParallel(float* projectiondata,Geometry geo,hipArray** d_cuArr
595595
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
596596
//cuda Array
597597
if (alloc){
598-
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent);
598+
hipMalloc3DArray(&d_cuArrTex[0], &channelDesc, extent, 0);
599599
cudaCheckErrors("Texture memory allocation fail");
600600
}
601601
hipMemcpy3DParms copyParams = {0};
@@ -625,4 +625,4 @@ void CreateTextureParallel(float* projectiondata,Geometry geo,hipArray** d_cuArr
625625
hipCreateTextureObject(&texImage[0], &texRes, &texDescr, NULL);
626626
cudaCheckErrors("Texture object creation fail");
627627

628-
}
628+
}

0 commit comments

Comments
 (0)