Skip to content

Commit

Permalink
ocl: minor fixes and adjustments
Browse files Browse the repository at this point in the history 
* Adopt best-matching device even for device-0 if ACC_OPENCL_ACTIVATE is not set or no device is activated.
* Fix: allow zero-sized allocation (c_dbcsr_acc_dev_mem_allocate).
* Improved finding OpenCL header file (CL/cl.h) in Makefile.
* Refactored initialization into init and pre-init.
* Minor rework of environment setup.
* Updated tuned parameters.
* Improved verbose output.
  • Loading branch information
@hfp
hfp committed Jan 6, 2025
1 parent 76ccc19 commit 78bfb48
Show file tree
Hide file tree
Showing 8 changed files with 289 additions and 237 deletions.
18 changes: 12 additions & 6 deletions src/acc/opencl/Makefile
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,10 +22,12 @@ endif
UNAME := $(shell uname)
HEADERONLY ?= 0
STATIC ?= 1
INTEL ?= 0
GNU ?= 0
DEV ?= 0

# Intel Compiler
ICX := $(shell which icx 2>/dev/null)
INTEL ?= $(if $(ICX),$(if $(filter-out 0,$(GNU)),0,2),0)

# select from set of predefined triplet specifications
SPECID ?= 0
# limit shape in tests (zero or negative for unlimited)
Expand Down Expand Up @@ -200,11 +202,15 @@ else
endif
else ifeq (,$(OPENCL_INC))
ifneq (,$(wildcard $(OPENCL_ROOT)/include/CL/cl.h))
LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
OPENCL_INC := $(OPENCL_ROOT)/include
LDFLAGS += -L$(OPENCL_ROOT)/lib64
else ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
OPENCL_INC := $(OPENCL_ROOT)/include/sycl
LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
else ifneq (,$(ICX))
OPENCL_ROOT := $(abspath $(dir $(ICX))/..)
ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
OPENCL_INC := $(OPENCL_ROOT)/include/sycl
endif
endif
endif
# OPENCL_INC: directory containing CL/cl.h.
Expand Down
267 changes: 149 additions & 118 deletions src/acc/opencl/acc_opencl.c
View file

Large diffs are not rendered by default.

