From 7c132821a8ba3afc4e0f80fe1e2882b98d31feff Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 4 Sep 2016 20:37:29 +0530 Subject: [PATCH 1/4] Removed NVCC check for hipCtxXXX functions in module_api/runKernel.cpp Change-Id: I2bdd4fadf41063ec60626f1850e16f8307ebe6b5 --- samples/0_Intro/module_api/runKernel.cpp | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index 1c0dd9f349..8c1ed1598f 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/samples/0_Intro/module_api/runKernel.cpp @@ -47,14 +47,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 +104,6 @@ int main(){ std::cout< Date: Sun, 4 Sep 2016 21:15:49 +0530 Subject: [PATCH 2/4] Update README with branching and tagging information Change-Id: I2d801ef4f0abe6eb0280f86443bc55725701252a --- README.md | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/README.md b/README.md index f9f653f2ba..6d6c1b2553 100644 --- a/README.md +++ b/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) From 0cc51f2a7239d95f874ab4fca9535dff2b9457b8 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 21:25:14 +0530 Subject: [PATCH 3/4] module_api sample: Remove unnecessary platform checks Change-Id: I1d531264d51ff952a3a68d554672b6d293e23379 --- samples/0_Intro/module_api/Makefile | 29 ++++++++---------------- samples/0_Intro/module_api/runKernel.cpp | 9 +------- 2 files changed, 10 insertions(+), 28 deletions(-) diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile index bda0a7dd93..81e876ddbf 100644 --- a/samples/0_Intro/module_api/Makefile +++ b/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 diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index 8c1ed1598f..c6bba63b06 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/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; From dc5f2b944161e2fb638a213c68f0327f6b4cf1c0 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 21:39:28 +0530 Subject: [PATCH 4/4] hip_kernel_language.md: Document difference in arguments for binary code object Change-Id: I96e347e8582cbd1dbc3776fbafcb9e61563538db --- docs/markdown/hip_kernel_language.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 628cd36c26..ed9dc2fbe9 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/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.