') + ')', 'gi'); if (regex.test(text)) { found = true; var frag = document.createDocumentFragment(); var parts = text.split(regex); parts.forEach(function(part, i) { if (i % 2 === 0) { frag.appendChild(document.createTextNode(part)); } else { var span = document.createElement('span'); span.className = 'userscript-highlight'; span.textContent = part; frag.appendChild(span); } }); node.parentNode.replaceChild(frag, node); } }); } else if (node.nodeType === 1 && node.childNodes) { // element var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT']; if (!skipTags.includes(node.tagName)) { Array.from(node.childNodes).forEach(highlight); } } } highlight(document.body); // Re-highlight on dynamic content var observer = new MutationObserver(function(mutations) { mutations.forEach(function(m) { m.addedNodes.forEach(function(node) { if (node.nodeType === 1 || node.nodeType === 3) highlight(node); }); }); }); observer.observe(document.body, { childList: true, subtree: true }); })(); } } catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); } })(); (function(){ try { var __m = "*"; var __re = new RegExp('^' + ".*" + ', 'i'); if (__m === '*' || __re.test(location.href)) { // Strip utm_, fbclid, gclid, etc. from all links on page (function() { var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content', 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid', 'ref', 'ref_src', 'source', 'medium', 'campaign']; function cleanUrl(url) { try { var u = new URL(url, window.location.origin); var changed = false; trackingParams.forEach(function(p) { if (u.searchParams.has(p)) { u.searchParams.delete(p); changed = true; } }); return changed ? u.toString() : url; } catch (e) { return url; } } function cleanLinks() { document.querySelectorAll('a[href]').forEach(function(a) { var clean = cleanUrl(a.href); if (clean !== a.href) a.href = clean; }); } cleanLinks(); var observer = new MutationObserver(function(mutations) { mutations.forEach(function(m) { m.addedNodes.forEach(function(node) { if (node.nodeType === 1) { if (node.tagName === 'A') cleanLinks(); node.querySelectorAll('a[href]').forEach(function(a) { var clean = cleanUrl(a.href); if (clean !== a.href) a.href = clean; }); } }); }); }); observer.observe(document.body, { childList: true, subtree: true }); })(); } } catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); } })(); (function(){ try { var __m = "youtube.com"; var __re = new RegExp('^' + "youtube\\.com" + ', 'i'); if (__m === '*' || __re.test(location.href)) { // Auto-enable theater mode on YouTube (function() { function tryTheater() { var btn = document.querySelector('button[aria-label="Theater mode"], ytd-player #player button[title="Theater mode"]'); if (btn && !btn.classList.contains('activated')) { btn.click(); } } // Try immediately tryTheater(); // Try after navigation (SPA) var lastUrl = location.href; setInterval(function() { if (location.href !== lastUrl) { lastUrl = location.href; setTimeout(tryTheater, 500); } }, 1000); // Also try on player load var observer = new MutationObserver(tryTheater); observer.observe(document.body, { childList: true, subtree: true }); })(); } } catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); } })(); (function(){ try { var __m = "*"; var __re = new RegExp('^' + ".*" + ', 'i'); if (__m === '*' || __re.test(location.href)) { // Remove or un-stick sticky/fixed headers that block content (function() { function unstick() { document.querySelectorAll('header, nav, [role="banner"], .header, .navbar, .sticky, .fixed-top, [style*="position: fixed"], [style*="position:sticky"]').forEach(function(el) { if (el.style.position === 'fixed' || el.style.position === 'sticky' || getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') { el.style.position = 'static'; el.style.top = 'auto'; el.style.zIndex = 'auto'; } }); } unstick(); var observer = new MutationObserver(unstick); observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] }); })(); } } catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); } })(); })(); Cuda and OpenCL fixes by LourensVeen · Pull Request #8 · treecode/sapporo2 · GitHub
Skip to content

