From d95642b85c7d33439e4db26aa17476d7550e5fb3 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 [ROCm/clr commit: bb5f15e199c7c2359283c249eb49959114c02056] --- .../samples/0_Intro/module_api/runKernel.cpp | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp b/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp index 1c0dd9f349..8c1ed1598f 100644 --- a/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/clr/hipamd/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 [ROCm/clr commit: fa83be60049044d81bde4f679cc895095b1f93ed] --- projects/clr/hipamd/README.md | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index f9f653f2ba..6d6c1b2553 100644 --- a/projects/clr/hipamd/README.md +++ b/projects/clr/hipamd/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 de5701b678646b5a57a730d1b6e6565ebd921745 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 [ROCm/clr commit: 40a392800af6f4bccbe297405537a884f1eee7b2] --- .../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/projects/clr/hipamd/samples/0_Intro/module_api/Makefile b/projects/clr/hipamd/samples/0_Intro/module_api/Makefile index bda0a7dd93..81e876ddbf 100644 --- a/projects/clr/hipamd/samples/0_Intro/module_api/Makefile +++ b/projects/clr/hipamd/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/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp b/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp index 8c1ed1598f..c6bba63b06 100644 --- a/projects/clr/hipamd/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/clr/hipamd/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 a2ea4a894409586870e2335e8d3b4bb50b64b693 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 [ROCm/clr commit: 462a1d97b2e625e5c380d4ebd5c0a0a8bb1df8dc] --- projects/clr/hipamd/docs/markdown/hip_kernel_language.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 628cd36c26..ed9dc2fbe9 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/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.