4 changes: 4 additions & 0 deletions src/acc/opencl/acc_opencl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -338,8 +338,12 @@ typedef struct c_dbcsr_acc_opencl_config_t {
c_dbcsr_acc_opencl_timer_t timer; /* c_dbcsr_acc_opencl_device_t? */
/** Kernel-parameters are matched against device's UID */
cl_uint devmatch;
/** Split devices into sub-devices (if possible) */
cl_int devsplit;
/** Verbosity level (output on stderr). */
cl_int verbosity;
/** Guessed number of ranks per node (local), and rank-ID. */
cl_int nranks, nrank;
/** Non-zero if library is initialized (negative: no device). */
cl_int ndevices;
/** Maximum number of threads (omp_get_max_threads). */
Expand Down
141 changes: 75 additions & 66 deletions src/acc/opencl/acc_opencl_mem.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -189,7 +189,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
}
# endif
# if defined(ACC_OPENCL_XHINTS)
if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
host_ptr = malloc(nbytes);
if (NULL != host_ptr) flags = CL_MEM_USE_HOST_PTR;
}
Expand All @@ -202,7 +202,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
: c_dbcsr_acc_opencl_stream_default());
mapped = clEnqueueMapBuffer(str->queue, memory, CL_TRUE /*always block*/,
# if defined(ACC_OPENCL_XHINTS) && (defined(CL_VERSION_1_2) || defined(CL_MAP_WRITE_INVALIDATE_REGION))
(4 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
(8 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
# endif
(CL_MAP_READ | CL_MAP_WRITE),
0 /*offset*/, nbytes, 0, NULL, NULL, &result);
Expand Down Expand Up @@ -250,7 +250,7 @@ int c_dbcsr_acc_host_mem_deallocate(void* host_mem, void* stream) {
void* host_ptr = NULL;
int result_release;
# if defined(ACC_OPENCL_XHINTS)
if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) &&
if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) &&
(0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS)) &&
EXIT_SUCCESS == clGetMemObjectInfo(info.memory, CL_MEM_HOST_PTR, sizeof(void*), &host_ptr, NULL) && NULL != host_ptr)
{
Expand Down Expand Up @@ -343,79 +343,88 @@ int c_dbcsr_acc_dev_mem_allocate(void** dev_mem, size_t nbytes) {
}
# endif
assert(NULL != dev_mem && NULL != context);
if (0 != nbytes) {
# if defined(ACC_OPENCL_MEM_DEVPTR)
if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
assert(0 == c_dbcsr_acc_opencl_config.device.unified);
*dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
if (EXIT_SUCCESS != result) *dev_mem = NULL;
}
else
# endif
{
const int devuid = c_dbcsr_acc_opencl_config.device.uid;
const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
(0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
? 0
: (1u << 22));
memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
assert(0 == c_dbcsr_acc_opencl_config.device.unified);
*dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
if (EXIT_SUCCESS != result) *dev_mem = NULL;
}
if (EXIT_SUCCESS == result) {
# if defined(ACC_OPENCL_MEM_DEVPTR)
const c_dbcsr_acc_opencl_stream_t* str = NULL;
static cl_kernel kernel = NULL;
const size_t size = 1;
ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
/* determine device-side value of device-memory object by running some kernel */
assert(NULL != memory && NULL != str && NULL != str->queue);
if (NULL == kernel) { /* generate kernel */
const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
" const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
" const size_t i = get_global_id(0);\n"
" ptr[i] = cast.u + i;\n"
"}\n";
assert(sizeof(size_t) == sizeof(cl_ulong));
result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
}
/* TODO: backup/restore memory */
if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
if (EXIT_SUCCESS == result) {
result = clEnqueueNDRangeKernel(
str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
}
if (EXIT_SUCCESS == result) {
result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
else
# endif
{
const int devuid = c_dbcsr_acc_opencl_config.device.uid;
const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
(0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
? 0
: (1u << 22));
memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
}
assert(EXIT_SUCCESS != result || NULL != memptr);
if (EXIT_SUCCESS == result) {
c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
assert(NULL != memptr);
if (NULL != info) {
info->memory = memory;
info->memptr = memptr;
*dev_mem = memptr;
# if defined(ACC_OPENCL_MEM_DEVPTR)
const c_dbcsr_acc_opencl_stream_t* str = NULL;
static cl_kernel kernel = NULL;
const size_t size = 1;
ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
/* determine device-side value of device-memory object by running some kernel */
assert(NULL != memory && NULL != str && NULL != str->queue);
if (NULL == kernel) { /* generate kernel */
const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
" const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
" const size_t i = get_global_id(0);\n"
" ptr[i] = cast.u + i;\n"
"}\n";
assert(sizeof(size_t) == sizeof(cl_ulong));
result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
}
else result = EXIT_FAILURE;
}
ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
/* TODO: backup/restore memory */
if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
if (EXIT_SUCCESS == result) {
result = clEnqueueNDRangeKernel(
str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
}
if (EXIT_SUCCESS == result) {
result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
}
assert(EXIT_SUCCESS != result || NULL != memptr);
if (EXIT_SUCCESS == result) {
c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
assert(NULL != memptr);
if (NULL != info) {
info->memory = memory;
info->memptr = memptr;
*dev_mem = memptr;
}
else result = EXIT_FAILURE;
}
ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
# else
*dev_mem = memptr = memory;
*dev_mem = memptr = memory;
# endif
}
if (EXIT_SUCCESS != result) {
if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
*dev_mem = NULL;
}
}
if (EXIT_SUCCESS != result) {
if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
*dev_mem = NULL;
if (0 != c_dbcsr_acc_opencl_config.verbosity) {
if (EXIT_SUCCESS == result && 0 != c_dbcsr_acc_opencl_config.debug) {
fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu successfully allocated\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
else if (EXIT_SUCCESS != result) {
fprintf(stderr, "ERROR ACC/OpenCL: memory=%p pointer=%p size=%llu failed to allocate\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
}
}
if (0 != c_dbcsr_acc_opencl_config.debug && 0 != c_dbcsr_acc_opencl_config.verbosity && EXIT_SUCCESS == result) {
fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu allocated\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
else *dev_mem = NULL;
assert(EXIT_SUCCESS == result || NULL == *dev_mem);
# if defined(__DBCSR_ACC) && defined(ACC_OPENCL_PROFILE)
c_dbcsr_timestop(&routine_handle);
Expand Down
2 changes: 1 addition & 1 deletion src/acc/opencl/acc_opencl_stream.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -118,7 +118,7 @@ int c_dbcsr_acc_stream_create(void** stream_p, const char* name, int priority) {
# endif
{
# if defined(ACC_OPENCL_XHINTS)
if ((2 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
if ((4 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
struct {
cl_command_queue_properties properties;
cl_bitfield capabilities;
Expand Down
24 changes: 13 additions & 11 deletions src/acc/opencl/smm/opencl_libsmm.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -451,20 +451,24 @@ int libsmm_acc_init(void) {
}
# if defined(OPENCL_KERNELS_PARAMS_SMM) && defined(OPENCL_KERNELS_DEVICES)
if (EXIT_SUCCESS == result && (0 == ntuned || 0 != key_direct_skip)) {
unsigned int default_uid = c_dbcsr_acc_opencl_config.device.uid;
const char *line = OPENCL_KERNELS_PARAMS_SMM, *next;
# if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
int active_match = -1;
if (EXIT_SUCCESS == c_dbcsr_acc_opencl_device_name(c_dbcsr_acc_opencl_config.device.id, bufname, ACC_OPENCL_BUFFERSIZE,
NULL /*platform*/, 0 /*platform_maxlen*/, /*cleanup*/ 1))
{ /* determine best-matching parameters based on name of device */
int i = 0, best = 0;
if (1 >= c_dbcsr_acc_opencl_config.devmatch) {
c_dbcsr_acc_opencl_device_uid(c_dbcsr_acc_opencl_config.device.id, bufname, &default_uid);
}
for (; i < ndevices_params; ++i) {
const int score = libxsmm_strimatch(bufname, OPENCL_KERNELS_DEVICES[i], NULL);
unsigned int uid;
if (best < score ||
((best == score) &&
EXIT_SUCCESS == c_dbcsr_acc_opencl_device_uid(NULL /*device*/, OPENCL_KERNELS_DEVICES[i], &uid) &&
uid == c_dbcsr_acc_opencl_config.device.uid))
uid == default_uid))
{
active_match = i;
best = score;
Expand All @@ -489,22 +493,20 @@ int libsmm_acc_init(void) {
{
key.devuid = 0;
}
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key)); /* duplicate? */
if (NULL == config_init) {
if (NULL == libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
if (NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) ++ntuned;
else { /* failed to register */
result = EXIT_FAILURE;
break;
}
else ++ntuned;
}
else if (config_init->gflops < config.gflops) { /* update */
memcpy(config_init, &config, sizeof(config));
}
# if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
if (active_match == i && 0 != c_dbcsr_acc_opencl_config.device.uid &&
c_dbcsr_acc_opencl_config.device.uid != key.devuid)
{
key.devuid = c_dbcsr_acc_opencl_config.device.uid;
if (active_match == i && 0 != default_uid && default_uid != key.devuid) {
key.devuid = default_uid;
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
if (NULL == config_init && NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
static int info = 0;
Expand Down Expand Up @@ -532,7 +534,7 @@ int libsmm_acc_init(void) {
}
# endif
# if defined(OPENCL_KERNELS_DEVICES)
if (EXIT_SUCCESS == result && 0 != ntuned &&
if (EXIT_SUCCESS == result && 0 != ntuned && 0 == c_dbcsr_acc_opencl_config.nrank &&
(2 <= c_dbcsr_acc_opencl_config.verbosity || 0 > c_dbcsr_acc_opencl_config.verbosity))
{
fprintf(stderr, "INFO ACC/LIBSMM: PARAMS in %i set%s loaded targeting ", ntuned, 1 != ntuned ? "s" : "");
Expand Down Expand Up @@ -1084,7 +1086,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
(NULL == env_nz || '\0' == *env_nz) ? (0 != defaults ? /*default*/ 0 : config->nz) : atoi(env_nz), 0, 1);
new_config.al = LIBXSMM_CLMP(/* bug: AL=1 */
(NULL == env_al || '\0' == *env_al)
? (0 == (64 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
? (0 == (32 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
: atoi(env_al),
0, 1);
new_config.tb = LIBXSMM_CLMP(
Expand All @@ -1095,7 +1097,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
(NULL == env_ap || '\0' == *env_ap) ? (0 != defaults ? /*default*/ 0 : config->ap) : atoi(env_ap), 0, 1);
new_config.aa = LIBXSMM_CLMP(/* bug: AA=2 XF=1 */
(NULL == env_aa || '\0' == *env_aa) ? (0 != defaults ? default_aa : config->aa) : atoi(env_aa), 0,
(0 == (32 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
(0 == (16 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
new_config.ab = LIBXSMM_CLMP(
(NULL == env_ab || '\0' == *env_ab) ? (0 != defaults ? default_ab : config->ab) : atoi(env_ab), 0, 2);
new_config.ac = LIBXSMM_CLMP(
Expand Down
8 changes: 4 additions & 4 deletions src/acc/opencl/smm/params/tune_multiply_Mi250.csv
View file
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ gfx90a [0x989f];3;8;5;8;30000;0;15;8;1;1;1;1;-2;0;0;0;1;0;0;0;0;0
gfx90a [0x989f];3;8;6;6;30000;0;12;8;1;8;1;-2;-2;0;0;0;1;0;2;0;0;0
gfx90a [0x989f];3;8;6;8;30000;0;15;8;1;1;1;-2;-1;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;7;7;30000;0;13;8;1;5;1;1;-2;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;4;1;-2;0;0;0;0;1;1;0;2;1;0
gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;6;1;-1;-2;0;0;0;1;0;0;0;0;0
gfx90a [0x989f];3;8;8;3;30000;0;15;8;1;1;1;1;-1;0;0;0;1;0;2;0;0;0
gfx90a [0x989f];3;8;8;4;30000;0;15;8;1;7;1;0;-1;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;8;5;30000;0;15;8;1;8;1;0;-1;0;0;0;1;0;2;2;0;0
Expand Down Expand Up @@ -212,7 +212,7 @@ gfx90a [0x989f];3;9;9;23;30000;0;15;9;1;1;1;-1;-1;0;1;1;1;1;2;2;0;0
gfx90a [0x989f];3;9;16;9;30000;0;24;9;1;8;1;0;-2;0;0;0;1;0;2;2;0;0
gfx90a [0x989f];3;9;16;16;30000;0;12;9;1;1;1;1;-2;0;0;0;1;0;2;2;0;0
gfx90a [0x989f];3;9;23;9;30000;0;26;9;1;8;1;0;-2;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;8;1;1;3;0;1;0;1;0;1;2;1;0
gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;5;1;0;2;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;10;3;3;30000;0;10;10;1;8;1;0;-2;0;1;0;1;1;2;0;0;0
gfx90a [0x989f];3;10;3;10;30000;0;12;10;1;1;1;-1;-2;0;0;1;1;1;2;0;0;0
gfx90a [0x989f];3;10;4;4;30000;0;13;10;1;3;1;-2;-1;1;1;0;1;1;0;2;0;0
Expand Down Expand Up @@ -249,7 +249,7 @@ gfx90a [0x989f];3;11;11;3;30000;0;22;11;1;2;1;-1;1;0;1;0;1;1;2;0;1;0
gfx90a [0x989f];3;11;11;4;30000;0;20;11;1;2;1;-1;-2;0;1;1;1;1;0;0;0;0
gfx90a [0x989f];3;11;11;5;30000;0;20;11;1;5;1;-2;1;0;1;1;1;1;2;0;1;0
gfx90a [0x989f];3;11;11;6;30000;0;24;11;1;1;1;-2;0;0;0;0;1;1;2;2;0;0
gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;7;1;0;1;0;1;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;9;1;1;1;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;11;30000;0;12;11;1;10;1;-2;-2;0;1;0;1;1;2;0;0;0
gfx90a [0x989f];3;11;11;16;30000;0;3;8;1;11;1;1;-1;1;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;23;30000;0;15;11;1;1;1;1;-1;0;1;1;1;1;2;2;0;0
Expand Down Expand Up @@ -351,7 +351,7 @@ gfx90a [0x989f];3;18;23;18;30000;0;4;18;1;4;1;1;3;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;18;23;23;30000;0;4;8;1;16;1;0;3;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;19;19;19;30000;0;40;8;1;10;1;1;3;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;19;19;23;30000;0;40;8;1;15;1;-1;-2;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;-1;3;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;1;0;0;1;0;1;0;1;0;0;0
gfx90a [0x989f];3;28;28;28;30000;0;3;28;1;28;28;-2;2;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;32;32;32;30000;0;25;32;1;20;1;-2;0;0;1;0;1;0;2;0;0;0
gfx90a [0x989f];3;35;17;17;30000;0;15;35;1;29;1;1;0;0;1;0;1;0;2;1;0;0
Expand Down
Loading

0 comments on commit 78bfb48

Please sign in to comment.