diff --git a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp index f492643e41..f0ae0f582f 100644 --- a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp @@ -216,11 +216,11 @@ __global__ void testKernel(hipLaunchParm lp,int *g_odata) // Atomic increment (modulo 17+1) //atomicInc((unsigned int *)&g_odata[5], 17); - atomicInc((unsigned int *)&g_odata[5]); + //atomicInc((unsigned int *)&g_odata[5]); // Atomic decrement // atomicDec((unsigned int *)&g_odata[6], 137); - atomicDec((unsigned int *)&g_odata[6]); + //atomicDec((unsigned int *)&g_odata[6]); // Atomic compare-and-swap atomicCAS(&g_odata[7], tid-1, tid); diff --git a/projects/hip/tests/src/hip_anyall.cpp b/projects/hip/tests/src/hip_anyall.cpp index c6ac36e2e3..52a2a13db9 100644 --- a/projects/hip/tests/src/hip_anyall.cpp +++ b/projects/hip/tests/src/hip_anyall.cpp @@ -28,22 +28,25 @@ THE SOFTWARE. #define HIP_ASSERT(x) (assert((x)==hipSuccess)) __global__ void - warpvote(hipLaunchParm lp, int* device_any, int* device_all , int Num_Warps_per_Block) + warpvote(hipLaunchParm lp, int* device_any, int* device_all , int Num_Warps_per_Block, int pshift) { int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; - device_any[hipThreadIdx_x>>6] = __any(tid -77); - device_all[hipThreadIdx_x>>6] = __all(tid -77); + device_any[hipThreadIdx_x>>pshift] = __any(tid -77); + device_all[hipThreadIdx_x>>pshift] = __all(tid -77); } int main(int argc, char *argv[]) -{ int warpSize; +{ int warpSize, pshift; hipDeviceProp_t devProp; hipDeviceGetProperties(&devProp, 0); - if(strncmp(devProp.name,"Fiji",1)==0) warpSize =64; - else warpSize =32; + if(strncmp(devProp.name,"Fiji",1)==0) +{ warpSize =64; + pshift =6; +} + else {warpSize =32; pshift=5;} int anycount =0; int allcount =0; int Num_Threads_per_Block = 1024; @@ -65,7 +68,7 @@ for (int i=0; i> 6; + const unsigned int warp_num = hipThreadIdx_x >> pshift; atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popcll(__ballot(tid - 245))); } int main(int argc, char *argv[]) -{ int warpSize; +{ int warpSize, pshift; hipDeviceProp_t devProp; hipDeviceGetProperties(&devProp, 0); - if(strncmp(devProp.name,"Fiji",1)==0) warpSize =64; - else warpSize =32; + if(strncmp(devProp.name,"Fiji",1)==0) {warpSize =64; pshift =6;} + else {warpSize =32; pshift =5;} unsigned int Num_Threads_per_Block = 512; unsigned int Num_Blocks_per_Grid = 1; unsigned int Num_Warps_per_Block = Num_Threads_per_Block/warpSize; @@ -33,7 +33,7 @@ int main(int argc, char *argv[]) HIP_ASSERT(hipMemcpy(device_ballot, host_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyHostToDevice)); - hipLaunchKernel(gpu_ballot, dim3(Num_Blocks_per_Grid),dim3(Num_Threads_per_Block),0,0, device_ballot,Num_Warps_per_Block); + hipLaunchKernel(gpu_ballot, dim3(Num_Blocks_per_Grid),dim3(Num_Threads_per_Block),0,0, device_ballot,Num_Warps_per_Block,pshift); HIP_ASSERT(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyDeviceToHost)); @@ -45,7 +45,7 @@ int main(int argc, char *argv[]) divergent_count++;} } -if (divergent_count==1) printf("PASSED"); else printf("FAILED"); +if (divergent_count==1) printf("PASSED\n"); else printf("FAILED\n"); return EXIT_SUCCESS; }