Cuda and OpenCL fixes#8

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
LourensVeen wants to merge 3 commits into
base:master
Choose a base branch
from
Open

Cuda and OpenCL fixes #8

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
005f1c2
Fix error when building for newer GPU architectures
LourensVeen Oct 31, 2023
0263a6d
Fix build for newer CUDA
LourensVeen Oct 30, 2023
2cc55c3
Fix errors in OpenCL compilation
LourensVeen Oct 31, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion lib/CUDAKernels/kernels.cu
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -206,7 +206,7 @@ __device__ __forceinline__ double RSQRT(double val) { return rsqrt(val); }
// template<> __device__ __forceinline__ double RSQRT(double val) { return 1.0/sqrt(val); }



#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull =
Expand All@@ -220,6 +220,7 @@ __device__ double atomicAdd(double* address, double val)
} while (assumed != old);
return __longlong_as_double(old);
}
#endif


__device__ __forceinline__ double atomicMin(double *address, double val)
Expand Down
2 changes: 1 addition & 1 deletion lib/Makefile
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -41,7 +41,7 @@ NVCCVERSION=$(shell "${NVCC}" --version | grep ^Cuda | sed 's/^.* //g')
ifeq "${NVCCVERSION}" "V5.5.22"
NVCCFLAGS ?= -arch sm_20
else
NVCCFLAGS ?= -arch sm_30
NVCCFLAGS ?= -arch sm_50
endif

#NVCCFLAGS = -arch sm_35
Expand Down
4 changes: 3 additions & 1 deletion lib/include/cudadev.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -710,12 +710,14 @@ namespace dev {
// jitOptionCount++;
// }



#if CUDA_VERSION < 6000
if(computeMode < CU_TARGET_COMPUTE_20)
{
fprintf(stderr,"Sapporo2 requires at least a Fermi or newer NVIDIA architecture.\n");
exit(-1);
}
#endif

//Set the architecture
// {
Expand Down
9 changes: 8 additions & 1 deletion lib/include/defines.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -55,8 +55,15 @@ inline const char* get_kernelName(const int integrator,
case SIXTH:
if(precision == DOUBLESINGLE)
{
#ifdef _OCL_
fprintf(stderr, "ERROR: Sixth order integrator with double single precision");
fprintf(stderr, "ERROR: is not implemented in OpenCL, only in CUDA. Please");
fprintf(stderr, "ERROR: file an issue on GitHub if you need this combination.");
exit(1);
#else

@jbedorfjbedorfNov 5, 2023

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I guess there's no need for the else? Given that there's the exit above? In that case maybe change the #else into #endif

@LourensVeenLourensVeenNov 5, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The problem here actually was that float3 was undefined when compiling with OpenCL, causing a compiler error. So we could move the perThreadSM = ... line outside of the #ifdef, but then it wouldn't compile for OpenCL.

I think I later saw a header somewhere that aliases some OpenCL types to CUDA-like names, so maybe float3 can be added there to fix it instead, I'll have a look.

perThreadSM = sizeof(float4)*2 + sizeof(float4) + sizeof(float3);
return "dev_evaluate_gravity_sixth_DS";
#endif
return "dev_evaluate_gravity_sixth_DS";
}
else if(precision == DOUBLE){
#ifdef _OCL_
Expand Down
4 changes: 2 additions & 2 deletions lib/include/ocldev.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -574,8 +574,8 @@ namespace dev {
void copy(const memory &src, const cl_bool OCL_BLOCKING = CL_TRUE) {
assert(ContextFlag);
if (n != src.n) {
ocl_free();
cmalloc(src.n, DeviceMemFlags);
ocl_free();
allocate(src.n, DeviceMemFlags);

@jbedorfjbedorfNov 5, 2023

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, this looks to be the right thing todo.

}
oclSafeCall(clEnqueueCopyBuffer(CommandQueue,
src.DeviceMem,
Expand Down