Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions."

This reverts commit 1c50968 [ROCm/hip commit: 32e11e7dc6]
2017-11-29 21:49:10 +00:00
@@ -29,7 +29,7 @@ THE SOFTWARE.
 const int NN = 1 << 21;

 __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
 			x[i] = sqrt(powf(3.14159,i));
@@ -39,7 +39,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 }

 __global__ void nKernel(hipLaunchParm lp, float *y){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	y[tid] = y[tid] + 1.0f;
 }

@@ -31,7 +31,7 @@ THE SOFTWARE.
 const int NN = 1 << 21;

 __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
 			x[i] = sqrt(powf(3.14159,i));
@@ -41,7 +41,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 }

 __global__ void nKernel(hipLaunchParm lp, float *y){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	y[tid] = y[tid] + 1.0f;
 }

@@ -41,8 +41,8 @@ vectorADDRepeat(hipLaunchParm lp,
            size_t NELEM,
            int repeat)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;

    for (int j=1; j<=repeat;j++) {
        for (size_t i=offset; i<NELEM; i+=stride) {
@@ -73,7 +73,7 @@ void D2H(T *Dst, T *Src, size_t size){

 template<typename T>
 __global__ void Inc(hipLaunchParm lp, T *In){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+int tx = threadIdx.x + blockIdx.x * blockDim.x;
 In[tx] = In[tx] + 1;
 }