AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

[dests] Fix hipTestClock, hipTestNew & hipTestGlobalVariable tests for nvcc

Telusuri Sumber 
nvcc does not support global kernels in struct/class

Change-Id: I2d7297e0c3725564215e20dbdd31c0bb8d7a07de


[ROCm/clr commit: f63ffaf6e5]
This commit is contained in:
Maneesh Gupta
2018-09-17 15:32:05 +05:30
orang tua 6e7f871105
melakukan 0bae7eaad3
3 mengubah file dengan 8 tambahan dan 19 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp
+1 -4
Melihat File
@@ -33,8 +33,6 @@ THE SOFTWARE.
#define LEN 512 #define LEN 512
#define SIZE 2048 #define SIZE 2048
struct TestClock {
static __global__ void kernel1(int* Ad) { static __global__ void kernel1(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x; int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = clock() + clock64() + __clock() + __clock64(); Ad[tid] = clock() + clock64() + __clock() + __clock64();
@@ -61,9 +59,8 @@ struct TestClock {
assert(0 != A[i]); assert(0 != A[i]);
} }
} }
};
int main() { int main() {
TestClock().run(); run();
passed(); passed();
} }
projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp
+1 -3
Melihat File
@@ -33,7 +33,6 @@ THE SOFTWARE.
#define LEN 512 #define LEN 512
#define SIZE 2048 #define SIZE 2048
struct TestPlacementNew {
class A { class A {
public: public:
__device__ A() { __device__ A() {
@@ -63,9 +62,8 @@ struct TestPlacementNew {
assert(i == A[i]); assert(i == A[i]);
} }
} }
};
int main() { int main() {
TestPlacementNew().run(); run();
passed(); passed();
} }
projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp
+6 -12
Melihat File
@@ -33,15 +33,14 @@ THE SOFTWARE.
#define LEN 512 #define LEN 512
#define SIZE 2048 #define SIZE 2048
struct TestConstantGlobalVar { __constant__ int ConstantGlobalVar = 123;
static __constant__ int ConstantGlobalVar;
static __global__ void kernel(int* Ad) { static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x; int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = ConstantGlobalVar; Ad[tid] = ConstantGlobalVar;
} }
void run() { void runTestConstantGlobalVar() {
int *A, *Ad; int *A, *Ad;
A = new int[LEN]; A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) { for (unsigned i = 0; i < LEN; i++) {
@@ -56,11 +55,8 @@ struct TestConstantGlobalVar {
assert(123 == A[i]); assert(123 == A[i]);
} }
} }
};
__constant__ int TestConstantGlobalVar::ConstantGlobalVar = 123;
struct TestGlobalArray { __device__ int GlobalArray[LEN];
static __device__ int GlobalArray[LEN];
static __global__ void kernelWrite() { static __global__ void kernelWrite() {
int tid = threadIdx.x + blockIdx.x * blockDim.x; int tid = threadIdx.x + blockIdx.x * blockDim.x;
@@ -71,7 +67,7 @@ struct TestGlobalArray {
Ad[tid] = GlobalArray[tid]; Ad[tid] = GlobalArray[tid];
} }
void run() { void runTestGlobalArray() {
int *A, *Ad; int *A, *Ad;
A = new int[LEN]; A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) { for (unsigned i = 0; i < LEN; i++) {
@@ -87,11 +83,9 @@ struct TestGlobalArray {
assert(i == A[i]); assert(i == A[i]);
} }
} }
};
__device__ int TestGlobalArray::GlobalArray[LEN];
int main() { int main() {
TestConstantGlobalVar().run(); runTestConstantGlobalVar();
TestGlobalArray().run(); runTestGlobalArray();
passed(); passed();
} }