Merge branch 'amd-develop' into amd-master

Change-Id: Ia1320caf8bf1f32f7377639b1b9db8fd26f23113

README.md

@@ -11,6 +11,20 @@ Key features include:
 
 New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms.  Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application.  HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port.
 
+## Repository branches:
+
+The HIP repository maintains several branches. The branches that are of importance are:
+
+* master branch: This is the stable branch. All stable releases are based on this branch.
+* developer-preview branch: This is the branch were the new features still under development are visible. While this maybe of interest to many, it should be noted that this branch and the features under development might not be stable.
+
+## Release tagging:
+
+HIP releases are typically of two types. The tag naming convention is different for both types of releases to help differentiate them.
+
+* release_x.yy.zzzz: These are the stable releases based on the master branch. This type of release is typically made once a month.
+* preview_x.yy.zzzz: These denote pre-release code and are based on the developer-preview branch. This type of release is typically made once a week.
+
 ## More Info:
 - [Installation](INSTALL.md) 
 - [HIP FAQ](docs/markdown/hip_faq.md)

docs/markdown/hip_kernel_language.md

@@ -678,5 +678,5 @@ The file format for binary is `.co` which means Code Object. The following comma
 [INPUT FILE] = Name of the file containing kernels
 [OUTPUT FILE] = Name of the generated code object file```
 
-
+Note that one important fact to remember when using binary code objects is that the number of arguments to the kernel are different on HCC and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel.
 

samples/0_Intro/module_api/Makefile

@@ -3,31 +3,20 @@ ifeq (,$(HIP_PATH))
 	HIP_PATH=../../..
 endif
 HIPCC=$(HIP_PATH)/bin/hipcc
-OPT=
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
-all: vcpy_kernel.compile runKernel.hip.out
+ifeq (${HIP_PLATFORM}, hcc)
+	GENCO_FLAGS=--target-isa=fiji
+endif
+
+all: vcpy_kernel.code runKernel.hip.out
 
 runKernel.hip.out: runKernel.cpp
-	$(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out
+	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
 
-ifeq (${HIP_PLATFORM}, hcc)
-
-vcpy_kernel.compile: vcpy_kernel.cpp
-	$(HIPCC) --genco --target-isa=fiji vcpy_kernel.cpp -o vcpy_kernel.co
+vcpy_kernel.code: vcpy_kernel.cpp
+	$(HIPCC) --genco $(GENCO_FLAGS) $< -o $@
 
 clean:
-	rm -f *.co *.out
-
-endif
-
-ifeq (${HIP_PLATFORM}, nvcc)
-
-vcpy_kernel.compile: vcpy_kernel.cpp
-	$(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx
-
-clean:
-	rm -f *.ptx *.out
-
-endif
+	rm -f *.code *.out
 

samples/0_Intro/module_api/runKernel.cpp

@@ -26,15 +26,8 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN<<2
 
-#ifdef __HIP_PLATFORM_HCC__ 
-#define fileName "vcpy_kernel.co"
+#define fileName "vcpy_kernel.code"
 #define kernel_name "hello_world"
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
-#define fileName "vcpy_kernel.ptx"
-#define kernel_name "hello_world"
-#endif
 
 int main(){
     float *A, *B;
@@ -47,14 +40,11 @@ int main(){
         B[i] = 0.0f;
     }
 
-
-#ifdef __HIP_PLATFORM_NVCC__
   	hipInit(0);
-	  hipDevice_t device;
-	  hipCtx_t context;
-	  hipDeviceGet(&device, 0);
-	  hipCtxCreate(&context, 0, device);
-#endif
+	hipDevice_t device;
+	hipCtx_t context;
+	hipDeviceGet(&device, 0);
+    hipCtxCreate(&context, 0, device);
 
     hipMalloc((void**)&Ad, SIZE);
     hipMalloc((void**)&Bd, SIZE);
@@ -107,9 +97,6 @@ int main(){
         std::cout<<A[i]<<" - "<<B[i]<<std::endl;
     }
 
-#ifdef __HIP_PLATFORM_NVCC__
-	  hipCtxDetach(context);
-#endif 
-
+    hipCtxDestroy(context);
     return 0;
 }